arrayfire · GitHub

Commit 7fc3856

author

Peter Andreas Entschev

committed

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

2 files changed

Lines changed: 8 additions & 4 deletions

File tree

  • src/backend

    • cuda/kernel

    • opencl/kernel

Lines changed: 4 additions & 2 deletions

Original file line numberDiff line numberDiff line change

@@ -333,8 +333,10 @@ __global__ void sub(

333333

{

334334

unsigned i = blockIdx.x * blockDim.x + threadIdx.x;

335335
336-

for (unsigned l = 0; l < n_layers; l++)

337-

out.ptr[l*nel + i] = in.ptr[(l+1)*nel + i] - in.ptr[l*nel + i];

336+

if (i < nel) {

337+

for (unsigned l = 0; l < n_layers; l++)

338+

out.ptr[l*nel + i] = in.ptr[(l+1)*nel + i] - in.ptr[l*nel + i];

339+

}

338340

}

339341
340342

#define SCPTR(Y, X) (s_center[(Y) * s_i + (X)])

Lines changed: 4 additions & 2 deletions

Original file line numberDiff line numberDiff line change

@@ -255,8 +255,10 @@ __kernel void sub(

255255

{

256256

unsigned i = get_global_id(0);

257257
258-

for (unsigned l = 0; l < n_layers; l++)

259-

out[l*nel + i] = in[l*nel + i] - in[(l+1)*nel + i];

258+

if (i < nel) {

259+

for (unsigned l = 0; l < n_layers; l++)

260+

out[l*nel + i] = in[l*nel + i] - in[(l+1)*nel + i];

261+

}

260262

}

261263
262264

#define LCPTR(Y, X) (l_center[(Y) * l_i + (X)])

Read the original on github.com ↗