fix(dnn/cuda): fix INTMAX overflow in warp_perspective_cuda

GitOrigin-RevId: d7354e74e269c7953bc954fe628f06e57ab11899

dnn/src/cuda/warp_perspective/forward.cu

@@ -28,7 +28,7 @@ struct DirectSrcVisitor {
     const ctype* ptr;
 
     __device__ __forceinline__ const ctype* get(int batch, int im_size) {
-        return ptr + batch * im_size;
+        return ptr + static_cast<int64_t>(batch) * static_cast<int64_t>(im_size);
     }
 
     void move_batch(size_t batch, size_t im_size) {
@@ -54,7 +54,7 @@ struct IndexedSrcVisitor {
                     orig_batch, batch, N_SRC);
             batch = 0;
         }
-        return ptr + batch * im_size;
+        return ptr + static_cast<int64_t>(batch) * static_cast<int64_t>(im_size);
     }
 
     void move_batch(size_t batch, size_t) {

dnn/test/cuda/warp_perspective.cpp

@@ -14,6 +14,7 @@
 #include "test/common/benchmarker.h"
 #include "test/common/warp_perspective.h"
 #include "test/common/opr_proxy.h"
+#include "test/cuda/utils.h"
 
 namespace {
 
@@ -217,6 +218,30 @@ TEST_F(CUDA, WARP_PERSPECTIVE_FORWARD)
     }
 }
 
+TEST_F(CUDA, WARP_PERSPECTIVE_FORWARD_INTMAX)
+{
+    require_compute_capability(6, 0);
+    using Param = WarpPerspective::Param;
+    Checker<WarpPerspectiveForward> checker(handle_cuda());
+    WarpPerspectiveMatRNG rng;
+    checker.set_rng(1, &rng);
+    for (auto bmode: {WarpPerspective::BorderMode::REPLICATE})
+    {
+        WarpPerspective::Param param;
+        param.border_val = 0.3f;
+        param.bmode = bmode;
+        param.imode = Param::InterpolationMode::LINEAR;
+
+        param.format = Param::Format::NHWC;
+        checker.set_param(param);
+        checker.set_epsilon(0.15).set_max_avg_error(4e-2);
+		size_t n = (INT_MAX) / (512 * 512 * 3);
+		checker.execs(
+				{{n + 1, 512, 512, 3}, {n + 1, 3, 3}, {n + 1, 25, 25, 3}});
+    }
+}
+
+
 TEST_F(CUDA, WARP_PERSPECTIVE_FORWARD_FP16)
 {
     using Param = WarpPerspective::Param;