[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?

[Junkai-Wu]

@HydraQYH After investigation, the timeout issue is caused by sm120 kernel calling sm90 scheduler where sm120 kernel didn't need to call is_last_tile before and after the removal of the macro, it will call this function which may cause some schedule issue. Maybe revert to previously implementation which using IsBlockScaled but with a proper name would be a safe option.

[HydraQYH]

@Junkai-Wu Okay. I rebaseed the code and reverted to previously implementation. I changed is_blockscaled to IsBlockScaledDispatchPolicy, which I believe is reasonable because the template parameter is just a DispatchPolicy. I also fixed the typo in Cooperative. And it's ready for review.

[Junkai-Wu]

@HydraQYH The internal pipeline still fails. After checking, I found all sm120 kernels should not call the is_last_tile function, not just sm120 blockscaled kernels. I refactored the change and reran the internal pipeline. I'll leave comments in this PR after the pipeline passes.

[Junkai-Wu]

Add if constexpr (!IsSm120Family) condition here.

[HydraQYH]

Done.

[Junkai-Wu]

@HydraQYH The internal pipeline passes. Leave suggested implementations in the comments.

[HydraQYH]

@Junkai-Wu Thank you very much for your help. I have picked all your suggestions.