From f68a326bd94b711a68fbd98370f01e0174e04245 Mon Sep 17 00:00:00 2001 From: Dmitri Smirnov Date: Thu, 30 Apr 2020 15:29:39 -0700 Subject: [PATCH] Implement Pow(12) for cpu and cuda (#3727) * Implement Pow(12) cpu and cuda. --- .../providers/cpu/cpu_execution_provider.cc | 10 +- .../providers/cpu/math/element_wise_ops.cc | 148 ++++++++++++++---- .../providers/cpu/math/element_wise_ops.h | 1 - .../cuda/cu_inc/binary_elementwise_impl.cuh | 50 +++--- .../core/providers/cuda/cu_inc/common.cuh | 7 +- .../providers/cuda/cuda_execution_provider.cc | 16 +- .../cuda/math/binary_elementwise_ops.cc | 142 ++++++++++++++++- .../cuda/math/binary_elementwise_ops.h | 7 + .../cuda/math/binary_elementwise_ops_impl.cu | 71 +++++++-- .../cuda/math/binary_elementwise_ops_impl.h | 40 +++-- .../cpu/math/element_wise_ops_test.cc | 104 ++++++++++++ .../test/python/onnx_backend_test_series.py | 18 +-- 12 files changed, 507 insertions(+), 107 deletions(-) diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index c28e966899..279f1f225e 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -81,8 +81,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, float, Reciprocal); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, float, Sqrt); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, double, Sqrt); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, float, Pow); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, double, Pow); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, 11, Pow); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, float, Exp); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, double, Exp); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, float, Log); @@ -423,6 +422,8 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Mi class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Max); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool); +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Pow); + class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t, ReduceMax); @@ -513,8 +514,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1059,6 +1059,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo()), \ KERNEL_CLASS); -// var args are type constraints +// var args are type constraints for T and T1 #define REG_ELEMENTWISE_KERNEL_NONT(OP_TYPE, VERSION, KERNEL_CLASS, ...) \ - ONNX_CPU_OPERATOR_KERNEL( \ - OP_TYPE, \ - VERSION, \ - KernelDefBuilder() \ + ONNX_CPU_OPERATOR_KERNEL( \ + OP_TYPE, \ + VERSION, \ + KernelDefBuilder() \ .TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>()) \ .TypeConstraint("T1", BuildKernelDefConstraints<__VA_ARGS__>()), \ KERNEL_CLASS); -// var args are type constraints +// var args are type constraints for T and T1 #define REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(OP_TYPE, VERSION_FROM, VERSION_TO, KERNEL_CLASS, ...) \ - ONNX_CPU_OPERATOR_VERSIONED_KERNEL( \ - OP_TYPE, \ - VERSION_FROM, \ - VERSION_TO, \ - KernelDefBuilder() \ + ONNX_CPU_OPERATOR_VERSIONED_KERNEL( \ + OP_TYPE, \ + VERSION_FROM, \ + VERSION_TO, \ + KernelDefBuilder() \ .TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>()) \ .TypeConstraint("T1", BuildKernelDefConstraints<__VA_ARGS__>()), \ KERNEL_CLASS); @@ -114,8 +114,10 @@ REG_ELEMENTWISE_TYPED_KERNEL(Reciprocal, 6, float, Reciprocal); REG_ELEMENTWISE_TYPED_KERNEL(Sqrt, 6, float, Sqrt); REG_ELEMENTWISE_TYPED_KERNEL(Sqrt, 6, double, Sqrt); -REG_ELEMENTWISE_TYPED_KERNEL(Pow, 7, float, Pow); -REG_ELEMENTWISE_TYPED_KERNEL(Pow, 7, double, Pow); +REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(Pow, 7, 11, Pow, float, double); +// To reduce templetization we choose to support the below types for both +// base and the exponent. This gives us 16 permutations +REG_ELEMENTWISE_KERNEL_NONT(Pow, 12, Pow, int32_t, int64_t, float, double); REG_ELEMENTWISE_TYPED_KERNEL(Exp, 6, float, Exp); REG_ELEMENTWISE_TYPED_KERNEL(Exp, 6, double, Exp); @@ -273,25 +275,115 @@ Status Sqrt::Compute(OpKernelContext* ctx) const { return Status::OK(); } -template -Status Pow::Compute(OpKernelContext* context) const { - const Tensor& Y = *context->Input(1); - std::function, ConstEigenVectorMap, T)> input1scalar = - [](EigenVectorMap output, ConstEigenVectorMap input0, T input1) { output = Eigen::pow(input0.array(), input1); }; +namespace pow_internal { + +template +void PowImpl(OpKernelContext* context, const Tensor& X, const Tensor& Y) { + TBroadcaster bc{X, Y}; + Tensor* const output_tensor = context->Output(0, bc.GetOutputShape()); + TBroadcastOutput output{bc.GetSpanSize(), *output_tensor}; + + // Scalar base + auto input0scalar = [](gsl::span output, T X, gsl::span Y) { + std::transform(Y.cbegin(), Y.cend(), output.begin(), + [X](E y) { + return static_cast(std::pow(X, y)); + }); + }; + + // Scalar exponent switch to possibly available optimizations + std::function, gsl::span X, E Y)> input1scalar = + [](gsl::span output, gsl::span X, E Y) { + std::transform(X.cbegin(), X.cend(), output.begin(), + [Y](T x) { + return static_cast(std::pow(x, Y)); + }); + }; + if (Y.Shape().Size() == 1) { - T value = *Y.Data(); - if (value == 2.0) { - input1scalar = [](EigenVectorMap output, ConstEigenVectorMap input0, T) { output = Eigen::square(input0.array()); }; - } else if (value == 3.0) { - input1scalar = [](EigenVectorMap output, ConstEigenVectorMap input0, T) { output = Eigen::cube(input0.array()); }; + auto exp = *Y.template Data(); + if (exp == E{2}) { + input1scalar = [](gsl::span output, gsl::span X, E) { + std::transform(X.cbegin(), X.cend(), output.begin(), + [](T x) { + return static_cast(x * x); + }); + }; + } else if (exp == E{3}) { + input1scalar = [](gsl::span output, gsl::span X, E) { + std::transform(X.cbegin(), X.cend(), output.begin(), + [](T x) { + return static_cast(x * x * x); + }); + }; } } - return BroadcastTwo( - *context, - [](EigenVectorMap output, T input0, ConstEigenVectorMap input1) { output = Eigen::pow(input0, input1.array()); }, - input1scalar, - [](EigenVectorMap output, ConstEigenVectorMap input0, ConstEigenVectorMap input1) { output = Eigen::pow(input0.array(), input1.array()); }); + auto general = [](gsl::span output, gsl::span X, gsl::span Y) { + std::transform( + X.cbegin(), X.cend(), Y.cbegin(), output.begin(), + [](T x, E y) { + return static_cast(std::pow(x, y)); + }); + }; + + BroadcastLoopSpan(bc, output, input0scalar, input1scalar, general); +} + +template +Status DispatchOnBase(OpKernelContext* context, const Tensor& X, const Tensor& Y) { + namespace on = ONNX_NAMESPACE; + Status s; + switch (Y.GetElementType()) { + case on::TensorProto_DataType_INT32: + PowImpl(context, X, Y); + break; + case on::TensorProto_DataType_INT64: + PowImpl(context, X, Y); + break; + case on::TensorProto_DataType_FLOAT: + PowImpl(context, X, Y); + break; + case on::TensorProto_DataType_DOUBLE: + PowImpl(context, X, Y); + break; + default: + s = ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Unsupported Y type: ", + DataTypeImpl::ToString(Y.DataType())); + } + return s; +} + +} // namespace pow_internal + +Status +Pow::Compute(OpKernelContext* context) const { + const Tensor& X = *context->Input(0); + const Tensor& Y = *context->Input(1); + + namespace on = ONNX_NAMESPACE; + using namespace pow_internal; + + Status s; + // Switch on base type first + switch (X.GetElementType()) { + case on::TensorProto_DataType_INT32: + s = DispatchOnBase(context, X, Y); + break; + case on::TensorProto_DataType_INT64: + s = DispatchOnBase(context, X, Y); + break; + case on::TensorProto_DataType_FLOAT: + s = DispatchOnBase(context, X, Y); + break; + case on::TensorProto_DataType_DOUBLE: + s = DispatchOnBase(context, X, Y); + break; + default: + s = ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Unsupported X type: ", + DataTypeImpl::ToString(X.DataType())); + } + return s; } template diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.h b/onnxruntime/core/providers/cpu/math/element_wise_ops.h index 79e53f5b9e..217104d462 100644 --- a/onnxruntime/core/providers/cpu/math/element_wise_ops.h +++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.h @@ -111,7 +111,6 @@ class Sqrt final : public OpKernel { Status Compute(OpKernelContext* context) const override; }; -template class Pow final : public OpKernel { public: Pow(const OpKernelInfo& info) : OpKernel(info) { diff --git a/onnxruntime/core/providers/cuda/cu_inc/binary_elementwise_impl.cuh b/onnxruntime/core/providers/cuda/cu_inc/binary_elementwise_impl.cuh index 2a00a04aa7..67c3012a6c 100644 --- a/onnxruntime/core/providers/cuda/cu_inc/binary_elementwise_impl.cuh +++ b/onnxruntime/core/providers/cuda/cu_inc/binary_elementwise_impl.cuh @@ -10,20 +10,20 @@ namespace onnxruntime { namespace cuda { // broadcast by computing output coordinate from offset, using fast_divmod -template +template __global__ void _BinaryElementWise( int32_t output_rank, const TArray lhs_padded_strides, const T* lhs_data, const TArray rhs_padded_strides, - const T* rhs_data, + const T1* rhs_data, const TArray fdm_output_strides, T* output_data, const FuncT& functor, CUDA_LONG N) { CUDA_LONG start = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x; T lvalue[NumElementsPerThread]; - T rvalue[NumElementsPerThread]; + T1 rvalue[NumElementsPerThread]; CUDA_LONG id = start; #pragma unroll @@ -68,16 +68,16 @@ __global__ void _BinaryElementWise( } // for scalar broadcast or non-broadcast case -template +template __global__ void _BinaryElementWiseSimple( const T* lhs_data, - const T* rhs_data, + const T1* rhs_data, T* output_data, const FuncT& func, CUDA_LONG N) { CUDA_LONG start = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x; T lvalue[NumElementsPerThread]; - T rvalue[NumElementsPerThread]; + T1 rvalue[NumElementsPerThread]; CUDA_LONG id = start; #pragma unroll @@ -102,17 +102,17 @@ __global__ void _BinaryElementWiseSimple( } // for rhs per-channel broadcast case -template +template __global__ void _BinaryElementWiseRhsPerChannelBatch1( const T* lhs_data, - const T* rhs_data, + const T1* rhs_data, const fast_divmod fdm_H, T* output_data, FuncT func, CUDA_LONG N) { CUDA_LONG start = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x; T lvalue[NumElementsPerThread]; - T rvalue[NumElementsPerThread]; + T1 rvalue[NumElementsPerThread]; CUDA_LONG id = start; #pragma unroll @@ -137,10 +137,10 @@ __global__ void _BinaryElementWiseRhsPerChannelBatch1( } } -template +template __global__ void _BinaryElementWiseRhsPerChannelBatchN( const T* lhs_data, - const T* rhs_data, + const T1* rhs_data, const fast_divmod fdm_H, const fast_divmod fdm_C, T* output_data, @@ -148,7 +148,7 @@ __global__ void _BinaryElementWiseRhsPerChannelBatchN( CUDA_LONG N) { CUDA_LONG start = NumElementsPerThread * NumThreadsPerBlock * blockIdx.x + threadIdx.x; T lvalue[NumElementsPerThread]; - T rvalue[NumElementsPerThread]; + T1 rvalue[NumElementsPerThread]; CUDA_LONG id = start; #pragma unroll @@ -177,10 +177,10 @@ __global__ void _BinaryElementWiseRhsPerChannelBatchN( } } -template +template void BinaryElementWiseNoBroadcastImpl( const T* lhs_data, - const T* rhs_data, + const T1* rhs_data, T* output_data, const FuncT& func, size_t count) { @@ -189,7 +189,7 @@ void BinaryElementWiseNoBroadcastImpl( int blocksPerGrid = static_cast(CeilDiv(count, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread)); CUDA_LONG N = static_cast(count); - _BinaryElementWiseSimple<<>>( + _BinaryElementWiseSimple<<>>( lhs_data, rhs_data, output_data, @@ -197,13 +197,13 @@ void BinaryElementWiseNoBroadcastImpl( N); } -template +template void BinaryElementWiseImpl( int32_t output_rank_or_simple_broadcast, const TArray* lhs_padded_strides, const T* lhs_data, const TArray* rhs_padded_strides, - const T* rhs_data, + const T1* rhs_data, const TArray* fdm_output_strides, const fast_divmod& fdm_H, const fast_divmod& fdm_C, @@ -216,21 +216,21 @@ void BinaryElementWiseImpl( int blocksPerGrid = static_cast(CeilDiv(count, GridDim::maxThreadsPerBlock * GridDim::maxElementsPerThread)); CUDA_LONG N = static_cast(count); if (output_rank_or_simple_broadcast == static_cast(SimpleBroadcast::NoBroadcast)) { - _BinaryElementWiseSimple<<>>( + _BinaryElementWiseSimple<<>>( lhs_data, rhs_data, output_data, func, N); } else if (output_rank_or_simple_broadcast == static_cast(SimpleBroadcast::LeftScalar)) { - _BinaryElementWiseSimple<<>>( + _BinaryElementWiseSimple<<>>( lhs_data, rhs_data, output_data, func, N); } else if (output_rank_or_simple_broadcast == static_cast(SimpleBroadcast::RightScalar)) { - _BinaryElementWiseSimple<<>>( lhs_data, rhs_data, @@ -238,7 +238,7 @@ void BinaryElementWiseImpl( func, N); } else if (output_rank_or_simple_broadcast == static_cast(SimpleBroadcast::RightPerChannelBatch1)) { - _BinaryElementWiseRhsPerChannelBatch1<<>>( + _BinaryElementWiseRhsPerChannelBatch1<<>>( lhs_data, rhs_data, fdm_H, @@ -246,7 +246,7 @@ void BinaryElementWiseImpl( func, N); } else if (output_rank_or_simple_broadcast == static_cast(SimpleBroadcast::RightPerChannelBatchN)) { - _BinaryElementWiseRhsPerChannelBatchN<<>>( + _BinaryElementWiseRhsPerChannelBatchN<<>>( lhs_data, rhs_data, fdm_H, @@ -256,7 +256,7 @@ void BinaryElementWiseImpl( N); } else { if (lhs_padded_strides && rhs_padded_strides && lhs_padded_strides->size_ && rhs_padded_strides->size_) - _BinaryElementWise<<>>( + _BinaryElementWise<<>>( output_rank_or_simple_broadcast, *lhs_padded_strides, lhs_data, @@ -267,7 +267,7 @@ void BinaryElementWiseImpl( func, N); else if (lhs_padded_strides && lhs_padded_strides->size_) - _BinaryElementWise<<>>( + _BinaryElementWise<<>>( output_rank_or_simple_broadcast, *lhs_padded_strides, lhs_data, @@ -278,7 +278,7 @@ void BinaryElementWiseImpl( func, N); else - _BinaryElementWise<<>>( + _BinaryElementWise<<>>( output_rank_or_simple_broadcast, *lhs_padded_strides, lhs_data, diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh index 96395a947d..b1794d4276 100644 --- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh +++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh @@ -181,8 +181,11 @@ __device__ __inline__ half2 _Tanh(half2 a) { return __float22half2_rn(tmp); } -template -__device__ __inline__ T _Pow(T a, T b); +// Capture permutations of int32/64/float/double +template +__device__ __inline__ T _Pow(T a, T1 b) { + return static_cast(pow(static_cast(a), static_cast(b))); +} template <> __device__ __inline__ float _Pow(float a, float b) { return powf(a, b); } diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 1611b6c5b4..b7c9fb22c1 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -313,9 +313,9 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, 10, float, Softmax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, 10, double, Softmax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, 10, MLFloat16, Softmax); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, float, Pow); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, double, Pow); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, MLFloat16, Pow); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 11, float, Pow); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 11, double, Pow); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 11, MLFloat16, Pow); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, float, PRelu); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, double, PRelu); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, MLFloat16, PRelu); @@ -755,6 +755,8 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, MaxPool); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, MaxPool); +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, Pow); + class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, ReduceMax); @@ -829,9 +831,9 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1270,6 +1272,8 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index dc57ab1f5c..e8eeb4bd5e 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -85,7 +85,7 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin return Status::OK(); } -#define BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(x, ver, T) \ +#define BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED_V(x, class_name, ver, T) \ ONNX_OPERATOR_TYPED_KERNEL_EX( \ x, \ kOnnxDomain, \ @@ -93,7 +93,19 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin T, \ kCudaExecutionProvider, \ KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ - x); + class_name); + +#define BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(x, ver, T) \ + BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED_V(x, x, ver, T) + +#define BINARY_ELEMENTWISE_REGISTER_KERNEL_NONTEMP(x, class_name, ver, ...) \ + ONNX_OPERATOR_TYPED_KERNEL_EX( \ + x, \ + kOnnxDomain, \ + ver, \ + kCudaExecutionProvider, \ + KernelDefBuilder().TypeConstraint("T", BuildKernelDefConstraints<>(__VAR_ARGS__)), \ + class_name); #define BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(x, ver, T) \ ONNX_OPERATOR_TYPED_KERNEL_EX( \ @@ -116,6 +128,17 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ x); +#define BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED_CLASS(x, class_name, startver, endver, T) \ + ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \ + x, \ + kOnnxDomain, \ + startver, \ + endver, \ + T, \ + kCudaExecutionProvider, \ + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + class_name); + #define BINARY_ELEMENTWISE_COMPUTE(x, T) \ template <> \ Status x::ComputeInternal(OpKernelContext* context) const { \ @@ -139,6 +162,10 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin BINARY_ELEMENTWISE_REGISTER_KERNEL_TYPED(name, ver, T) \ BINARY_ELEMENTWISE_COMPUTE(name, T) +#define BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, T) \ + BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED_CLASS(name, class_name, startver, endver, T) \ + BINARY_ELEMENTWISE_COMPUTE(class_name, T) + #define BINARY_LOGICALOP_TYPED(name, ver, T) \ BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(name, ver, T) \ BINARY_ELEMENTWISE_COMPUTE(name, T) @@ -211,6 +238,11 @@ Status BinaryElementwise::Prepare(OpKernelContext* context, Bin BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, float) \ BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, double) +#define BINARY_OP_REGISTER_VERSIONED_CLASS_HFD(name, class_name, startver, endver) \ + BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, MLFloat16) \ + BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, float) \ + BINARY_OP_TYPED_VERSIONED_V(name, class_name, startver, endver, double) + #define BINARY_OP_REGISTER_VERSIONED_UZILHFD(name, startver, endver) \ BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, uint32_t) \ BINARY_ELEMENTWISE_REGISTER_KERNEL_VERSIONED_TYPED(name, startver, endver, uint64_t) \ @@ -222,12 +254,116 @@ BINARY_OP_UZILHFD(Add, 7) BINARY_OP_UZILHFD(Sub, 7) BINARY_OP_UZILHFD(Mul, 7) BINARY_OP_UZILHFD(Div, 7) -BINARY_OP_HFD(Pow, 7) +BINARY_OP_REGISTER_VERSIONED_CLASS_HFD(Pow, Pow_7, 7, 11) BINARY_LOGICALOP_TYPED(And, 7, bool) BINARY_LOGICALOP_TYPED(Or, 7, bool) BINARY_LOGICALOP_TYPED(Xor, 7, bool) BINARY_OP_HFD(PRelu, 7) +// Pow version 12 +ONNX_OPERATOR_KERNEL_EX( + Pow, + kOnnxDomain, + 12, + kCudaExecutionProvider, + KernelDefBuilder().TypeConstraint("T", BuildKernelDefConstraints()) + .TypeConstraint("T1", BuildKernelDefConstraints()), + Pow); + +namespace pow12_internal { +template +Status DispatchOnFirstArg(const BinaryElementwisePreparation& prepare) { + namespace on = ONNX_NAMESPACE; + Status s; + switch (prepare.rhs_tensor->GetElementType()) { + case on::TensorProto_DataType_INT32: + ImplT1_Pow::MappedType, typename ToCudaType::MappedType>( + prepare.output_rank_or_simple_broadcast, + &prepare.lhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.lhs_tensor->template Data()), + &prepare.rhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.rhs_tensor->template Data()), + &prepare.fdm_output_strides, + prepare.fdm_H, + prepare.fdm_C, + reinterpret_cast::MappedType*>(prepare.output_tensor->template MutableData()), + prepare.output_tensor->Shape().Size()); + break; + case on::TensorProto_DataType_INT64: + ImplT1_Pow::MappedType, typename ToCudaType::MappedType>( + prepare.output_rank_or_simple_broadcast, + &prepare.lhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.lhs_tensor->template Data()), + &prepare.rhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.rhs_tensor->template Data()), + &prepare.fdm_output_strides, + prepare.fdm_H, + prepare.fdm_C, + reinterpret_cast::MappedType*>(prepare.output_tensor->template MutableData()), + prepare.output_tensor->Shape().Size()); + break; + case on::TensorProto_DataType_FLOAT: + ImplT1_Pow::MappedType, typename ToCudaType::MappedType>( + prepare.output_rank_or_simple_broadcast, + &prepare.lhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.lhs_tensor->template Data()), + &prepare.rhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.rhs_tensor->template Data()), + &prepare.fdm_output_strides, + prepare.fdm_H, + prepare.fdm_C, + reinterpret_cast::MappedType*>(prepare.output_tensor->template MutableData()), + prepare.output_tensor->Shape().Size()); + break; + case on::TensorProto_DataType_DOUBLE: + ImplT1_Pow::MappedType, typename ToCudaType::MappedType>( + prepare.output_rank_or_simple_broadcast, + &prepare.lhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.lhs_tensor->template Data()), + &prepare.rhs_padded_strides, + reinterpret_cast::MappedType*>(prepare.rhs_tensor->template Data()), + &prepare.fdm_output_strides, + prepare.fdm_H, + prepare.fdm_C, + reinterpret_cast::MappedType*>(prepare.output_tensor->template MutableData()), + prepare.output_tensor->Shape().Size()); + break; + default: + s = ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Unsupported Y type: ", + DataTypeImpl::ToString(prepare.rhs_tensor->DataType())); + } + return s; +} +} // namespace pow12_internal + +Status Pow::ComputeInternal(OpKernelContext* context) const { + BinaryElementwisePreparation prepare; + Prepare(context, &prepare); + namespace on = ONNX_NAMESPACE; + using namespace pow12_internal; + + Status s; + + switch (prepare.lhs_tensor->GetElementType()) { + case on::TensorProto_DataType_INT32: + s = DispatchOnFirstArg(prepare); + break; + case on::TensorProto_DataType_INT64: + s = DispatchOnFirstArg(prepare); + break; + case on::TensorProto_DataType_FLOAT: + s = DispatchOnFirstArg(prepare); + break; + case on::TensorProto_DataType_DOUBLE: + s = DispatchOnFirstArg(prepare); + break; + default: + s = ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Unsupported X type: ", + DataTypeImpl::ToString(prepare.lhs_tensor->DataType())); + } + return s; +} + template Status VariadicInputBase::ComputeMethod(OpKernelContext* context, ImplCompute Impl_Compute) const { const auto& node = Node(); diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h index f9ff9e4cfb..44a43c26ff 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h @@ -164,6 +164,13 @@ class Div final : public BinaryElementwise { }; template +class Pow_7 final : public BinaryElementwise { + public: + Pow_7(const OpKernelInfo& info) : BinaryElementwise(info) {} + Status ComputeInternal(OpKernelContext* context) const override; +}; + +// Since version 12 class Pow final : public BinaryElementwise { public: Pow(const OpKernelInfo& info) : BinaryElementwise(info) {} diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu index ad4d428313..35f483f040 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu @@ -9,12 +9,12 @@ namespace onnxruntime { namespace cuda { -#define OP(name, expr) \ - template \ - struct OP_##name { \ - __device__ __inline__ T operator()(T a, T b) const { \ - return (expr); \ - } \ +#define OP(name, expr) \ + template \ + struct OP_##name { \ + __device__ __inline__ T operator()(T a, T1 b) const { \ + return (expr); \ + } \ }; #define BINARY_ELEMENTWISE_IMPL(name) \ @@ -28,15 +28,36 @@ namespace cuda { fdm_H, \ fdm_C, \ output_data, \ - OP_##name(), \ + OP_##name(), \ count); \ } -#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, T) \ - template void Impl_##x(int32_t output_rank, \ - const TArray* lhs_padded_strides, const T* lhs_data, \ - const TArray* rhs_padded_strides, const T* rhs_data, \ - const TArray* fdm_output_strides, const fast_divmod& fdm_H, const fast_divmod& fdm_C, T* output_data, size_t count); +#define BINARY_ELEMENTWISE_IMPL_T1(name) \ + BINARY_ELEMENTWISE_IMPL_DECLARATION_T1(name) { \ + BinaryElementWiseImpl(output_rank_or_simple_broadcast, \ + lhs_padded_strides, \ + lhs_data, \ + rhs_padded_strides, \ + rhs_data, \ + fdm_output_strides, \ + fdm_H, \ + fdm_C, \ + output_data, \ + OP_##name(), \ + count); \ + } + +#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, T) \ + template void Impl_##x(int32_t output_rank, \ + const TArray* lhs_padded_strides, const T* lhs_data, \ + const TArray* rhs_padded_strides, const T* rhs_data, \ + const TArray* fdm_output_strides, const fast_divmod& fdm_H, const fast_divmod& fdm_C, T* output_data, size_t count); + +#define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(x, T, T1) \ + template void ImplT1_##x(int32_t output_rank, \ + const TArray* lhs_padded_strides, const T* lhs_data, \ + const TArray* rhs_padded_strides, const T1* rhs_data, \ + const TArray* fdm_output_strides, const fast_divmod& fdm_H, const fast_divmod& fdm_C, T* output_data, size_t count); #define SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(x) \ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(x, uint32_t) \ @@ -85,7 +106,7 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Add, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Sub) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Mul) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Div) -SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Pow) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Pow_7) SPECIALIZED_BINARY_ELEMENTWISE_IMPL(And, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Or, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool) @@ -96,5 +117,29 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Max) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Min) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Less) +// create declarations for op and impl for Pow +OP(Pow, _Pow(a, b)) +BINARY_ELEMENTWISE_IMPL_T1(Pow) + +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int32_t, int32_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int32_t, int64_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int32_t, float) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int32_t, double) + +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int64_t, int32_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int64_t, int64_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int64_t, float) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, int64_t, double) + +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, float, int32_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, float, int64_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, float, float) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, float, double) + +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, double, int32_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, double, int64_t) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, double, float) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T1(Pow, double, double) + } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h index ac717524b1..ec47fadb86 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h @@ -20,7 +20,7 @@ namespace cuda { BINARY_OP_NAME_EXPR(Sub, (a - b)) \ BINARY_OP_NAME_EXPR(Mul, (a * b)) \ BINARY_OP_NAME_EXPR(Div, (a / b)) \ - BINARY_OP_NAME_EXPR(Pow, _Pow(a, b)) \ + BINARY_OP_NAME_EXPR(Pow_7, _Pow(a, b)) \ BINARY_OP_NAME_EXPR(And, (a & b)) \ BINARY_OP_NAME_EXPR(Or, (a | b)) \ BINARY_OP_NAME_EXPR(Xor, (a ^ b)) \ @@ -34,23 +34,39 @@ namespace cuda { // NOTE that cu files are compiled with nvcc and should not refer to any onnxruntime headers // so struct BinaryElementwisePreparation cannot be used here -#define BINARY_ELEMENTWISE_IMPL_DECLARATION(name) \ - template \ - void Impl_##name( \ - int32_t output_rank_or_simple_broadcast, \ - const TArray* lhs_padded_strides, \ - const T* lhs_data, \ - const TArray* rhs_padded_strides, \ - const T* rhs_data, \ +#define BINARY_ELEMENTWISE_IMPL_DECLARATION(name) \ + template \ + void Impl_##name( \ + int32_t output_rank_or_simple_broadcast, \ + const TArray* lhs_padded_strides, \ + const T* lhs_data, \ + const TArray* rhs_padded_strides, \ + const T* rhs_data, \ const TArray* fdm_output_strides, \ - const fast_divmod& fdm_H, \ - const fast_divmod& fdm_C, \ - T* output_data, \ + const fast_divmod& fdm_H, \ + const fast_divmod& fdm_C, \ + T* output_data, \ size_t count) #define BINARY_OP_NAME_EXPR(name, expr) BINARY_ELEMENTWISE_IMPL_DECLARATION(name); BINARY_OPS() #undef BINARY_OP_NAME_EXPR +#define BINARY_ELEMENTWISE_IMPL_DECLARATION_T1(name) \ + template \ + void ImplT1_##name( \ + int32_t output_rank_or_simple_broadcast, \ + const TArray* lhs_padded_strides, \ + const T* lhs_data, \ + const TArray* rhs_padded_strides, \ + const T1* rhs_data, \ + const TArray* fdm_output_strides, \ + const fast_divmod& fdm_H, \ + const fast_divmod& fdm_C, \ + T* output_data, \ + size_t count) + +BINARY_ELEMENTWISE_IMPL_DECLARATION_T1(Pow); + } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index ed6b8ea735..2199ca9083 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -621,6 +621,110 @@ TEST(MathOpTest, Pow_Broadcast_Scalar1) { test.Run(); } +TEST(MathOpTest, Pow_Float_12) { + OpTester test("Pow", 12); + std::vector dims{2, 2}; + test.AddInput("X", dims, + {2.0f, 2.0f, + std::sqrt(2.0f), 1.0f}); + test.AddInput("Y", dims, + {0.0f, 8.0f, + 2.0f, 9.0f}); + test.AddOutput("Z", dims, + {1.0f, 256.0f, + 2.0f, 1.0f}); + test.Run(); +} + +TEST(MathOpTest, Pow_Double_12) { + OpTester test("Pow", 12); + std::vector dims{2, 2}; + test.AddInput("X", dims, + {2.0, 2.0, + std::sqrt(2.0), 1.0}); + test.AddInput("Y", dims, + {0.0, 8.0, + 2.0, 9.0}); + test.AddOutput("Z", dims, + {1.0, 256.0, + 2.0, 1.0}); + test.Run(); +} + +TEST(MathOpTest, Pow_Broadcast_Scalar0_12) { + OpTester test("Pow", 12); + + std::vector dims{3}; + test.AddInput("X", {}, {2.0f}); + test.AddInput("Y", dims, {1.0f, 2.0f, 3.0f}); + test.AddOutput("Z", dims, {2.0f, 4.0f, 8.0f}); + test.Run(); +} + +TEST(MathOpTest, Pow_Broadcast_Scalar1_12) { + OpTester test("Pow", 12); + + std::vector dims{3}; + test.AddInput("X", dims, {1.0f, 2.0f, 3.0f}); + test.AddInput("Y", {}, {2.0f}); + test.AddOutput("Z", dims, {1.0f, 4.0f, 9.0f}); + test.Run(); +} + +TEST(MathOpTest, Pow_float_int64) { + OpTester test("Pow", 12); + std::vector dims{3}; + test.AddInput("X", dims, {1.0f, 2.0f, 3.0f}); + test.AddInput("Y", dims, {4, 5, 6}); + test.AddOutput("Z", dims, {1.f, 32.f, 729.f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); +} + +TEST(MathOpTest, Pow_int64_float) { + OpTester test("Pow", 12); + std::vector dims{3}; + test.AddInput("X", dims, {1, 2, 3}); + test.AddInput("Y", dims, {4.f, 5.f, 6.f}); + test.AddOutput("Z", dims, {1, 32, 729}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); +} + +TEST(MathOpTest, Pow_float_int32) { + OpTester test("Pow", 12); + std::vector dims{3}; + test.AddInput("X", dims, {1.0f, 2.0f, 3.0f}); + test.AddInput("Y", dims, {4, 5, 6}); + test.AddOutput("Z", dims, {1.f, 32.f, 729.f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); +} + +TEST(MathOpTest, Pow_int32_float) { + OpTester test("Pow", 12); + std::vector dims{3}; + test.AddInput("X", dims, {1, 2, 3}); + test.AddInput("Y", dims, {4.f, 5.f, 6.f}); + test.AddOutput("Z", dims, {1, 32, 729}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); +} + +TEST(MathOpTest, Pow_int64_double) { + OpTester test("Pow", 12); + std::vector dims{3}; + test.AddInput("X", dims, {1, 2, 3}); + test.AddInput("Y", dims, {4.f, 5.f, 6.f}); + test.AddOutput("Z", dims, {1, 32, 729}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); +} + +TEST(MathOpTest, Pow_double_int64) { + OpTester test("Pow", 12); + std::vector dims{3}; + test.AddInput("X", dims, {1., 2., 3.}); + test.AddInput("Y", dims, {4, 5, 6}); + test.AddOutput("Z", dims, {1., 32., 729.}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider}); +} + TEST(MathOpTest, Exp_float) { OpTester test("Exp"); std::vector dims{2, 2}; diff --git a/onnxruntime/test/python/onnx_backend_test_series.py b/onnxruntime/test/python/onnx_backend_test_series.py index 5b9d8760ae..3cd8c88460 100644 --- a/onnxruntime/test/python/onnx_backend_test_series.py +++ b/onnxruntime/test/python/onnx_backend_test_series.py @@ -132,20 +132,8 @@ def create_backend_test(testname=None): '^test_negative_log_likelihood_loss_input_shape_is_NCd1d2_with_weight_reduction_sum_cpu', '^test_negative_log_likelihood_loss_input_shape_is_NCd1d2_with_weight_reduction_sum_ignore_index_cpu', '^test_nesterov_momentum_cpu', - '^test_pow_bcast_array_cpu', - '^test_pow_bcast_scalar_cpu', - '^test_pow_cpu', - '^test_pow_example_cpu', - '^test_pow_types_float32_int32_cpu', - '^test_pow_types_float32_int64_cpu', '^test_pow_types_float32_uint32_cpu', '^test_pow_types_float32_uint64_cpu', - '^test_pow_types_float_cpu', - '^test_pow_types_int32_float32_cpu', - '^test_pow_types_int32_int32_cpu', - '^test_pow_types_int64_float32_cpu', - '^test_pow_types_int64_int64_cpu', - '^test_pow_types_int_cpu', '^test_softmax_cross_entropy_mean_3d_cpu', '^test_softmax_cross_entropy_mean_3d_expanded_cpu', '^test_softmax_cross_entropy_mean_cpu', @@ -182,7 +170,11 @@ def create_backend_test(testname=None): '^test_argmin_negative_axis.*', '^test_hardmax_negative_axis.*', '^test_gemm_default_no_bias_cpu', '^test_flatten_negative_axis.*', '^test_reduce_[a-z1-9_]*_negative_axes_.*', 'test_squeeze_negative_axes_cpu', 'test_unsqueeze_negative_axes_cpu', 'test_constant_pad_cpu', - 'test_edge_pad_cpu', 'test_reflect_pad_cpu', '^test_split_zero_size_splits_.*','^test_argmax_keepdims_example_select_last_index_.*', '^test_argmax_no_keepdims_example_select_last_index_.*','^test_argmin_no_keepdims_example_select_last_index_.*','^test_argmin_keepdims_example_select_last_index_.*' + 'test_edge_pad_cpu', 'test_reflect_pad_cpu', '^test_split_zero_size_splits_.*', + '^test_argmax_keepdims_example_select_last_index_cpu', '^test_argmax_no_keepdims_example_select_last_index_cpu', + '^test_argmin_no_keepdims_example_select_last_index_cpu','^test_argmin_keepdims_example_select_last_index_cpu', + '^test_pow_types_float32_int32_*', '^test_pow_types_float32_int64_*', '^test_pow_types_float_*', + '^test_pow_types_int32_float32_*', '^test_pow_types_int_*', '^test_pow_types_int64_float32_*' ] if c2.supports_device('DNNL'):