Skip to content

GitLab

Commit a5e3869f authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed thread synchronization bug

parent b1a22214
Hide whitespace changes
Inline Side-by-side
Showing with 4 additions and 2 deletions
platforms/opencl/src/kernels/sort.cl
View file @ a5e3869f
...@@ -23,17 +23,19 @@ __kernel void computeRange(__global TYPE* data, int length, __global float2* ran ...@@ -23,17 +23,19 @@ __kernel void computeRange(__global TYPE* data, int length, __global float2* ran
// Now reduce them. // Now reduce them.
buffer[get_local_id(0)] = minimum; buffer[get_local_id(0)] = minimum;
barrier(CLK_LOCAL_MEM_FENCE);
for (int step = 1; step < get_local_size(0); step *= 2) { for (int step = 1; step < get_local_size(0); step *= 2) {
if (get_local_id(0)+step < get_local_size(0) && get_local_id(0)%(2*step) == 0) if (get_local_id(0)+step < get_local_size(0) && get_local_id(0)%(2*step) == 0)
buffer[get_local_id(0)] = min(buffer[get_local_id(0)], buffer[get_local_id(0)+step]); buffer[get_local_id(0)] = min(buffer[get_local_id(0)], buffer[get_local_id(0)+step]);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
minimum = buffer[0]; minimum = buffer[0];
buffer[get_local_id(0)] = maximum; buffer[get_local_id(0)] = maximum;
barrier(CLK_LOCAL_MEM_FENCE);
for (int step = 1; step < get_local_size(0); step *= 2) { for (int step = 1; step < get_local_size(0); step *= 2) {
if (get_local_id(0)+step < get_local_size(0) && get_local_id(0)%(2*step) == 0) if (get_local_id(0)+step < get_local_size(0) && get_local_id(0)%(2*step) == 0)
buffer[get_local_id(0)] = max(buffer[get_local_id(0)], buffer[get_local_id(0)+step]); buffer[get_local_id(0)] = max(buffer[get_local_id(0)], buffer[get_local_id(0)+step]);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
maximum = buffer[0]; maximum = buffer[0];
if (get_local_id(0) == 0) if (get_local_id(0) == 0)
......
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