Module tensormap 

TMA descriptors (CUtensorMap) for cp.async.bulk.tensor loads.

A TMA descriptor moves address generation, tiling and BOUNDS CHECKING for a global→shared copy into hardware: the kernel supplies tile coordinates and the copy engine does the rest, with out-of-range elements zero-filled instead of masked by hand. That removes the per-row index arithmetic, the tail predication and the register round-trip that cp.async still pays.

Why this exists

kernels/gb10/common/gated_delta_rule_fla.cu claimed “GB10 sm_121 has cp.async.cg (NO TMA)”. That is false: cp.async.bulk.tensor + mbarrier compile for sm_121a under CUDA 13.0, and the FlashQLA GDN spine measured on a GB10 at 15.0 ms uses exactly that against our 105.7 ms. Nothing in kernels/ used TMA before this module.

Contract (the parts that bite)

• The descriptor is 128 bytes aligned to 64 — hence repr(C, align(64)). It is passed BY VALUE to a __grid_constant__ const CUtensorMap parameter, which is why KernelLaunch::arg_tensormap has to occupy 16 slots.
• global_strides carries rank - 1 entries in bytes. The innermost stride is implicit and must equal the element size, so a tensor whose fast axis is not contiguous cannot be described.
• Every stride must be 16-byte aligned, and so must global_address.
• Dimensions are in ELEMENTS and are ordered fastest-varying FIRST — the opposite of the row-major [rows][cols] we write everywhere else. Getting this backwards does not error; it silently transposes the load.

Structs

TensorMap

A 128-byte TMA descriptor, aligned as the driver requires.