Fix race condition in kEstimateQuantiles (bitsandbytes-foundation#1061)

SunMarc · Feb 14, 2024 · 5b28fd3 · 5b28fd3
1 parent 344e851
commit 5b28fd3
Showing 1 changed file with 2 additions and 0 deletions.
diff --git a/csrc/kernels.cu b/csrc/kernels.cu
@@ -654,6 +654,8 @@ __global__ void kEstimateQuantiles(T *__restrict__ const A, float *code, const f
       for(int j = threadIdx.x; j < BLOCK_ESTIMATE; j+=blockDim.x)
           temp_storage.smem_qidx[j] = -1;
 
+      __syncthreads();
+
       if(threadIdx.x < 256)
       {
           float q_interval = (1.0f-(2.0f*offset))/255.0f;