Skip to content

Commit cf09543

Browse files
committed
Delete unnecessary #ifdef / #endif.
1 parent b0f28c1 commit cf09543

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
@@ -604,9 +604,10 @@ class GemmUniversal<
604604
// Get pipeline stage increments from tensor shapes
605605
auto k_tile_count = size<3>(gA_mkl);
606606
607-
#ifdef CUTLASS_ENABLE_GDC_FOR_SM90
607+
// Ensure that the kernel does not touch
608+
// unflushed global memory prior to this instruction
608609
cutlass::arch::wait_on_dependent_grids();
609-
#endif
610+
610611
if (warp_group_role == WarpGroupRole::Producer) {
611612
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
612613

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
@@ -642,9 +642,10 @@ class GemmUniversal<
642642
// Get pipeline stage increments from tensor shapes
643643
auto k_tile_count = size<3>(gA_mkl);
644644
645-
#ifdef CUTLASS_ENABLE_GDC_FOR_SM90
645+
// Ensure that the kernel does not touch
646+
// unflushed global memory prior to this instruction
646647
cutlass::arch::wait_on_dependent_grids();
647-
#endif
648+
648649
if (warp_group_role == WarpGroupRole::Producer) {
649650
cutlass::arch::warpgroup_reg_dealloc<LoadRegisterRequirement>();
650651

0 commit comments

Comments
 (0)