tvm::backend::cuda Namespace Reference

Namespaces
namespace	attr
namespace	op_attr

Classes
struct	TCGen05InstrDescriptorAttrs
	Static options for the dense tcgen05 instruction descriptor. More...
struct	TCGen05InstrDescriptorBlockScaledAttrs
	Static options for the block-scaled tcgen05 instruction descriptor. More...
struct	TensorMapEncodeTiledAttr
	Fixed encoding options for tensormap_encode_tiled. More...

Functions
TVM_TIRX_DLL const Op &	tensormap_encode_tiled_op ()
	Encode a tiled tensor map at invocation time.
Frontend-only NVIDIA IKET annotations
These CUDA device intrinsics have opaque call effects and no registered attrs schema or attrs defaults. Marker, range-start, range-end and range-push annotations have a registered variadic operand tail; the Python convenience builders and IKET lowering use it for one optional scalar payload. C denotes tvm.backend.cuda.op and T denotes the TIRx script namespace.
TVM_TIRX_DLL const Op &	iket_mark_op ()
	Get the NVIDIA IKET marker operator, tirx.cuda.iket_mark. Python: C.cuda_iket_mark(name, payload=None). Script: T.cuda.iket.mark(name, payload=None). Operands: [0] name, the marker name; [1...] args, the optional payload tail. Result: void. Effects: opaque marker annotation.
TVM_TIRX_DLL const Op &	iket_range_start_op ()
	Get the NVIDIA IKET token-range start operator, tirx.cuda.iket_range_start. Python: C.cuda_iket_range_start(name, payload=None). Script: T.cuda.iket.range_start(name, payload=None). Operands: [0] name, the range name; [1...] args, the optional payload tail. Result: uint32 range token. Effects: opaque range-start annotation.
TVM_TIRX_DLL const Op &	iket_range_end_op ()
	Get the NVIDIA IKET token-range end operator, tirx.cuda.iket_range_end. Python: C.cuda_iket_range_end(token, payload=None). Script: T.cuda.iket.range_end(token, payload=None). Operands: [0] token, an IntExpr identifying the range to end; [1...] args, the optional payload tail. IKET lowering requires a uint32 token. Result: void. Effects: opaque range-end annotation.
TVM_TIRX_DLL const Op &	iket_range_push_op ()
	Get the NVIDIA IKET stack-range push operator, tirx.cuda.iket_range_push. Python: C.cuda_iket_range_push(name, payload=None). Script: T.cuda.iket.range_push(name, payload=None). Operands: [0] name, the range name; [1...] args, the optional payload tail. Result: void. Effects: opaque stack-range push annotation.
TVM_TIRX_DLL const Op &	iket_range_pop_op ()
	Get the NVIDIA IKET stack-range pop operator, tirx.cuda.iket_range_pop. Python: C.cuda_iket_range_pop(). Script: T.cuda.iket.range_pop(). Operands: none; IKET lowering enforces the empty operand list. Result: void. Effects: opaque stack-range pop annotation.
TVM_TIRX_DLL const Op &	iket_sentinel_token_op ()
	Get the NVIDIA IKET sentinel-token operator, tirx.cuda.iket_sentinel_token. Python: C.cuda_iket_sentinel_token(name). Script: T.cuda.iket.sentinel_token(name). Operands: [0] name, the event name. Result: uint32 sentinel token. Registered effects: opaque. The sentinel carries token-flow identity; IKET lowering emits no runtime event for it.

Function Documentation

iket_mark_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::iket_mark_op

(

)

Get the NVIDIA IKET marker operator, tirx.cuda.iket_mark. Python: C.cuda_iket_mark(name, payload=None). Script: T.cuda.iket.mark(name, payload=None). Operands: [0] name, the marker name; [1...] args, the optional payload tail. Result: void. Effects: opaque marker annotation.

iket_range_end_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::iket_range_end_op

(

)

Get the NVIDIA IKET token-range end operator, tirx.cuda.iket_range_end. Python: C.cuda_iket_range_end(token, payload=None). Script: T.cuda.iket.range_end(token, payload=None). Operands: [0] token, an IntExpr identifying the range to end; [1...] args, the optional payload tail. IKET lowering requires a uint32 token. Result: void. Effects: opaque range-end annotation.

iket_range_pop_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::iket_range_pop_op

(

)

Get the NVIDIA IKET stack-range pop operator, tirx.cuda.iket_range_pop. Python: C.cuda_iket_range_pop(). Script: T.cuda.iket.range_pop(). Operands: none; IKET lowering enforces the empty operand list. Result: void. Effects: opaque stack-range pop annotation.

iket_range_push_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::iket_range_push_op

(

)

Get the NVIDIA IKET stack-range push operator, tirx.cuda.iket_range_push. Python: C.cuda_iket_range_push(name, payload=None). Script: T.cuda.iket.range_push(name, payload=None). Operands: [0] name, the range name; [1...] args, the optional payload tail. Result: void. Effects: opaque stack-range push annotation.

iket_range_start_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::iket_range_start_op

(

)

Get the NVIDIA IKET token-range start operator, tirx.cuda.iket_range_start. Python: C.cuda_iket_range_start(name, payload=None). Script: T.cuda.iket.range_start(name, payload=None). Operands: [0] name, the range name; [1...] args, the optional payload tail. Result: uint32 range token. Effects: opaque range-start annotation.

iket_sentinel_token_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::iket_sentinel_token_op

(

)

Get the NVIDIA IKET sentinel-token operator, tirx.cuda.iket_sentinel_token. Python: C.cuda_iket_sentinel_token(name). Script: T.cuda.iket.sentinel_token(name). Operands: [0] name, the event name. Result: uint32 sentinel token. Registered effects: opaque. The sentinel carries token-flow identity; IKET lowering emits no runtime event for it.

tensormap_encode_tiled_op()

TVM_TIRX_DLL const Op & tvm::backend::cuda::tensormap_encode_tiled_op

(

)

Encode a tiled tensor map at invocation time.

TensorMapEncodeTiledAttr stores the descriptor dtype, rank and fixed options. Arguments are descriptor and data pointers, global dimensions (rank), byte strides (rank - 1), box dimensions (rank), then element strides (rank).