From bd8966b9d011802d7810a6b1967a977158594cb2 Mon Sep 17 00:00:00 2001 From: "akihiro.takagi" Date: Thu, 10 Aug 2023 11:31:27 +0900 Subject: [PATCH 1/3] Avoid address overflow (HostImage) --- test/host_image.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/test/host_image.h b/test/host_image.h index 83a068f..5423a06 100644 --- a/test/host_image.h +++ b/test/host_image.h @@ -94,8 +94,8 @@ class HostImage memcpy(rhs.data, data, elemSize(type) * rows * step); } - template T* ptr(int y = 0) { return (T*)data + y * step; } - template const T* ptr(int y = 0) const { return (T*)data + y * step; } + template T* ptr(int y = 0) { return (T*)data + y * (size_t)step; } + template const T* ptr(int y = 0) const { return (T*)data + y * (size_t)step; } void* data; int rows, cols, step; From 401f27b34f53a4e4300b6fb00b5acc925a990733 Mon Sep 17 00:00:00 2001 From: "akihiro.takagi" Date: Mon, 21 Aug 2023 11:07:14 +0900 Subject: [PATCH 2/3] Avoid address overflow (random_fill) --- test/test_utility.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/test/test_utility.h b/test/test_utility.h index 1b96140..960a970 100644 --- a/test/test_utility.h +++ b/test/test_utility.h @@ -14,25 +14,25 @@ template constexpr T max_of() { return std::numeric_limits::max( template <> constexpr uint8_t max_of() { return 255; } template -static void random_fill_(T* dst, int n, T minv = min_of(), T maxv = max_of()) +static void random_fill_(T* dst, size_t n, T minv = min_of(), T maxv = max_of()) { std::uniform_int_distribution dist(minv, maxv); - for (int i = 0; i < n; ++i) + for (size_t i = 0; i < n; ++i) dst[i] = dist(g_engine); } template <> -void random_fill_(uint8_t* dst, int n, uint8_t minv, uint8_t maxv) +void random_fill_(uint8_t* dst, size_t n, uint8_t minv, uint8_t maxv) { std::uniform_int_distribution dist(minv, maxv); - for (int i = 0; i < n; ++i) + for (size_t i = 0; i < n; ++i) dst[i] = static_cast(dist(g_engine)); } template static void random_fill_(sgm::HostImage& image, T minv = min_of(), T maxv = max_of()) { - random_fill_(image.ptr(), image.rows * image.step, minv, maxv); + random_fill_(image.ptr(), image.rows * (size_t)image.step, minv, maxv); } static void random_fill(sgm::HostImage& image) From 99be2decf9c416422bfa11866facbe6c4ed8d4a0 Mon Sep 17 00:00:00 2001 From: "akihiro.takagi" Date: Mon, 21 Aug 2023 11:26:38 +0900 Subject: [PATCH 3/3] Avoid address overflow (winner_takes_all) --- src/winner_takes_all.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/winner_takes_all.cu b/src/winner_takes_all.cu index 08c2a37..2c0cc9f 100644 --- a/src/winner_takes_all.cu +++ b/src/winner_takes_all.cu @@ -91,7 +91,7 @@ __global__ void winner_takes_all_kernel( ? REDUCTION_PER_THREAD : ACCUMULATION_INTERVAL; - const unsigned int cost_step = MAX_DISPARITY * width * height; + const size_t cost_step = static_cast(MAX_DISPARITY) * width * height; const unsigned int warp_id = threadIdx.x / WARP_SIZE; const unsigned int lane_id = threadIdx.x % WARP_SIZE;