Skip to content

Commit e4dd222

Browse files
committed
Fix int32 overflow in max_pool3d forward kernel
Widen parameters and variables in MaxPool3dKernelFunctor and max_pool3d_with_indices_out_template from int to int64_t to prevent overflow when computing strides for large tensors. This mirrors the backward kernel fix from #3558 The forward kernel was partially protected by int64_t local variables and explicit casts at usage sites, so the existing test shape [70, 32, 100, 100, 100] did not trigger failures. However, the stride values themselves (e.g. in_batch_stride = features * T * H * W) were computed as int, which silently overflows for shapes where a single stride exceeds INT32_MAX (e.g. [2, 2200, 100, 100, 100] gives in_batch_stride = 2.2B). Changes: - Fix typo: MaxPool3dKerenlFunctor -> MaxPool3dKernelFunctor - Widen dimension and stride params/members to int64_t (kernel, stride, padding, and dilation params remain int) - Remove now-redundant (int64_t) casts in operator() - Add test/regressions/test_max_pool3d_fwd_int64.py with shapes that overflow int32 batch strides in the forward pass
1 parent 42e64e2 commit e4dd222

2 files changed

Lines changed: 140 additions & 67 deletions

File tree

src/ATen/native/xpu/sycl/DilatedMaxPool3d.cpp

