Fixed out-of-bounds memory access in CUDA/OpenCL SIFT

arrayfire · Oct 14, 2015 · 7fc3856 · 7fc3856
1 parent f8462bc
commit 7fc3856
Show file tree

Hide file tree

Showing 2 changed files with 8 additions and 4 deletions.
diff --git 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
@@ -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)])