Fix the race condition in cumsum operator (#42205)

* Fix the race condition in cumsum operator * Optimize cumsum operator
PaddlePaddle · May 5, 2022 · ff455c7 · ff455c7 · paddle-bot-old · May 5, 2022
1 parent a391762
commit ff455c7
Showing 1 changed file with 5 additions and 7 deletions.
diff --git a/paddle/phi/kernels/gpu/cumsum_kernel.cu b/paddle/phi/kernels/gpu/cumsum_kernel.cu
@@ -39,14 +39,12 @@ __device__ void BlockReverse(
   int tx = threadIdx.x;
 
   int offset = tx;
-  int in_index = src_base + offset;
-  if (offset >= valid_item) {
-    sh_mem[offset] = 0;
-  } else {
-    int sh_mem_index = BLOCK_SIZE - offset - 1;
-    T data = idata[in_index];
-    sh_mem[sh_mem_index] = data;
+  T src_data = 0;
+  int src_offset = BLOCK_SIZE - offset - 1;
+  if (src_offset < valid_item) {
+    src_data = idata[src_base + src_offset];
   }
+  sh_mem[offset] = src_data;
 
   __syncthreads();
   int out_index = dst_base - offset;