diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 21c4c2cb42..23e665244b 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -323,6 +323,12 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Sum); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Sum); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Sum); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Max); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Max); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Max); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Max); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Max); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Max); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, float, Greater); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, double, Greater); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, MLFloat16, Greater); @@ -606,6 +612,12 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + 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 c58c3c782b..8ab48ef15a 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -245,6 +245,62 @@ Status Sum::ComputeInternal(OpKernelContext* context) const { return Status::OK(); } +template +Status Max::ComputeInternal(OpKernelContext* context) const { + typedef typename ToCudaType::MappedType CudaT; + const auto& node = Node(); + const auto& node_name = node.Name(); + auto input_count = node.InputArgCount().front(); + ORT_RETURN_IF_NOT(input_count >= 1, "Must have 1 or more inputs"); + + if (input_count == 1) { + auto input_tensor = context->Input(0); + const auto& input_shape = input_tensor->Shape(); + auto output_tensor = context->Output(0, input_shape); + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(output_tensor->MutableDataRaw(), input_tensor->DataRaw(), sizeof(CudaT) * input_shape.Size(), cudaMemcpyDeviceToDevice)); + } else { + // compute output shape first, using broadcast rule + TensorShape output_shape; + ORT_RETURN_IF_ERROR(ComputeOutputShape(node_name, context->Input(0)->Shape(), context->Input(1)->Shape(), output_shape)); + for (int index = 2; index < input_count; index++) { + TensorShape previous_output_shape = output_shape; + ORT_RETURN_IF_ERROR(ComputeOutputShape(node_name, previous_output_shape, context->Input(index)->Shape(), output_shape)); + } + Tensor* output_tensor = context->Output(0, output_shape); + BinaryElementwisePreparation prepare(this); + + // More than 2 inputs, set output to 0, add input0 to output, so that input0 can be broadcast with output shape correctly + CUDA_RETURN_IF_ERROR(cudaMemset(output_tensor->MutableDataRaw(), 0, output_shape.Size() * sizeof(CudaT))); + ORT_RETURN_IF_ERROR(BinaryElementwiseBroadcastPrepare(0, output_tensor, context->Input(0), output_tensor, &prepare)); + Impl_Add( + prepare.output_rank_or_simple_broadcast, + prepare.lhs_padded_strides.GpuPtr(), + reinterpret_cast(prepare.lhs_tensor->template Data()), + prepare.rhs_padded_strides.GpuPtr(), + reinterpret_cast(prepare.rhs_tensor->template Data()), + prepare.fdm_output_strides.GpuPtr(), + prepare.fdm_H, + prepare.fdm_C, + reinterpret_cast(prepare.output_tensor->template MutableData()), + prepare.output_tensor->Shape().Size()); + for (int index = 1; index < input_count; index++) { + ORT_RETURN_IF_ERROR(BinaryElementwiseBroadcastPrepare(0, output_tensor, context->Input(index), output_tensor, &prepare)); + Impl_Max( + prepare.output_rank_or_simple_broadcast, + prepare.lhs_padded_strides.GpuPtr(), + reinterpret_cast(prepare.lhs_tensor->template Data()), + prepare.rhs_padded_strides.GpuPtr(), + reinterpret_cast(prepare.rhs_tensor->template Data()), + prepare.fdm_output_strides.GpuPtr(), + prepare.fdm_H, + prepare.fdm_C, + reinterpret_cast(prepare.output_tensor->template MutableData()), + prepare.output_tensor->Shape().Size()); + } + } + return Status::OK(); +} + //Greater op output tensor type is bool, so it cannot directly fit in the macros //for other elementwise ops template @@ -261,7 +317,7 @@ Status Greater::ComputeInternal(OpKernelContext* context) const { BinaryElementwisePreparation prepare(this); ORT_RETURN_IF_ERROR(BinaryElementwiseBroadcastPrepare(0, input0, input1, output_tensor, &prepare)); - Impl_Compare( + Impl_Greater( prepare.output_rank_or_simple_broadcast, prepare.lhs_padded_strides.GpuPtr(), reinterpret_cast(prepare.lhs_tensor->template Data()), @@ -280,6 +336,8 @@ BINARY_OP_REGISTER_UZILHFD(Sum, 8) BINARY_OP_REGISTER_VERSIONED_UZILHFD(Sum, 6, 7) BINARY_OP_REGISTER_UZILHFD(Greater, 9) BINARY_OP_REGISTER_VERSIONED_HFD(Greater, 7, 8) +BINARY_OP_REGISTER_HFD(Max, 8) +BINARY_OP_REGISTER_VERSIONED_HFD(Max, 6, 7) } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h index 9149548d1f..2f3e34c813 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h @@ -203,5 +203,14 @@ class Greater final : public CudaKernel { Status ComputeInternal(OpKernelContext* context) const override; }; + +template +class Max final : public CudaKernel { + public: + Max(const OpKernelInfo& info) : CudaKernel(info) { + } + + Status ComputeInternal(OpKernelContext* context) const override; +}; } // namespace cuda } // namespace onnxruntime 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 c276dc4761..f7e61cbce6 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu @@ -81,7 +81,8 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(And, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Or, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(PRelu) -SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Compare) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Greater) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Max) } // 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 f032a6a6ee..2c5c64b363 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h @@ -25,7 +25,8 @@ namespace cuda { BINARY_OP_NAME_EXPR(Or, (a | b)) \ BINARY_OP_NAME_EXPR(Xor, (a ^ b)) \ BINARY_OP_NAME_EXPR(PRelu, (a > (T)0 ? a : a * b)) \ - BINARY_OP_NAME_EXPR(Compare, (a > b) ? 1 : 0) + BINARY_OP_NAME_EXPR(Greater, (a > b) ? 1 : 0) \ + BINARY_OP_NAME_EXPR(Max, _Max(a, b)) // NOTE that cu files are compiled with nvcc and should not refer to any onnxruntime headers // so struct BinaryElementwisePreparation cannot be used here 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 f56da177ef..e8bc5db12f 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -585,6 +585,21 @@ TEST(MathOpTest, Max_8) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //Input batch size is inconsistent } +TEST(MathOpTest, Max_8_2inputbroadcast) { + OpTester test("Max", 8); + test.AddInput("data_0", {1, 3}, + {1.0f, 2.0f, 3.0f}); + test.AddInput("data_1", {3, 3}, + {10.0f, 20.0f, 30.0f, + 40.0f, 50.0f, 60.0f, + 70.0f, 80.0f, 90.0f}); + test.AddOutput("max", {3, 3}, + {10.0f, 20.0f, 30.0f, + 40.0f, 50.0f, 60.0f, + 70.0f, 80.0f, 90.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //Input batch size is inconsistent +} + TEST(MathOpTest, Not) { OpTester test("Not"); std::vector dims{2};