From 13ba9328d2e38af8f45c2626bf6d94d65117a339 Mon Sep 17 00:00:00 2001 From: YOUNGJOON KIM Date: Sat, 21 Feb 2026 19:17:26 +0900 Subject: [PATCH 1/3] ROI pool CUDA: compare in acc_type to fix half/MSVC build --- torchvision/csrc/ops/cuda/roi_pool_kernel.cu | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/torchvision/csrc/ops/cuda/roi_pool_kernel.cu b/torchvision/csrc/ops/cuda/roi_pool_kernel.cu index 3a9374bb438..ccb8888cacb 100644 --- a/torchvision/csrc/ops/cuda/roi_pool_kernel.cu +++ b/torchvision/csrc/ops/cuda/roi_pool_kernel.cu @@ -63,10 +63,14 @@ __global__ void roi_pool_forward_kernel_impl( int maxidx = -1; const T* offset_input = input + (roi_batch_ind * channels + c) * height * width; + using acc_t = at::acc_type; for (int h = hstart; h < hend; ++h) { for (int w = wstart; w < wend; ++w) { int input_index = h * width + w; - if (offset_input[input_index] > maxval) { + acc_t v = static_cast(offset_input[input_index]); + acc_t mv = static_cast(maxval); + + if (v > mv) { maxval = offset_input[input_index]; maxidx = input_index; } From c03784b5f66ca6584d71a7eba779dcc110da740c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Patrick=20Kim=20=28=EA=B9=80=EC=98=81=EC=A4=80=29?= <43011129+kimchioverfit@users.noreply.github.com> Date: Mon, 23 Feb 2026 18:11:28 +0900 Subject: [PATCH 2/3] Add AccumulateType header to roi_pool_kernel.cu Add --- torchvision/csrc/ops/cuda/roi_pool_kernel.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/torchvision/csrc/ops/cuda/roi_pool_kernel.cu b/torchvision/csrc/ops/cuda/roi_pool_kernel.cu index ccb8888cacb..4b78dc2e705 100644 --- a/torchvision/csrc/ops/cuda/roi_pool_kernel.cu +++ b/torchvision/csrc/ops/cuda/roi_pool_kernel.cu @@ -4,7 +4,7 @@ #include #include #include - +#include #include "cuda_helpers.h" namespace vision { From 788bf6a94235d599ccd01d0ad8846e72df3a2a03 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Patrick=20Kim=20=28=EA=B9=80=EC=98=81=EC=A4=80=29?= <43011129+kimchioverfit@users.noreply.github.com> Date: Tue, 24 Feb 2026 15:14:38 +0900 Subject: [PATCH 3/3] Apply clang-format to roi_pool_kernel.cu No functional changes. Only formatting adjustments required by clang-format (include order and spacing). --- torchvision/csrc/ops/cuda/roi_pool_kernel.cu | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/torchvision/csrc/ops/cuda/roi_pool_kernel.cu b/torchvision/csrc/ops/cuda/roi_pool_kernel.cu index 4b78dc2e705..894531901ea 100644 --- a/torchvision/csrc/ops/cuda/roi_pool_kernel.cu +++ b/torchvision/csrc/ops/cuda/roi_pool_kernel.cu @@ -1,10 +1,11 @@ #include +#include #include #include #include #include #include -#include + #include "cuda_helpers.h" namespace vision { @@ -67,7 +68,7 @@ __global__ void roi_pool_forward_kernel_impl( for (int h = hstart; h < hend; ++h) { for (int w = wstart; w < wend; ++w) { int input_index = h * width + w; - acc_t v = static_cast(offset_input[input_index]); + acc_t v = static_cast(offset_input[input_index]); acc_t mv = static_cast(maxval); if (v > mv) {