diff --git a/onnxruntime/core/providers/cuda/nn/avg_pool_impl.cu b/onnxruntime/core/providers/cuda/nn/avg_pool_impl.cu new file mode 100644 index 0000000000000..028f3f2bde9a4 --- /dev/null +++ b/onnxruntime/core/providers/cuda/nn/avg_pool_impl.cu @@ -0,0 +1,247 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "avg_pool_impl.h" + +#include "core/providers/cuda/cu_inc/common.cuh" +#include "core/providers/cuda/shared_inc/fast_divmod.h" +#include "core/providers/cuda/shared_inc/cuda_utils.h" + +namespace onnxruntime { +namespace cuda { + +// Accumulate half in float for precision; keep native type otherwise. +template +struct AveragePoolAccumulator { + using type = T; +}; +template <> +struct AveragePoolAccumulator { + using type = float; +}; +template <> +struct AveragePoolAccumulator { + using type = float; +}; + +template +__global__ void AveragePoolWithPadKernel( + int64_t channels, + int64_t height, + int64_t width, + int64_t depth, + int64_t pooled_height, + int64_t pooled_width, + int64_t pooled_depth, + int64_t kernel_h, + int64_t kernel_w, + int64_t kernel_d, + int64_t stride_h, + int64_t stride_w, + int64_t stride_d, + int64_t pad_h_head, + int64_t pad_w_head, + int64_t pad_d_head, + int64_t pad_h_tail, + int64_t pad_w_tail, + int64_t pad_d_tail, + int64_t dilation_h, + int64_t dilation_w, + int64_t dilation_d, + fast_divmod fdm_c, + fast_divmod fdm_h, + fast_divmod fdm_w, + fast_divmod fdm_d, + bool count_include_pad, + const T* p_input, + int64_t output_size, + T* p_output) { + int id = blockIdx.x * blockDim.x + threadIdx.x; + if (id >= output_size) return; + + auto compute_offset = + [height, width, depth, channels](int n_index, int c_index, int h_index, int w_index, int d_index) -> int64_t { + if constexpr (Layout == LAYOUT_NCHW) { + return (((n_index * channels + c_index) * height + h_index) * width + w_index) * depth + d_index; + } else if constexpr (Layout == LAYOUT_NHWC) { + return (((n_index * height + h_index) * width + w_index) * depth + d_index) * channels + c_index; + } + }; + + int d_index, w_index, h_index, c_index, n_index, id_tmp; + if constexpr (Layout == LAYOUT_NCHW) { + fdm_d.divmod(id, id_tmp, d_index); + fdm_w.divmod(id_tmp, id_tmp, w_index); + fdm_h.divmod(id_tmp, id_tmp, h_index); + fdm_c.divmod(id_tmp, n_index, c_index); + } else if constexpr (Layout == LAYOUT_NHWC) { + fdm_c.divmod(id, id_tmp, c_index); + fdm_d.divmod(id_tmp, id_tmp, d_index); + fdm_w.divmod(id_tmp, id_tmp, w_index); + fdm_h.divmod(id_tmp, n_index, h_index); + } + + // Window bounds mirror the CPU AveragePool{1,2,3}DTask reference exactly. + int64_t h_start = h_index * stride_h - pad_h_head; + int64_t w_start = w_index * stride_w - pad_w_head; + int64_t d_start = d_index * stride_d - pad_d_head; + + int64_t h_end = _Min(h_start + kernel_h * dilation_h, height + pad_h_tail); + int64_t w_end = _Min(w_start + kernel_w * dilation_w, width + pad_w_tail); + int64_t d_end = _Min(d_start + kernel_d * dilation_d, depth + pad_d_tail); + + using AccT = typename AveragePoolAccumulator::type; + AccT acc = static_cast(0); + int64_t counted = 0; + + int64_t offset = compute_offset(n_index, c_index, 0, 0, 0); + const T* p_slice = p_input + offset; + for (int64_t h = h_start; h < h_end; h += dilation_h) { + if (h < 0 || h >= height) continue; + for (int64_t w = w_start; w < w_end; w += dilation_w) { + if (w < 0 || w >= width) continue; + for (int64_t d = d_start; d < d_end; d += dilation_d) { + if (d < 0 || d >= depth) continue; + acc += static_cast(p_slice[compute_offset(0, 0, h, w, d)]); + ++counted; + } + } + } + + AccT result = static_cast(0); + if (counted > 0) { + if (count_include_pad) { + int64_t divisor = (1 + (h_end - h_start - 1) / dilation_h) * + (1 + (w_end - w_start - 1) / dilation_w) * + (1 + (d_end - d_start - 1) / dilation_d); + result = acc / static_cast(divisor); + } else { + result = acc / static_cast(counted); + } + } + p_output[id] = static_cast(result); +} + +template +void AveragePoolWithPad( + cudaStream_t stream, + const TensorShape& input_shape, + const TensorShape& output_shape, + const gsl::span& kernel_shape, + const gsl::span& stride_shape, + const gsl::span& pads, + const gsl::span& dilations, + bool count_include_pad, + const T* p_input, + T* p_output) { + int64_t channels, height, width, depth; + int64_t pooled_height, pooled_width, pooled_depth; + if constexpr (Layout == LAYOUT_NCHW) { + channels = input_shape[1]; + height = input_shape[2]; + width = kernel_shape.size() > 1 ? input_shape[3] : 1; + depth = kernel_shape.size() > 2 ? input_shape[4] : 1; + + pooled_height = output_shape[2]; + pooled_width = kernel_shape.size() > 1 ? output_shape[3] : 1; + pooled_depth = kernel_shape.size() > 2 ? output_shape[4] : 1; + } else if constexpr (Layout == LAYOUT_NHWC) { + height = input_shape[1]; + width = kernel_shape.size() > 1 ? input_shape[2] : 1; + depth = kernel_shape.size() > 2 ? input_shape[3] : 1; + channels = input_shape[input_shape.NumDimensions() - 1]; + + pooled_height = output_shape[1]; + pooled_width = kernel_shape.size() > 1 ? output_shape[2] : 1; + pooled_depth = kernel_shape.size() > 2 ? output_shape[3] : 1; + } + + const int64_t rank = static_cast(kernel_shape.size()); + int64_t kernel_h = kernel_shape[0]; + int64_t kernel_w = rank > 1 ? kernel_shape[1] : 1; + int64_t kernel_d = rank > 2 ? kernel_shape[2] : 1; + int64_t stride_h = stride_shape[0]; + int64_t stride_w = rank > 1 ? stride_shape[1] : 1; + int64_t stride_d = rank > 2 ? stride_shape[2] : 1; + + // pads: [x1_begin,...,xN_begin, x1_end,...,xN_end]. Begin at [i], end at [rank + i]. + int64_t pad_h_head = pads[0]; + int64_t pad_w_head = rank > 1 ? pads[1] : 0; + int64_t pad_d_head = rank > 2 ? pads[2] : 0; + int64_t pad_h_tail = pads[rank + 0]; + int64_t pad_w_tail = rank > 1 ? pads[rank + 1] : 0; + int64_t pad_d_tail = rank > 2 ? pads[rank + 2] : 0; + + int64_t dilation_h = dilations[0]; + int64_t dilation_w = rank > 1 ? dilations[1] : 1; + int64_t dilation_d = rank > 2 ? dilations[2] : 1; + + int64_t output_size = output_shape.Size(); + if (output_size == 0) return; + + fast_divmod fdm_c(static_cast(channels)); + fast_divmod fdm_h(static_cast(pooled_height)); + fast_divmod fdm_w(static_cast(pooled_width)); + fast_divmod fdm_d(static_cast(pooled_depth)); + + int blocksPerGrid = (int)((output_size + GridDim::maxThreadsPerBlock - 1) / GridDim::maxThreadsPerBlock); + AveragePoolWithPadKernel<<>>( + channels, + height, + width, + depth, + pooled_height, + pooled_width, + pooled_depth, + kernel_h, + kernel_w, + kernel_d, + stride_h, + stride_w, + stride_d, + pad_h_head, + pad_w_head, + pad_d_head, + pad_h_tail, + pad_w_tail, + pad_d_tail, + dilation_h, + dilation_w, + dilation_d, + fdm_c, + fdm_h, + fdm_w, + fdm_d, + count_include_pad, + p_input, + output_size, + p_output); +} + +#define INSTANTIATE_AVERAGEPOOLWITHPAD(T, Layout) \ + template void AveragePoolWithPad( \ + cudaStream_t stream, \ + const TensorShape& input_shape, \ + const TensorShape& output_shape, \ + const gsl::span& kernel_shape, \ + const gsl::span& stride_shape, \ + const gsl::span& pads, \ + const gsl::span& dilations, \ + bool count_include_pad, \ + const T* p_input, \ + T* p_output); + +INSTANTIATE_AVERAGEPOOLWITHPAD(float, LAYOUT_NCHW) +INSTANTIATE_AVERAGEPOOLWITHPAD(double, LAYOUT_NCHW) +INSTANTIATE_AVERAGEPOOLWITHPAD(half, LAYOUT_NCHW) +INSTANTIATE_AVERAGEPOOLWITHPAD(BFloat16, LAYOUT_NCHW) + +#ifdef ENABLE_CUDA_NHWC_OPS +INSTANTIATE_AVERAGEPOOLWITHPAD(float, LAYOUT_NHWC) +INSTANTIATE_AVERAGEPOOLWITHPAD(double, LAYOUT_NHWC) +INSTANTIATE_AVERAGEPOOLWITHPAD(half, LAYOUT_NHWC) +INSTANTIATE_AVERAGEPOOLWITHPAD(BFloat16, LAYOUT_NHWC) +#endif + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/nn/avg_pool_impl.h b/onnxruntime/core/providers/cuda/nn/avg_pool_impl.h new file mode 100644 index 0000000000000..9cf51a5442f9f --- /dev/null +++ b/onnxruntime/core/providers/cuda/nn/avg_pool_impl.h @@ -0,0 +1,39 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/framework/tensor_shape.h" + +namespace onnxruntime { +namespace cuda { + +// Custom average-pooling CUDA kernel that honors per-side (asymmetric) padding. +// +// cuDNN's pooling descriptor stores a single symmetric pad value per axis, so it cannot +// represent ONNX asymmetric padding (pad_begin != pad_end), which arises from explicit +// asymmetric `pads` or `auto_pad = SAME_UPPER/SAME_LOWER` resolving to asymmetric pads. This +// kernel is the CUDA fallback for that case and mirrors the CPU reference functor +// (AveragePool{1,2,3}DTask) exactly: +// start = out_idx * stride - pad_begin +// end = min(start + kernel * dilation, in_size + pad_end) +// sum over cells [start, end) with dilation step that are in [0, in_size) +// count_include_pad == 1: divisor = product of (1 + (end - start - 1) / dilation) +// count_include_pad == 0: divisor = number of summed in-bounds cells +// +// `pads` is the full ONNX layout [x1_begin,...,xN_begin, x1_end,...,xN_end] (2 * rank). +template +void AveragePoolWithPad( + cudaStream_t stream, + const TensorShape& input_shape, + const TensorShape& output_shape, + const gsl::span& kernel_shape, + const gsl::span& stride_shape, + const gsl::span& pads, + const gsl::span& dilations, + bool count_include_pad, + const T* p_input, + T* p_output); + +} // namespace cuda +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/nn/pool.cc b/onnxruntime/core/providers/cuda/nn/pool.cc index 89c6047769cf0..b955e68268a2f 100644 --- a/onnxruntime/core/providers/cuda/nn/pool.cc +++ b/onnxruntime/core/providers/cuda/nn/pool.cc @@ -6,6 +6,7 @@ #include "core/providers/cuda/nn/pool.h" #include "core/providers/cuda/cudnn_common.h" #include "core/providers/cuda/nn/max_pool_with_index.h" +#include "core/providers/cuda/nn/avg_pool_impl.h" #include "core/providers/cuda/math/unary_elementwise_ops_impl.h" using namespace onnxruntime::common; @@ -205,6 +206,37 @@ Status Pool::ComputeInternal(OpKernelContext* context) cons auto x_data = reinterpret_cast(X->Data()); auto y_data = reinterpret_cast(Y->MutableData()); + // cuDNN's pooling descriptor cannot represent two ONNX features: + // (1) Asymmetric padding: it stores a single symmetric pad value per axis and applies it to + // both sides, so it silently drops ONNX end pads when pad_begin != pad_end (explicit + // asymmetric pads, or auto_pad=SAME_UPPER/SAME_LOWER resolving to asymmetric pads). + // (2) Dilation: the pooling descriptor has no dilation parameter at all, so any dilation > 1 + // is silently ignored. + // Either case produces wrong sums and divisors on cuDNN, so route it to the custom kernel, + // which honors per-side pads AND dilation and matches the CPU reference divisor exactly. + // Symmetric, non-dilated pooling (the common case, including all global pooling) keeps the + // fast cuDNN path unchanged, so there is zero perf regression. (MaxPool<8> guards dilation the + // same way via !default_dilations.) Global pooling is always symmetric and never dilated, and + // its kernel_shape/strides/dilations are left unpopulated (PoolAttributes returns early), so it + // is excluded here and stays on cuDNN. + if constexpr (PoolType::type == onnxruntime::PoolType::kAveragePool) { + if (!pool_attrs_.global_pooling) { + const size_t spatial_rank = kernel_shape.size(); + bool asymmetric_pads = false; + for (size_t i = 0; i < spatial_rank; ++i) { + if (pads[i] != pads[spatial_rank + i]) { + asymmetric_pads = true; + break; + } + } + if (asymmetric_pads || !pool_attrs_.default_dilations) { + AveragePoolWithPad(Stream(context), x_shape, y_shape, kernel_shape, strides, pads, + pool_attrs_.dilations, pool_attrs_.count_include_pad, x_data, y_data); + return Status::OK(); + } + } + } + TensorShapeVector x_dims_cudnn(x_dims.begin(), x_dims.end()); TensorShapeVector y_dims_cudnn(y_dims); if (kernel_shape.size() < 2) { diff --git a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc index e7f870fc43d7f..95c0c9d35c7f0 100644 --- a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc +++ b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc @@ -959,10 +959,13 @@ TEST(PoolTest, AveragePool_CountIncludePad_AsymmetricPads) { test.AddInput("X", x_dims, x_vals); test.AddOutput("Y", expected_dims, expected_vals); - // This test targets the CPU fix only. Exclude EPs whose external libraries - // (cuDNN, CoreML, etc.) also produce wrong results for this case. + // The CUDA custom AveragePoolWithPad kernel now honors per-side (asymmetric) pads, so the + // CUDA (NCHW) leg is un-excluded here to lock in that fix. kCudaNHWCExecutionProvider is + // excluded here only to avoid redundant coverage: the asymmetric NHWC-CUDA decode branch is + // now exercised (and passing) by the 1D/2D AveragePool_CUDA_* parity tests below. The remaining + // exclusions are EPs whose external libraries (CoreML, etc.) still produce wrong results here. test.Run(OpTester::ExpectResult::kExpectSuccess, "", - {kCudaExecutionProvider, kCudaNHWCExecutionProvider, + {kCudaNHWCExecutionProvider, kTensorrtExecutionProvider, kAclExecutionProvider, kOpenVINOExecutionProvider, kDnnlExecutionProvider, kCoreMLExecutionProvider, kQnnExecutionProvider, kDmlExecutionProvider}); @@ -994,10 +997,14 @@ TEST(PoolTest, AveragePool3D_CountIncludePad_AsymmetricPads) { 0.5f, 0.25f}; test.AddInput("X", x3d_dims, x3d_vals); test.AddOutput("Y", expected3d_dims, expected3d_vals); - // This test targets the CPU fix only. Exclude EPs whose external libraries - // (cuDNN, CoreML, etc.) also produce wrong results for this case. + // The CUDA custom AveragePoolWithPad kernel now honors per-side (asymmetric) pads in 3D, so the + // CUDA (NCHW) leg is un-excluded here to lock in that fix. kCudaNHWCExecutionProvider stays + // excluded here: the asymmetric NHWC-CUDA decode branch is now exercised (and passing) by the + // 1D/2D AveragePool_CUDA_* parity tests below, but 3D (NDHWC) NHWC pooling is not among them, so + // it remains a follow-up. The remaining exclusions are EPs whose external libraries (CoreML, + // etc.) still produce wrong results here. test.Run(OpTester::ExpectResult::kExpectSuccess, "", - {kCudaExecutionProvider, kCudaNHWCExecutionProvider, + {kCudaNHWCExecutionProvider, kTensorrtExecutionProvider, kAclExecutionProvider, kOpenVINOExecutionProvider, kDnnlExecutionProvider, kCoreMLExecutionProvider, kQnnExecutionProvider, kDmlExecutionProvider}); @@ -1095,6 +1102,282 @@ TEST(PoolTest, AveragePool_19_ceil_count_include_pad_1d) { {kTensorrtExecutionProvider, kAclExecutionProvider, kOpenVINOExecutionProvider, kDmlExecutionProvider}); } +// --------------------------------------------------------------------------- +// CUDA AveragePool asymmetric-padding parity tests. +// +// cuDNN's pooling descriptor stores one symmetric pad per axis, so it silently drops the +// ONNX end pad when pad_begin != pad_end, producing wrong averages on CUDA (e.g. for a 1D +// pad of (0,3) cuDNN yields [4, 6.5, 8] while the CPU reference yields [4, 5.571, 4], and a +// 2D pad of (0,0,3,3) diverges by up to 53.25). The custom AveragePoolWithPad CUDA kernel +// fixes this. These cases keep the CUDA EP UN-excluded so the CUDA leg actually runs and must +// match the CPU reference oracle. Expected values are the CPU reference outputs. +// +// ceil_mode + count_include_pad cases use opset 19 so the CPU leg runs the already-correct v19 +// reference functor and validates the CUDA kernel independently of the separate CPU opset-7..18 +// MLAS fix (PR #29629); the CUDA routing is opset-independent, so this still exercises the fix. +// +// Exception: the fp16 case below runs CUDA-only (CPU has no fp16 AveragePool kernel on x64 and +// the Arm64 NEON fp16 pooling kernel mishandles the ceil_mode + count_include_pad divisor), so +// the "CPU leg runs the v19 reference oracle" statement above does not apply to it — see its own +// comment for how it validates the CUDA half accumulate-in-float path without a CPU oracle. +// +// The float cases share a single exclusion set (kPoolingEpsExcludedFromCeilCipTests) so the list cannot +// drift test-to-test. It names every EP whose pooling does NOT implement ONNX's asymmetric-pad / +// dilated / ceil_mode + count_include_pad clamped-divisor semantics (they would produce wrong +// values and cannot serve as an oracle). The CPU EP (correct v19 reference) stays un-excluded as +// the float oracle, and the CUDA + CUDA-NHWC EPs stay un-excluded as the tested targets — the +// asymmetric NHWC-CUDA path is intentionally exercised here (USE_CUDA_NHWC_OPS defaults ON) and +// passes. kWebGpuExecutionProvider is listed defensively: it auto-skips in a CUDA-only build +// (DefaultWebGpuExecutionProvider returns nullptr), but naming it keeps a future WebGPU build leg +// from re-triggering the CI failure this list fixes. +// --------------------------------------------------------------------------- +const std::unordered_set kPoolingEpsExcludedFromCeilCipTests = { + kTensorrtExecutionProvider, kNvTensorRTRTXExecutionProvider, kDnnlExecutionProvider, + kOpenVINOExecutionProvider, kAclExecutionProvider, kCoreMLExecutionProvider, + kQnnExecutionProvider, kDmlExecutionProvider, kWebGpuExecutionProvider}; + +TEST(PoolTest, AveragePool_CUDA_asymmetric_tail_pad_1d) { + OpTester test("AveragePool", 19); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3}); + test.AddAttribute("pads", std::vector{0, 3}); + test.AddAttribute("kernel_shape", std::vector{7}); + test.AddAttribute("ceil_mode", (int64_t)1); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f}; + std::vector x_dims = {1, 1, 9}; + std::vector expected_dims = {1, 1, 3}; + std::vector expected_vals = {4.0f, 5.5714283f, 4.0f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +TEST(PoolTest, AveragePool_CUDA_asymmetric_tail_pad_1d_exclude_pad) { + OpTester test("AveragePool", 18); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3}); + test.AddAttribute("pads", std::vector{0, 3}); + test.AddAttribute("kernel_shape", std::vector{7}); + test.AddAttribute("ceil_mode", (int64_t)1); + test.AddAttribute("count_include_pad", (int64_t)0); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f}; + std::vector x_dims = {1, 1, 9}; + std::vector expected_dims = {1, 1, 3}; + // exclude-pad divides by in-bounds cells only. + std::vector expected_vals = {4.0f, 6.5f, 8.0f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +TEST(PoolTest, AveragePool_CUDA_asymmetric_tail_pad_2d) { + OpTester test("AveragePool", 19); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3, 3}); + test.AddAttribute("pads", std::vector{0, 0, 3, 3}); + test.AddAttribute("kernel_shape", std::vector{7, 7}); + test.AddAttribute("ceil_mode", (int64_t)1); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals(81); + for (int i = 0; i < 81; ++i) { + x_vals[i] = static_cast(i + 1); + } + std::vector x_dims = {1, 1, 9, 9}; + std::vector expected_dims = {1, 1, 3, 3}; + std::vector expected_vals = {31.0f, 28.714287f, 17.5f, + 45.857143f, 41.142857f, 24.642858f, + 33.5f, 29.785715f, 17.75f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +TEST(PoolTest, AveragePool_CUDA_asymmetric_tail_pad_2d_exclude_pad) { + OpTester test("AveragePool", 18); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3, 3}); + test.AddAttribute("pads", std::vector{0, 0, 3, 3}); + test.AddAttribute("kernel_shape", std::vector{7, 7}); + test.AddAttribute("ceil_mode", (int64_t)1); + test.AddAttribute("count_include_pad", (int64_t)0); + + std::vector x_vals(81); + for (int i = 0; i < 81; ++i) { + x_vals[i] = static_cast(i + 1); + } + std::vector x_dims = {1, 1, 9, 9}; + std::vector expected_dims = {1, 1, 3, 3}; + std::vector expected_vals = {31.0f, 33.5f, 35.0f, + 53.5f, 56.0f, 57.5f, + 67.0f, 69.5f, 71.0f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +// auto_pad=SAME_UPPER produces naturally asymmetric pads (here pad(0,1)); proves the latent +// SAME-pad bug on CUDA is also fixed by the same kernel. +TEST(PoolTest, AveragePool_CUDA_same_upper_asymmetric_1d) { + OpTester test("AveragePool", 18); + + test.AddAttribute("auto_pad", "SAME_UPPER"); + test.AddAttribute("strides", std::vector{2}); + test.AddAttribute("kernel_shape", std::vector{3}); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f, 10.0f}; + std::vector x_dims = {1, 1, 10}; + std::vector expected_dims = {1, 1, 5}; + std::vector expected_vals = {2.0f, 4.0f, 6.0f, 8.0f, 6.3333335f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +// Regression guard: symmetric pads must STAY on the fast cuDNN path and remain correct. +TEST(PoolTest, AveragePool_CUDA_symmetric_pad_regression_1d) { + OpTester test("AveragePool", 19); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3}); + test.AddAttribute("pads", std::vector{3, 3}); + test.AddAttribute("kernel_shape", std::vector{7}); + test.AddAttribute("ceil_mode", (int64_t)1); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f}; + std::vector x_dims = {1, 1, 9}; + std::vector expected_dims = {1, 1, 4}; + std::vector expected_vals = {1.4285715f, 4.0f, 5.5714283f, 4.0f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +// MaxPool asymmetric-pad probe/regression on CUDA (verify-then-decide per design). MaxPool +// ignores pad cells (no divisor); asymmetric tail pad only changes output size, computed +// correctly upstream. CUDA un-excluded to confirm parity with the CPU reference. +TEST(PoolTest, MaxPool_CUDA_asymmetric_tail_pad_1d) { + OpTester test("MaxPool", 12); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3}); + test.AddAttribute("pads", std::vector{0, 3}); + test.AddAttribute("kernel_shape", std::vector{7}); + test.AddAttribute("ceil_mode", (int64_t)1); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f}; + std::vector x_dims = {1, 1, 9}; + std::vector expected_dims = {1, 1, 3}; + std::vector expected_vals = {7.0f, 9.0f, 9.0f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +// fp16 asymmetric-pad AveragePool. Runs on the CUDA EP ONLY (via an explicit provider list): +// the CPU AveragePool has no fp16 kernel on x64, and the Arm64 NEON fp16 pooling kernel does +// not honor the ceil_mode + count_include_pad divisor rule (a separate, pre-existing CPU +// limitation), so it cannot serve as the fp16 oracle. This test validates that the CUDA +// AveragePoolWithPad kernel's half accumulate-in-float path matches the reference values. +TEST(PoolTest, AveragePool_CUDA_asymmetric_tail_pad_1d_fp16) { + auto cuda_ep = DefaultCudaExecutionProvider(); + if (!cuda_ep) { + return; + } + + OpTester test("AveragePool", 19); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{3}); + test.AddAttribute("pads", std::vector{0, 3}); + test.AddAttribute("kernel_shape", std::vector{7}); + test.AddAttribute("ceil_mode", (int64_t)1); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals = {MLFloat16(1.0f), MLFloat16(2.0f), MLFloat16(3.0f), + MLFloat16(4.0f), MLFloat16(5.0f), MLFloat16(6.0f), + MLFloat16(7.0f), MLFloat16(8.0f), MLFloat16(9.0f)}; + std::vector x_dims = {1, 1, 9}; + std::vector expected_dims = {1, 1, 3}; + std::vector expected_vals = {MLFloat16(4.0f), MLFloat16(5.5714283f), MLFloat16(4.0f)}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.SetOutputTolerance(0.005f); + + std::vector> execution_providers; + execution_providers.push_back(std::move(cuda_ep)); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); +} + +// Symmetric pads BUT dilation > 1. cuDNN's pooling descriptor has no dilation parameter, so the +// old (asymmetric-pads-only) guard let this fall through to cuDNN, which silently ignored the +// dilation and produced the wrong result. The dilation guard (!default_dilations) now routes this +// to the custom kernel. opset 19 so the CPU AveragePoolV19 reference (which honors dilation) also +// runs and must match. Expected values come from that CPU reference. +TEST(PoolTest, AveragePool_CUDA_symmetric_pad_dilation_1d) { + OpTester test("AveragePool", 19); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{1}); + test.AddAttribute("pads", std::vector{2, 2}); + test.AddAttribute("kernel_shape", std::vector{3}); + test.AddAttribute("dilations", std::vector{2}); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f}; + std::vector x_dims = {1, 1, 9}; + std::vector expected_dims = {1, 1, 9}; + std::vector expected_vals = {1.3333334f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, + 4.6666665f, 5.3333335f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +// auto_pad=SAME_LOWER produces naturally asymmetric pads with the extra pad on the LOW side +// (here pad(1,0)); companion to the SAME_UPPER case. opset 19 so the CPU reference also runs. +TEST(PoolTest, AveragePool_CUDA_same_lower_asymmetric_1d) { + OpTester test("AveragePool", 19); + + test.AddAttribute("auto_pad", "SAME_LOWER"); + test.AddAttribute("strides", std::vector{2}); + test.AddAttribute("kernel_shape", std::vector{3}); + test.AddAttribute("count_include_pad", (int64_t)1); + + std::vector x_vals = {1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f, 10.0f}; + std::vector x_dims = {1, 1, 10}; + std::vector expected_dims = {1, 1, 5}; + std::vector expected_vals = {1.0f, 3.0f, 5.0f, 7.0f, 9.0f}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", kPoolingEpsExcludedFromCeilCipTests); +} + +// bf16 AveragePool is intentionally NOT tested here: although the CUDA kernel instantiates +// BFloat16 (compile-checked) and AveragePoolWithPad accumulates it in float like fp16, the ONNX +// AveragePool schema type constraint does not include tensor(bfloat16), so OpTester's model +// type-checker rejects such a graph at load. The fp16 case above already exercises the +// accumulate-in-float path. + TEST(PoolTest, GlobalAveragePool) { OpTester test("GlobalAveragePool");