Skip to content

GitLab

"csrc/smxx/decode/vscode:/vscode.git/clone" did not exist on "945ced44aa42465620848cf1d5a40af2ce63f097"
Commit 42cbcb89 authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed a bug that caused a crash on Fermi.

parent 9f4bd981
Hide whitespace changes
Inline Side-by-side
Showing with 4 additions and 5 deletions
platforms/opencl/src/kernels/sort.cl
View file @ 42cbcb89
...@@ -74,12 +74,11 @@ __kernel void computeBucketPositions(int numBuckets, __global int* bucketOffset, ...@@ -74,12 +74,11 @@ __kernel void computeBucketPositions(int numBuckets, __global int* bucketOffset,
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
// Perform a parallel prefix sum. // Perform a parallel prefix sum.
for (int step = 1; step < get_local_size(0); step *= 2) { for (int step = 1; step < get_local_size(0); step *= 2) {
int add = buffer[get_local_id(0)-step]; int add = (get_local_id(0) >= step ? buffer[get_local_id(0)-step] : 0);
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
if (get_local_id(0) >= step) buffer[get_local_id(0)] += add;
buffer[get_local_id(0)] += add;
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
...@@ -170,4 +169,4 @@ __kernel void sortBuckets(__global TYPE* data, __global TYPE* buckets, int numBu ...@@ -170,4 +169,4 @@ __kernel void sortBuckets(__global TYPE* data, __global TYPE* buckets, int numBu
} }
} }
} }
} }
\ No newline at end of file
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