Mojo function

build_mma_ts_ws

build_mma_ts_ws[b_dtype: DType, *, b_BMN: Int, b_BK: Int, b_swizzle: TensorMapSwizzle, b_is_k_major: Bool](kind: String, *, operand_size: Int, num_k_mmas: Int, tcgen05_mma_type: String, mma_k: Int = 16) -> String

Build PTX for TS (TMEM-A, SMEM-B) .ws MMA over multiple K-tiles.

Each K-tile requires a separate TMEM address operand for A passed via inline assembly ($7,$8, ...). B is an SMEM descriptor whose low-half is bumped by the tile-layout offset for each K-tile.

Returns:

String

Was this page helpful?

Thank you! We'll create more content like this.

Thank you for helping us improve!

😔 What went wrong?

Some code doesn’t work

It includes inaccurate information

It's missing information I need

It was difficult to understand

Other

Submit