tvm_tirx.backend.cuda.script

tvm_tirx.backend.cuda.script#

CUDA TVMScript namespaces.

Classes#

STIRNamespace()

Nodes the s_tir pipeline's own passes build.

PTXLegacyNamespace()

Apache-compatible spellings of instructions the dialect already covers.

CudaWgmmaNamespace()

WGMMA companions that are not PTX instructions (pure-C / empty-asm helpers).

CudaTcgen05Namespace()

tcgen05 companions that are not PTX instructions (pure-C descriptor packers).

IketNamespace()

Frontend-only NVIDIA IKET annotations.

NVSHMEMNamespace()

The NVSHMEM intrinsics submodule.

NVSHMEMGetMemNBINamespace()

The NVSHMEM GetMemNBI intrinsics submodule.

NVSHMEMPutMemNBINamespace()

The NVSHMEM PutMemNBI intrinsics submodule.

NVSHMEMPutMemSignalNBINamespace()

The NVSHMEM PutMemSignalNBI intrinsics submodule.

PTXNamespace([table])

Script-facing PTX chains, also valid as tile instruction operands.

Functions#

mma_store(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.mma_store.

mma_fill(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.mma_fill.

__shfl_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.__shfl_sync.

__shfl_up_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.__shfl_up_sync.

__shfl_down_sync(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.__shfl_down_sync.

__shfl_xor_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.__shfl_xor_sync.

__activemask(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.__activemask.

any_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.any_sync.

ballot_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.ballot_sync.

bfloat1622float2(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.bfloat1622float2.

bfloat162float(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.bfloat162float.

block_dim(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.block_dim.

block_idx(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.block_idx.

clock64(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.clock64.

cluster_cta_id(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.cluster_cta_id.

cluster_dim(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.cluster_dim.

cluster_id(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.cluster_id.

cluster_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.cluster_sync.

cta_pair_id(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.cta_pair_id.

cta_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.cta_sync.

cvta_generic_to_shared(*args[, attrs, ...])

Construct a call to tirx.cuda.cvta_generic_to_shared.

decl_tmem(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.decl_tmem.

elect_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.elect_sync.

fadd2_rn(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.fadd2_rn.

fdividef(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.fdividef.

ffs_u32(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.ffs_u32.

float22bfloat162_rn(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.float22bfloat162_rn.

float22bfloat162_rn_from_float2(*args[, ...])

Construct a call to tirx.cuda.float22bfloat162_rn_from_float2.

float22half2(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.float22half2.

float2_x(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.float2_x.

float2_y(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.float2_y.

float8tohalf8(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.float8tohalf8.

float_as_uint(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.float_as_uint.

fmul2_rn(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.fmul2_rn.

fp8x4_e4m3_from_float4(*args[, attrs, ...])

Construct a call to tirx.cuda.fp8x4_e4m3_from_float4.

get_tmem_addr(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.get_tmem_addr.

grid_dim(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.grid_dim.

grid_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.grid_sync.

half2float(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.half2float.

half8tofloat8(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.half8tofloat8.

hmax2(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.hmax2.

hmin2(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.hmin2.

iket_mark(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.iket_mark.

iket_official_event(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.iket_official_event.

iket_range_end(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.iket_range_end.

iket_range_pop(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.iket_range_pop.

iket_range_push(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.iket_range_push.

iket_range_start(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.iket_range_start.

iket_sentinel_token(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.iket_sentinel_token.

lane_id(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.lane_id.

linear_thread_id(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.linear_thread_id.

make_float2(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.make_float2.

mbarrier_wait(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.mbarrier_wait.

mbarrier_wait_acquire_cluster(*args[, ...])

Construct a call to tirx.cuda.mbarrier_wait_acquire_cluster.

nano_sleep(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.nano_sleep.

printf(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.printf.

reduce_add_sync_u32(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.reduce_add_sync_u32.

reduce_min_sync_u32(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.reduce_min_sync_u32.

runtime_instr_desc(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.runtime_instr_desc.

smem_addr_from_uint64(*args[, attrs, ...])

Construct a call to tirx.cuda.smem_addr_from_uint64.

syncthreads_and(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.syncthreads_and.

syncthreads_or(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.syncthreads_or.

tcgen05_encode_instr_descriptor(*args[, ...])

Construct a call to tirx.cuda.tcgen05_encode_instr_descriptor.

tcgen05_encode_instr_descriptor_block_scaled(*args)

Construct a call to tirx.cuda.tcgen05_encode_instr_descriptor_block_scaled.

tcgen05_encode_matrix_descriptor(*args[, ...])

Construct a call to tirx.cuda.tcgen05_encode_matrix_descriptor.

tensormap_encode_tiled(*args[, attrs, ...])

Construct a call to tirx.cuda.tensormap_encode_tiled.

thread_fence(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.thread_fence.

thread_idx(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.thread_idx.

thread_in_warpgroup(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.thread_in_warpgroup.

thread_rank(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.thread_rank.

trap_when_assert_failed(*args[, attrs, ...])

Construct a call to tirx.cuda.trap_when_assert_failed.

uint_as_float(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.uint_as_float.

warp_id(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.warp_id.

warp_in_warpgroup(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.warp_in_warpgroup.

warp_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.warp_sync.

warpgroup_id(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.warpgroup_id.

warpgroup_sync(*args[, attrs, ty_args, loc, ty])

Construct a call to tirx.cuda.warpgroup_sync.

wgmma_encode_matrix_descriptor(*args[, ...])

Construct a call to tirx.cuda.wgmma_encode_matrix_descriptor.

wgmma_noop_barrier(*args[, attrs, ty_args, ...])

Construct a call to tirx.cuda.wgmma_noop_barrier.

Namespaces#