Added missing synchronization to avoid WAR hazards between tiles. (NV…

…IDIA#386)
LinaKim93 · Dec 20, 2021 · 288af36 · 288af36
1 parent 0dc3ba6
commit 288af36
Showing 1 changed file with 3 additions and 0 deletions.
diff --git a/include/cutlass/gemm/kernel/gemm_grouped.h b/include/cutlass/gemm/kernel/gemm_grouped.h
@@ -546,6 +546,9 @@ struct GemmGrouped {
       // Compute threadblock-scoped matrix multiply-add
       int gemm_k_iterations = (problem_size.k() + Mma::Shape::kK - 1) / Mma::Shape::kK;
 
+      // Wait for all threads to finish their epilogue phases from the previous tile.
+      __syncthreads();
+
       // Compute threadblock-scoped matrix multiply-add
       mma(
         gemm_k_iterations,