Skip to content

Commit db9708f

Browse files
committed
Delete unnecessary #ifdef / #endif.
1 parent 9a902f3 commit db9708f

File tree

2 files changed

+6
-4
lines changed

2 files changed

+6
-4
lines changed

include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_cooperative.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -608,9 +608,10 @@ class GemmUniversal<
608608
// Get pipeline stage increments from tensor shapes
609609
auto k_tile_count = size<3>(gA_mkl);
610610
611-
#ifdef CUTLASS_ENABLE_GDC_FOR_SM90
611+
// Ensure that the kernel does not touch
612+
// unflushed global memory prior to this instruction
612613
cutlass::arch::wait_on_dependent_grids();
613-
#endif
614+
614615
if (warp_group_role == WarpGroupRole::Producer) {
615616
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
616617

include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -646,9 +646,10 @@ class GemmUniversal<
646646
// Get pipeline stage increments from tensor shapes
647647
auto k_tile_count = size<3>(gA_mkl);
648648
649-
#ifdef CUTLASS_ENABLE_GDC_FOR_SM90
649+
// Ensure that the kernel does not touch
650+
// unflushed global memory prior to this instruction
650651
cutlass::arch::wait_on_dependent_grids();
651-
#endif
652+
652653
if (warp_group_role == WarpGroupRole::Producer) {
653654
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
654655

0 commit comments

Comments
 (0)