diff --git a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc index 827aeca791..fbb52db0f4 100644 --- a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc @@ -615,7 +615,7 @@ Status ReduceComputeCore(CUDAExecutionProvider& cuda_ep, const Tensor& input, Pr // cudnnReduceTensor has issue if input and output has same size, which will happen if the axis to be reduced has dim value of 1. // the output is zeros of the output size if (input_count == output_count) { - CUDA_RETURN_IF_ERROR(cudaMemsetAsync(output.template MutableData(), static_cast(0), output_count * sizeof(int64_t))); + CUDA_RETURN_IF_ERROR(cudaMemsetAsync(output.template MutableData(), static_cast(0), output_count * sizeof(int64_t), stream)); } else { if (temp_X) { auto temp_output = cuda_ep.GetScratchBuffer(output_count); @@ -748,7 +748,7 @@ Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, cudnnRe IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); \ Impl_Cast(Stream(), reinterpret_cast(X->template Data()), temp_X.get(), X->Shape().Size()); \ \ - ORT_RETURN_IF_ERROR(reduce_desc.Set(cudnn_reduce_op, cudnn_type_X, CUDNN_REDUCE_TENSOR_FLATTENED_INDICES)); \ + ORT_RETURN_IF_ERROR(reduce_desc.Set(cudnn_reduce_op, cudnn_type_X, CUDNN_REDUCE_TENSOR_NO_INDICES)); \ ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_cudnn, cudnn_type_X)); \ ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_cudnn, cudnn_type_X)); \ CUDNN_RETURN_IF_ERROR( \ @@ -859,7 +859,7 @@ Status ReduceKernel::ComputeImpl Impl_Cast(Stream(), reinterpret_cast(X->template Data()), temp_X.get(), X->Shape().Size()); - ORT_RETURN_IF_ERROR(reduce_desc.Set(cudnn_reduce_op, cudnn_type_X, CUDNN_REDUCE_TENSOR_FLATTENED_INDICES)); + ORT_RETURN_IF_ERROR(reduce_desc.Set(cudnn_reduce_op, cudnn_type_X, CUDNN_REDUCE_TENSOR_NO_INDICES)); ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_cudnn, cudnn_type_X)); ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_cudnn, cudnn_type_X)); CUDNN_RETURN_IF_ERROR( diff --git a/onnxruntime/core/providers/rocm/reduction/reduction_ops.cc b/onnxruntime/core/providers/rocm/reduction/reduction_ops.cc index 55aac97d7a..886a8ab2c6 100644 --- a/onnxruntime/core/providers/rocm/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/rocm/reduction/reduction_ops.cc @@ -155,7 +155,7 @@ Status ReduceKernel::ReduceKernelShared( m, n, false); } case ApplicableMatrixReduction::Columns: - // don't call reduce_matrix_columns() since it will reset initial output data + // don't call reduce_matrix_columns() since it will reset initial output data default: break; } @@ -506,7 +506,7 @@ Status ReduceComputeCore(ROCMExecutionProvider& rocm_ep, const Tensor& input, Pr input_data_buffer = rocm_ep.GetScratchBuffer(input_count); input_data = reinterpret_cast(input_data_buffer.get()); fast_divmod tmp_div; - Impl_Mul(stream, + Impl_Mul(stream, static_cast(SimpleBroadcast::NoBroadcast), nullptr, reinterpret_cast(input.template Data()), nullptr, reinterpret_cast(input.template Data()), nullptr, @@ -617,7 +617,8 @@ Status ReduceComputeCore(ROCMExecutionProvider& rocm_ep, const Tensor& input, Pr &zero, output_tensor, reinterpret_cast(output.template MutableData()))); } } - } else { // For ArgMax & ArgMin ops, use the indicies as the output with int64 type + } else { + // For ArgMax & ArgMin ops, use the indicies as the output with int64 type // miopenReduceTensor has issue if input and output has same size, which will happen if the axis to be reduced has dim value of 1. // the output is zeros of the output size if (input_count == output_count) { @@ -692,183 +693,94 @@ Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, miopenR calculate_log_, calculate_sqt_, log_sum_exp_, fast_reduction); } -template <> -template <> -Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, miopenReduceTensorOp_t miopen_reduce_op) const { - typedef typename ToHipType::MappedType HipT; - - const Tensor* X = ctx->Input(0); - std::vector axes; - - size_t num_inputs = ctx->InputCount(); - if (num_inputs == 2) { - //override the attribute value with the input value for reduction_axes - const Tensor* axes_tensor = ctx->Input(1); - ORT_ENFORCE(axes_tensor != nullptr, "Axes input is null"); - ORT_ENFORCE(axes_tensor->Shape().NumDimensions() == 1, "An axes tensor must be a vector tensor."); - auto nDims = static_cast(axes_tensor->Shape()[0]); - const auto* data = axes_tensor->template Data(); - axes.assign(data, data + nDims); - } else { - axes.assign(axes_.begin(), axes_.end()); +#define SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(T) \ + template <> \ + template <> \ + Status ReduceKernel::ComputeImpl( \ + OpKernelContext * ctx, miopenReduceTensorOp_t miopen_reduce_op) const { \ + typedef typename ToHipType::MappedType HipT; \ + const Tensor* X = ctx->Input(0); \ + std::vector axes; \ + size_t num_inputs = ctx->InputCount(); \ + if (num_inputs == 2) { \ + const Tensor* axes_tensor = ctx->Input(1); \ + ORT_ENFORCE(axes_tensor != nullptr, "Axes input is null"); \ + ORT_ENFORCE(axes_tensor->Shape().NumDimensions() == 1, "An axes tensor must be a vector tensor."); \ + auto nDims = static_cast(axes_tensor->Shape()[0]); \ + const auto* data = axes_tensor->template Data(); \ + axes.assign(data, data + nDims); \ + } else { \ + axes.assign(axes_.begin(), axes_.end()); \ + } \ + \ + if (axes.empty() && noop_with_empty_axes_) { \ + auto* Y = ctx->Output(0, X->Shape()); \ + HIP_RETURN_IF_ERROR(hipMemcpyAsync(Y->template MutableData(), X->template Data(), X->SizeInBytes(), \ + hipMemcpyDeviceToDevice, Stream())); \ + return Status::OK(); \ + } \ + \ + PrepareReduceMetadata prepare_reduce_metadata; \ + ORT_RETURN_IF_ERROR(PrepareForReduce(X, keepdims_, axes, prepare_reduce_metadata)); \ + \ + Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims); \ + \ + int64_t input_count = prepare_reduce_metadata.input_count; \ + int64_t output_count = prepare_reduce_metadata.output_count; \ + std::vector& input_dims_miopen = prepare_reduce_metadata.input_dims_miopen; \ + std::vector& output_dims_miopen = prepare_reduce_metadata.output_dims_miopen; \ + \ + if (input_count == 0) { \ + assert(Y->Shape().Size() == 0); \ + return Status::OK(); \ + } \ + \ + if (input_count == output_count) { \ + if (Y->template MutableData() != X->template Data()) { \ + HIP_RETURN_IF_ERROR(hipMemcpyAsync(Y->template MutableData(), X->template Data(), \ + input_count * sizeof(T), hipMemcpyDeviceToDevice, Stream())); \ + } \ + return Status::OK(); \ + } \ + \ + HIP_RETURN_IF_ERROR(hipMemsetAsync(Y->MutableDataRaw(), 0, Y->SizeInBytes(), Stream())); \ + \ + size_t indices_bytes = 0; \ + size_t workspace_bytes = 0; \ + MiopenTensor input_tensor; \ + MiopenTensor output_tensor; \ + MiopenReduceDescriptor reduce_desc; \ + \ + miopenDataType_t miopen_type_X = miopenFloat; \ + IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); \ + Impl_Cast(Stream(), reinterpret_cast(X->template Data()), temp_X.get(), X->Shape().Size()); \ + \ + ORT_RETURN_IF_ERROR(reduce_desc.Set(miopen_reduce_op, miopen_type_X, MIOPEN_REDUCE_TENSOR_NO_INDICES)); \ + ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_miopen, miopen_type_X)); \ + ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_miopen, miopen_type_X)); \ + MIOPEN_RETURN_IF_ERROR( \ + miopenGetReductionIndicesSize(MiopenHandle(), reduce_desc, input_tensor, output_tensor, &indices_bytes)); \ + MIOPEN_RETURN_IF_ERROR( \ + miopenGetReductionWorkspaceSize(MiopenHandle(), reduce_desc, input_tensor, output_tensor, &workspace_bytes)); \ + IAllocatorUniquePtr indices_rocm = GetScratchBuffer(indices_bytes); \ + IAllocatorUniquePtr workspace_rocm = GetScratchBuffer(workspace_bytes); \ + \ + const auto one = Consts::One; \ + const auto zero = Consts::Zero; \ + auto temp_Y = GetScratchBuffer(output_count); \ + MIOPEN_RETURN_IF_ERROR(miopenReduceTensor(MiopenHandle(), reduce_desc, indices_rocm.get(), indices_bytes, \ + workspace_rocm.get(), workspace_bytes, &one, input_tensor, temp_X.get(), \ + &zero, output_tensor, temp_Y.get())); \ + \ + Impl_Cast(Stream(), temp_Y.get(), reinterpret_cast(Y->template MutableData()), output_count); \ + \ + return Status::OK(); \ } - // empty axes and no-op - if (axes.empty() && noop_with_empty_axes_) { - auto* Y = ctx->Output(0, X->Shape()); - HIP_RETURN_IF_ERROR(hipMemcpyAsync(Y->template MutableData(), X->template Data(), X->SizeInBytes(), hipMemcpyDeviceToDevice, Stream())); - return Status::OK(); - } - - PrepareReduceMetadata prepare_reduce_metadata; - - ORT_RETURN_IF_ERROR(PrepareForReduce(X, - keepdims_, - axes, - prepare_reduce_metadata)); - - Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims); - - int64_t input_count = prepare_reduce_metadata.input_count; - int64_t output_count = prepare_reduce_metadata.output_count; - std::vector& input_dims_miopen = prepare_reduce_metadata.input_dims_miopen; - std::vector& output_dims_miopen = prepare_reduce_metadata.output_dims_miopen; - - // special case when there is a dim value of 0 in the shape. - if (input_count == 0) { - assert(Y->Shape().Size() == 0); - return Status::OK(); - } - - // miopenReduceTensor for ReduceSum has issue if input and output has same size, we just need to copy the data for this case - if (input_count == output_count) { - if (Y->template MutableData() != X->template Data()) { - HIP_RETURN_IF_ERROR(hipMemcpyAsync(Y->template MutableData(), X->template Data(), input_count * sizeof(int32_t), hipMemcpyDeviceToDevice, Stream())); - } - return Status::OK(); - } - - // This reduction keep adding values to this buffer. If a non-zero value, say 1000, is here, the sum will start with 1000. - // Therefore zeroing out the memory is required - HIP_RETURN_IF_ERROR(hipMemsetAsync(Y->MutableDataRaw(), 0, Y->SizeInBytes(), Stream())); - - size_t indices_bytes = 0; - size_t workspace_bytes = 0; - MiopenTensor input_tensor; - MiopenTensor output_tensor; - MiopenReduceDescriptor reduce_desc; - - miopenDataType_t miopen_type_X = miopenFloat; - IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); - Impl_Cast(Stream(), reinterpret_cast(X->template Data()), temp_X.get(), X->Shape().Size()); - - ORT_RETURN_IF_ERROR(reduce_desc.Set(miopen_reduce_op, miopen_type_X, MIOPEN_REDUCE_TENSOR_FLATTENED_INDICES)); - ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_miopen, miopen_type_X)); - ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_miopen, miopen_type_X)); - MIOPEN_RETURN_IF_ERROR(miopenGetReductionIndicesSize(MiopenHandle(), reduce_desc, input_tensor, output_tensor, &indices_bytes)); - MIOPEN_RETURN_IF_ERROR(miopenGetReductionWorkspaceSize(MiopenHandle(), reduce_desc, input_tensor, output_tensor, &workspace_bytes)); - IAllocatorUniquePtr indices_rocm = GetScratchBuffer(indices_bytes); - IAllocatorUniquePtr workspace_rocm = GetScratchBuffer(workspace_bytes); - - const auto one = Consts::One; - const auto zero = Consts::Zero; - auto temp_Y = GetScratchBuffer(output_count); - MIOPEN_RETURN_IF_ERROR(miopenReduceTensor(MiopenHandle(), - reduce_desc, - indices_rocm.get(), - indices_bytes, - workspace_rocm.get(), - workspace_bytes, - &one, - input_tensor, - temp_X.get(), - &zero, - output_tensor, - temp_Y.get())); - - Impl_Cast(Stream(), temp_Y.get(), Y->template MutableData(), output_count); - - return Status::OK(); -} - -template <> -template <> -Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, miopenReduceTensorOp_t miopen_reduce_op) const { - typedef typename ToHipType::MappedType HipT; - - const Tensor* X = ctx->Input(0); - PrepareReduceMetadata prepare_reduce_metadata; - - ORT_RETURN_IF_ERROR(PrepareForReduce(X, - keepdims_, - axes_, - prepare_reduce_metadata)); - - Tensor* Y = ctx->Output(0, prepare_reduce_metadata.squeezed_output_dims); - - int64_t input_count = prepare_reduce_metadata.input_count; - int64_t output_count = prepare_reduce_metadata.output_count; - std::vector& input_dims_miopen = prepare_reduce_metadata.input_dims_miopen; - std::vector& output_dims_miopen = prepare_reduce_metadata.output_dims_miopen; - - // special case when there is a dim value of 0 in the shape. - if (input_count == 0) { - assert(Y->Shape().Size() == 0); - return Status::OK(); - } - - // miopenReduceTensor has issue if input and output has same size, we just need to copy the data for this case - auto* const dst = Y->template MutableData(); - const auto* const src = X->template Data(); - if (input_count == output_count) { - if (src != dst) { - HIP_RETURN_IF_ERROR(hipMemcpyAsync(dst, src, input_count * sizeof(int8_t), hipMemcpyDeviceToDevice, Stream())); - } - return Status::OK(); - } - - // This reduction keep adding values to this buffer. If a non-zero value, say 1000, is here, the sum will start with 1000. - // Therefore zeroing out the memory is required - HIP_RETURN_IF_ERROR(hipMemsetAsync(Y->MutableDataRaw(), 0, Y->SizeInBytes(), Stream())); - - size_t indices_bytes = 0; - size_t workspace_bytes = 0; - MiopenTensor input_tensor; - MiopenTensor output_tensor; - MiopenReduceDescriptor reduce_desc; - - miopenDataType_t miopen_type_X = miopenFloat; - IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); - Impl_Cast(Stream(), reinterpret_cast(src), temp_X.get(), X->Shape().Size()); - - ORT_RETURN_IF_ERROR(reduce_desc.Set(miopen_reduce_op, miopen_type_X, MIOPEN_REDUCE_TENSOR_FLATTENED_INDICES)); - ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_miopen, miopen_type_X)); - ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_miopen, miopen_type_X)); - MIOPEN_RETURN_IF_ERROR(miopenGetReductionIndicesSize(MiopenHandle(), reduce_desc, input_tensor, output_tensor, &indices_bytes)); - MIOPEN_RETURN_IF_ERROR(miopenGetReductionWorkspaceSize(MiopenHandle(), reduce_desc, input_tensor, output_tensor, &workspace_bytes)); - IAllocatorUniquePtr indices_rocm = GetScratchBuffer(indices_bytes); - IAllocatorUniquePtr workspace_rocm = GetScratchBuffer(workspace_bytes); - - const auto one = Consts::One; - const auto zero = Consts::Zero; - auto temp_Y = GetScratchBuffer(output_count); - MIOPEN_RETURN_IF_ERROR(miopenReduceTensor(MiopenHandle(), - reduce_desc, - indices_rocm.get(), - indices_bytes, - workspace_rocm.get(), - workspace_bytes, - &one, - input_tensor, - temp_X.get(), - &zero, - output_tensor, - temp_Y.get())); - - Impl_Cast(Stream(), temp_Y.get(), dst, output_count); - - return Status::OK(); -} +SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(int32_t) +SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(int64_t) +SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(int8_t) +SPECIALIZED_REDUCEKERNEL_COMPUTEIMPL(uint8_t) namespace ReductionOps { @@ -916,6 +828,13 @@ template Tensor ReduceCompute( // bool keep_dims, bool calculate_log, bool calculate_sqt, bool log_sum_exp, // bool fast_reduction, const TensorShape* input_shape_override); +template Tensor ReduceCompute( + ROCMExecutionProvider& rocm_ep, miopenReduceTensorOp_t miopen_reduce_op, + AllocatorPtr allocator, + const Tensor& input, const std::vector& axes, + bool keep_dims, bool calculate_log, bool calculate_sqt, bool log_sum_exp, + bool fast_reduction, const TensorShape* input_shape_override); + } // namespace ReductionOps #define REGISTER_KERNEL_HFD(name) \ @@ -937,8 +856,9 @@ REGISTER_KERNEL_TYPED_12(ReduceMax, MLFloat16) REGISTER_KERNEL_TYPED_12(ReduceMax, float) // REGISTER_KERNEL_TYPED_12(ReduceMax, double) REGISTER_KERNEL_TYPED_12(ReduceMax, int32_t) +REGISTER_KERNEL_TYPED_12(ReduceMax, int64_t) REGISTER_KERNEL_TYPED_12(ReduceMax, int8_t) -// REGISTER_KERNEL_TYPED_12(ReduceMax, uint8_t) +REGISTER_KERNEL_TYPED_12(ReduceMax, uint8_t) REGISTER_KERNEL_HFD(ReduceMean) @@ -947,7 +867,7 @@ REGISTER_KERNEL_TYPED_12(ReduceMin, float) // REGISTER_KERNEL_TYPED_12(ReduceMin, double) REGISTER_KERNEL_TYPED_12(ReduceMin, int32_t) REGISTER_KERNEL_TYPED_12(ReduceMin, int8_t) -// REGISTER_KERNEL_TYPED_12(ReduceMin, uint8_t) +REGISTER_KERNEL_TYPED_12(ReduceMin, uint8_t) REGISTER_KERNEL_HFD(ReduceProd) @@ -955,6 +875,7 @@ REGISTER_KERNEL_TYPED_13(ReduceSum, MLFloat16) REGISTER_KERNEL_TYPED_13(ReduceSum, float) // REGISTER_KERNEL_TYPED_13(ReduceSum, double) REGISTER_KERNEL_TYPED_13(ReduceSum, int32_t) +REGISTER_KERNEL_TYPED_13(ReduceSum, int64_t) REGISTER_KERNEL_HFD(ReduceLogSum) REGISTER_KERNEL_HFD(ReduceSumSquare) diff --git a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc index 5a56bbe657..2805cb4d4e 100644 --- a/onnxruntime/core/providers/rocm/rocm_execution_provider.cc +++ b/onnxruntime/core/providers/rocm/rocm_execution_provider.cc @@ -529,6 +529,7 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kO class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, double, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, MLFloat16, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, int32_t, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, int64_t, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, float, ReduceMean); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, double, ReduceMean); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 1, 10, MLFloat16, ReduceMean); @@ -713,6 +714,7 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kO class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 11, double, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 11, int32_t, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 11, int64_t, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, float, ReduceMean); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, double, ReduceMean); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 11, 12, MLFloat16, ReduceMean); @@ -800,6 +802,7 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kO class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, double, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, MLFloat16, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, int32_t, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, int64_t, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, int8_t, ReduceMax); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 12, 12, uint8_t, ReduceMax); @@ -975,6 +978,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, double, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, MLFloat16, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int32_t, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int64_t, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, int8_t, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, uint8_t, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kRocmExecutionProvider, kOnnxDomain, 13, float, ReduceMean); @@ -1241,6 +1245,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { // BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, // BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1438,6 +1443,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { // BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, // BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1521,15 +1527,16 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { // BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, // BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1696,8 +1703,9 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { // BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, // BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1707,7 +1715,7 @@ static Status RegisterRocmKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - // BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, // BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/tools/ci_build/github/pai/pai-excluded-tests.txt b/tools/ci_build/github/pai/pai-excluded-tests.txt index 7bee65b29a..294674c2e5 100644 --- a/tools/ci_build/github/pai/pai-excluded-tests.txt +++ b/tools/ci_build/github/pai/pai-excluded-tests.txt @@ -24,8 +24,6 @@ ReductionOpTest.ReduceLogSumExp_do_not_keepdims_2_double ReductionOpTest.ReduceLogSumExp_keepdims_double ReductionOpTest.ReduceLogSumExp_double ReductionOpTest.ReduceMax_double -ReductionOpTest.ReduceMax_int32 -ReductionOpTest.ReduceMax_int8 ReductionOpTest.ReduceMean_default_axes_keepdims_double ReductionOpTest.ReduceMean_default_axes_do_not_keep_dims_double ReductionOpTest.ReduceMean_do_not_keepdims_double @@ -34,8 +32,6 @@ ReductionOpTest.ReduceMean_keepdims_double ReductionOpTest.ReduceMean_double ReductionOpTest.ReduceMean0DTensor_double ReductionOpTest.ReduceMin_double -ReductionOpTest.ReduceMin_int32 -ReductionOpTest.ReduceMin_int8 ReductionOpTest.ReduceSum_double ReductionOpTest.ReduceSumSquare_double ReductionOpTest.ReduceInfMax_double