Lines changed: 69 additions & 67 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ DISABLE_RETURN_TYPE_WARNING_BEGIN
3333
namespace at::native::xpu {
3434

3535
template <typename scalar_t, bool channels_last_>
36-
struct MaxPool3dKerenlFunctor {
36+
struct MaxPool3dKernelFunctor {
3737
void operator()(sycl::nd_item<1> item) const {
3838
auto outputIndex = item.get_global_id(0);
3939
if (outputIndex < OutputSize_) {
@@ -62,9 +62,12 @@ struct MaxPool3dKerenlFunctor {
6262
int tStart = oTime * dT_ - pT_;
6363
int hStart = oRow * dH_ - pH_;
6464
int wStart = oColumn * dW_ - pW_;
65-
int tEnd = std::min(tStart + (kT_ - 1) * dilationT_ + 1, itime_);
66-
int hEnd = std::min(hStart + (kH_ - 1) * dilationH_ + 1, iheight_);
67-
int wEnd = std::min(wStart + (kW_ - 1) * dilationW_ + 1, iwidth_);
65+
int tEnd = std::min(
66+
tStart + (kT_ - 1) * dilationT_ + 1, static_cast<int>(itime_));
67+
int hEnd = std::min(
68+
hStart + (kH_ - 1) * dilationH_ + 1, static_cast<int>(iheight_));
69+
int wEnd = std::min(
70+
wStart + (kW_ - 1) * dilationW_ + 1, static_cast<int>(iwidth_));
6871

6972
while (tStart < 0)
7073
tStart += dilationT_;
@@ -78,9 +81,9 @@ struct MaxPool3dKerenlFunctor {
7881
int64_t ioffset;
7982

8083
if constexpr (!channels_last_) {
81-
ioffset = (int64_t)slice * in_cf_c_stride_;
84+
ioffset = slice * in_cf_c_stride_;
8285
} else {
83-
ioffset = ((int64_t)batch * in_batch_stride_) + channel;
86+
ioffset = batch * in_batch_stride_ + channel;
8487
}
8588

8689
scalar_t max = at::numeric_limits<scalar_t>::lower_bound();
@@ -89,7 +92,7 @@ struct MaxPool3dKerenlFunctor {
8992
for (int h = hStart; h < hEnd; h += dilationH_) {
9093
for (int w = wStart; w < wEnd; w += dilationW_) {
9194
scalar_t val;
92-
int index = t * in_hw_stride_ + h * iwidth_ + w;
95+
int64_t index = t * in_hw_stride_ + h * iwidth_ + w;
9396
if constexpr (!channels_last_) {
9497
val = inputData_[ioffset + index];
9598
} else {
@@ -107,29 +110,28 @@ struct MaxPool3dKerenlFunctor {
107110

108111
int64_t out_index;
109112
if constexpr (!channels_last_) {
110-
out_index = (int64_t)slice * out_cf_c_stride_ +
111-
oTime * out_cf_d_stride_ + oRow * owidth_ + oColumn;
113+
out_index = slice * out_cf_c_stride_ + oTime * out_cf_d_stride_ +
114+
oRow * owidth_ + oColumn;
112115
} else {
113-
out_index = (int64_t)batch * out_batch_stride_ +
114-
oTime * out_cl_d_stride_ + oRow * out_cl_h_stride_ +
115-
oColumn * features_ + channel;
116+
out_index = batch * out_batch_stride_ + oTime * out_cl_d_stride_ +
117+
oRow * out_cl_h_stride_ + oColumn * features_ + channel;
116118
}
117119
outputData_[out_index] = max;
118120
indicesData_[out_index] = maxIndex;
119121
}
120122
}
121-
MaxPool3dKerenlFunctor(
123+
MaxPool3dKernelFunctor(
122124
const scalar_t* inputData,
123125
scalar_t* outputData,
124126
int64_t* indicesData,
125-
int features,
126-
int itime,
127-
int iheight,
128-
int iwidth,
129-
int obatch,
130-
int otime,
131-
int oheight,
132-
int owidth,
127+
int64_t features,
128+
int64_t itime,
129+
int64_t iheight,
130+
int64_t iwidth,
131+
int64_t obatch,
132+
int64_t otime,
133+
int64_t oheight,
134+
int64_t owidth,
133135
int kT,
134136
int kH,
135137
int kW,
@@ -143,17 +145,17 @@ struct MaxPool3dKerenlFunctor {
143145
int dilationH,
144146
int dilationW,
145147
int64_t OutputSize,
146-
int out_cf_d_stride,
147-
int out_cf_c_stride,
148-
int in_cf_d_stride,
149-
int in_cf_c_stride,
150-
int out_cl_h_stride,
151-
int out_cl_d_stride,
152-
int in_cl_h_stride,
153-
int in_cl_d_stride,
154-
int in_batch_stride,
155-
int out_batch_stride,
156-
int in_hw_stride)
148+
int64_t out_cf_d_stride,
149+
int64_t out_cf_c_stride,
150+
int64_t in_cf_d_stride,
151+
int64_t in_cf_c_stride,
152+
int64_t out_cl_h_stride,
153+
int64_t out_cl_d_stride,
154+
int64_t in_cl_h_stride,
155+
int64_t in_cl_d_stride,
156+
int64_t in_batch_stride,
157+
int64_t out_batch_stride,
158+
int64_t in_hw_stride)
157159
: inputData_(inputData),
158160
outputData_(outputData),
159161
indicesData_(indicesData),
@@ -194,14 +196,14 @@ struct MaxPool3dKerenlFunctor {
194196
const scalar_t* inputData_;
195197
scalar_t* outputData_;
196198
int64_t* indicesData_;
197-
int features_;
198-
int itime_;
199-
int iheight_;
200-
int iwidth_;
201-
int obatch_;
202-
int otime_;
203-
int oheight_;
204-
int owidth_;
199+
int64_t features_;
200+
int64_t itime_;
201+
int64_t iheight_;
202+
int64_t iwidth_;
203+
int64_t obatch_;
204+
int64_t otime_;
205+
int64_t oheight_;
206+
int64_t owidth_;
205207
int kT_;
206208
int kH_;
207209
int kW_;
@@ -215,32 +217,32 @@ struct MaxPool3dKerenlFunctor {
215217
int dilationH_;
216218
int dilationW_;
217219
int64_t OutputSize_;
218-
int out_cf_d_stride_;
219-
int out_cf_c_stride_;
220-
int in_cf_d_stride_;
221-
int in_cf_c_stride_;
222-
int out_cl_h_stride_;
223-
int out_cl_d_stride_;
224-
int in_cl_h_stride_;
225-
int in_cl_d_stride_;
226-
int in_batch_stride_;
227-
int out_batch_stride_;
228-
int in_hw_stride_;
220+
int64_t out_cf_d_stride_;
221+
int64_t out_cf_c_stride_;
222+
int64_t in_cf_d_stride_;
223+
int64_t in_cf_c_stride_;
224+
int64_t out_cl_h_stride_;
225+
int64_t out_cl_d_stride_;
226+
int64_t in_cl_h_stride_;
227+
int64_t in_cl_d_stride_;
228+
int64_t in_batch_stride_;
229+
int64_t out_batch_stride_;
230+
int64_t in_hw_stride_;
229231
};
230232

231233
template <typename scalar_t, bool channels_last>
232234
void max_pool3d_with_indices_out_template(
233235
const scalar_t* inputData,
234236
scalar_t* outputData,
235237
int64_t* indicesData,
236-
int features,
237-
int itime,
238-
int iheight,
239-
int iwidth,
240-
int obatch,
241-
int otime,
242-
int oheight,
243-
int owidth,
238+
int64_t features,
239+
int64_t itime,
240+
int64_t iheight,
241+
int64_t iwidth,
242+
int64_t obatch,
243+
int64_t otime,
244+
int64_t oheight,
245+
int64_t owidth,
244246
int kT,
245247
int kH,
246248
int kW,
@@ -255,10 +257,10 @@ void max_pool3d_with_indices_out_template(
255257
int dilationW) {
256258
int64_t OutputSize = obatch * features * otime * oheight * owidth;
257259

258-
int out_cf_d_stride = 0, out_cf_c_stride = 0, in_cf_d_stride = 0,
259-
in_cf_c_stride = 0;
260-
int out_cl_h_stride = 0, out_cl_d_stride = 0, in_cl_h_stride = 0,
261-
in_cl_d_stride = 0;
260+
int64_t out_cf_d_stride = 0, out_cf_c_stride = 0, in_cf_d_stride = 0,
261+
in_cf_c_stride = 0;
262+
int64_t out_cl_h_stride = 0, out_cl_d_stride = 0, in_cl_h_stride = 0,
263+
in_cl_d_stride = 0;
262264
if constexpr (!channels_last) {
263265
out_cf_d_stride = owidth * oheight;
264266
out_cf_c_stride = otime * out_cf_d_stride;
@@ -268,10 +270,10 @@ void max_pool3d_with_indices_out_template(
268270
out_cl_h_stride = owidth * features;
269271
out_cl_d_stride = oheight * out_cl_h_stride;
270272
}
271-
auto in_batch_stride = itime * iheight * iwidth * features;
272-
auto out_batch_stride = otime * oheight * owidth * features;
273-
auto in_hw_stride = iwidth * iheight;
274-
MaxPool3dKerenlFunctor<scalar_t, channels_last> kfn(
273+
int64_t in_batch_stride = itime * iheight * iwidth * features;
274+
int64_t out_batch_stride = otime * oheight * owidth * features;
275+
int64_t in_hw_stride = iwidth * iheight;
276+
MaxPool3dKernelFunctor<scalar_t, channels_last> kfn(
275277
inputData,
276278
outputData,
277279
indicesData,
Lines changed: 71 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,71 @@
1+
# Copyright 2020-2026 Intel Corporation
2+
#
3+
# Licensed under the Apache License, Version 2.0 (the "License");
4+
# you may not use this file except in compliance with the License.
5+
# You may obtain a copy of the License at
6+
#
7+
# http://www.apache.org/licenses/LICENSE-2.0
8+
9+
# Owner(s): ["module: intel"]
10+
"""
11+
Regression test for int32 overflow in max_pool3d forward kernel.
12+
13+
The forward kernel computes strides (e.g. in_batch_stride = C * D * H * W)
14+
that can exceed INT32_MAX for large tensors. This test exercises shapes
15+
where batch strides overflow int32 in both channels_last_3d and contiguous
16+
memory formats.
17+
"""
18+
19+
import torch
20+
from torch.testing._internal.common_device_type import largeTensorTest
21+
from torch.testing._internal.common_utils import (
22+
instantiate_parametrized_tests,
23+
parametrize,
24+
run_tests,
25+
subtest,
26+
TestCase,
27+
)
28+
29+
30+
class TestMaxPool3dFwdInt64(TestCase):
31+
"""Test that the forward pass of max_pool3d handles int64 strides correctly.
32+
33+
Exercises two overflow scenarios:
34+
- channels_last_3d: in_batch_stride = C * D * H * W > INT32_MAX
35+
(shape [2, 2200, 100, 100, 100] -> stride 2.2B).
36+
- contiguous: in_cf_c_stride = D * H * W > INT32_MAX
37+
(shape [1, 2, 1300, 1300, 1300] -> stride 2.197B).
38+
"""
39+
40+
@largeTensorTest("20GB")
41+
@parametrize(
42+
"shape,memory_format",
43+
[
44+
subtest(
45+
((2, 2200, 100, 100, 100), torch.channels_last_3d),
46+
name="channels_last_3d",
47+
),
48+
subtest(
49+
((1, 2, 1300, 1300, 1300), torch.contiguous_format), name="contiguous"
50+
),
51+
],
52+
)
53+
def test_pool3d_fwd_large_size_int64(self, shape, memory_format):
54+
x = torch.randn(*shape, dtype=torch.half, device="xpu").to(
55+
memory_format=memory_format
56+
)
57+
y = torch.nn.functional.max_pool3d(x, 5)
58+
torch.xpu.synchronize()
59+
60+
ref_x = x.detach().cpu().float().contiguous()
61+
ref_y = torch.nn.functional.max_pool3d(ref_x, 5)
62+
63+
self.assertEqual(y.cpu(), ref_y, exact_dtype=False)
64+
del x, y, ref_x, ref_y
65+
66+
67+
instantiate_parametrized_tests(TestMaxPool3dFwdInt64)
68+
69+
70+
if __name__ == "__main__":
71+
run_tests()

0 commit comments

Comments
 (0)