[Bug Fix]Bypass launch grids for SM120 Kernel with SM90 Mainloop & SM100 TileScheduler

[HydraQYH]

Follow: #2719. Some SM120 kernels will use SM90 Cooperative/Pingpong Mainloop with SM100 TileScheduler. SM100 TileScheduler do not have is_last_tile function. This can cause compilation errors when we remove unnecessary conditional compilation.
In this PR, If SM120 blockscaled is detected at compile time, then bypass launch_dependent_grids.

[HydraQYH]

@Junkai-Wu I believe the current solution is only temporary, and in the future we should add is_last_tile to the SM100 TileScheduler. cc @hwu36 @IonThruster @d-k-b

[Junkai-Wu]

@Junkai-Wu I believe the current solution is only temporary, and in the future we should add is_last_tile to the SM100 TileScheduler. cc @hwu36 @IonThruster @d-k-b

Agreed. @HydraQYH Do you think if you can do the corresponding changes in this PR? I think just an empty const function with false return should be OK.

[HydraQYH]

@Junkai-Wu I believe the current solution is only temporary, and in the future we should add is_last_tile to the SM100 TileScheduler. cc @hwu36 @IonThruster @d-k-b

Agreed. @HydraQYH Do you think if you can do the corresponding changes in this PR? I think just an empty const function with false return should be OK.

@Junkai-Wu Thank you for your suggestion, I will do it as soon as possible.

[HydraQYH]

@Junkai-Wu I rebase code and just add is_last_tile for SM100 TileScheduler. After testing with following command:

cmake .. -DCUTLASS_NVCC_ARCHS=120a -DCUTLASS_BUILD_FOR_PROFILER_REGRESSIONS=ON
make VERBOSE=1 cutlass_profiler -j16

There are no compilation errors anymore. And this PR is ready for review.

[Junkai-Wu]

@HydraQYH The changes look much cleaner now. Thanks for the quick action. Leave minor comments above.

[Junkai-Wu]

@HydraQYH I ran the internal pipeline with these changes and got a timeout issue of this unit test: https://github.com/NVIDIA/cutlass/blob/main/test/unit/gemm/device/sm120_tensorop_gemm/CMakeLists.txt#L51

The issue disappeared when I added the macro back on the is_last_tile function. I've asked corresponding developer to help identify the issue. You can also investigate it if possible.

[HydraQYH]

@Junkai-Wu Sorry, I don't have an SM120 device, so all I can do is compile. However, I noticed that removing this macro in Pingpong causes the program to return immediately sometime:

cutlass/include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized_pingpong.hpp

Lines 803 to 814 in d4e16f5

	#ifdef CUTLASS_ENABLE_GDC_FOR_SM90
	// It is possible to have work tiles start off invalid,
	// so we have to check that first.
	if (not work_tile_info.is_valid()) {
	// Hint on an early release of global memory resources.
	// The timing of calling this function only influences performance,
	// not functional correctness.
	cutlass::arch::launch_dependent_grids();
	return;
	}
	#endif

I think this radical approach may be risky, and may result in some semaphores not being released yet. In contrast, other changes will not affect the program's execution. I've reverted the changes to the early stop part. Could you try this?