diff --git a/dnn/src/cuda/warp_perspective/forward.cu b/dnn/src/cuda/warp_perspective/forward.cu index 3cf1333d0be0679cfddfa263ea4e1a4c8af24367..a16c78303f189f2528bdbe252e5c3bd74568e066 100644 --- a/dnn/src/cuda/warp_perspective/forward.cu +++ b/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(batch) * static_cast(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(batch) * static_cast(im_size); } void move_batch(size_t batch, size_t) { diff --git a/dnn/test/cuda/warp_perspective.cpp b/dnn/test/cuda/warp_perspective.cpp index 6997676f0b9359d56f25076c3793b15ace682ec7..6d72943721e3e3ecc5570600fea2c61cd23cf9b7 100644 --- a/dnn/test/cuda/warp_perspective.cpp +++ b/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 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;