diff options
| author | Yang Rong <rong.r.yang@intel.com> | 2013-06-27 16:47:58 +0800 |
|---|---|---|
| committer | Zhigang Gong <zhigang.gong@linux.intel.com> | 2013-06-27 18:54:19 +0800 |
| commit | 2dd6e7f82536aa59c59025c061cf06c6b84f43fe (patch) | |
| tree | 22c2fb03b44205bb008862bc392fb4d11b201f7c /kernels | |
| parent | 1f854302ef8fb59cd5b4a5921e86dd84799ea3b0 (diff) | |
Fix some math function error in simd16.
INT DIV splite to simd8 but forget to set quarter_control.
Will fail when predication enable.
Change the atomic test case to trigger this bug.
Signed-off-by: Yang Rong <rong.r.yang@intel.com>
Reviewed-by: Zhigang Gong <zhigang.gong@linux.intel.com>
Diffstat (limited to 'kernels')
| -rw-r--r-- | kernels/compiler_atomic_functions.cl | 8 |
1 files changed, 4 insertions, 4 deletions
diff --git a/kernels/compiler_atomic_functions.cl b/kernels/compiler_atomic_functions.cl index 24f17c29..61ce2f44 100644 --- a/kernels/compiler_atomic_functions.cl +++ b/kernels/compiler_atomic_functions.cl @@ -7,8 +7,8 @@ __kernel void compiler_atomic_functions(__global int *dst, __local int *tmp, __g case 1: atomic_dec(&tmp[i]); break; case 2: atomic_add(&tmp[i], src[lid]); break; case 3: atomic_sub(&tmp[i], src[lid]); break; - case 4: atomic_and(&tmp[i], ~(src[lid]<<(lid>>2))); break; - case 5: atomic_or (&tmp[i], src[lid]<<(lid>>2)); break; + case 4: atomic_and(&tmp[i], ~(src[lid]<<(lid / 4))); break; + case 5: atomic_or (&tmp[i], src[lid]<<(lid / 4)); break; case 6: atomic_xor(&tmp[i], src[lid]); break; case 7: atomic_min(&tmp[i], -src[lid]); break; case 8: atomic_max(&tmp[i], src[lid]); break; @@ -23,8 +23,8 @@ __kernel void compiler_atomic_functions(__global int *dst, __local int *tmp, __g case 1: atomic_dec(&dst[i]); break; case 2: atomic_add(&dst[i], src[lid]); break; case 3: atomic_sub(&dst[i], src[lid]); break; - case 4: atomic_and(&dst[i], ~(src[lid]<<(lid>>2))); break; - case 5: atomic_or (&dst[i], src[lid]<<(lid>>2)); break; + case 4: atomic_and(&dst[i], ~(src[lid]<<(lid / 4))); break; + case 5: atomic_or (&dst[i], src[lid]<<(lid / 4)); break; case 6: atomic_xor(&dst[i], src[lid]); break; case 7: atomic_min(&dst[i], -src[lid]); break; case 8: atomic_max(&dst[i], src[lid]); break; |