From 20644043e528fd1eeb196ea461d59695267c84f4 Mon Sep 17 00:00:00 2001 From: Ryan Hill Date: Mon, 12 Apr 2021 17:08:11 -0700 Subject: [PATCH] Fix merge breaks --- .../core/framework/provider_bridge_ort.cc | 4 + .../core/providers/cuda/cuda_common.cc | 17 - onnxruntime/core/providers/cuda/cuda_common.h | 20 +- .../providers/cuda/cuda_provider_factory.cc | 1 - .../openvino/ov_versions/data_ops.cc | 966 +++++++++--------- .../provider_bridge_provider.cc | 3 + .../shared_library/provider_interfaces.h | 10 + onnxruntime/test/framework/test_utils.cc | 1 - onnxruntime/test/framework/test_utils.h | 1 - .../test/providers/cuda/gemm_options_test.cc | 4 + orttraining/orttraining/models/bert/main.cc | 3 +- .../test/graph/gradient_graph_builder_test.cc | 3 +- .../session/training_session_test_utils.cc | 1 - .../training_ops/cuda/controlflow/yield.cc | 3 +- .../training_ops/cuda/nn/conv_grad.cc | 18 +- .../training_ops/cuda/nn/dropout_grad.cc | 22 +- .../training_ops/cuda/tensor/concat.cc | 3 +- 17 files changed, 533 insertions(+), 547 deletions(-) diff --git a/onnxruntime/core/framework/provider_bridge_ort.cc b/onnxruntime/core/framework/provider_bridge_ort.cc index a63d01fae6..8ebe2ef096 100644 --- a/onnxruntime/core/framework/provider_bridge_ort.cc +++ b/onnxruntime/core/framework/provider_bridge_ort.cc @@ -53,6 +53,7 @@ void Foo() {} #include "orttraining/training_ops/cpu/controlflow/group.h" #include "orttraining/training_ops/cpu/controlflow/record.h" #include "orttraining/training_ops/cpu/controlflow/wait.h" +#include "orttraining/training_ops/cpu/controlflow/yield.h" #include "orttraining/training_ops/cpu/loss/softmax_cross_entropy_loss.h" #include "orttraining/training_ops/cpu/tensor/split.h" #endif @@ -490,6 +491,7 @@ struct ProviderHostImpl : ProviderHost { void KernelDefBuilder__Alias(KernelDefBuilder* p, int input_index, int output_index) override { p->Alias(input_index, output_index); } void KernelDefBuilder__Alias(KernelDefBuilder* p, const std::vector>& aliases) override { p->Alias(aliases); } void KernelDefBuilder__VariadicAlias(KernelDefBuilder* p, int input_offset, int output_offset) override { p->VariadicAlias(input_offset, output_offset); } + void KernelDefBuilder__ExternalOutputs(KernelDefBuilder* p) override { p->ExternalOutputs(); } std::unique_ptr KernelDefBuilder__Build(KernelDefBuilder* p) override { return p->Build(); } @@ -654,6 +656,7 @@ struct ProviderHostImpl : ProviderHost { bool OpKernelInfo__TryGetConstantInput(const OpKernelInfo* p, int input_index, const Tensor** constant_input_value) override { return p->TryGetConstantInput(input_index, constant_input_value); } uint32_t OpKernelInfo__GetInputCount(const OpKernelInfo* p) override { return p->GetInputCount(); } + uint32_t OpKernelInfo__GetOutputCount(const OpKernelInfo* p) override { return p->GetOutputCount(); } const Node& OpKernelInfo__node(const OpKernelInfo* p) override { return p->node(); } // SessionState @@ -836,6 +839,7 @@ struct ProviderHostImpl : ProviderHost { void contrib__GetNDCFromLogitAndLabelShape(const TensorShape& logit_shape, const TensorShape& label_shape, int64_t& N_D, int64_t& C) override { contrib::GetNDCFromLogitAndLabelShape(logit_shape, label_shape, N_D, C); } void contrib__GetPermutationAndShape(bool ncd_to_ndc, const TensorShape& tensor_shape, std::vector& new_shape, std::vector& permutations) override { contrib::GetPermutationAndShape(ncd_to_ndc, tensor_shape, new_shape, permutations); } Status contrib__PrepareForTrainingCompute(const TensorShape& input_shape, int num_outputs, int64_t& axis, int& before_dims, int& after_dims_including_split_axis, int& after_dims_excluding_split, std::vector& split_sizes) override { return contrib::PrepareForTrainingCompute(input_shape, num_outputs, axis, before_dims, after_dims_including_split_axis, after_dims_excluding_split, split_sizes); } + Status YieldOp__Compute(const contrib::YieldOp* p, OpKernelContext* context) override { return p->Compute(context); } #endif } provider_host_; diff --git a/onnxruntime/core/providers/cuda/cuda_common.cc b/onnxruntime/core/providers/cuda/cuda_common.cc index 771c0338f1..57477f167c 100644 --- a/onnxruntime/core/providers/cuda/cuda_common.cc +++ b/onnxruntime/core/providers/cuda/cuda_common.cc @@ -2,8 +2,6 @@ // Licensed under the MIT License. #include "core/providers/cuda/cuda_common.h" -#include "core/common/logging/logging.h" -#include "core/common/logging/severity.h" #include "core/platform/env_var_utils.h" namespace onnxruntime { @@ -29,20 +27,5 @@ const HalfGemmOptions* HalfGemmOptions::GetInstance() { return &instance; } -void HalfGemmOptions::Initialize(int value) -{ - compute_16f_ = (value & 0x01) > 0; -#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000 - disallow_reduced_precision_reduction_ = (value & 0x02) > 0; - pedantic_ = (value & 0x04) > 0; - LOGS_DEFAULT(INFO) << "ORT_CUDA_GEMM_OPTIONS: compute_16f=" << instance.compute_16f_ - << " disallow_reduced_precision_reduction=" << instance.disallow_reduced_precision_reduction_ - << " pedantic=" << instance.pedantic_; -#else - LOGS_DEFAULT(INFO) << "ORT_CUDA_GEMM_OPTIONS: compute_16f=" << instance.compute_16f_; -#endif - initialized_ = true; -} - } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_common.h b/onnxruntime/core/providers/cuda/cuda_common.h index f785785c81..5298ee5a04 100644 --- a/onnxruntime/core/providers/cuda/cuda_common.h +++ b/onnxruntime/core/providers/cuda/cuda_common.h @@ -103,12 +103,12 @@ inline bool CalculateFdmStrides(gsl::span p, const std::vector= 7 : true ); +#if defined(CUDA_VERSION) && CUDA_VERSION < 11000 + enable_ = (mode == CUBLAS_TENSOR_OP_MATH ? prop.major >= 7 : true); #else enable_ = (mode == CUBLAS_TF32_TENSOR_OP_MATH ? prop.major >= 8 : true); #endif - + if (enable_) { cublasGetMathMode(handle, &mode_); enable_ = (mode_ != mode); @@ -164,7 +164,19 @@ class HalfGemmOptions { bool IsCompute16F() const { return compute_16f_; } - void Initialize(int value); + void Initialize(int value) { + compute_16f_ = (value & 0x01) > 0; +#if defined(CUDA_VERSION) && CUDA_VERSION >= 11000 + disallow_reduced_precision_reduction_ = (value & 0x02) > 0; + pedantic_ = (value & 0x04) > 0; + LOGS_DEFAULT(INFO) << "ORT_CUDA_GEMM_OPTIONS: compute_16f=" << instance.compute_16f_ + << " disallow_reduced_precision_reduction=" << instance.disallow_reduced_precision_reduction_ + << " pedantic=" << instance.pedantic_; +#else + LOGS_DEFAULT(INFO) << "ORT_CUDA_GEMM_OPTIONS: compute_16f=" << instance.compute_16f_; +#endif + initialized_ = true; + } private: // Default is FP32. Aggregate in FP16 might be faster but the cost is loss in precision. diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index ac2ae55576..9ccc7209ad 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -149,4 +149,3 @@ ORT_API(onnxruntime::Provider*, GetProvider) { return &onnxruntime::g_provider; } } - diff --git a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc index 9c1c090016..c5dd205da0 100644 --- a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc +++ b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc @@ -29,141 +29,140 @@ namespace onnxruntime { namespace openvino_ep { std::set ops_supported_only_in_model = { - "Cast", - "Concat", - "ConstantOfShape", - "Dropout", - "Expand", - "EyeLike", - "Exp", - "GatherND", - "Identity", - "Loop", - "NonMaxSuppression", - "NonZero", - "Not", - "OneHot", - "Pad", - "Range", - "ReduceMin", - "Resize", - "Round", - "Shape", - "Split", - "Tile", - "TopK" -}; + "Cast", + "Concat", + "ConstantOfShape", + "Dropout", + "Expand", + "EyeLike", + "Exp", + "GatherND", + "Identity", + "Loop", + "NonMaxSuppression", + "NonZero", + "Not", + "OneHot", + "Pad", + "Range", + "ReduceMin", + "Resize", + "Round", + "Shape", + "Split", + "Tile", + "TopK"}; std::vector supported_op_mode = { - {"Abs", V_2020_4,{"CPU", "GPU"}}, - {"Acos", V_2020_4,{"CPU"}}, - {"Acosh", V_2020_4,{"CPU"}}, - {"Add", V_2020_4,{"All"}}, - {"And", V_2020_4,{"All"}}, - {"ArgMax", V_2020_4,{"CPU"}}, - {"ArgMax", V_2021_1,{"All"}}, - {"ArgMin", V_2020_4,{"CPU"}}, - {"ArgMin", V_2021_2,{"CPU","MYRIAD"}}, - {"Asin", V_2020_4,{"CPU", "GPU"}}, - {"Asinh", V_2020_4,{"CPU", "GPU"}}, - {"Atan", V_2020_4,{"CPU", "GPU"}}, - {"Atanh", V_2020_4,{"CPU"}}, - {"AveragePool", V_2020_4,{"All"}}, - {"BatchNormalization", V_2020_4,{"All"}}, - {"Cast", V_2020_4,{"All"}}, - {"Ceil", V_2020_4,{"GPU"}}, - {"Ceil", V_2021_3,{"MYRIAD"}}, - {"Ceil", V_2021_2,{"GPU","MYRIAD"}}, - {"Clip", V_2020_4,{"All"}}, - {"Concat", V_2020_4,{"All"}}, - {"Constant", V_2020_4,{"All"}}, - {"ConstantOfShape", V_2020_4,{"All"}}, - {"Conv", V_2020_4,{"All"}}, - {"ConvTranspose", V_2020_4,{"All"}}, - {"Cos", V_2020_4,{"CPU"}}, - {"Cosh", V_2020_4,{"CPU"}}, - {"DepthToSpace", V_2020_4,{"All"}}, - {"Div", V_2020_4,{"All"}}, - {"Dropout", V_2020_4,{"All"}}, - {"Elu", V_2020_4,{"All"}}, - {"Equal", V_2020_4,{"All"}}, - {"Erf", V_2020_4,{"All"}}, - {"Exp", V_2020_4,{"All"}}, - {"Expand", V_2021_1,{"MYRIAD"}}, - {"Flatten", V_2020_4,{"All"}}, - {"Floor", V_2020_4,{"All"}}, - {"Gather", V_2020_4,{"All"}}, - {"GatherElements", V_2021_3,{"MYRIAD"}}, - {"GatherND", V_2021_2,{"MYRIAD"}}, - {"Gemm", V_2020_4,{"All"}}, - {"GlobalAveragePool", V_2020_4,{"All"}}, - {"GlobalLpPool", V_2020_4,{"CPU", "GPU"}}, - {"Greater", V_2020_4,{"All"}}, - {"Identity", V_2020_4,{"All"}}, - {"InstanceNormalization", V_2020_4,{"All"}}, - {"HardSigmoid", V_2020_4,{"CPU", "GPU"}}, - {"LeakyRelu", V_2020_4,{"All"}}, - {"Less", V_2020_4,{"All"}}, - {"Log", V_2020_4,{"All"}}, - {"Loop", V_2021_3,{"MYRIAD"}}, - {"LRN", V_2020_4,{"All"}}, - {"LSTM", V_2020_4,{"All"}}, - {"MatMul", V_2020_4,{"All"}}, - {"Max", V_2020_4,{"All"}}, - {"MaxPool", V_2020_4,{"All"}}, - {"Mean", V_2020_4,{"All"}}, - {"Min", V_2020_4,{"All"}}, - {"Mul", V_2020_4,{"All"}}, - {"Neg", V_2020_4,{"All"}}, - {"NonMaxSuppression", V_2021_1,{"All"}}, - {"NonZero", V_2021_1,{"CPU", "MYRIAD"}}, - {"Not", V_2021_1,{"All"}}, - {"Not", V_2020_4,{"CPU", "GPU"}}, - {"OneHot", V_2020_4,{"All"}}, - {"Pad", V_2020_4,{"All"}}, - {"Pow", V_2020_4,{"All"}}, - {"PRelu", V_2020_4,{"All"}}, - {"Range", V_2021_2,{"MYRIAD"}}, - {"Reciprocal", V_2020_4,{"All"}}, - {"ReduceLogSum", V_2020_4,{"CPU", "MYRIAD"}}, - {"ReduceMax", V_2020_4 ,{"All"}}, - {"ReduceMean", V_2020_4,{"All"}}, - {"ReduceMin", V_2020_4,{"All"}}, - {"ReduceProd", V_2020_4,{"CPU"}}, - {"ReduceSum", V_2020_4,{"All"}}, - {"ReduceSumSquare", V_2020_4,{"CPU", "MYRIAD"}}, - {"Relu", V_2020_4,{"All"}}, - {"Resize", V_2020_4,{"CPU"}}, - {"Resize", V_2021_3,{"MYRIAD"}}, - {"Reshape", V_2020_4,{"All"}}, - {"RoiAlign", V_2021_1,{"All"}}, - {"Round", V_2021_2,{"MYRIAD"}}, - {"Scatter", V_2021_1,{"MYRIAD"}}, - {"ScatterElements", V_2021_2,{"MYRIAD"}}, - {"Selu", V_2020_4,{"CPU", "GPU"}}, - {"Shape", V_2020_4,{"All"}}, - {"Sigmoid", V_2020_4,{"All"}}, - {"Sign", V_2020_4,{"CPU"}}, - {"Sign", V_2020_4,{"CPU"}}, - {"Sinh", V_2020_4,{"CPU"}}, - {"SinFloat", V_2020_4,{"MYRIAD"}}, - {"Slice", V_2020_4,{"All"}}, - {"Softmax", V_2020_4,{"All"}}, - {"SpaceToDepth", V_2020_4,{"All"}}, - {"Split", V_2020_4,{"All"}}, - {"Sqrt", V_2020_4,{"All"}}, - {"Squeeze", V_2020_4,{"All"}}, - {"Softsign", V_2020_4,{"CPU"}}, - {"Sub", V_2020_4,{"All"}}, - {"Sum", V_2020_4,{"All"}}, - {"Tan", V_2020_4,{"CPU", "GPU"}}, - {"Tanh", V_2020_4,{"All"}}, - {"Tile", V_2021_3,{"MYRIAD"}}, - {"Transpose", V_2020_4,{"All"}}, - {"TopK", V_2020_4,{"All"}}, - {"Unsqueeze", V_2020_4,{"All"}}, - {"Upsample", V_2021_1,{"CPU"}}, - {"Where", V_2021_2,{"MYRIAD"}}, + {"Abs", V_2020_4, {"CPU", "GPU"}}, + {"Acos", V_2020_4, {"CPU"}}, + {"Acosh", V_2020_4, {"CPU"}}, + {"Add", V_2020_4, {"All"}}, + {"And", V_2020_4, {"All"}}, + {"ArgMax", V_2020_4, {"CPU"}}, + {"ArgMax", V_2021_1, {"All"}}, + {"ArgMin", V_2020_4, {"CPU"}}, + {"ArgMin", V_2021_2, {"CPU", "MYRIAD"}}, + {"Asin", V_2020_4, {"CPU", "GPU"}}, + {"Asinh", V_2020_4, {"CPU", "GPU"}}, + {"Atan", V_2020_4, {"CPU", "GPU"}}, + {"Atanh", V_2020_4, {"CPU"}}, + {"AveragePool", V_2020_4, {"All"}}, + {"BatchNormalization", V_2020_4, {"All"}}, + {"Cast", V_2020_4, {"All"}}, + {"Ceil", V_2020_4, {"GPU"}}, + {"Ceil", V_2021_3, {"MYRIAD"}}, + {"Ceil", V_2021_2, {"GPU", "MYRIAD"}}, + {"Clip", V_2020_4, {"All"}}, + {"Concat", V_2020_4, {"All"}}, + {"Constant", V_2020_4, {"All"}}, + {"ConstantOfShape", V_2020_4, {"All"}}, + {"Conv", V_2020_4, {"All"}}, + {"ConvTranspose", V_2020_4, {"All"}}, + {"Cos", V_2020_4, {"CPU"}}, + {"Cosh", V_2020_4, {"CPU"}}, + {"DepthToSpace", V_2020_4, {"All"}}, + {"Div", V_2020_4, {"All"}}, + {"Dropout", V_2020_4, {"All"}}, + {"Elu", V_2020_4, {"All"}}, + {"Equal", V_2020_4, {"All"}}, + {"Erf", V_2020_4, {"All"}}, + {"Exp", V_2020_4, {"All"}}, + {"Expand", V_2021_1, {"MYRIAD"}}, + {"Flatten", V_2020_4, {"All"}}, + {"Floor", V_2020_4, {"All"}}, + {"Gather", V_2020_4, {"All"}}, + {"GatherElements", V_2021_3, {"MYRIAD"}}, + {"GatherND", V_2021_2, {"MYRIAD"}}, + {"Gemm", V_2020_4, {"All"}}, + {"GlobalAveragePool", V_2020_4, {"All"}}, + {"GlobalLpPool", V_2020_4, {"CPU", "GPU"}}, + {"Greater", V_2020_4, {"All"}}, + {"Identity", V_2020_4, {"All"}}, + {"InstanceNormalization", V_2020_4, {"All"}}, + {"HardSigmoid", V_2020_4, {"CPU", "GPU"}}, + {"LeakyRelu", V_2020_4, {"All"}}, + {"Less", V_2020_4, {"All"}}, + {"Log", V_2020_4, {"All"}}, + {"Loop", V_2021_3, {"MYRIAD"}}, + {"LRN", V_2020_4, {"All"}}, + {"LSTM", V_2020_4, {"All"}}, + {"MatMul", V_2020_4, {"All"}}, + {"Max", V_2020_4, {"All"}}, + {"MaxPool", V_2020_4, {"All"}}, + {"Mean", V_2020_4, {"All"}}, + {"Min", V_2020_4, {"All"}}, + {"Mul", V_2020_4, {"All"}}, + {"Neg", V_2020_4, {"All"}}, + {"NonMaxSuppression", V_2021_1, {"All"}}, + {"NonZero", V_2021_1, {"CPU", "MYRIAD"}}, + {"Not", V_2021_1, {"All"}}, + {"Not", V_2020_4, {"CPU", "GPU"}}, + {"OneHot", V_2020_4, {"All"}}, + {"Pad", V_2020_4, {"All"}}, + {"Pow", V_2020_4, {"All"}}, + {"PRelu", V_2020_4, {"All"}}, + {"Range", V_2021_2, {"MYRIAD"}}, + {"Reciprocal", V_2020_4, {"All"}}, + {"ReduceLogSum", V_2020_4, {"CPU", "MYRIAD"}}, + {"ReduceMax", V_2020_4, {"All"}}, + {"ReduceMean", V_2020_4, {"All"}}, + {"ReduceMin", V_2020_4, {"All"}}, + {"ReduceProd", V_2020_4, {"CPU"}}, + {"ReduceSum", V_2020_4, {"All"}}, + {"ReduceSumSquare", V_2020_4, {"CPU", "MYRIAD"}}, + {"Relu", V_2020_4, {"All"}}, + {"Resize", V_2020_4, {"CPU"}}, + {"Resize", V_2021_3, {"MYRIAD"}}, + {"Reshape", V_2020_4, {"All"}}, + {"RoiAlign", V_2021_1, {"All"}}, + {"Round", V_2021_2, {"MYRIAD"}}, + {"Scatter", V_2021_1, {"MYRIAD"}}, + {"ScatterElements", V_2021_2, {"MYRIAD"}}, + {"Selu", V_2020_4, {"CPU", "GPU"}}, + {"Shape", V_2020_4, {"All"}}, + {"Sigmoid", V_2020_4, {"All"}}, + {"Sign", V_2020_4, {"CPU"}}, + {"Sign", V_2020_4, {"CPU"}}, + {"Sinh", V_2020_4, {"CPU"}}, + {"SinFloat", V_2020_4, {"MYRIAD"}}, + {"Slice", V_2020_4, {"All"}}, + {"Softmax", V_2020_4, {"All"}}, + {"SpaceToDepth", V_2020_4, {"All"}}, + {"Split", V_2020_4, {"All"}}, + {"Sqrt", V_2020_4, {"All"}}, + {"Squeeze", V_2020_4, {"All"}}, + {"Softsign", V_2020_4, {"CPU"}}, + {"Sub", V_2020_4, {"All"}}, + {"Sum", V_2020_4, {"All"}}, + {"Tan", V_2020_4, {"CPU", "GPU"}}, + {"Tanh", V_2020_4, {"All"}}, + {"Tile", V_2021_3, {"MYRIAD"}}, + {"Transpose", V_2020_4, {"All"}}, + {"TopK", V_2020_4, {"All"}}, + {"Unsqueeze", V_2020_4, {"All"}}, + {"Upsample", V_2021_1, {"CPU"}}, + {"Where", V_2021_2, {"MYRIAD"}}, }; void DataOps::populate_types_supported() { @@ -219,72 +218,67 @@ void DataOps::populate_op_mode_supported() { no_dimension_supported_.push_back({"Resize", V_2021_2, {"Myriad"}}); no_dimension_supported_.push_back({"Equal", V_2021_2, {"Myriad"}}); - no_dimension_supported_.push_back({"Unsqueeze", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Squeeze", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Cast", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Gather", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Mul", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Sub", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Min", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Div", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Floor", V_2020_4,{"All"}}); - no_dimension_supported_.push_back({"Where", V_2021_2,{"All"}}); - no_dimension_supported_.push_back({"Range", V_2021_2,{"All"}}); - no_dimension_supported_.push_back({"ArgMin", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Max", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Add", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Less", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Greater", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Clip", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Resize", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Equal", V_2021_2,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Reshape", V_2021_3,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Ceil", V_2021_3,{"MYRIAD"}}); - no_dimension_supported_.push_back({"Loop", V_2021_3,{"MYRIAD"}}); - no_dimension_supported_.push_back({"ReduceMin", V_2021_3,{"MYRIAD"}}); + no_dimension_supported_.push_back({"Unsqueeze", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Squeeze", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Cast", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Gather", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Mul", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Sub", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Min", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Div", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Floor", V_2020_4, {"All"}}); + no_dimension_supported_.push_back({"Where", V_2021_2, {"All"}}); + no_dimension_supported_.push_back({"Range", V_2021_2, {"All"}}); + no_dimension_supported_.push_back({"ArgMin", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Max", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Add", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Less", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Greater", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Clip", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Resize", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Equal", V_2021_2, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Reshape", V_2021_3, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Ceil", V_2021_3, {"MYRIAD"}}); + no_dimension_supported_.push_back({"Loop", V_2021_3, {"MYRIAD"}}); + no_dimension_supported_.push_back({"ReduceMin", V_2021_3, {"MYRIAD"}}); + subgraph_supported_.push_back({"Mul", V_2020_4, {"All"}}); + subgraph_supported_.push_back({"Transpose", V_2020_4, {"All"}}); + subgraph_supported_.push_back({"Unsqueeze", V_2020_4, {"All"}}); + subgraph_supported_.push_back({"Cast", V_2020_4, {"All"}}); + subgraph_supported_.push_back({"Concat", V_2020_4, {"All"}}); + subgraph_supported_.push_back({"Gather", V_2020_4, {"All"}}); + subgraph_supported_.push_back({"Div", V_2020_4, {"MYRIAD"}}); + subgraph_supported_.push_back({"Sub", V_2020_4, {"MYRIAD"}}); + subgraph_supported_.push_back({"Identity", V_2021_1, {"CPU"}}); + subgraph_supported_.push_back({"Div", V_2021_1, {"CPU"}}); + subgraph_supported_.push_back({"Sub", V_2021_1, {"CPU"}}); - - subgraph_supported_.push_back({"Mul", V_2020_4,{"All"}}); - subgraph_supported_.push_back({"Transpose", V_2020_4,{"All"}}); - subgraph_supported_.push_back({"Unsqueeze", V_2020_4,{"All"}}); - subgraph_supported_.push_back({"Cast", V_2020_4,{"All"}}); - subgraph_supported_.push_back({"Concat", V_2020_4,{"All"}}); - subgraph_supported_.push_back({"Gather", V_2020_4,{"All"}}); - subgraph_supported_.push_back({"Div", V_2020_4,{"MYRIAD"}}); - subgraph_supported_.push_back({"Sub", V_2020_4,{"MYRIAD"}}); - subgraph_supported_.push_back({"Identity", V_2021_1,{"CPU"}}); - subgraph_supported_.push_back({"Div", V_2021_1,{"CPU"}}); - subgraph_supported_.push_back({"Sub", V_2021_1,{"CPU"}}); - - //populate unsupportedmode_t { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - for (size_t i = 0; i < node->InputDefs().size(); i++) { - if (node->InputDefs()[i]->TypeAsProto()->tensor_type().elem_type() != ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT) - return true; - } - return false; - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + for (size_t i = 0; i < node->InputDefs().size(); i++) { + if (node->InputDefs()[i]->TypeAsProto()->tensor_type().elem_type() != ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT) + return true; + } + return false; + }}; op_list_.insert({"Abs", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //tensor type does not support select last index - auto& attributes = node->GetAttributes(); - auto last_index_arg = attributes.count("select_last_index") > 0 ? attributes.at("select_last_index").i() : 0; - if (last_index_arg != 0) - return true; - // tensor type supports float as input for argmax and argmin - if (node->InputDefs()[0]->TypeAsProto()->tensor_type().elem_type() != ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT) - return true; - return false; - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //tensor type does not support select last index + auto& attributes = node->GetAttributes(); + auto last_index_arg = attributes.count("select_last_index") > 0 ? attributes.at("select_last_index").i() : 0; + if (last_index_arg != 0) + return true; + // tensor type supports float as input for argmax and argmin + if (node->InputDefs()[0]->TypeAsProto()->tensor_type().elem_type() != ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT) + return true; + return false; + }}; op_list_.insert({"ArgMax", obj}); op_list_.insert({"ArgMin", obj}); } @@ -301,72 +295,70 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"AveragePool", obj}); } { - UnsupportedOpMode obj = {{V_2021_1,V_2021_2,V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //auto pad null value is not supported - const auto& attributes = node->GetAttributes(); - auto auto_attr = attributes.find("auto_pad"); - if (auto_attr->second().s() == "") { - return true; - } - // default value of ceil_mode (0) is supported. - auto ceil_attr = attributes.find("ceil_mode"); - if (ceil_attr != attributes.end() && ceil_attr->second().i() != 0) return true; - return (!dimension_unsupported(node)); - } - }; + UnsupportedOpMode obj = {{V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //auto pad null value is not supported + const auto& attributes = node->GetAttributes(); + auto auto_attr = attributes.find("auto_pad"); + if (auto_attr->second().s() == "") { + return true; + } + // default value of ceil_mode (0) is supported. + auto ceil_attr = attributes.find("ceil_mode"); + if (ceil_attr != attributes.end() && ceil_attr->second().i() != 0) return true; + return (!dimension_unsupported(node)); + }}; op_list_.insert({"AveragePool", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //Only float 16, float and double data types are supported - const bool data_is_float = node->InputDefs()[0]->Type()->find("float") != std::string::npos; - const bool data_is_float16 = node->InputDefs()[0]->Type()->find("float16") != std::string::npos; - const bool data_is_double = node->InputDefs()[0]->Type()->find("double") != std::string::npos; - return !(data_is_float || data_is_float16 || data_is_double); - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //Only float 16, float and double data types are supported + const bool data_is_float = node->InputDefs()[0]->Type()->find("float") != std::string::npos; + const bool data_is_float16 = node->InputDefs()[0]->Type()->find("float16") != std::string::npos; + const bool data_is_double = node->InputDefs()[0]->Type()->find("double") != std::string::npos; + return !(data_is_float || data_is_float16 || data_is_double); + }}; op_list_.insert({"Clip", obj}); } - { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - if (GetInputCount(node, initializers) > 1) - return true; - return false; - } }; + { + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + if (GetInputCount(node, initializers) > 1) + return true; + return false; + }}; op_list_.insert({"Conv", obj}); } - { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - if (GetInputCount(node, initializers) > 1) - return true; - return false; - } }; + { + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + if (GetInputCount(node, initializers) > 1) + return true; + return false; + }}; op_list_.insert({"ConvTranspose", obj}); } - { - UnsupportedOpMode obj = {{V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& ) { - auto& attributes = node->GetAttributes(); - if (attributes.count("auto_pad") == 0 || attributes.at("auto_pad").s() == "") { - return true; - } - return false; - } }; + { + UnsupportedOpMode obj = {{V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + auto& attributes = node->GetAttributes(); + if (attributes.count("auto_pad") == 0 || attributes.at("auto_pad").s() == "") { + return true; + } + return false; + }}; op_list_.insert({"Conv", obj}); } - { - UnsupportedOpMode obj = {{V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& ) { - auto& attributes = node->GetAttributes(); - if (attributes.count("auto_pad") == 0 || attributes.at("auto_pad").s() == "") { - return true; - } - return false; - } }; + { + UnsupportedOpMode obj = {{V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + auto& attributes = node->GetAttributes(); + if (attributes.count("auto_pad") == 0 || attributes.at("auto_pad").s() == "") { + return true; + } + return false; + }}; op_list_.insert({"ConvTranspose", obj}); } { @@ -384,19 +376,18 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"ConvTranspose", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2,V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - // all ConvInteger zero points need to be constants - if (node->InputDefs().size() == 3) { - return (initializers.find(node->InputDefs()[2]->Name()) == initializers.end()); - } else if (node->InputDefs().size() == 4) { - return initializers.find(node->InputDefs()[2]->Name()) == initializers.end() || - initializers.find(node->InputDefs()[3]->Name()) == initializers.end(); - } - return false; - } - }; - op_list_.insert({"ConvInteger", obj}); + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + // all ConvInteger zero points need to be constants + if (node->InputDefs().size() == 3) { + return (initializers.find(node->InputDefs()[2]->Name()) == initializers.end()); + } else if (node->InputDefs().size() == 4) { + return initializers.find(node->InputDefs()[2]->Name()) == initializers.end() || + initializers.find(node->InputDefs()[3]->Name()) == initializers.end(); + } + return false; + }}; + op_list_.insert({"ConvInteger", obj}); } { UnsupportedOpMode obj = {{V_2020_4, V_2021_1}, @@ -434,121 +425,115 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"And", obj}); } { - UnsupportedOpMode obj = {{V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& ) { - if (device_id_.find("GPU") != std::string::npos) { - const auto& input = node->InputDefs()[0]; - auto graph_inputs = graph_viewer_.GetInputs(); - auto it = find(graph_inputs.begin(), graph_inputs.end(), input); - if (it != graph_inputs.end()) { - const auto& indices_arg = node->InputDefs()[1]; - if (indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_INT64) - return true; - } - } - return false; - } - }; + UnsupportedOpMode obj = {{V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + if (device_id_.find("GPU") != std::string::npos) { + const auto& input = node->InputDefs()[0]; + auto graph_inputs = graph_viewer_.GetInputs(); + auto it = find(graph_inputs.begin(), graph_inputs.end(), input); + if (it != graph_inputs.end()) { + const auto& indices_arg = node->InputDefs()[1]; + if (indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_INT64) + return true; + } + } + return false; + }}; op_list_.insert({"Gather", obj}); } { UnsupportedOpMode obj = {{V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& ) { - const auto& indices_arg = node->InputDefs()[0]; - const auto& output_arg = node->OutputDefs()[0]; - if (indices_arg->TypeAsProto()->tensor_type().elem_type() != output_arg->TypeAsProto()->tensor_type().elem_type()) - return true; - if ((indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT16) || - (indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_UINT8) || - (indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT)) { - return false; - } - return true; - }}; + [this](const Node* node, const InitializedTensorSet&) { + const auto& indices_arg = node->InputDefs()[0]; + const auto& output_arg = node->OutputDefs()[0]; + if (indices_arg->TypeAsProto()->tensor_type().elem_type() != output_arg->TypeAsProto()->tensor_type().elem_type()) + return true; + if ((indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT16) || + (indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_UINT8) || + (indices_arg->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT)) { + return false; + } + return true; + }}; op_list_.insert({"GatherElements", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - const auto& input = node->InputDefs()[0]; - const auto& output = node->OutputDefs()[0]; - auto graph_inputs = this->graph_viewer_.GetInputs(); - auto graph_outputs = this->graph_viewer_.GetOutputs(); - auto input_it = find(graph_inputs.begin(), graph_inputs.end(), input); - auto output_it = find(graph_outputs.begin(), graph_outputs.end(), output); - if (input_it != graph_inputs.end() && output_it != graph_outputs.end()) - return true; - return false; - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + const auto& input = node->InputDefs()[0]; + const auto& output = node->OutputDefs()[0]; + auto graph_inputs = this->graph_viewer_.GetInputs(); + auto graph_outputs = this->graph_viewer_.GetOutputs(); + auto input_it = find(graph_inputs.begin(), graph_inputs.end(), input); + auto output_it = find(graph_outputs.begin(), graph_outputs.end(), output); + if (input_it != graph_inputs.end() && output_it != graph_outputs.end()) + return true; + return false; + }}; op_list_.insert({"Identity", obj}); } - { + { UnsupportedOpMode obj = {{V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - //Loop has to be initializer - const auto& cond = node->InputDefs()[1]; - return (initializers.find(cond->Name()) == initializers.end()); - } - }; + [this](const Node* node, const InitializedTensorSet& initializers) { + //Loop has to be initializer + const auto& cond = node->InputDefs()[1]; + return (initializers.find(cond->Name()) == initializers.end()); + }}; op_list_.insert({"Loop", obj}); } { - UnsupportedOpMode obj = {{V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //MaxPool "indices" output is not currently supported. - //if (node->OutputDefs().size() > 1) - // return true; - const auto& attributes = node->GetAttributes(); - /* default value of ceil_mode (0) is supported. + UnsupportedOpMode obj = {{V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //MaxPool "indices" output is not currently supported. + //if (node->OutputDefs().size() > 1) + // return true; + const auto& attributes = node->GetAttributes(); + /* default value of ceil_mode (0) is supported. auto ceil_attr = attributes.find("ceil_mode"); if (ceil_attr != attributes.end() && ceil_attr->second().i() != 0) return true;*/ - auto auto_attr = attributes.find("auto_pad"); - //auto pad null value is not supported - if (auto_attr->second().s() == "") - return true; - // dilations attrs are not supported in nGraph - if (attributes.find("dilations") != attributes.end()) - return true; - return(!this->dimension_unsupported(node)); - } - }; + auto auto_attr = attributes.find("auto_pad"); + //auto pad null value is not supported + if (auto_attr->second().s() == "") + return true; + // dilations attrs are not supported in nGraph + if (attributes.find("dilations") != attributes.end()) + return true; + return (!this->dimension_unsupported(node)); + }}; op_list_.insert({"MaxPool", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //MaxPool "indices" output is not currently supported. - if (node->OutputDefs().size() > 1) - return true; - const auto& attributes = node->GetAttributes(); - // default value of ceil_mode (0) is supported. - auto ceil_attr = attributes.find("ceil_mode"); - if (ceil_attr != attributes.end() && ceil_attr->second().i() != 0) - return true; - auto auto_attr = attributes.find("auto_pad"); - //auto pad null value is not supported - if (auto_attr->second().s() == "") - return true; - // dilations attrs are not supported in nGraph - if (attributes.find("dilations") != attributes.end()) - return true; - return(!this->dimension_unsupported(node)); - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2}, + [this](const Node* node, const InitializedTensorSet&) { + //MaxPool "indices" output is not currently supported. + if (node->OutputDefs().size() > 1) + return true; + const auto& attributes = node->GetAttributes(); + // default value of ceil_mode (0) is supported. + auto ceil_attr = attributes.find("ceil_mode"); + if (ceil_attr != attributes.end() && ceil_attr->second().i() != 0) + return true; + auto auto_attr = attributes.find("auto_pad"); + //auto pad null value is not supported + if (auto_attr->second().s() == "") + return true; + // dilations attrs are not supported in nGraph + if (attributes.find("dilations") != attributes.end()) + return true; + return (!this->dimension_unsupported(node)); + }}; op_list_.insert({"MaxPool", obj}); } { - UnsupportedOpMode obj = {{V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - if (device_id_.find("MYRIAD") == std::string::npos) { - if (GetInputCount(node, initializers) == 1) - return true; - } - return false; - } - }; + UnsupportedOpMode obj = {{V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + if (device_id_.find("MYRIAD") == std::string::npos) { + if (GetInputCount(node, initializers) == 1) + return true; + } + return false; + }}; op_list_.insert({"Max", obj}); op_list_.insert({"Min", obj}); op_list_.insert({"Mean", obj}); @@ -581,31 +566,29 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"Min", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //All matmuls except float have computation missmatch - const bool A_is_float = node->InputDefs()[0]->Type()->find("float") != std::string::npos; - const bool B_is_float = node->InputDefs()[1]->Type()->find("float") != std::string::npos; - return (A_is_float && B_is_float) ? false : true; - } - }; - op_list_.insert({"MatMul", obj}); + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //All matmuls except float have computation missmatch + const bool A_is_float = node->InputDefs()[0]->Type()->find("float") != std::string::npos; + const bool B_is_float = node->InputDefs()[1]->Type()->find("float") != std::string::npos; + return (A_is_float && B_is_float) ? false : true; + }}; + op_list_.insert({"MatMul", obj}); } - { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - // all MatMulInteger zero points need to be constants - if (node->InputDefs().size() == 3) { - // not found in initializers -> not const - return initializers.find(node->InputDefs()[2]->Name()) == initializers.end(); - } else if (node->InputDefs().size() == 4) { - // not found in initializers -> not const - return ((initializers.find(node->InputDefs()[2]->Name()) == initializers.end()) || - (initializers.find(node->InputDefs()[2]->Name()) == initializers.end())) ; - } - return false; - } - }; + { + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + // all MatMulInteger zero points need to be constants + if (node->InputDefs().size() == 3) { + // not found in initializers -> not const + return initializers.find(node->InputDefs()[2]->Name()) == initializers.end(); + } else if (node->InputDefs().size() == 4) { + // not found in initializers -> not const + return ((initializers.find(node->InputDefs()[2]->Name()) == initializers.end()) || + (initializers.find(node->InputDefs()[2]->Name()) == initializers.end())); + } + return false; + }}; op_list_.insert({"MatMulInteger", obj}); } { @@ -626,15 +609,14 @@ void DataOps::populate_op_mode_supported() { } { UnsupportedOpMode obj = {{V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - auto graph_outputs = graph_viewer_.GetOutputs(); - const auto& output = node->OutputDefs()[0]; - auto output_it = find(graph_outputs.begin(), graph_outputs.end(), output); - if (output_it != graph_outputs.end()) - return true; - return false; - } - }; + [this](const Node* node, const InitializedTensorSet&) { + auto graph_outputs = graph_viewer_.GetOutputs(); + const auto& output = node->OutputDefs()[0]; + auto output_it = find(graph_outputs.begin(), graph_outputs.end(), output); + if (output_it != graph_outputs.end()) + return true; + return false; + }}; op_list_.insert({"NonMaxSuppression", obj}); } { @@ -649,53 +631,51 @@ void DataOps::populate_op_mode_supported() { } { UnsupportedOpMode obj = {{V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - if (device_id_.find("GPU") != std::string::npos) { - auto x_data_type = node->InputDefs()[0]->TypeAsProto()->tensor_type().elem_type(); - auto y_data_type = node->InputDefs()[1]->TypeAsProto()->tensor_type().elem_type(); - return x_data_type != y_data_type; - } - //currently both inputs with int32 or int64 datatype are not supported - const bool A_is_int32 = node->InputDefs()[0]->Type()->find("int32") != std::string::npos; - const bool B_is_int32 = node->InputDefs()[1]->Type()->find("int32") != std::string::npos; - const bool A_is_int64 = node->InputDefs()[0]->Type()->find("int64") != std::string::npos; - const bool B_is_int64 = node->InputDefs()[1]->Type()->find("int64") != std::string::npos; - if((A_is_int32 && B_is_int32) || (A_is_int64 && B_is_int64)) - return true; - return false; - } - }; + [this](const Node* node, const InitializedTensorSet&) { + if (device_id_.find("GPU") != std::string::npos) { + auto x_data_type = node->InputDefs()[0]->TypeAsProto()->tensor_type().elem_type(); + auto y_data_type = node->InputDefs()[1]->TypeAsProto()->tensor_type().elem_type(); + return x_data_type != y_data_type; + } + //currently both inputs with int32 or int64 datatype are not supported + const bool A_is_int32 = node->InputDefs()[0]->Type()->find("int32") != std::string::npos; + const bool B_is_int32 = node->InputDefs()[1]->Type()->find("int32") != std::string::npos; + const bool A_is_int64 = node->InputDefs()[0]->Type()->find("int64") != std::string::npos; + const bool B_is_int64 = node->InputDefs()[1]->Type()->find("int64") != std::string::npos; + if ((A_is_int32 && B_is_int32) || (A_is_int64 && B_is_int64)) + return true; + return false; + }}; op_list_.insert({"Pow", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - auto slope = node->InputDefs()[1]; - //PRelu slope has to be an initializer or needs to come from a constant node - if (initializers.count(slope->Name())) - return false; - else { - for (auto input_node = node->InputNodesBegin(); input_node != node->InputNodesEnd(); ++input_node) { - if (GetInputCount(this->graph_viewer_.GetNode((*input_node).Index()), initializers) == 0) - return false; - } } - return true; - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + auto slope = node->InputDefs()[1]; + //PRelu slope has to be an initializer or needs to come from a constant node + if (initializers.count(slope->Name())) + return false; + else { + for (auto input_node = node->InputNodesBegin(); input_node != node->InputNodesEnd(); ++input_node) { + if (GetInputCount(this->graph_viewer_.GetNode((*input_node).Index()), initializers) == 0) + return false; + } + } + return true; + }}; op_list_.insert({"PRelu", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - bool non_const_zero_point = false; - // check if any of the zero points is NOT in the initializers list - non_const_zero_point |= initializers.find(node->InputDefs()[2]->Name()) == initializers.end(); - non_const_zero_point |= initializers.find(node->InputDefs()[5]->Name()) == initializers.end(); - non_const_zero_point |= initializers.find(node->InputDefs()[7]->Name()) == initializers.end(); - // QLinearMatMul is not supported if any of the zero points is a dynamic input - return non_const_zero_point; - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + bool non_const_zero_point = false; + // check if any of the zero points is NOT in the initializers list + non_const_zero_point |= initializers.find(node->InputDefs()[2]->Name()) == initializers.end(); + non_const_zero_point |= initializers.find(node->InputDefs()[5]->Name()) == initializers.end(); + non_const_zero_point |= initializers.find(node->InputDefs()[7]->Name()) == initializers.end(); + // QLinearMatMul is not supported if any of the zero points is a dynamic input + return non_const_zero_point; + }}; op_list_.insert({"QLinearMatMul", obj}); } { @@ -720,26 +700,25 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"Resize", obj}); } { - UnsupportedOpMode obj = {{V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - const auto& attributes = node->GetAttributes(); - auto axis_attr = attributes.find("axis"); - //Negative axis is not supported - if (axis_attr->second().i() < 0) - return true; - return false; - } - }; + UnsupportedOpMode obj = {{V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + const auto& attributes = node->GetAttributes(); + auto axis_attr = attributes.find("axis"); + //Negative axis is not supported + if (axis_attr->second().i() < 0) + return true; + return false; + }}; op_list_.insert({"Scatter", obj}); op_list_.insert({"ScatterElements", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1, V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& initializers) { - //start, end, axes need to be a initializer - bool cond_for_slice = false; - const auto& data_arg = node->InputDefs()[0]; - auto graph_inputs = graph_viewer_.GetInputs(); + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet& initializers) { + //start, end, axes need to be a initializer + bool cond_for_slice = false; + const auto& data_arg = node->InputDefs()[0]; + auto graph_inputs = graph_viewer_.GetInputs(); auto it = find(graph_inputs.begin(), graph_inputs.end(), data_arg); if (it != graph_inputs.end()) { @@ -760,15 +739,14 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"Slice", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - //Shape can't have empty axes attribute - const auto& attributes = node->GetAttributes(); - if (attributes.count("axes") == 0) - return true; - return false; - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //Shape can't have empty axes attribute + const auto& attributes = node->GetAttributes(); + if (attributes.count("axes") == 0) + return true; + return false; + }}; op_list_.insert({"Squeeze", obj}); } { @@ -779,25 +757,24 @@ void DataOps::populate_op_mode_supported() { op_list_.insert({"TopK", obj}); } { - UnsupportedOpMode obj = {{V_2020_4,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - return (!this->dimension_unsupported(node)); - } - }; + UnsupportedOpMode obj = {{V_2020_4, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + return (!this->dimension_unsupported(node)); + }}; op_list_.insert({"Unsqueeze", obj}); } { - UnsupportedOpMode obj = {{V_2021_1,V_2021_2, V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet& ) { - //check for attributes - auto& upsample_attr = node->GetAttributes(); - if (upsample_attr.count("scales") > 0) { - auto& upsample_arg = upsample_attr.at("scales"); - auto float_size = upsample_arg.floats_size(); - if (float_size > 2 && (upsample_arg.floats(0) != 1.f || upsample_arg.floats(1) != 1.f)) { - return true; - } - } + UnsupportedOpMode obj = {{V_2021_1, V_2021_2, V_2021_3}, + [this](const Node* node, const InitializedTensorSet&) { + //check for attributes + auto& upsample_attr = node->GetAttributes(); + if (upsample_attr.count("scales") > 0) { + auto& upsample_arg = upsample_attr.at("scales"); + auto float_size = upsample_arg.floats_size(); + if (float_size > 2 && (upsample_arg.floats(0) != 1.f || upsample_arg.floats(1) != 1.f)) { + return true; + } + } //check for input dimensions const auto& x_arg = node->InputDefs()[0]; @@ -821,22 +798,21 @@ void DataOps::populate_op_mode_supported() { } { UnsupportedOpMode obj = {{V_2021_2}, - [this](const Node* node, const Provider_InitializedTensorSet& ) { - //float data type is not supported - const bool data_is_float = node->InputDefs()[1]->Type()->find("float") != std::string::npos; - return data_is_float; - } }; + [this](const Node* node, const InitializedTensorSet&) { + //float data type is not supported + const bool data_is_float = node->InputDefs()[1]->Type()->find("float") != std::string::npos; + return data_is_float; + }}; op_list_.insert({"Where", obj}); } { UnsupportedOpMode obj = {{V_2021_3}, - [this](const Node* node, const Provider_InitializedTensorSet&) { - return (!this->dimension_unsupported(node)); - } - }; + [this](const Node* node, const InitializedTensorSet&) { + return (!this->dimension_unsupported(node)); + }}; op_list_.insert({"ReduceSum", obj}); } -} +} bool DataOps::op_is_supported(std::string name, std::vector& op_list) { for (size_t i = 0; i < op_list.size(); i++) { @@ -1113,23 +1089,23 @@ bool DataOps::IsOpSupportedOnlyInModel(std::string name) { } bool DataOps::SpecialConditionForClusterSizeOne(std::unordered_set& ng_required_initializers, const Node* node) { - if (node->OpType() == "Where") { - if (device_id_.find("MYRIAD") != std::string::npos) { - if (node->InputDefs()[1]->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT) - return true; - } - } else if (node->OpType() == "Reshape") { - const auto& shape_arg = node->InputDefs()[1]; - if (ng_required_initializers.find(shape_arg->Name()) == ng_required_initializers.end()) { - return true; - } - } else if (node->OpType() == "Expand") { - // nGraph only supports constant shape input values - const auto& output = node->OutputDefs()[0]; - if (output->TypeAsProto()->tensor_type().elem_type() != ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT16) - return true; - } else if (node->OpType() == "RoiAlign") { - using onnx_dtype = ONNX_NAMESPACE::TensorProto_DataType; + if (node->OpType() == "Where") { + if (device_id_.find("MYRIAD") != std::string::npos) { + if (node->InputDefs()[1]->TypeAsProto()->tensor_type().elem_type() == ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT) + return true; + } + } else if (node->OpType() == "Reshape") { + const auto& shape_arg = node->InputDefs()[1]; + if (ng_required_initializers.find(shape_arg->Name()) == ng_required_initializers.end()) { + return true; + } + } else if (node->OpType() == "Expand") { + // nGraph only supports constant shape input values + const auto& output = node->OutputDefs()[0]; + if (output->TypeAsProto()->tensor_type().elem_type() != ONNX_NAMESPACE::TensorProto_DataType::TensorProto_DataType_FLOAT16) + return true; + } else if (node->OpType() == "RoiAlign") { + using onnx_dtype = ONNX_NAMESPACE::TensorProto_DataType; onnx_dtype input_0_data_type = (ONNX_NAMESPACE::TensorProto_DataType)node->InputDefs()[0]->TypeAsProto()->tensor_type().elem_type(); onnx_dtype input_1_data_type = (ONNX_NAMESPACE::TensorProto_DataType)node->InputDefs()[1]->TypeAsProto()->tensor_type().elem_type(); diff --git a/onnxruntime/core/providers/shared_library/provider_bridge_provider.cc b/onnxruntime/core/providers/shared_library/provider_bridge_provider.cc index 7f5ceee439..7a6ad42177 100644 --- a/onnxruntime/core/providers/shared_library/provider_bridge_provider.cc +++ b/onnxruntime/core/providers/shared_library/provider_bridge_provider.cc @@ -352,6 +352,7 @@ std::unique_ptr CopyOpKernelInfo(const OpKernelInfo& info) { #ifdef ENABLE_TRAINING #include "orttraining/training_ops/cpu/controlflow/group.h" +#include "orttraining/training_ops/cpu/controlflow/yield.h" #endif namespace onnxruntime { @@ -453,6 +454,8 @@ Tensor* AttentionBase::GetPresent(OpKernelContext* context, const Tensor* past, return g_host->AttentionBase__GetPresent(this, context, past, batch_size, head_size, sequence_length, past_sequence_length); } +Status YieldOp::Compute(OpKernelContext* context) const { return g_host->YieldOp__Compute(this, context); } + } // namespace contrib #endif diff --git a/onnxruntime/core/providers/shared_library/provider_interfaces.h b/onnxruntime/core/providers/shared_library/provider_interfaces.h index 7848995c6e..14cc0d9276 100644 --- a/onnxruntime/core/providers/shared_library/provider_interfaces.h +++ b/onnxruntime/core/providers/shared_library/provider_interfaces.h @@ -38,6 +38,7 @@ class LongformerAttentionBase; class AttentionBase; class Group; class PassThrough; +class YieldOp; } // namespace contrib template @@ -423,6 +424,7 @@ struct ProviderHost { virtual void KernelDefBuilder__Alias(KernelDefBuilder* p, int input_index, int output_index) = 0; virtual void KernelDefBuilder__Alias(KernelDefBuilder* p, const std::vector>& aliases) = 0; virtual void KernelDefBuilder__VariadicAlias(KernelDefBuilder* p, int input_offset, int output_offset) = 0; + virtual void KernelDefBuilder__ExternalOutputs(KernelDefBuilder* p) = 0; virtual std::unique_ptr KernelDefBuilder__Build(KernelDefBuilder* p) = 0; @@ -572,6 +574,7 @@ struct ProviderHost { virtual bool OpKernelInfo__TryGetConstantInput(const OpKernelInfo* p, int input_index, const Tensor** constant_input_value) = 0; virtual uint32_t OpKernelInfo__GetInputCount(const OpKernelInfo* p) = 0; + virtual uint32_t OpKernelInfo__GetOutputCount(const OpKernelInfo* p) = 0; virtual const Node& OpKernelInfo__node(const OpKernelInfo* p) = 0; // SessionState @@ -753,6 +756,7 @@ struct ProviderHost { virtual void contrib__GetNDCFromLogitAndLabelShape(const TensorShape& logit_shape, const TensorShape& label_shape, int64_t& N_D, int64_t& C) = 0; virtual void contrib__GetPermutationAndShape(bool ncd_to_ndc, const TensorShape& tensor_shape, std::vector& new_shape, std::vector& permutations) = 0; virtual Status contrib__PrepareForTrainingCompute(const TensorShape& input_shape, int num_outputs, int64_t& axis, int& before_dims, int& after_dims_including_split_axis, int& after_dims_excluding_split, std::vector& split_sizes) = 0; + virtual Status YieldOp__Compute(const contrib::YieldOp* p, OpKernelContext* context) = 0; #endif }; @@ -1154,6 +1158,11 @@ struct KernelDefBuilder { return *this; } + KernelDefBuilder& ExternalOutputs() { + g_host->KernelDefBuilder__ExternalOutputs(this); + return *this; + } + std::unique_ptr Build() { return g_host->KernelDefBuilder__Build(this); } @@ -1445,6 +1454,7 @@ struct OpKernelInfo { const KernelDef& GetKernelDef() const { return g_host->OpKernelInfo__GetKernelDef(this); } uint32_t GetInputCount() const { return g_host->OpKernelInfo__GetInputCount(this); } + uint32_t GetOutputCount() const { return g_host->OpKernelInfo__GetOutputCount(this); } const Node& node() const noexcept { return g_host->OpKernelInfo__node(this); } diff --git a/onnxruntime/test/framework/test_utils.cc b/onnxruntime/test/framework/test_utils.cc index 83f6ed6c76..8c5414a4a8 100644 --- a/onnxruntime/test/framework/test_utils.cc +++ b/onnxruntime/test/framework/test_utils.cc @@ -70,4 +70,3 @@ std::map CountOpsInGraph(const Graph& graph, bool recurse_into } // namespace test } // namespace onnxruntime - diff --git a/onnxruntime/test/framework/test_utils.h b/onnxruntime/test/framework/test_utils.h index 48f72cc2d5..c332a435de 100644 --- a/onnxruntime/test/framework/test_utils.h +++ b/onnxruntime/test/framework/test_utils.h @@ -114,4 +114,3 @@ std::map CountOpsInGraph(const Graph& graph, bool recurse_into } // namespace test } // namespace onnxruntime - diff --git a/onnxruntime/test/providers/cuda/gemm_options_test.cc b/onnxruntime/test/providers/cuda/gemm_options_test.cc index 3dec06da25..426b1425e0 100644 --- a/onnxruntime/test/providers/cuda/gemm_options_test.cc +++ b/onnxruntime/test/providers/cuda/gemm_options_test.cc @@ -8,6 +8,10 @@ namespace onnxruntime { namespace cuda { + +// Initialize the singleton instance +HalfGemmOptions HalfGemmOptions::instance; + namespace test { TEST(CudaGemmOptionsTest, DefaultOptions) { diff --git a/orttraining/orttraining/models/bert/main.cc b/orttraining/orttraining/models/bert/main.cc index 97ddd4bf62..685a0972f0 100644 --- a/orttraining/orttraining/models/bert/main.cc +++ b/orttraining/orttraining/models/bert/main.cc @@ -610,7 +610,7 @@ void setup_training_params(BertParameters& params) { 0, nullptr}; - if (params.cuda_mem_limit_in_gb > 0) { + if (params.gpu_mem_limit_in_gb > 0) { info.gpu_mem_limit = gsl::narrow(params.gpu_mem_limit_in_gb * 1024 * 1024 * 1024); } info.cudnn_conv_algo_search = OrtCudnnConvAlgoSearch::EXHAUSTIVE; @@ -884,4 +884,3 @@ int main(int argc, char* argv[]) { #endif return 0; } - diff --git a/orttraining/orttraining/test/graph/gradient_graph_builder_test.cc b/orttraining/orttraining/test/graph/gradient_graph_builder_test.cc index e1873fac7a..e57a3e0dcd 100644 --- a/orttraining/orttraining/test/graph/gradient_graph_builder_test.cc +++ b/orttraining/orttraining/test/graph/gradient_graph_builder_test.cc @@ -18,7 +18,7 @@ #if defined(USE_CUDA) || defined(USE_ROCM) #include "bert_toy_fetches.h" -#elif USE_ROCM +#ifdef USE_ROCM #include "core/providers/rocm/rocm_execution_provider.h" #endif #endif @@ -1896,4 +1896,3 @@ TEST(GradientGraphBuilderTest, TrainingSession_WithPipeline) { } // namespace test } // namespace onnxruntime - diff --git a/orttraining/orttraining/test/session/training_session_test_utils.cc b/orttraining/orttraining/test/session/training_session_test_utils.cc index 129a918dff..eaaea01349 100644 --- a/orttraining/orttraining/test/session/training_session_test_utils.cc +++ b/orttraining/orttraining/test/session/training_session_test_utils.cc @@ -246,4 +246,3 @@ std::unique_ptr BuildAndRunTrainingSessionWithChecks( } // namespace training_session_test_utils } // namespace test } // namespace onnxruntime - diff --git a/orttraining/orttraining/training_ops/cuda/controlflow/yield.cc b/orttraining/orttraining/training_ops/cuda/controlflow/yield.cc index 32756afa8a..5ebbf0d1d6 100644 --- a/orttraining/orttraining/training_ops/cuda/controlflow/yield.cc +++ b/orttraining/orttraining/training_ops/cuda/controlflow/yield.cc @@ -1,6 +1,7 @@ // Copyright (c) Microsoft Corporation. All rights reserved. // Licensed under the MIT License. +#include "core/providers/shared_library/provider_api.h" #include "orttraining/training_ops/cpu/controlflow/yield.h" #include "core/providers/cuda/cuda_fwd.h" @@ -12,7 +13,7 @@ ONNX_OPERATOR_KERNEL_EX( kMSDomain, 1, kCudaExecutionProvider, - KernelDefBuilder() + (*KernelDefBuilder::Create()) .TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()) .ExternalOutputs(), onnxruntime::contrib::YieldOp); diff --git a/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc b/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc index ca742ad897..bff551e548 100644 --- a/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc +++ b/orttraining/orttraining/training_ops/cuda/nn/conv_grad.cc @@ -10,14 +10,14 @@ namespace onnxruntime { namespace cuda { -#define REGISTER_GRADIENT_KERNEL_TYPED(T) \ - ONNX_OPERATOR_TYPED_KERNEL_EX( \ - ConvGrad, \ - kMSDomain, \ - 1, \ - T, \ - kCudaExecutionProvider, \ - KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ +#define REGISTER_GRADIENT_KERNEL_TYPED(T) \ + ONNX_OPERATOR_TYPED_KERNEL_EX( \ + ConvGrad, \ + kMSDomain, \ + 1, \ + T, \ + kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()).TypeConstraint("T", DataTypeImpl::GetTensorType()), \ ConvGrad); REGISTER_GRADIENT_KERNEL_TYPED(float) @@ -226,4 +226,4 @@ Status ConvGrad::ComputeBiasGradient(Tensor* dB, const Tensor* dY) const { } } // namespace cuda -} // namespace onnxruntime \ No newline at end of file +} // namespace onnxruntime diff --git a/orttraining/orttraining/training_ops/cuda/nn/dropout_grad.cc b/orttraining/orttraining/training_ops/cuda/nn/dropout_grad.cc index 5bbfefb80e..3ce65f6a4d 100644 --- a/orttraining/orttraining/training_ops/cuda/nn/dropout_grad.cc +++ b/orttraining/orttraining/training_ops/cuda/nn/dropout_grad.cc @@ -9,17 +9,17 @@ namespace onnxruntime { namespace cuda { -#define REGISTER_GRADIENT_KERNEL(OpName) \ - ONNX_OPERATOR_KERNEL_EX( \ - OpName, \ - kMSDomain, \ - 1, \ - kCudaExecutionProvider, \ - KernelDefBuilder() \ - .TypeConstraint("T", ALL_IEEE_FLOAT_TENSOR_TYPES) \ - .TypeConstraint("T1", ALL_IEEE_FLOAT_TENSOR_TYPES) \ - .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ - .InputMemoryType(2), \ +#define REGISTER_GRADIENT_KERNEL(OpName) \ + ONNX_OPERATOR_KERNEL_EX( \ + OpName, \ + kMSDomain, \ + 1, \ + kCudaExecutionProvider, \ + (*KernelDefBuilder::Create()) \ + .TypeConstraint("T", ALL_IEEE_FLOAT_TENSOR_TYPES) \ + .TypeConstraint("T1", ALL_IEEE_FLOAT_TENSOR_TYPES) \ + .TypeConstraint("T2", DataTypeImpl::GetTensorType()) \ + .InputMemoryType(OrtMemTypeCPUInput, 2), \ DropoutGrad); REGISTER_GRADIENT_KERNEL(DropoutGrad) diff --git a/orttraining/orttraining/training_ops/cuda/tensor/concat.cc b/orttraining/orttraining/training_ops/cuda/tensor/concat.cc index a1b91e149d..103cc91892 100644 --- a/orttraining/orttraining/training_ops/cuda/tensor/concat.cc +++ b/orttraining/orttraining/training_ops/cuda/tensor/concat.cc @@ -10,7 +10,7 @@ ONNX_OPERATOR_KERNEL_EX(ConcatTraining, 1, kCudaExecutionProvider, (*KernelDefBuilder::Create()) - .OutputMemoryType(1) + .OutputMemoryType(OrtMemTypeCPUInput, 1) .TypeConstraint("T", DataTypeImpl::AllFixedSizeTensorTypes()), ConcatTraining); @@ -81,4 +81,3 @@ Status ConcatTraining::ComputeInternal(OpKernelContext* ctx) const { } } // namespace cuda } // namespace onnxruntime -