[BUG] Sm90 & Sm100 Array gemm kernels read ahead of `wait_on_dependent_grids()`

[Algy]

Which component has the problem?

CUTLASS C++

Bug Report

Describe the bug

First, the first element of the array containing problem shapes are accessed in the constructor of group tile scheduler.

cutlass/include/cutlass/gemm/kernel/sm90_tile_scheduler_group.hpp

Line 267 in 8debf77

ProblemShape problem_shape = params_.problem_shapes_.get_problem_shape(0);

cutlass/include/cutlass/gemm/group_array_problem_shape.hpp

Line 68 in 8debf77

return problem_shapes[group_idx];

TileScheduler is constructed before the wait_on_dependent_grids(), which could risk a chance reading the pointer arrays before dependent data gets flushed into global memory by a preceding kernel.

cutlass/include/cutlass/gemm/kernel/sm90_gemm_array_tma_warpspecialized_pingpong.hpp

Lines 466 to 475 in 8debf77

	auto scheduler = [&] () {
	// Group scheduler requires a different constructor that takes a response ptr
	if constexpr (cute::is_same_v<SchedulerTag, GroupScheduler>) {
	return TileScheduler{params.scheduler, shared_storage.scheduler_response};
	}
	else {
	return TileScheduler{params.scheduler};
	}
	} ();

cutlass/include/cutlass/gemm/kernel/sm100_gemm_array_tma_warpspecialized.hpp

Lines 836 to 842 in 8debf77

	TileScheduler scheduler(
	(!IsTensorMapUpdateAsync || is_participant.sched || is_participant.tensor_map_updater)
	? &shared_storage.clc_response[0][0]
	: &shared_storage.clc_response[1][0],
	params.scheduler,
	block_id_in_cluster
	);

This would cause a race condition when PDL is enabled.

Steps/Code to reproduce bug

Hard to reproduce as it's a race condition.

I wish there was some compiler support to spot PDL related bugs, e.g. ,print a warning log when codes in a kernel access to global address before wait_on_dependent_grids(). It was frustrating to figure out this kind of issues.

[hwu36]

@depaulmillz

[depaulmillz]

This should be fixed now on main.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.