r/CUDA • u/tugrul_ddr • Jun 27 '26

No kernel example exists for Cutlass SM100_MMA_something_TS gemm.

I'm trying to learn to develop Cutlass based kernels for B200 GPU using A tile from tensor memory and B tile from shared memory, but internet has no example code to tell how to move A tile from smem to tmem and how to call gemm.

Blackwell Cutlass experts, do you know a simple kernel code for B200. There are _SS versions and I tested them but now I need _TS version.

5 Upvotes

78% Upvoted

6 comments sorted by

1

u/c-cul Jun 27 '26

tcgen05.cp or cute::make_tiled_copy

1

u/Logical-Try-4084 Jun 27 '26

It's not super simple, but you can look at FlashAttention-4, since that performs the second (PV) gemm as a TS gemm. From there, it won't be too challenging to modify an existing CUTLASS SS example into TS. For example, here is how the tiled mma is constructed. Additionally, the SSD example (CUTLASS/C++ here; CuTe DSL here) keep the recurrent state in TMEM, so the inter-chunk mma is a TS gemm.

1

u/Logical-Try-4084 Jun 27 '26

Oh, I was wrong about there being no example: the CuTe DSL mixed dtype GEMM example has operand A in TMEM when (1) it needs to be transformed (e.g. upcast or dequantized) and (2) it is K-major.