[arrayfire] 129/248: Fixed out-of-bounds memory access in CUDA/OpenCL SIFT

[Ghislain Vaillant]

This is an automated email from the git hooks/post-receive script.

ghisvail-guest pushed a commit to branch dfsg-clean
in repository arrayfire.

commit 7fc3856721b4852c63e87fea33adf84078fb103b
Author: Peter Andreas Entschev <peter at arrayfire.com>
Date:   Wed Oct 14 18:00:54 2015 -0400

    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 56dc681..e94aeb1 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 f62ff37..dc968d4 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)])

-- 
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/debian-science/packages/arrayfire.git