From 7fc3856721b4852c63e87fea33adf84078fb103b Mon Sep 17 00:00:00 2001 From: Peter Andreas Entschev Date: Wed, 14 Oct 2015 18:00:54 -0400 Subject: [PATCH] Fixed out-of-bounds memory access in CUDA/OpenCL SIFT --- src/backend/cuda/kernel/sift_nonfree.hpp | 6 ++++-- src/backend/opencl/kernel/sift_nonfree.cl | 6 ++++-- 2 files changed, 8 insertions(+), 4 deletions(-) diff --git a/src/backend/cuda/kernel/sift_nonfree.hpp b/src/backend/cuda/kernel/sift_nonfree.hpp index 56dc6817b0..e94aeb1377 100644 --- a/src/backend/cuda/kernel/sift_nonfree.hpp +++ b/src/backend/cuda/kernel/sift_nonfree.hpp @@ -333,8 +333,10 @@ __global__ void sub( { unsigned i = blockIdx.x * blockDim.x + threadIdx.x; - for (unsigned l = 0; l < n_layers; l++) - out.ptr[l*nel + i] = in.ptr[(l+1)*nel + i] - in.ptr[l*nel + i]; + if (i < nel) { + for (unsigned l = 0; l < n_layers; l++) + out.ptr[l*nel + i] = in.ptr[(l+1)*nel + i] - in.ptr[l*nel + i]; + } } #define SCPTR(Y, X) (s_center[(Y) * s_i + (X)]) diff --git a/src/backend/opencl/kernel/sift_nonfree.cl b/src/backend/opencl/kernel/sift_nonfree.cl index f62ff37612..dc968d4f4d 100644 --- a/src/backend/opencl/kernel/sift_nonfree.cl +++ b/src/backend/opencl/kernel/sift_nonfree.cl @@ -255,8 +255,10 @@ __kernel void sub( { unsigned i = get_global_id(0); - for (unsigned l = 0; l < n_layers; l++) - out[l*nel + i] = in[l*nel + i] - in[(l+1)*nel + i]; + if (i < nel) { + for (unsigned l = 0; l < n_layers; l++) + out[l*nel + i] = in[l*nel + i] - in[(l+1)*nel + i]; + } } #define LCPTR(Y, X) (l_center[(Y) * l_i + (X)])