/third_party/mindspore/mindspore/ccsrc/runtime/framework/actor/ |
D | output_actor.cc | 23 TensorPtr CreateOutputTensor(const AnfNodePtr &output_node, size_t output_index, size_t output_posi… in CreateOutputTensor() argument 26 << ", output index: " << output_index << ", output position: " << output_position; in CreateOutputTensor() 30 auto type_id = AnfAlgo::GetOutputInferDataType(output_node, output_index); in CreateOutputTensor() 32 auto shape = AnfAlgo::GetOutputInferShape(output_node, output_index); in CreateOutputTensor() 35 tensor->set_padding_type(AnfAlgo::GetOutputReshapeType(output_node, output_index)); in CreateOutputTensor() 38 const auto &device_tensor = AnfAlgo::GetMutableOutputAddr(output_node, output_index, false); in CreateOutputTensor() 86 auto output_index = output_nodes_[i].second; in UpdateOutputDeviceAddress() local 88 const auto &device_tensor = AnfAlgo::GetMutableOutputAddr(output_node, output_index, false); in UpdateOutputDeviceAddress() 100 AnfAlgo::SetOutputAddr(new_device_tensor, output_index, output_node.get()); in UpdateOutputDeviceAddress() 108 void OutputActor::CollectOutput(const AnfNodePtr &output_node, size_t output_index, size_t output_p… in CollectOutput() argument [all …]
|
/third_party/mindspore/mindspore/ccsrc/runtime/device/ascend/executor/ |
D | aicpu_ext_info_handle.cc | 167 bool AicpuExtInfoHandler::UpdateOutputShapeAndType(uint32_t output_index, const NotNull<AnfNodePtr>… in UpdateOutputShapeAndType() argument 168 if (output_index >= output_num_) { in UpdateOutputShapeAndType() 169 …MS_LOG(ERROR) << "output_index:" << output_index << " >= output_num_:" << output_num_ << ", node: … in UpdateOutputShapeAndType() 173 auto shape = AnfAlgo::GetOutputDeviceShape(anf_node, output_index); in UpdateOutputShapeAndType() 174 auto max_shape = AnfAlgo::GetOutputMaxShape(anf_node, output_index); in UpdateOutputShapeAndType() 190 if (output_index >= output_shape_and_type_.size()) { in UpdateOutputShapeAndType() 191 MS_LOG(ERROR) << "Invalid output_index: " << output_index in UpdateOutputShapeAndType() 198 return UpdateShapeAndType(tmp_shape, NOT_NULL(output_shape_and_type_[output_index])); in UpdateOutputShapeAndType() 201 bool AicpuExtInfoHandler::GetOutputShapeAndType(uint32_t output_index, NotNull<std::vector<int64_t>… in GetOutputShapeAndType() argument 203 MS_LOG(DEBUG) << "Get " << node_name_ << " Output:" << output_index << " Shape And Type"; in GetOutputShapeAndType() [all …]
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/ |
D | gpu_kernel_factory.cc | 100 for (size_t output_index = 0; output_index < kernel_info->GetOutputNum(); output_index++) { in ReducePrecision() local 101 if (kernel_info->GetOutputDeviceType(output_index) == kNumberTypeInt64 && in ReducePrecision() 102 …(iter->second)[attr_index].first.GetOutputAttr(output_index % attr_size).first == kNumberTypeInt32… in ReducePrecision() 103 builder->SetOutputDeviceType(kNumberTypeInt32, output_index); in ReducePrecision() 104 …G(WARNING) << "Kernel [" << kernel_name << "] does not support int64, cast output " << output_index in ReducePrecision() 163 for (size_t output_index = 0; output_index < kernel_info->GetOutputNum(); output_index++) { in GpuKernelAttrCheck() local 164 if (kernel_info->GetOutputDeviceType(output_index) != in GpuKernelAttrCheck() 165 (iter->second)[attr_index].first.GetOutputAttr(output_index % attr_size).first) { in GpuKernelAttrCheck()
|
/third_party/mindspore/mindspore/ccsrc/backend/optimizer/ascend/format_type/ |
D | deal_ref_and_split_unsupported_transdata.cc | 65 … const CNodePtr &cnode, const size_t output_index, in AddRefNodePairToKernelGraph() argument 70 session::AnfWithOutIndex final_pair = std::make_pair(cnode, output_index); in AddRefNodePairToKernelGraph() 100 … const CNodePtr &cnode, size_t output_index, in AddAdditionalToRefOutput() argument 105 size_t final_index = output_index; in AddAdditionalToRefOutput() 116 auto cur_format = AnfAlgo::GetOutputFormat(cnode, output_index); in AddAdditionalToRefOutput() 117 auto cur_type = AnfAlgo::GetOutputDeviceDataType(cnode, output_index); in AddAdditionalToRefOutput() 118 auto cur_shape = AnfAlgo::GetOutputInferShape(cnode, output_index); in AddAdditionalToRefOutput() 119 auto detail_shape = AnfAlgo::GetOutputDetailShape(cnode, output_index); in AddAdditionalToRefOutput() 144 AddRefNodePairToKernelGraph(func_graph, cnode, output_index, input_index); in AddAdditionalToRefOutput() 188 for (size_t output_index = 0; output_index < output_num; ++output_index) { in DealRefForMultipleOutput() local [all …]
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/ |
D | kernel_build_info.cc | 30 std::string KernelBuildInfo::GetOutputFormat(size_t output_index) const { in GetOutputFormat() 31 if (output_index >= outputs_format_.size()) { in GetOutputFormat() 32 MS_LOG(ERROR) << "The index [" << output_index << "] is exceed the number of output"; in GetOutputFormat() 35 return outputs_format_[output_index]; in GetOutputFormat() 46 TypeId KernelBuildInfo::GetOutputDeviceType(size_t output_index) const { in GetOutputDeviceType() 47 if (output_index >= outputs_device_type_.size()) { in GetOutputDeviceType() 48 MS_LOG(ERROR) << "The index [" << output_index << "] is exceed the number of output"; in GetOutputDeviceType() 51 return outputs_device_type_[output_index]; in GetOutputDeviceType() 90 std::string KernelBuildInfo::GetOutputReshapeType(size_t output_index) const { in GetOutputReshapeType() 94 if (output_index >= output_reshape_type_.size()) { in GetOutputReshapeType() [all …]
|
/third_party/mindspore/mindspore/ccsrc/backend/optimizer/pass/ |
D | replace_node_by_proxy.cc | 39 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GenerateKernelBuildInfo() local 40 outputs_device_format.push_back(AnfAlgo::GetOutputFormat(cnode, output_index)); in GenerateKernelBuildInfo() 41 outputs_device_type.push_back(AnfAlgo::GetOutputDeviceDataType(cnode, output_index)); in GenerateKernelBuildInfo() 42 outputs_shape.push_back(AnfAlgo::GetOutputInferShape(cnode, output_index)); in GenerateKernelBuildInfo()
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/gpu/cuda_impl/ |
D | rcwm_small_impl.cu | 24 <<<1, BLOCK, 0, stream>>>(seedc, input_size, input, output_mask, output_index, k); \ 43 __global__ void L2Rcwm(int seedc, int input_size, const K *input, K *output_mask, S *output_index, … in L2Rcwm() argument 121 output_index[i] = shared_V[i]; in L2Rcwm() 126 void RCWMScaleK(int seedc, int input_size, K *input, int k, S *output_index, K *output_mask, cudaSt… in RCWMScaleK() argument 146 void CalRandomChoiceWithMaskSmall(int input_size, int seedc, int count, K *input, S *output_index, … in CalRandomChoiceWithMaskSmall() argument 148 RCWMScaleK<T, S, K>(seedc, input_size, input, count, output_index, output_mask, stream); in CalRandomChoiceWithMaskSmall() 152 … int *output_index, bool *output_mask, cudaStream_t stream);
|
D | random_choice_with_mask_impl.cu | 103 … const int d3, const int d4, const int d5, const T *input, S *output_index) { in Reshape2Index() argument 116 output_index[index_pos++] = 0; in Reshape2Index() 120 output_index[index_pos++] = pos_array[i]; in Reshape2Index() 184 …l__ void MoveToOutput(const int input_shape_size, const int count, const T *input, S *output_index, in MoveToOutput() argument 197 output_index[pos] = index_buff[idx]; in MoveToOutput() 211 output_index[pos] = index_buff[idx]; in MoveToOutput() 220 output_index[pos] = static_cast<S>(0); in MoveToOutput() 231 … const T *input, S *output_index, T *output_mask, S *index_buff, S *mask_buff, S *rank_buff, in CalRandomChoiceWithMask() argument 256 …BLOCKS(count), GET_THREADS, 0, stream>>>(input_shape_size, count, input, output_index, output_mask, in CalRandomChoiceWithMask() 262 … const bool *input, int *output_index, bool *output_mask, int *index_buff,
|
D | topk_impl.cu | 27 …<<<block_num_limit, BLOCK, 0, stream>>>(outer_size, inner_size, input, output, output_index, k_cut… 101 … S *output_index, S k_cut, const T &init_K, const int &outer_id, T *shared_K, in TopKStep() argument 171 output_index[outer_id * k_cut + (*k_prime) + i] = shared_V[i] % inner_size; in TopKStep() 178 …void TopKBlock(int outer_size, int inner_size, const T *input, T *output, S *output_index, S k_cut, in TopKBlock() argument 200 …outer_size, inner_size, input, output, output_index, k_cut, init_K, outer_id, shared_K, shared_V, … in TopKBlock() 207 …pK(const int outer_size, const int inner_size, const T *input, S k_cut, T *output, S *output_index, in FastTopK() argument 225 int *output_index, const half init_K, cudaStream_t stream); 227 int *output_index, const float init_K, cudaStream_t stream);
|
D | general_reduction_impl.cu | 48 T *output, S *output_index, bool fp16_flag, T init_K) { in ThreadReduction() argument 73 output_index[outer_id * inner_size + inner_id] = threadV; in ThreadReduction() 79 S *output_index, bool fp16_flag, T init_K) { in WarpReduction() argument 119 output_index[outer_id * inner_size + inner_id] = threadV; in WarpReduction() 127 T *output, S *output_index, bool fp16_flag, T init_K) { in Warp4Reduction() argument 208 output_index[outer_id * inner_size + inner_id] = shared_V[groupId * kWarpGroup]; in Warp4Reduction() 216 T *output, S *output_index, bool fp16_flag, T init_K) { in BlockReduction() argument 293 output_index[outer_id * inner_size + inner_id] = threadV; in BlockReduction() 300 S *output_index, cudaStream_t stream) { in GeneralReduction() argument 310 small, outer_size, bound, inner_size, input, output, output_index, fp16_flag, init_K); in GeneralReduction() [all …]
|
/third_party/mindspore/mindspore/ccsrc/backend/optimizer/ascend/enhancer/ |
D | insert_tensor_move_for_getnext.cc | 40 for (size_t output_index = 0; output_index < output_num; ++output_index) { in InsertTensorMoveForGetNextOutputs() local 41 auto tuple_get_item = CreatTupleGetItemNode(func_graph, node, output_index); in InsertTensorMoveForGetNextOutputs()
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/host/ |
D | host_kernel_metadata.cc | 51 for (size_t output_index = 0; output_index < output_num; ++output_index) { in HostMetadataInfo() local 53 outputs_type.push_back(AnfAlgo::GetOutputInferDataType(kernel_node, output_index)); in HostMetadataInfo()
|
/third_party/mindspore/mindspore/ccsrc/runtime/device/cpu/ |
D | kernel_select_cpu.cc | 65 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GetOutputDtypes() local 67 dtype = AnfAlgo::GetOutputInferDataType(kernel_node, output_index); in GetOutputDtypes() 74 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GetOutputFormat() local 104 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GetOutputFormatsAndDtypes() local 105 output_formats->emplace_back(kernel_attr.GetOutputAttr(output_index).second); in GetOutputFormatsAndDtypes() 106 auto dtype = kernel_attr.GetOutputAttr(output_index).first; in GetOutputFormatsAndDtypes()
|
/third_party/mindspore/tests/ut/python/parallel/ |
D | test_semi_auto_two_subgraphs.py | 57 def __init__(self, network, output_index): argument 60 self.output_index = output_index 63 predict = self.network(x1)[self.output_index] 91 self.loss_net_w = IthOutputCell(network, output_index=0) 92 self.loss_net_d = IthOutputCell(network, output_index=1)
|
D | test_auto_parallel_for_loop_multi_subgraph.py | 78 def __init__(self, network, output_index): argument 81 self.output_index = output_index 84 predict = self.network(x)[self.output_index] 109 self.loss_net_w = IthOutputCell(network, output_index=0) 110 self.loss_net_d = IthOutputCell(network, output_index=1)
|
D | test_auto_parallel_double_subgraphs.py | 61 def __init__(self, network, output_index): argument 64 self.output_index = output_index 67 predict = self.network(x)[self.output_index] 92 self.loss_net_w = IthOutputCell(network, output_index=0) 93 self.loss_net_d = IthOutputCell(network, output_index=1)
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/aicpu/ |
D | aicpu_kernel_build.cc | 231 for (size_t output_index = 0; output_index < output_num; output_index++) { in SetNodeOutputs() local 234 std::vector<size_t> output_shape = AnfAlgo::GetOutputDeviceShape(anf_node, output_index); in SetNodeOutputs() 242 TypeId output_type = AnfAlgo::GetOutputDeviceDataType(anf_node, output_index); in SetNodeOutputs() 343 for (size_t output_index = 0; output_index < output_num; output_index++) { in SetExtInfoOutputShapeType() local 344 std::vector<size_t> output_shape = AnfAlgo::GetOutputDeviceShape(anf_node, output_index); in SetExtInfoOutputShapeType() 345 TypeId output_type = AnfAlgo::GetOutputDeviceDataType(anf_node, output_index); in SetExtInfoOutputShapeType() 347 outputs[output_index].type = output_data_type; in SetExtInfoOutputShapeType() 351 outputs[output_index].dims[output_shape_index] = SizeToLong(output_shape[output_shape_index]); in SetExtInfoOutputShapeType() 354 outputs[output_index].dims[output_shape_index] = LLONG_MIN; in SetExtInfoOutputShapeType()
|
D | aicpu_kernel_metadata.cc | 67 for (size_t output_index = 0; output_index < output_num; ++output_index) { in AicpuMetadataInfoForSpecialNodes() local 69 (void)outputs_type.emplace_back(AnfAlgo::GetOutputInferDataType(kernel_node, output_index)); in AicpuMetadataInfoForSpecialNodes()
|
/third_party/mindspore/mindspore/ccsrc/runtime/device/ascend/ |
D | kernel_select_ascend.cc | 68 for (size_t output_index = 0; output_index < kernel_build_info.GetOutputNum(); ++output_index) { in MatchInferOutputDataType() local 69 …if (kernel_build_info.GetOutputDeviceType(output_index) != AnfAlgo::GetOutputInferDataType(cnode, … in MatchInferOutputDataType() 174 for (size_t output_index = 0; output_index < output_num; ++output_index) { in UpdateCurMatchCounts() local 176 if (kernel_build_info.GetOutputDeviceType(output_index) == in UpdateCurMatchCounts() 177 AnfAlgo::GetOutputInferDataType(kernel_node, output_index)) { in UpdateCurMatchCounts() 180 if (kernel_build_info.GetOutputFormat(output_index) == pri_match_format) { in UpdateCurMatchCounts() 271 for (size_t output_index = 0; output_index < kernel_build_info->GetOutputNum(); ++output_index) { in TagRaiseReduce() local 272 auto in_dtype = AnfAlgo::GetOutputInferDataType(cnode, output_index); in TagRaiseReduce() 273 auto device_dtype = kernel_build_info->GetOutputDeviceType(output_index); in TagRaiseReduce()
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/cpu/ |
D | one_hot_cpu_kernel.cc | 64 size_t output_index = stride_num * depth_ * stride_ + i % stride_; in Launch() local 68 output[output_index] = on_value; in Launch() 70 output[output_index] = off_value; in Launch() 72 output_index += stride_; in Launch()
|
/third_party/mindspore/mindspore/ccsrc/fl/server/kernel/ |
D | optimizer_kernel.h | 78 for (size_t output_index = 0; output_index < output_num; ++output_index) { in InitServerKernelInputOutputSize() local 79 std::vector<size_t> shape = AnfAlgo::GetOutputInferShape(kernel_node, output_index); in InitServerKernelInputOutputSize()
|
/third_party/mindspore/mindspore/ccsrc/backend/optimizer/gpu/ |
D | add_relu_v2_fusion.cc | 43 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GenerateKernelBuildInfo() local 44 outputs_type.push_back(AnfAlgo::GetOutputInferDataType(node, output_index)); in GenerateKernelBuildInfo()
|
D | matmul_biasadd_fusion.cc | 43 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GenerateKernelBuildInfo() local 44 outputs_type.push_back(AnfAlgo::GetOutputInferDataType(node, output_index)); in GenerateKernelBuildInfo()
|
D | add_relu_grad_v2_fusion.cc | 43 for (size_t output_index = 0; output_index < output_num; ++output_index) { in GenerateKernelBuildInfo() local 44 outputs_type.push_back(AnfAlgo::GetOutputInferDataType(node, output_index)); in GenerateKernelBuildInfo()
|
/third_party/mindspore/mindspore/ccsrc/backend/kernel_compiler/hccl/ |
D | hccl_kernel_metadata.cc | 97 for (size_t output_index = 0; output_index < output_num; ++output_index) { in HcclMetadataInfo() local 98 outputs_format.emplace_back(GetKernelFormat(kernel_node, output_index)); in HcclMetadataInfo()
|