tilelang.contrib.cutedsl.ldst¶

Explicit-width and read-only-cache loads and stores for the CuTeDSL backend.

Mirrors tl::load_global_* / tl::store_global_* / tl::load_shared_* / tl::store_shared_* in src/tl_templates/cuda/copy.h and copy_sm100.h, and CUDA’s __ldg.

Functions¶

`ldg32`(ptr[, pred])	`tl::load_global_32[_conditional]`: `@p ld.global.u32`; 0 when `pred` is false.
`ldg64`(ptr[, pred])	`tl::load_global_64[_conditional]`: `@p ld.global.v2.u32`; zeros when `pred` is false.
`ldg128`(ptr[, pred])	`tl::load_global_128[_conditional]`: `@p ld.global.v4.u32`; zeros when `pred` is false.
`ldg256`(ptr[, pred])	`tl::load_global_256[_conditional]`: `@p ld.global.v4.u64` on sm_100+, else two
`stg32`(ptr, value[, pred, loc, ip])	`tl::store_global_32[_conditional]`: `@p st.global.u32`.
`stg64`(ptr, value[, pred, loc, ip])	`tl::store_global_64[_conditional]`: `@p st.global.v2.u32`.
`stg128`(ptr, value[, pred, loc, ip])	`tl::store_global_128[_conditional]`: `@p st.global.v4.u32`.
`stg256`(ptr, value[, pred, loc, ip])	`tl::store_global_256[_conditional]`: `@p st.global.v4.u64` on sm_100+, else two
`lds32`(ptr)	`tl::load_shared_32`: a plain 32-bit load.
`lds64`(ptr)	`tl::load_shared_64`: a plain 64-bit (`uint2`) load.
`lds128`(ptr)	`tl::load_shared_128`: a plain 128-bit (`uint4`) load.
`sts32`(ptr, value)	`tl::store_shared_32`: a plain 32-bit store.
`sts64`(ptr, value)	`tl::store_shared_64`: a plain 64-bit (`uint2`) store.
`sts128`(ptr, value)	`tl::store_shared_128`: a plain 128-bit (`uint4`) store.

Module Contents¶

tilelang.contrib.cutedsl.ldst.ldg32(ptr, pred=True)¶: tl::load_global_32[_conditional]: @p ld.global.u32; 0 when pred is false.

tilelang.contrib.cutedsl.ldst.ldg64(ptr, pred=True)¶: tl::load_global_64[_conditional]: @p ld.global.v2.u32; zeros when pred is false.

tilelang.contrib.cutedsl.ldst.ldg128(ptr, pred=True)¶: tl::load_global_128[_conditional]: @p ld.global.v4.u32; zeros when pred is false.

tilelang.contrib.cutedsl.ldst.ldg256(ptr, pred=True)¶: tl::load_global_256[_conditional]: @p ld.global.v4.u64 on sm_100+, else two @p ld.global.v4.u32; zeros when pred is false.

tilelang.contrib.cutedsl.ldst.stg32(ptr, value, pred=True, *, loc=None, ip=None)¶: tl::store_global_32[_conditional]: @p st.global.u32.

tilelang.contrib.cutedsl.ldst.stg64(ptr, value, pred=True, *, loc=None, ip=None)¶: tl::store_global_64[_conditional]: @p st.global.v2.u32.

tilelang.contrib.cutedsl.ldst.stg128(ptr, value, pred=True, *, loc=None, ip=None)¶: tl::store_global_128[_conditional]: @p st.global.v4.u32.

tilelang.contrib.cutedsl.ldst.stg256(ptr, value, pred=True, *, loc=None, ip=None)¶: tl::store_global_256[_conditional]: @p st.global.v4.u64 on sm_100+, else two @p st.global.v4.u32.

tilelang.contrib.cutedsl.ldst.lds32(ptr)¶: tl::load_shared_32: a plain 32-bit load.

tilelang.contrib.cutedsl.ldst.lds64(ptr)¶: tl::load_shared_64: a plain 64-bit (uint2) load.

tilelang.contrib.cutedsl.ldst.lds128(ptr)¶: tl::load_shared_128: a plain 128-bit (uint4) load.

tilelang.contrib.cutedsl.ldst.sts32(ptr, value)¶: tl::store_shared_32: a plain 32-bit store.

tilelang.contrib.cutedsl.ldst.sts64(ptr, value)¶: tl::store_shared_64: a plain 64-bit (uint2) store.

tilelang.contrib.cutedsl.ldst.sts128(ptr, value)¶: tl::store_shared_128: a plain 128-bit (uint4) store.