diff options
| author | Cedric Nugteren <web@cedricnugteren.nl> | 2019-09-06 19:30:52 +0200 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2019-09-06 19:30:52 +0200 |
| commit | 701ac9bf76da372c19a147dfd3ec10d31eaa5a52 (patch) | |
| tree | edb452a4e5b42ff6a8a3ad1c45925fe5d4b8bab4 | |
| parent | ec501055f93f6f46fa0d0ea1208eeda0e9190518 (diff) | |
| parent | 9560193a9e4de27889e3980c8353a050ff2e00da (diff) | |
Merge pull request #368 from etomzak/master
Fix out-of-bounds read/write in XhadFaster
| -rw-r--r-- | src/kernels/level1/xhad.opencl | 6 |
1 files changed, 4 insertions, 2 deletions
diff --git a/src/kernels/level1/xhad.opencl b/src/kernels/level1/xhad.opencl index 3880b7a4..24e0c76c 100644 --- a/src/kernels/level1/xhad.opencl +++ b/src/kernels/level1/xhad.opencl @@ -97,10 +97,12 @@ void XhadFaster(const int n, const real_arg arg_alpha, const real_arg arg_beta, const real alpha = GetRealArg(arg_alpha); const real beta = GetRealArg(arg_beta); - if (get_global_id(0) < n / (VW)) { + const int num_desired_threads = n / (VW * WPT); + + if (get_global_id(0) < num_desired_threads) { #pragma unroll for (int _w = 0; _w < WPT; _w += 1) { - const int id = _w*get_global_size(0) + get_global_id(0); + const int id = _w * num_desired_threads + get_global_id(0); realV xvalue = xgm[id]; realV yvalue = ygm[id]; realV zvalue = zgm[id]; |