Fixed thread synchronization bug

9c1a2669 · Peter Eastman · a5e3869f · 9c1a2669
Commit 9c1a2669 authored Jan 27, 2010 by Peter Eastman
Show whitespace changes
Inline Side-by-side

Showing with 3 additions and 1 deletion

platforms/opencl/src/kernels/sort.cl platforms/opencl/src/kernels/sort.cl +3 -1

No files found.
--- a/platforms/opencl/src/kernels/sort.cl
+++ b/platforms/opencl/src/kernels/sort.cl
@@ -76,8 +76,10 @@ __kernel void computeBucketPositions(int numBuckets, __global int* bucketOffset,
        // Perform a parallel prefix sum.
        
        for (int step = 1; step < get_local_size(0); step *= 2) {
+            int add = buffer[get_local_id(0)-step];
+            barrier(CLK_LOCAL_MEM_FENCE);
            if (get_local_id(0) >= step)
-                buffer[get_local_id(0)] += buffer[get_local_id(0)-step];
+                buffer[get_local_id(0)] += add;
            barrier(CLK_LOCAL_MEM_FENCE);
        }