Commit de666e30 authored by peastman's avatar peastman
Browse files

Fixed race condition

parent b849c856
......@@ -6,6 +6,7 @@
*/
__device__ real reduceValue(real value, volatile real* temp) {
const int thread = threadIdx.x;
__syncthreads();
temp[thread] = value;
__syncthreads();
for (uint step = 1; step < 32; step *= 2) {
......
......@@ -6,6 +6,7 @@
*/
real reduceValue(real value, __local volatile real* temp) {
const int thread = get_local_id(0);
barrier(CLK_LOCAL_MEM_FENCE);
temp[thread] = value;
barrier(CLK_LOCAL_MEM_FENCE);
for (uint step = 1; step < 32; step *= 2) {
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment