can_implement check question

[luote1996]

cutlass/include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized.hpp

Line 246 in 8dbce01

[[maybe_unused]] Arguments const& args) {

why use StrideA{} check but not actual parameters at runtime args.dA to check?

[BulletsHo]

Great question! The reason can_implement uses StrideA{} instead of runtime args.dA boils down to compile-time static polymorphism and zero-overhead dispatch in the CUTLASS 3.x architecture.

Here is why it is designed this way:

Compile-Time Evaluation & TMA Descriptor Generation:
In CUTLASS 3.x, collective mainloops like MainloopSm90TmaGmmaWarpSpecialized rely heavily on compile-time layout and stride types to construct TMA hardware descriptors (make_tma_copy_A_sm90) and shared memory layouts (SmemLayoutA). StrideA{} provides the compile-time type traits needed to evaluate whether the layout satisfies the strict 128-bit alignment requirements enforced by Hopper's TMA engine.

Avoiding Runtime Overhead:
can_implement acts as an early-exit guard before kernel launch and is often evaluated in constexpr contexts or during plan/instantiation phases. Relying on args.dA would turn a static structural check into a runtime evaluation, going against CUTLASS's design philosophy of resolving hardware compatibility upfront.

Note that while runtime leading dimensions can vary, the layout structure (e.g., whether it's strictly row-major or column-major with standard strides) is captured by the StrideA template parameter, making the static check on StrideA{} both necessary and sufficient for structural capability validation.

[Hosseinghorbani0]

Thanks for the explanation. I looked a bit closer at the implementation, and I think there is an important distinction between the compile-time stride type and the runtime stride value here.

In can_implement(), the check is specifically:

check_alignment<min_tma_aligned_elements_A>(cute::make_shape(M,K,L), StrideA{})

while to_underlying_arguments() later constructs the actual tensor using args.dA:

make_tensor(ptr_A, make_layout(make_shape(M,K,L), args.dA))

So I understand why StrideA{} is useful for validating properties that are determined by the stride type/layout.

My remaining question is about cases where StrideA is a dynamic stride type and the actual runtime args.dA does not satisfy the required TMA alignment. In that situation, would can_implement() still return true even though the TMA descriptor is ultimately constructed from the runtime stride?

If so, is the assumption that the caller is responsible for guaranteeing that args.dA is compatible with the StrideA type and satisfies the required alignment, rather than can_implement() validating the concrete runtime stride?

That seems to be the key distinction between checking the layout type and checking the actual runtime tensor layout.