diff --git a/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc b/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc index bb7a50bec6f99..633dfe26c2532 100644 --- a/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc @@ -879,25 +879,35 @@ bool check_and_reduce_empty_set_input(OpKernelContext* ctx, const gsl::span input_axes; if (ctx->InputCount() == 2) { ORT_ENFORCE(axes.empty(), "Axes input and attribute should not both be present for reduction."); - // second input holds the axes. const Tensor* axes_tensor = ctx->Input(1); - auto nDims = static_cast(axes_tensor->Shape()[0]); - const auto* data = axes_tensor->Data(); - input_axes.insert(input_axes.begin(), data, data + nDims); + if (axes_tensor != nullptr) { + 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->Data(); + input_axes.insert(input_axes.begin(), data, data + nDims); + } + // axes_tensor == nullptr means no axes provided → reduce all dims } else { input_axes.resize(axes.size()); std::copy(axes.begin(), axes.end(), input_axes.begin()); } - gsl::span shape_dims = input_shape.GetDims(); - const int64_t input_shape_size = narrow(shape_dims.size()); + // Normalize negative axes + const int64_t rank = narrow(input_shape.NumDimensions()); + for (auto& axis : input_axes) { + axis = HandleNegativeAxis(axis, rank); + } + + // Build reduced output shape (linear scan — rank is always < 8) TensorShapeVector output_shape_vector; - for (int64_t i = 0; i < input_shape_size; ++i) { - if (input_axes.empty() || std::find(input_axes.begin(), input_axes.end(), i) != input_axes.end()) { + for (int64_t i = 0; i < rank; ++i) { + bool is_reduced = input_axes.empty() || + std::find(input_axes.begin(), input_axes.end(), i) != input_axes.end(); + if (is_reduced) { if (keepdims) { output_shape_vector.push_back(1); } @@ -968,17 +978,22 @@ template void CommonReduce1Loop(OpKernelContext* ctx, const gsl::span& axes_, int64_t keepdims_, bool noop_with_empty_axes) { - if (check_and_reduce_empty_set_input(ctx, axes_, keepdims_ != 0)) { - return; - } - + // Resolve effective axes first (from input tensor or attribute). TensorShapeVector tmp_axes; auto effective_axes = GetEffectiveAxes(ctx, axes_, tmp_axes); + + // noop_with_empty_axes takes precedence: if no axes, copy input as-is + // (applying element-wise transforms if any). This applies even to empty + // tensors — a {1,0} input should stay {1,0}, not be reduced to scalar. if (effective_axes.empty() && noop_with_empty_axes) { ApplyNoopEmptyAxesElementwise(ctx); return; } + if (check_and_reduce_empty_set_input(ctx, axes_, keepdims_ != 0)) { + return; + } + FastReduceKind fast_kind; TensorShapeVector fast_shape; TensorShapeVector output_shape; diff --git a/onnxruntime/core/providers/cpu/reduction/reduction_ops.h b/onnxruntime/core/providers/cpu/reduction/reduction_ops.h index 03be8f3e6686c..3867eef22cdb2 100644 --- a/onnxruntime/core/providers/cpu/reduction/reduction_ops.h +++ b/onnxruntime/core/providers/cpu/reduction/reduction_ops.h @@ -386,8 +386,10 @@ class ReduceAggregatorMax : public ReduceAggregator { inline void update(const T& v) { this->accumulator_ = v > this->accumulator_ ? v : this->accumulator_; } static void fill_for_empty_set(Tensor& output) { - if constexpr (std::is_same_v) { /* bool specific impl */ - ORT_NOT_IMPLEMENTED(); + if constexpr (std::is_same_v) { + // ONNX spec: ReduceMax on empty bool set → false (boolean zero/identity) + auto* data = output.MutableData(); + std::fill(data, data + output.Shape().Size(), false); } else { EigenMap(output).array() = -std::numeric_limits::infinity(); } @@ -596,8 +598,10 @@ class ReduceAggregatorMin : public ReduceAggregator { inline void update(const T& v) { this->accumulator_ = v < this->accumulator_ ? v : this->accumulator_; } static void fill_for_empty_set(Tensor& output) { - if constexpr (std::is_same_v) { /* bool specific impl */ - ORT_NOT_IMPLEMENTED(); + if constexpr (std::is_same_v) { + // ONNX spec: ReduceMin on empty bool set → true (boolean max/identity) + auto* data = output.MutableData(); + std::fill(data, data + output.Shape().Size(), true); } else { EigenMap(output).array() = std::numeric_limits::infinity(); } diff --git a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc index a0a2f377d0c80..cef299fc41a8f 100644 --- a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc @@ -298,22 +298,17 @@ Status PrepareForReduce(const Tensor* X, prepare_reduce_metadata.output_dims = input_shape.AsShapeVector(); for (auto axis : axes) { axis = HandleNegativeAxis(axis, rank); - ORT_ENFORCE(input_dims[axis] != 0, - "Can't reduce on dim with value of 0 if 'keepdims' is false. " - "Invalid output shape would be produced. input_shape:", - input_shape); prepare_reduce_metadata.output_dims[axis] = 1; reduced[axis] = true; } } else { // no axes provided (i.e.) default axes => reduce on all dims + // Each reduced dim becomes 1 (even if the original dim was 0 — the + // reduction collapses the axis regardless of its size). prepare_reduce_metadata.output_dims.reserve(input_dims.size()); - for (auto dim : input_dims) { - ORT_ENFORCE(keepdims || dim != 0, - "Can't reduce on dim with value of 0 if 'keepdims' is false. " - "Invalid output shape would be produced. input_shape:", - input_shape); - prepare_reduce_metadata.output_dims.push_back(dim == 0 ? 0 : 1); + for (size_t i = 0; i < input_dims.size(); ++i) { + prepare_reduce_metadata.output_dims.push_back(1); + reduced[i] = true; } } @@ -377,7 +372,39 @@ Status ReduceComputeCore(const AllocatorPtr& gpu_allocator, const CudaKernel* ke auto& output_dims_cudnn = prepare_reduce_metadata.output_dims_cudnn; // special case when there is a dim value of 0 in the shape. if (input_count == 0) { - assert(output.Shape().Size() == 0); + // Empty input reduction: output may still be non-empty when only some + // axes are reduced. Per ONNX spec, fill with the reduction identity. + if (output_count > 0) { + // For types that don't support std::numeric_limits natively (MLFloat16, + // BFloat16), use float intermediary and convert via CudaT. + if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_AVG) { + // ReduceMean on empty set is undefined (0/0) per ONNX spec. + // Fill with 0 for consistency with CPU (ReduceAggregatorMean inherits + // ReduceAggregatorSum::fill_for_empty_set which fills 0). + CUDA_RETURN_IF_ERROR(cudaMemsetAsync(output.MutableDataRaw(), 0, + output.SizeInBytes(), stream)); + } else if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_MUL) { + // ReduceProd identity is 1. + CudaT one_val = ToCudaType::FromFloat(1.0f); + std::vector ones(output_count, one_val); + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(output.MutableDataRaw(), ones.data(), + output.SizeInBytes(), cudaMemcpyHostToDevice, stream)); + } else if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_MIN) { + CudaT inf_val = ToCudaType::FromFloat(std::numeric_limits::infinity()); + std::vector vals(output_count, inf_val); + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(output.MutableDataRaw(), vals.data(), + output.SizeInBytes(), cudaMemcpyHostToDevice, stream)); + } else if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_MAX) { + CudaT neg_inf_val = ToCudaType::FromFloat(-std::numeric_limits::infinity()); + std::vector vals(output_count, neg_inf_val); + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(output.MutableDataRaw(), vals.data(), + output.SizeInBytes(), cudaMemcpyHostToDevice, stream)); + } else { + // Sum, SumSquare, L1, L2: identity is 0. + CUDA_RETURN_IF_ERROR(cudaMemsetAsync(output.MutableDataRaw(), 0, + output.SizeInBytes(), stream)); + } + } return Status::OK(); } @@ -770,7 +797,34 @@ Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, cudnnRe auto& output_dims_cudnn = prepare_reduce_metadata.output_dims_cudnn; \ \ if (input_count == 0) { \ - assert(Y->Shape().Size() == 0); \ + /* Empty input reduction: fill output with the reduction identity. */ \ + /* ONNX spec: Sum→0, Prod→1, Min→+inf, Max→-inf, Mean→0. */ \ + if (Y->Shape().Size() > 0) { \ + typedef typename ToCudaType::MappedType CudaT_local; \ + if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_MUL) { \ + /* Identity is 1 for product */ \ + CudaT_local one_val = ToCudaType::FromFloat(1.0f); \ + std::vector ones(Y->Shape().Size(), one_val); \ + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(Y->MutableDataRaw(), ones.data(), \ + Y->SizeInBytes(), cudaMemcpyHostToDevice, Stream(ctx))); \ + } else if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_MIN) { \ + /* ONNX spec: "yields plus infinity (if supported) or max value" */ \ + CudaT_local inf_val = ToCudaType::FromFloat(std::numeric_limits::infinity()); \ + std::vector vals(Y->Shape().Size(), inf_val); \ + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(Y->MutableDataRaw(), vals.data(), \ + Y->SizeInBytes(), cudaMemcpyHostToDevice, Stream(ctx))); \ + } else if (cudnn_reduce_op == CUDNN_REDUCE_TENSOR_MAX) { \ + /* ONNX spec: "yields minus infinity (if supported) or minimum value" */ \ + CudaT_local neg_inf_val = ToCudaType::FromFloat(-std::numeric_limits::infinity()); \ + std::vector vals(Y->Shape().Size(), neg_inf_val); \ + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(Y->MutableDataRaw(), vals.data(), \ + Y->SizeInBytes(), cudaMemcpyHostToDevice, Stream(ctx))); \ + } else { \ + /* Sum, SumSquare, Mean, L1, L2, Amax: identity is 0 */ \ + CUDA_RETURN_IF_ERROR(cudaMemsetAsync(Y->MutableDataRaw(), 0, \ + Y->SizeInBytes(), Stream(ctx))); \ + } \ + } \ return Status::OK(); \ } \ \ diff --git a/onnxruntime/core/providers/webgpu/reduction/reduction_ops.cc b/onnxruntime/core/providers/webgpu/reduction/reduction_ops.cc index db51675e81513..2fac346601a13 100644 --- a/onnxruntime/core/providers/webgpu/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/webgpu/reduction/reduction_ops.cc @@ -169,6 +169,32 @@ std::unordered_map reduce_op_output_values_map = { {ReduceOpType::LogSum, "log(f32(bestValue))"}, }; +// WGSL expressions for the ONNX empty-set identity value of each reduction op. +// Per ONNX spec: Sum→0, Prod→1, Max→-inf, Min→+inf, Mean→0 (undefined, ORT uses 0). +// ArgMax/ArgMin on empty input is undefined; output 0. +// For ±infinity we bitcast from the IEEE 754 f32 bit pattern. +// 0xFF800000u = -inf, 0x7F800000u = +inf. +// bitcast() is always valid in WGSL (u32→f32 reinterpretation). +// output_value_t() then converts f32 inf to f16 inf when needed (f16 has inf). +// Using division-by-zero (1.0/0.0) would also produce ±inf per IEEE 754, but +// some WGSL shader validators reject constant division by zero. +std::unordered_map reduce_op_empty_identity_map = { + {ReduceOpType::Max, "output_value_t(bitcast(0xFF800000u))"}, // -inf + {ReduceOpType::Min, "output_value_t(bitcast(0x7F800000u))"}, // +inf + {ReduceOpType::Mean, "output_value_t(0)"}, + {ReduceOpType::Sum, "output_value_t(0)"}, + {ReduceOpType::Prod, "output_value_t(1)"}, + {ReduceOpType::SumSquare, "output_value_t(0)"}, + {ReduceOpType::LogSumExp, "output_value_t(bitcast(0xFF800000u))"}, // -inf (log(0) = -inf) + {ReduceOpType::L1, "output_value_t(0)"}, + {ReduceOpType::L2, "output_value_t(0)"}, + {ReduceOpType::LogSum, "output_value_t(bitcast(0xFF800000u))"}, // -inf (log(0) = -inf) + {ReduceOpType::ArgMax, "output_value_t(0)"}, + {ReduceOpType::ArgMin, "output_value_t(0)"}, + {ReduceOpType::ArgMax_select_last_index, "output_value_t(0)"}, + {ReduceOpType::ArgMin_select_last_index, "output_value_t(0)"}, +}; + std::unordered_map reduce_op_naive_code_map = { {ReduceOpType::Max, {"var max_element = first_element;", "max_element = max(max_element, current_element);", "let output_value = output_value_t(max_element);"}}, {ReduceOpType::Min, {"var min_element = first_element;", "min_element = min(min_element, current_element);", "let output_value = output_value_t(min_element);"}}, @@ -195,9 +221,15 @@ Status ReduceNaiveProgram::GenerateShaderCode(ShaderHelper& shader) const { const auto& code = reduce_op_naive_code_map.at(reduce_op_type_); const auto& output = shader.AddOutput("output", ShaderUsage::UseUniform | ShaderUsage::UseIndicesTypeAlias | ShaderUsage::UseValueTypeAlias); if (is_input_empty_) { + // Empty input: output the ONNX identity value for this reduction op. + // Don't use loop_header_/loop_footer_ — they may reference undefined + // variables (e.g., first_element) or divide by zero (ReduceMean). + const auto& identity = reduce_op_empty_identity_map.at(reduce_op_type_); shader.MainFunctionBody() << shader.GuardAgainstOutOfBoundsWorkgroupSizes("uniforms.output_size") - << code.loop_header_ - << code.loop_footer_ + // Use var to prevent WGSL constant-folding of the identity + // expression. Some validators reject inf-producing constant + // expressions (e.g. division by zero or f32→f16 inf conversion). + << "var output_value: output_value_t = " << identity << ";\n" << output.SetByOffset("global_idx", "output_value"); return Status::OK(); } diff --git a/onnxruntime/core/providers/webgpu/reduction/reduction_ops.h b/onnxruntime/core/providers/webgpu/reduction/reduction_ops.h index 6958b52a6c880..32a6e9269b0de 100644 --- a/onnxruntime/core/providers/webgpu/reduction/reduction_ops.h +++ b/onnxruntime/core/providers/webgpu/reduction/reduction_ops.h @@ -107,12 +107,12 @@ class ReduceMean final : public ReduceKernel { class ReduceMax final : public ReduceKernel { public: - ReduceMax(const OpKernelInfo& info) : ReduceKernel(info, "ReduceMax") {} + ReduceMax(const OpKernelInfo& info) : ReduceKernel(info, "ReduceMax", true) {} }; class ReduceMin final : public ReduceKernel { public: - ReduceMin(const OpKernelInfo& info) : ReduceKernel(info, "ReduceMin") {} + ReduceMin(const OpKernelInfo& info) : ReduceKernel(info, "ReduceMin", true) {} }; class ReduceSum final : public ReduceKernel { diff --git a/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc b/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc index 79617dc16e1f5..de133d4dac4f5 100644 --- a/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc +++ b/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc @@ -6050,7 +6050,6 @@ void test_empty_set(const std::string& op, int opset, bool axes_as_input, float kOpenVINOExecutionProvider, kQnnExecutionProvider, kTensorrtExecutionProvider, - kWebGpuExecutionProvider, }); } @@ -6318,209 +6317,324 @@ TEST(ReductionOpTest, ReduceSumSquare_NoopWithAxesNotProvided_ElementwiseSquare) test.ConfigEp(DefaultCpuExecutionProvider()).RunWithConfig(); } -// Opset 20 tests for ReduceMax and ReduceMin on CUDA. -// Verifies CUDA kernel registration at opset 20 works for all supported types. -#if defined(USE_CUDA) +// ============================================================================= +// Empty tensor reduction tests — ONNX spec compliance +// +// Per ONNX spec, reducing over an empty set produces the operator identity: +// ReduceSum → 0, ReduceProd → 1, ReduceMin → +inf, ReduceMax → -inf, +// ReduceMean → 0 (undefined per ONNX spec; ORT fills with 0 for consistency) +// +// These tests verify both CPU and CUDA providers handle empty tensors +// correctly for various shapes, keepdims settings, and axis configurations. +// +// Most other EPs (WebGPU, QNN, TensorRT, CoreML, DML, etc.) do not support +// empty tensor reductions, so they are excluded. +// ============================================================================= + +// EPs that failed CI on empty tensor reductions. +// Only exclude EPs with confirmed failures; leave others enabled so CI +// catches regressions if they arise. +// WebGPU: fixed in this PR (identity values for empty inputs). +const std::unordered_set kEmptyTensorExcludedEps = { + kDmlExecutionProvider, + kNnapiExecutionProvider, + kQnnExecutionProvider, + kTensorrtExecutionProvider, +}; -TEST(ReductionOpTest, ReduceMax_float_Opset20_Cuda) { - OpTester test("ReduceMax", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1.0f, 2.0f, 3.0f, 4.0f, - 5.0f, 6.0f, 7.0f, 8.0f, - 9.0f, 10.0f, 11.0f, 12.0f}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {4.0f, 8.0f, 12.0f}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +// --- ReduceSum empty tensor tests --- + +TEST(ReductionOpTest, ReduceSum_EmptyTensor_ExplicitAxis) { + // {1, 0} reduce axis 1 → {1}, identity = 0 + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMax_double_Opset20_Cuda) { - OpTester test("ReduceMax", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1.0, 2.0, 3.0, 4.0, - 5.0, 6.0, 7.0, 8.0, - 9.0, 10.0, 11.0, 12.0}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {4.0, 8.0, 12.0}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceSum_EmptyTensor_ExplicitAxis_KeepDims) { + // {1, 0} reduce axis 1 keepdims → {1, 1}, identity = 0 + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1, 1}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMax_half_Opset20_Cuda) { - OpTester test("ReduceMax", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - FloatsToMLFloat16s({1.0f, 2.0f, 3.0f, 4.0f, - 5.0f, 6.0f, 7.0f, 8.0f, - 9.0f, 10.0f, 11.0f, 12.0f})); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, FloatsToMLFloat16s({4.0f, 8.0f, 12.0f})); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceSum_EmptyTensor_MultiDim) { + // {2, 0, 3} reduce axis 1 → {2, 3}, all zeros + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {2, 3}, {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMax_int32_Opset20_Cuda) { - OpTester test("ReduceMax", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1, 2, 3, 4, - 5, 6, 7, 8, - 9, 10, 11, 12}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {4, 8, 12}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceSum_EmptyTensor_MultiDim_KeepDims) { + // {2, 0, 3} reduce axis 1 keepdims → {2, 1, 3}, all zeros + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {2, 1, 3}, {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMax_int64_Opset20_Cuda) { - OpTester test("ReduceMax", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1, 2, 3, 4, - 5, 6, 7, 8, - 9, 10, 11, 12}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {4, 8, 12}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceSum_EmptyTensor_DefaultAxes) { + // {1, 0} reduce all axes → scalar, identity = 0 + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(0)); + test.AddAttribute("noop_with_empty_axes", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_float_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1.0f, 2.0f, 3.0f, 4.0f, - 5.0f, 6.0f, 7.0f, 8.0f, - 9.0f, 10.0f, 11.0f, 12.0f}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {1.0f, 5.0f, 9.0f}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceSum_EmptyTensor_Shape_0_5) { + // {0, 5} reduce axis 0 → {5}, all zeros + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {0, 5}, {}); + test.AddInput("axes", {1}, {0}); + test.AddOutput("reduced", {5}, {0.0f, 0.0f, 0.0f, 0.0f, 0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_double_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1.0, 2.0, 3.0, 4.0, - 5.0, 6.0, 7.0, 8.0, - 9.0, 10.0, 11.0, 12.0}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {1.0, 5.0, 9.0}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +// --- ReduceMax empty tensor tests --- + +TEST(ReductionOpTest, ReduceMax_EmptyTensor_ExplicitAxis) { + // {1, 0} reduce axis 1 → {1}, identity = -inf + OpTester test("ReduceMax", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1}, {-std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_half_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - FloatsToMLFloat16s({1.0f, 2.0f, 3.0f, 4.0f, - 5.0f, 6.0f, 7.0f, 8.0f, - 9.0f, 10.0f, 11.0f, 12.0f})); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, FloatsToMLFloat16s({1.0f, 5.0f, 9.0f})); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceMax_EmptyTensor_ExplicitAxis_KeepDims) { + // {1, 0} reduce axis 1 keepdims → {1, 1}, identity = -inf + OpTester test("ReduceMax", 18); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1, 1}, {-std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_int32_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1, 2, 3, 4, - 5, 6, 7, 8, - 9, 10, 11, 12}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {1, 5, 9}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceMax_EmptyTensor_MultiDim) { + // {2, 0, 3} reduce axis 1 → {2, 3}, all -inf + OpTester test("ReduceMax", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {1}); + float neg_inf = -std::numeric_limits::infinity(); + test.AddOutput("reduced", {2, 3}, {neg_inf, neg_inf, neg_inf, neg_inf, neg_inf, neg_inf}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_int64_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1, 2, 3, 4, - 5, 6, 7, 8, - 9, 10, 11, 12}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {1, 5, 9}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceMax_EmptyTensor_DefaultAxes) { + // {1, 0} reduce all → scalar, identity = -inf + OpTester test("ReduceMax", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {}, {-std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_int8_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1, 2, 3, 4, - 5, 6, 7, 8, - 9, 10, 11, 12}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {1, 5, 9}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +// --- ReduceMin empty tensor tests --- + +TEST(ReductionOpTest, ReduceMin_EmptyTensor_ExplicitAxis) { + // {1, 0} reduce axis 1 → {1}, identity = +inf + OpTester test("ReduceMin", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1}, {std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -TEST(ReductionOpTest, ReduceMin_uint8_Opset20_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)1); - test.AddInput("data", {3, 2, 2}, - {1, 2, 3, 4, - 5, 6, 7, 8, - 9, 10, 11, 12}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3, 1, 1}, {1, 5, 9}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceMin_EmptyTensor_ExplicitAxis_KeepDims) { + // {1, 0} reduce axis 1 keepdims → {1, 1}, identity = +inf + OpTester test("ReduceMin", 18); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1, 1}, {std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -// Test ReduceMax at opset 20 with keepdims=0 on CUDA -TEST(ReductionOpTest, ReduceMax_float_Opset20_NoKeepdims_Cuda) { - OpTester test("ReduceMax", 20); - test.AddAttribute("keepdims", (int64_t)0); - test.AddInput("data", {3, 2, 2}, - {1.0f, 2.0f, 3.0f, 4.0f, - 5.0f, 6.0f, 7.0f, 8.0f, - 9.0f, 10.0f, 11.0f, 12.0f}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3}, {4.0f, 8.0f, 12.0f}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceMin_EmptyTensor_MultiDim) { + // {2, 0, 3} reduce axis 1 → {2, 3}, all +inf + OpTester test("ReduceMin", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {1}); + float pos_inf = std::numeric_limits::infinity(); + test.AddOutput("reduced", {2, 3}, {pos_inf, pos_inf, pos_inf, pos_inf, pos_inf, pos_inf}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -// Test ReduceMin at opset 20 with keepdims=0 on CUDA -TEST(ReductionOpTest, ReduceMin_float_Opset20_NoKeepdims_Cuda) { - OpTester test("ReduceMin", 20); - test.AddAttribute("keepdims", (int64_t)0); - test.AddInput("data", {3, 2, 2}, - {1.0f, 2.0f, 3.0f, 4.0f, - 5.0f, 6.0f, 7.0f, 8.0f, - 9.0f, 10.0f, 11.0f, 12.0f}); - test.AddInput("axes", {2}, {1, 2}); - test.AddOutput("reduced", {3}, {1.0f, 5.0f, 9.0f}); - std::vector> execution_providers; - execution_providers.push_back(DefaultCudaExecutionProvider()); - test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +TEST(ReductionOpTest, ReduceMin_EmptyTensor_DefaultAxes) { + // {1, 0} reduce all → scalar, identity = +inf + OpTester test("ReduceMin", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {}, {std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); } -#endif // defined(USE_CUDA) +// --- ReduceProd empty tensor tests --- + +TEST(ReductionOpTest, ReduceProd_EmptyTensor_ExplicitAxis) { + // {1, 0} reduce axis 1 → {1}, identity = 1 + OpTester test("ReduceProd", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1}, {1.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceProd_EmptyTensor_ExplicitAxis_KeepDims) { + // {1, 0} reduce axis 1 keepdims → {1, 1}, identity = 1 + OpTester test("ReduceProd", 18); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1, 1}, {1.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceProd_EmptyTensor_MultiDim) { + // {2, 0, 3} reduce axis 1 → {2, 3}, all ones + OpTester test("ReduceProd", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {2, 3}, {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceProd_EmptyTensor_DefaultAxes) { + // {1, 0} reduce all → scalar, identity = 1 + OpTester test("ReduceProd", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {}, {1.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +// --- ReduceMean empty tensor tests --- + +TEST(ReductionOpTest, ReduceMean_EmptyTensor_ExplicitAxis) { + // {1, 0} reduce axis 1 → {1}, mean of empty set = 0/0. + // Undefined per ONNX spec; ORT fills with 0 for consistency with CPU/CUDA. + OpTester test("ReduceMean", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {1}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceMean_EmptyTensor_MultiDim) { + // {2, 0, 3} reduce axis 1 → {2, 3}, all zeros + OpTester test("ReduceMean", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {1}); + test.AddOutput("reduced", {2, 3}, {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceMean_EmptyTensor_DefaultAxes) { + // {1, 0} reduce all → scalar, empty mean = 0 + OpTester test("ReduceMean", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +// --- Negative axes + empty tensor --- + +TEST(ReductionOpTest, ReduceSum_EmptyTensor_NegativeAxis) { + // {1, 0} reduce axis -1 (= axis 1) → {1}, identity = 0 + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddInput("axes", {1}, {-1}); + test.AddOutput("reduced", {1}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceMax_EmptyTensor_NegativeAxis) { + // {2, 0, 3} reduce axis -2 (= axis 1) → {2, 3}, identity = -inf + OpTester test("ReduceMax", 18); + test.AddAttribute("keepdims", int64_t(0)); + test.AddInput("data", {2, 0, 3}, {}); + test.AddInput("axes", {1}, {-2}); + float neg_inf = -std::numeric_limits::infinity(); + test.AddOutput("reduced", {2, 3}, {neg_inf, neg_inf, neg_inf, neg_inf, neg_inf, neg_inf}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +// --- noop_with_empty_axes + empty tensor --- + +TEST(ReductionOpTest, ReduceSum_EmptyTensor_NoopWithEmptyAxes) { + // {1, 0} with noop_with_empty_axes=1 and no axes → no-op, shape preserved + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(0)); + test.AddAttribute("noop_with_empty_axes", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + // No axes input → noop: output = input (same shape {1, 0}, empty) + test.AddOutput("reduced", {1, 0}, {}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +// --- Default axes + keepdims + empty tensor --- + +TEST(ReductionOpTest, ReduceSum_EmptyTensor_DefaultAxes_KeepDims) { + // {1, 0} reduce all axes with keepdims → {1, 1}, identity = 0 + OpTester test("ReduceSum", 13); + test.AddAttribute("keepdims", int64_t(1)); + test.AddAttribute("noop_with_empty_axes", int64_t(0)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {1, 1}, {0.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceMax_EmptyTensor_DefaultAxes_KeepDims) { + // {1, 0} reduce all axes with keepdims → {1, 1}, identity = -inf + OpTester test("ReduceMax", 18); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {1, 1}, {-std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +TEST(ReductionOpTest, ReduceMin_EmptyTensor_DefaultAxes_KeepDims) { + // {1, 0} reduce all axes with keepdims → {1, 1}, identity = +inf + OpTester test("ReduceMin", 18); + test.AddAttribute("keepdims", int64_t(1)); + test.AddInput("data", {1, 0}, {}); + test.AddOutput("reduced", {1, 1}, {std::numeric_limits::infinity()}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kEmptyTensorExcludedEps); +} + +// --- Bool reduction + empty tensor --- + +// NOTE: Bool ReduceMax/ReduceMin empty tensor tests omitted. +// ORT registers bool CPU kernels for ReduceMax/ReduceMin, but the ONNX +// schema at opset 18+ does not list bool as a valid input type. OpTester +// validation rejects them before execution. The fill_for_empty_set +// implementation is correct (Max→false, Min→true per spec) but cannot +// be tested through the standard graph-based test path. } // namespace test } // namespace onnxruntime