Align AvgPool ceil_mode on last value to torch by titaiwangms · Pull Request #16752 · microsoft/onnxruntime

titaiwangms · 2023-07-18T22:13:29Z

Previous to this PR, if ceil_mode is on, the calculation of a value would divide the kernel size, even if remaining pixels is less than the kernel size, which causes the difference in this operator between ORT and torch.

However, this fix only applies to the change in #15597, which only supports AvgPool since 19. The older opset version is remain the same, as it's using mlas files.

Also, the PR fixes the shape mismatch caused by sliding window starting from padding. More detail: onnx/onnx#6650 (And this PR is also validated with the tests added in onnx/onnx#6650)

1D and 3D [Fix](microsoft/onnxruntime#16752) can save the op after opset 19, but not before it. So an xfail is created when ceil_mode=True --------- Co-authored-by: Justin Chu <justinchuby@users.noreply.github.com>

justinchuby · 2023-07-20T15:04:33Z

Was this discrepancy clarified by the ONNX spec? Was it due to an implementation oversight of the spec? Hopefully the spec itself was not different than torch?

… shapes" In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `ceil_mode` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. [ghstack-poisoned]

In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `ceil_mode` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. [ghstack-poisoned]

… shapes" In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `ceil_mode` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. [ghstack-poisoned]

In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `ceil_mode` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. [ghstack-poisoned]

… shapes" In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `ceil_mode` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. [ghstack-poisoned]

In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `ceil_mode` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. [ghstack-poisoned]

In #87892, to pick up the corner cases found in #71549, the PR falls back the implementation of AvgPool to the way opset 9 implementing. However, it introduces a regression on dynamic shape cases found in #101397. This PR refactors the AvgPool op with the same implementation we have in onnxscript: microsoft/onnxscript#754. However, the corner case with `count_include_pad` remains unsolved in onnxruntime: microsoft/onnxruntime#16203. The calculuation on the last value of each dimension is different between ORT and PyTorch. But the fix can be proved in: microsoft/onnxruntime#16752, and it supports AvgPool since opset19. Pull Request resolved: #105683 Approved by: https://github.com/thiagocrepaldi

baijumeswani · 2023-09-26T17:09:53Z

Is this PR still relevant?

titaiwangms · 2023-11-07T20:30:45Z

Hi @baijumeswani Yes. But this PR only works on Pooling op after opset18. Could you help us to add this supports on older versions as well? Basically, ceil_mode goes off whenever there is right padding.

titaiwangms · 2025-01-23T23:00:16Z

@liqunfu @gramalingam @xadupre @justinchuby PTAL

### Description Pooling in reference currently has two bugs: (1) it has ["pads required"](https://github.com/onnx/onnx/blob/e292b4ae6d016c3231a801bfeb26f802ba95d82a/onnx/reference/ops/op_pool_common.py#L53) to make sure sliding window does not go out-of-bound, but it does not exclude pads required from pooling caculations. This causes microsoft/onnxruntime#16203 (2) In `ceil_mode`, the reference pooling pads the input image, but does not check if sliding window starts on pads. When we reach the end of the image, the window should stop working. This causes pytorch/pytorch#131272. Not directly, but related fix on MaxPooling: #5741. Detail: pytorch/pytorch#116420 (comment) ### Motivation and Context This PR fixes the two bugs and update their tests accordingly. I also drafted a PR to onnxruntime to have this update in CPU provider: microsoft/onnxruntime#16752 --------- Signed-off-by: titaiwangms <titaiwang@microsoft.com> Signed-off-by: Justin Chu <justinchuby@users.noreply.github.com> Signed-off-by: Andreas Fehlner <fehlner@arcor.de> Co-authored-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Andreas Fehlner <fehlner@arcor.de>

### Description Pooling in reference currently has two bugs: (1) it has ["pads required"](https://github.com/onnx/onnx/blob/e292b4ae6d016c3231a801bfeb26f802ba95d82a/onnx/reference/ops/op_pool_common.py#L53) to make sure sliding window does not go out-of-bound, but it does not exclude pads required from pooling caculations. This causes microsoft/onnxruntime#16203 (2) In `ceil_mode`, the reference pooling pads the input image, but does not check if sliding window starts on pads. When we reach the end of the image, the window should stop working. This causes pytorch/pytorch#131272. Not directly, but related fix on MaxPooling: #5741. Detail: pytorch/pytorch#116420 (comment) ### Motivation and Context This PR fixes the two bugs and update their tests accordingly. I also drafted a PR to onnxruntime to have this update in CPU provider: microsoft/onnxruntime#16752 --------- Signed-off-by: titaiwangms <titaiwang@microsoft.com> Signed-off-by: Justin Chu <justinchuby@users.noreply.github.com> Signed-off-by: Andreas Fehlner <fehlner@arcor.de> Co-authored-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Andreas Fehlner <fehlner@arcor.de> Signed-off-by: Andreas Fehlner <fehlner@arcor.de>

Fix #16203 Previous to this PR, if `ceil_mode` is on, the calculation of a value would divide the kernel size, even if remaining pixels is less than the kernel size, which causes the difference in this operator between ORT and torch. However, this fix only applies to the change in #15597, which only supports AvgPool since 19. The older opset version is remain the same, as it's using mlas files. Also, the PR fixes the shape mismatch caused by sliding window starting from padding. More detail: onnx/onnx#6650 (And this PR is also validated with the tests added in onnx/onnx#6650)

### Description This PR is to update the win-ort-main branch to the tip main branch as of 2025-02-11. ### PR List 74c778e [WebNN EP] Automatically move input CPU tensors to ml-tensor (#23073) 3775057 use correct total length to fix static kv_cache performance (#23615) 3901e96 remove --use_vcpkg flag for Python-CUDA-Packaging-Pipeline (#23631) c610df5 Add python_requires to package metadata (#23604) 2d27d68 [QNN EP] Add QNN EP to ARM64X build targets (#23635) e666503 [webgpu] no longer need pass-in gpu adapter for custom context (#23593) af679a0 Fix logic for selecting alternate name for blob (#23617) e206950 [ARM CPU] Add fp16 mlas kernels for exp, tanh, softmax, logsoftmax, softcap (#23597) 9ba5619 Update pybind and json to the latest (#23589) c54736c Migrate iOS release pipeline to 1 ES (#23606) 3981326 Increase timeout for Windows TensorRT CI (#23625) 0274b7b fix on trtCudaVersion (#23616) 740e9ab update run CI script (#23621) 5ef1832 [WebGPU] Support PIX Capture for WebGPU EP (#23192) 0114551 Fix for C4267 warning (#23610) 002916a Validate the context_file_path before EP compile graphs (#23611) 0887e36 [webgpu] Use pushErrorScope()/popErrorScope() once for an inference run (#23438) 65008cb Auto-generated baselines by 1ES Pipeline Templates (#23603) 09e5724 [CUDA] Fix beam search of num_beams > 32 (#23599) 82840f6 Implement Flash Attention 2 for webgpu EP (#23576) a6ea57b OpenVINO EP Weights Sharing Feature (#23553) 2c2ff4a [CUDA] Fix BeamSearchTest.DummyT5WithSequenceInputIds test failure in Windows (#23596) d981b15 [webgpu/js] Optimize resize webgpu op & fix precision issues (#23591) 328a13c Enable VCPKG in more pipelines (#23590) 6728d60 [TensorRT EP] support TensorRT 10.8-GA (#23592) d1fb58b Quantization tool: Allow user to override calibrator's session EP (#23559) 649ced4 Enable user loading model with external data from memory buffer (#23557) 544bdd6 Fix ConvTranspose for certain attribute combinations (#23488) 8f6ddf3 Delete extra cgmanifest entries and files (#23583) 5f6a315 Enable VCPKG in CI build (#23426) e1e3f62 Bump lintrunner from 0.12.5 to 0.12.7 (#23326) cd8775f Fix Node JS Samples (#23581) 6b4f9c4 [WebGPU EP] Batch Norm Implementation (#23525) 1fce51b Fix all instances of 4244 and 4267 warnings in OV EP code (#23567) c29ca1c Update QNN default version to 2.31 (#23573) 2fc75a4 [mobile] Add Android BrowserStack test project back (#23551) 9e18b6a [CUDA] Update nvcc flags (#23572) b47e1e6 [QNN EP] Make offloading graph input/output quantization (to CPU) the default (#23368) 75a9b40 [ROCm] Update CI to use rocm 6.3.2 (#23577) 26ff2b6 Bump ruff from 0.9.3 to 0.9.4 (#23563) b2560a7 Update react-native to 0.72 (#23509) faee912 [js] update JavaScript API to support QNN EP options (#23486) 816e8cb [EP Perf] Update env to ubuntu 22.04 (#23570) cddc271 Use Eigen in Round implementation (#23571) e8b0bdb Shape inference: ReduceMean dispatcher, quant_pre_process: skip_symbolic_shape bugfix (#23558) 267b493 delete the supported domain version upper bounds (#23237) bb7f961 remove log spam from cpuinfo (#23548) 169917b Use latest vcpkg commit in configuration, sync manifest with deps.txt (#23554) a9d4d08 Add of ReduceMax Gradient (#23501) 6bbf1bd [js/web] upgrade version of flatbuffers (#23545) 271c509 DP4AMatMul perf refinements (#23539) cb69c59 Add fusions for SigLIP and Conformer-Encoder (#23528) 61fae9b Remove "--enable_pybind" from webgpu pipeline (#23550) 0bb4ea6 Update BiasGelu fusion and related ops (#23518) 4dde74a Add more details to BrowserStack script failure (#23520) ead9d5c Set ANDROID_USE_LEGACY_TOOLCHAIN_FILE to false (#23544) 7e24088 Enable dlpack by default (#23110) dc2f7a9 Add overload of `TryParseStringWithClassicLocale()` that uses `std::from_chars()` (#23541) 5407c69 Fix the issue that the new generated EP context model not able to find external data (#23537) fbae88f [js/web] use the recommended workaround for Vite (#23531) d5338da Fix tensor external data info length parsing issue. (#23526) e3e4173 [ROCm EP] Fix transpose helper for gfx gridsize constraints (#23527) 80bc1d2 Enable Ep context with external data for CPU nodes (#23498) bf023ab [js/web] allow import .mjs/.wasm file (#23487) 655a23f [onnxruntime/build] Add new flag enable_generic_interface to build primary EPs by default (#23342) a770a8d Update RN to 0.71.19 (#23381) 1cf0ebd Delete Prefast workflow until the build failure is fixed (#23510) d2c5e24 Add of GlobalMaxPool Gradient (#23502) ded8730 Remove thrust::unary_function (#23506) 8db97a6 [webgpu] Bump version of Dawn to b9b4a370 (#23494) fdde2e2 Fix for gcc 13.3.1: Avoid creating a copy (#23500) 96ec1dd Bump ruff from 0.9.2 to 0.9.3 (#23496) 42f0c00 Adds the new System.Numerics.Tensors as an input/output type when using dotnet 8.0 and up. (#23261) 97c2bbe Fix shape infer of onnx GroupNorm (#23477) 1fc9c48 Enable coremltools for Linux build (#23481) 13348c5 [ARM CPU] hgemm optimized for gqa (#23107) c89a798 Enable opti on Microsoft.ML.OnnxRuntime with RelWithDebInfo config (#23463) d00ae32 Revert "[Mobile] Add BrowserStack Android MAUI Test (#23383)" (#23474) 8b1d3b3 Align AvgPool ceil_mode on last value to torch (#16752) 06fc73b [TRT EP Perf Tool] Add annotations import to python script to support annotations on Python 3.8 (#23466) ### Motivation and Context This update includes the change to add QNN EP to ARM64X build targets. --------- Signed-off-by: dependabot[bot] <support@github.com> Co-authored-by: Adrian Lizarraga <adlizarraga@microsoft.com> Co-authored-by: Ti-Tai Wang <titaiwang@microsoft.com> Co-authored-by: Caroline Zhu <wolfivyaura@gmail.com> Co-authored-by: Grégoire <gregoire.verdier@gmail.com> Co-authored-by: Jing Fang <126209182+fajin-corp@users.noreply.github.com> Co-authored-by: Changming Sun <chasun@microsoft.com> Co-authored-by: Yateng Hong <yatengh@microsoft.com> Co-authored-by: Michael Sharp <51342856+michaelgsharp@users.noreply.github.com> Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com> Co-authored-by: Malik Shahzad Muzaffar <shahzad.malik.muzaffar@cern.ch> Co-authored-by: Yulong Wang <7679871+fs-eire@users.noreply.github.com> Co-authored-by: Dmitri Smirnov <yuslepukhin@users.noreply.github.com> Co-authored-by: Corentin Maravat <101636442+cocotdf@users.noreply.github.com> Co-authored-by: Jian Chen <cjian@microsoft.com> Co-authored-by: Karim Vadsariya <karim.vadsariya@microsoft.com> Co-authored-by: Lei Cao <jslhcl@gmail.com> Co-authored-by: Karim Vadsariya <kvadsariya@microsoft.com> Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com> Co-authored-by: Hector Li <hecli@microsoft.com> Co-authored-by: Ted Themistokleous <107195283+TedThemistokleous@users.noreply.github.com> Co-authored-by: Ted Themistokleous <tedthemistokleous@amd.com> Co-authored-by: Edward Chen <18449977+edgchen1@users.noreply.github.com> Co-authored-by: Takeshi Watanabe <take-cheeze@users.noreply.github.com> Co-authored-by: Xavier Dupré <xadupre@users.noreply.github.com> Co-authored-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Tianlei Wu <tlwu@microsoft.com> Co-authored-by: kunal-vaishnavi <115581922+kunal-vaishnavi@users.noreply.github.com> Co-authored-by: Sushanth Rajasankar <44513542+sushraja-msft@users.noreply.github.com> Co-authored-by: PARK DongHa <luncliff@gmail.com> Co-authored-by: George Wu <jywu@microsoft.com> Co-authored-by: Xinpeng Dou <15529241576@163.com> Co-authored-by: Jambay Kinley <jambaykinley@microsoft.com> Co-authored-by: Yifan Li <109183385+yf711@users.noreply.github.com> Co-authored-by: Gavin Kinsey <98115505+ms-gavinkinsey@users.noreply.github.com> Co-authored-by: Prathik Rao <prathik.rao@gmail.com> Co-authored-by: Jon Campbell <jcampbell@cephable.com> Co-authored-by: Satya Kumar Jandhyala <satya.k.jandhyala@gmail.com> Co-authored-by: Joshua Lochner <admin@xenova.com> Co-authored-by: Ankit Maheshkar <ankit.maheshkar@intel.com> Co-authored-by: jatinwadhwa921 <jatin.wadhwa@intel.com> Co-authored-by: jatinwadhwa921 <110383850+jatinwadhwa921@users.noreply.github.com> Co-authored-by: saurabh <saurabh1.kale@intel.com> Co-authored-by: TejalKhade28 <tejal.khade@intel.com> Co-authored-by: sfatimar <sahar.fatima@intel.com> Co-authored-by: Javier E. Martinez <javier.e.martinez@intel.com> Co-authored-by: Preetha Veeramalai <preetha.veeramalai@intel.com> Co-authored-by: Eric Crawford <eric.r.crawford@intel.com> Co-authored-by: microsoft-github-policy-service[bot] <77245923+microsoft-github-policy-service[bot]@users.noreply.github.com> Co-authored-by: Jie Chen <jie.a.chen@intel.com> Co-authored-by: shaoboyan091 <shaoboyan@microsoft.com> Co-authored-by: David Hotham <david.hotham@microsoft.com> Co-authored-by: Guenther Schmuelling <guschmue@microsoft.com> Co-authored-by: Enrico Galli <enrico.galli@intel.com>

### Description Pooling in reference currently has two bugs: (1) it has ["pads required"](https://github.com/onnx/onnx/blob/e292b4ae6d016c3231a801bfeb26f802ba95d82a/onnx/reference/ops/op_pool_common.py#L53) to make sure sliding window does not go out-of-bound, but it does not exclude pads required from pooling caculations. This causes microsoft/onnxruntime#16203 (2) In `ceil_mode`, the reference pooling pads the input image, but does not check if sliding window starts on pads. When we reach the end of the image, the window should stop working. This causes pytorch/pytorch#131272. Not directly, but related fix on MaxPooling: onnx#5741. Detail: pytorch/pytorch#116420 (comment) ### Motivation and Context This PR fixes the two bugs and update their tests accordingly. I also drafted a PR to onnxruntime to have this update in CPU provider: microsoft/onnxruntime#16752 --------- Signed-off-by: titaiwangms <titaiwang@microsoft.com> Signed-off-by: Justin Chu <justinchuby@users.noreply.github.com> Signed-off-by: Andreas Fehlner <fehlner@arcor.de> Co-authored-by: Justin Chu <justinchuby@users.noreply.github.com> Co-authored-by: Andreas Fehlner <fehlner@arcor.de> Signed-off-by: seungwoo-ji <seungwoo.ji@nuvilab.com>

Fix #16203 Previous to this PR, if `ceil_mode` is on, the calculation of a value would divide the kernel size, even if remaining pixels is less than the kernel size, which causes the difference in this operator between ORT and torch. However, this fix only applies to the change in #15597, which only supports AvgPool since 19. The older opset version is remain the same, as it's using mlas files. Also, the PR fixes the shape mismatch caused by sliding window starting from padding. More detail: onnx/onnx#6650 (And this PR is also validated with the tests added in onnx/onnx#6650)

### Summary The ORT-Web JSEP pooling output-shape helper (`PoolConvUtil.computePoolOutputShape` → `computeShapeHelper` → `adjustPadAndReturnShape`) was **floor-only**: it had no `ceilMode` parameter and silently ignored `ceil_mode`. Under `ceil_mode=1` this allocated the output tensor one element too small along each affected spatial axis, producing a wrong output shape (and downstream shape mismatches) — independent of any pooling divisor concern. The same floor-only helper is duplicated in the legacy `js/web/lib/onnxjs/util.ts`; both copies are fixed. ### Fix - Thread a `ceilMode` parameter (default `0`, i.e. floor) through `computePoolOutputShape` → `computeShapeHelper` → `adjustPadAndReturnShape` in **both** `js/web/lib/wasm/jsep/util.ts` and `js/web/lib/onnxjs/util.ts`. - Route the NOTSET / VALID / SAME_UPPER / SAME_LOWER branches through a new `computeOutputSize()` helper that produces results identical to the C++ reference `PoolAttributes::ComputeOutputSize` (`onnxruntime/core/providers/cpu/nn/pool_attributes.h`), including the `ceil_mode` **"shrink the last window if it starts entirely in the trailing padding"** rule (ref: onnx/onnx#5741). - The floor path (`ceilMode` default) is algebraically unchanged, so Conv and `auto_pad` pad-adjustment are unaffected. - `js/web/lib/wasm/jsep/webgpu/ops/pool.ts` now passes `attributes.ceilMode` into the shape computation. ### Scope: SHAPE-only This PR fixes the **output-shape** computation only. End-to-end `ceil_mode` execution remains gated by the existing `throw` guards in `parseAveragePoolAttributes` / `parseMaxPoolAttributes`, because the WebGPU pooling kernel does not yet implement `ceil_mode` trailing-padding handling (and, for AveragePool, the `count_include_pad` divisor). Removing those throws + adding kernel support is a **tracked follow-up**; the now-correct shape path is exercised directly by the added unit tests until then. Comments at the throw sites and in the test header document this. ### Tests Adds `js/web/test/unittests/pool-output-shape.ts` (registered in `unittests/index.ts`), asserting output shapes for **both** implementations (jsep + onnxjs) against the C++ CPU reference ground truth: - `test_maxpool_2d_ceil` → `[1,1,2,2]`, `AveragePool_10_ceil1_2d` → `[1,1,2,3]` - AvgPool `ceil` 1D `[1,2,4]` / 2D `[1,1,3,3]` / 3D `[1,1,2,2,2]` (matching the CPU `pool_op_test.cc` cases) - `ceil_mode` with dilation, `VALID` / `SAME_UPPER` / `SAME_LOWER` auto_pad - a shrink-rule discriminator (naive `ceil()` would give 3; correct = 2) - floor-mode no-regression case ### Related - PyTorch context: pytorch/pytorch#183528 - Related ORT PR: #16752 ### Post-review fixups - **SAME_UPPER/SAME_LOWER integer division (external-review Major).** The `legacyTargetSize = (inSize + stride - 1) / stride` computation now uses `Math.floor(...)` to match C++ `pool_attributes.h` `ComputeSizePadDilations`, which uses integer division. This fixes a latent float-division divergence that mis-rounded the auto_pad pad distribution. The correction applies to **all** ceil modes (not only `ceil_mode=1`), aligning JSEP/onnxjs with the CPU/CUDA EPs. The floor-path (`ceil_mode=0`) **output shape is unchanged** — only the SAME_* pad split is corrected — so nothing previously-correct regresses. Added `ceil_mode=0` SAME_UPPER/SAME_LOWER non-divisible regression tests that assert the corrected pad out-param (SAME_UPPER → `[0,1]`, SAME_LOWER → `[1,0]`). - **Throw-message / TODO clarity.** The retained `ceil_mode` `throw` statements and the pool op TODO banner were reworded to make explicit that the output **shape** is now computed correctly, while `ceil_mode` **kernel execution** (padding/divisor) remains the pending WebGPU follow-up. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

…18 / MLAS path) (#29629) ### Summary Fixes wrong `AveragePool` output on the CPU EP when `ceil_mode=1` **and** `count_include_pad=1` for **opset 7–18** (the float MLAS path). This is the direct fix for the CPU-EP repro in pytorch/pytorch#183528. ### Root cause `Pool<float, AveragePool>::Compute` routes float `AveragePool` opset 7–18 to MLAS. With `count_include_pad=1`, MLAS divides every window by the **full kernel size**. Under `ceil_mode=1` the output grid gains a trailing window whose extent runs past `input + pad_tail` (phantom cells). MLAS's full-kernel divisor counts those phantom cells, producing an average that is too small on the boundary windows. The opset-19 reference functor (`AveragePool{1,2,3}DTask` in `pool_functors.h`) already handles this correctly: it clamps the window end to `input + pad_tail` and divides by `1 + (end - start - 1)/dilation`, excluding the phantom cells. ### Fix Dispatch-fallback, **zero MLAS edits**: 1. Extract the opset-19 reference average-pool loop into a shared free function `ComputeAveragePoolReference<T>(context, pool_attrs, tp)` (reads strides from `PoolAttributes`, `p=0`), and make `AveragePoolV19<T>::Compute` a thin wrapper over it — one canonical loop, no copy-paste drift. 2. In `Pool<float, AveragePool>::Compute`, route **only** the buggy combo `ceil_mode == 1 && count_include_pad && !global_pooling` to the reference loop. Every other case (`ceil_mode=0`, `count_include_pad=0`/exclude-pad, global pooling, all MaxPool) keeps the fast MLAS path unchanged. `global_pooling` is excluded because it has no ceil/pads (already correct) and leaves `strides[]` unpopulated; `ComputeAveragePoolReference` `ORT_ENFORCE`s that precondition. ### Perf Zero impact on the common path. The reference loop runs only for `ceil_mode=1 && count_include_pad=1` (rare); the guard is a cheap early branch. ### Relationship to #16752 ORT PR #16752 fixed this behavior for **opset ≥ 19** only (via the new v19 reference functor). Opset 7–18 float still went through MLAS and remained wrong. This PR closes that gap by reusing the same already-correct functor for the 7–18 dispatch fallback. ### Tests Added to `pool_op_test.cc` (CPU-focused; GPU/other EPs excluded so CI stays green): - `AveragePool_18_ceil_count_include_pad_1d` — opset-18 clone of the existing v19 test; the direct #183528 repro. - `AveragePool_18_ceil_count_include_pad_2d` — the `arange(1,17).reshape(1,1,4,4)`, k3/s2/pad1 case → `[1.556, 3.333, 2.0, 6.333, 11.0, 6.0, 4.5, 7.5, 4.0]`. - `AveragePool_18_ceil_count_include_pad_3d` — exercises the 3D functor path. - `AveragePool_18_ceil_count_exclude_pad_2d` — `count_include_pad=0` no-regression guard (stays on MLAS). Expected values are derived from the CPU v19 reference (equivalently, hand-computed per the ONNX spec). Full `PoolTest` suite passes locally. --------- Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>

align ceil_mode on last value to torch

d2880f3

titaiwangms added converter related to ONNX converters core runtime issues related to core runtime labels Jul 18, 2023

titaiwangms requested review from gramalingam, justinchuby, liqunfu and xadupre July 18, 2023 22:13

justinchuby reviewed Jul 18, 2023

View reviewed changes

Comment thread onnxruntime/test/providers/cpu/nn/pool_op_test.cc

justinchuby reviewed Jul 18, 2023

View reviewed changes

Comment thread onnxruntime/core/providers/cpu/nn/pool_functors.h

justinchuby previously approved these changes Jul 18, 2023

View reviewed changes

titaiwangms mentioned this pull request Jul 18, 2023

Add op (AvgPool) | feat (torchlib) microsoft/onnxscript#754

Merged

titaiwangms changed the title ~~align ceil_mode on last value to torch~~ Align ceil_mode on last value to torch Jul 18, 2023

titaiwangms changed the title ~~Align ceil_mode on last value to torch~~ Align AvgPool ceil_mode on last value to torch Jul 18, 2023

Merge branch 'main' into titaiwang/fix_avgpool_ceil_mode_and_pads

da2feec

liqunfu reviewed Jul 19, 2023

View reviewed changes

Comment thread onnxruntime/test/providers/cpu/nn/pool_op_test.cc Outdated

add f

f650a32

titaiwangms dismissed justinchuby’s stale review via f650a32 July 19, 2023 00:05

xadupre reviewed Jul 19, 2023

View reviewed changes

Comment thread onnxruntime/core/providers/cpu/nn/pool_functors.h

add the change to 2d and 3d

369c635

titaiwangms requested a review from justinchuby July 20, 2023 14:59

This was referenced Jul 20, 2023

[ONNX] Support of AvgPool2D when ceil_mode is True has disappeared with Torch 2.0 pytorch/pytorch#101397

Closed

[ONNX] Refactor AvgPool to support dynamic shapes pytorch/pytorch#105683

Closed

justinchuby previously approved these changes Aug 8, 2023

View reviewed changes

regisss mentioned this pull request Aug 12, 2023

adding onnx support for Funnel Transformer huggingface/optimum#1231

Closed

3 tasks

titaiwangms mentioned this pull request Aug 28, 2023

Ambiguity in Pool spec (AveragePool and MaxPool) onnx/onnx#1922

Open

titaiwangms added 3 commits January 17, 2025 17:05

Merge branch 'main' into titaiwang/fix_avgpool_ceil_mode_and_pads

d9664d9

Merge branch 'main' into titaiwang/fix_avgpool_ceil_mode_and_pads

1ec9cad

update ceil_mode

c84a858

titaiwangms dismissed justinchuby’s stale review via c84a858 January 23, 2025 20:24

titaiwangms mentioned this pull request Jan 23, 2025

Fix pooling pads issues onnx/onnx#6650

Merged

titaiwangms removed the converter related to ONNX converters label Jan 23, 2025

move the broken test to general category

488321a

justinchuby approved these changes Jan 23, 2025

View reviewed changes

titaiwangms merged commit 8b1d3b3 into microsoft:main Jan 24, 2025

ashrit-ms mentioned this pull request Feb 11, 2025

Update win-ort-main to tip main 250211 #23646

Merged

tengli-alaska mentioned this pull request Feb 12, 2026

Fix for AveragePool performance regression by guarding ceil_mode clamp #27331

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Align AvgPool ceil_mode on last value to torch - #16752

Align AvgPool ceil_mode on last value to torch#16752
titaiwangms merged 8 commits into
microsoft:mainfrom
titaiwangms:titaiwang/fix_avgpool_ceil_mode_and_pads

titaiwangms commented Jul 18, 2023 •

edited

Loading

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

justinchuby commented Jul 20, 2023

Uh oh!

baijumeswani commented Sep 26, 2023

Uh oh!

titaiwangms commented Nov 7, 2023 •

edited

Loading

Uh oh!

titaiwangms commented Jan 23, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Uh oh!

Conversation

titaiwangms commented Jul 18, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

justinchuby commented Jul 20, 2023

Uh oh!

baijumeswani commented Sep 26, 2023

Uh oh!

titaiwangms commented Nov 7, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

titaiwangms commented Jan 23, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

titaiwangms commented Jul 18, 2023 •

edited

Loading

titaiwangms commented Nov 7, 2023 •

edited

Loading