Module tensormap 

2026-09-25: TMA descriptors (CUtensorMap) for cp.async.bulk.tensor loads, encoded with the driver’s cuTensorMapEncodeTiled.

Owner: gpu-runtime. Invariants:

• TensorMap is 128 bytes aligned to 64, the size of CUtensorMap and the alignment cuTensorMapEncodeTiled requires (cuda.h); a test pins both.
• tiled_2d_bf16 rejects a global address or row stride that is not 16-byte aligned, a row stride narrower than cols, and a zero box dimension, before it calls the driver.

The driver takes dimensions fastest-varying first and rank - 1 strides in bytes, the innermost stride being the element size. tiled_2d_bf16 takes row-major [rows][cols] in elements and does that conversion. A descriptor is passed by value to a __grid_constant__ const CUtensorMap kernel parameter through KernelLaunch::arg_tensormap.

Structs

TensorMap

2026-09-25: A 128-byte TMA descriptor, aligned to 64 bytes.