Merge pull request #2102 from peastman/sort

Fixed race condition in sorting

platforms/cuda/src/kernels/sort.cu

@@ -144,6 +144,7 @@ __global__ void computeBucketPositions(unsigned int numBuckets, unsigned int* __
         // Load the bucket sizes into local memory.
 
         unsigned int globalIndex = startBucket+threadIdx.x;
+        __syncthreads();
         posBuffer[threadIdx.x] = (globalIndex < numBuckets ? bucketOffset[globalIndex] : 0);
         __syncthreads();
 

platforms/opencl/src/kernels/sort.cl

@@ -147,6 +147,7 @@ __kernel void computeBucketPositions(uint numBuckets, __global uint* restrict bu
         // Load the bucket sizes into local memory.
 
         uint globalIndex = startBucket+get_local_id(0);
+        barrier(CLK_LOCAL_MEM_FENCE);
         buffer[get_local_id(0)] = (globalIndex < numBuckets ? bucketOffset[globalIndex] : 0);
         barrier(CLK_LOCAL_MEM_FENCE);