Update layer_norm_cuda_kernel.cu

f1295380 · Eddie Yan · 1ec6b0e9 · f1295380
Commit f1295380 authored Oct 12, 2021 by Eddie Yan
Hide whitespace changes
Inline Side-by-side

Showing with 2 additions and 0 deletions

megatron/fused_kernels/layer_norm_cuda_kernel.cu megatron/fused_kernels/layer_norm_cuda_kernel.cu +2 -0

No files found.
--- a/megatron/fused_kernels/layer_norm_cuda_kernel.cu
+++ b/megatron/fused_kernels/layer_norm_cuda_kernel.cu
@@ -645,6 +645,8 @@ void cuComputeGradInput(
        k_grad_input[l] = static_cast<T>(f_grad_input);
      }
    }
+    // prevent race where buf is written again before reads are done
+    __syncthreads();
  }
 }