PrimFunc specific attribute names. More...

Functions
bool	IsPragmaKey (const std::string &attr_key)
	Check if attr_key is a pragma key extension.

Variables
constexpr const char *	kKernelLaunchParams = "tirx.kernel_launch_params"
	List of thread IterVar that a DeviceLaunch function corresponds to.

constexpr const char *	kLaunchBoundsMinBlocksPerSM = "tirx.launch_bounds_min_blocks_per_sm"
	CUDA launch bound minimum CTAs per SM.

constexpr const char *	kLaunchBoundsMaxBlocksPerCluster
	CUDA launch bound maximum CTAs per cluster.

constexpr const char *	kMaxRegisters = "tirx.max_registers"
	CUDA maximum registers per thread.

constexpr const char *	kRequiredBlockSize = "tirx.required_block_size"
	Require CUDA to use the statically-declared block and cluster dimensions.

constexpr const char *	kNoAlias = "tirx.noalias"
	Whether to set noalias rule on the function arguments.

constexpr const char *	kIsEntryFunc = "tirx.is_entry_func"
	Mark the function as the entry function of the final generated runtime module.

constexpr const char *	kIsGlobalFunc = "tirx.is_global_func"
	Mark the function as the global function called from the host.

constexpr const char *	kIsHostFunc = "tirx.is_host_func"
	Mark the function as run on the host, mutually exclusive with kTarget.

constexpr const char *	kIsScheduled = "tirx.is_scheduled"
	Mark the function as scheduled, so the default schedule will pass will skip it.

constexpr const char *	compute_scope = "compute_scope"
	Mark the scope as when computation start to happen. This can hint some code generator to create a new function for compute.

constexpr const char *	device_id = "device_id"
	The allocation device for global malloc in host.

constexpr const char *	device_scope = "device_scope"
	Mark that it is in the device scope.

constexpr const char *	device_type = "device_type"
	The device type.

constexpr const char *	extern_scope = "extern_scope"
	Mark the scope as generated by extern primitive. Such scope can contain arbitrary ir program and we need to be careful when making certain assumptions about the structure of the program.

constexpr const char *	pragma_auto_unroll_max_step = "pragma_auto_unroll_max_step"
	Pragma: auto-unroll, max_step.

constexpr const char *	pragma_import_c = "pragma_import_c"
	Import C source or file into the final code gen module.

constexpr const char *	pragma_import_llvm = "pragma_import_llvm"
	Import llvm source or file into the final code gen module.

constexpr const char *	pragma_unroll_explicit = "pragma_unroll_explicit"
	Pragma: unroll explicit.

constexpr const char *	storage_alignment = "storage_alignment"
	Mark storage alignment requirement of buffers.

constexpr const char *	thread_extent = "thread_extent"
	Mark launching extent of thread, used by device API.

constexpr const char *	kVolatile = "tirx.volatile"
	Annotation key on AllocBuffer marking the allocation as volatile.

constexpr const char *	buffer_data_alignment = "buffer_data_alignment"
	Mark buffer initial addr alignment in bytes.

constexpr const char *	buffer_allocated_addr = "buffer_allocated_addr"
	Mark buffer allocated addr in bytes.

constexpr const char *	tensorized_nki_instruction = "tensorized_nki_instruction"

constexpr const char *	kPersistentKernel = "tirx.persistent_kernel"
	Mark the kernel as persistent.

constexpr const char *	kDeviceEntry = "tirx.device_entry"
	Mark the device-region entry within a PrimFunc body. The `AttrStmt` so-keyed has a body that is the device-side region; anything before the marker (within the PrimFunc body) is host code. Value is `IntImm("bool", 1)` – a boolean marker, similar to `kPersistentKernel`.

Detailed Description

PrimFunc specific attribute names.

namespace of possible attributes in AttrStmt.attr_key

See also: tvm::attr

Function Documentation

◆ IsPragmaKey()

bool tvm::tirx::attr::IsPragmaKey ( const std::string & attr_key )

inline

Check if attr_key is a pragma key extension.

Parameters

attr_key The attr key to be compared

Returns: true if it is a pragma key

Variable Documentation

◆ buffer_allocated_addr

constexpr const char* tvm::tirx::attr::buffer_allocated_addr = "buffer_allocated_addr"

constexpr

Mark buffer allocated addr in bytes.

◆ buffer_data_alignment

constexpr const char* tvm::tirx::attr::buffer_data_alignment = "buffer_data_alignment"

constexpr

Mark buffer initial addr alignment in bytes.

◆ compute_scope

constexpr const char* tvm::tirx::attr::compute_scope = "compute_scope"

constexpr

Mark the scope as when computation start to happen. This can hint some code generator to create a new function for compute.

◆ device_id

constexpr const char* tvm::tirx::attr::device_id = "device_id"

constexpr

The allocation device for global malloc in host.

◆ device_scope

constexpr const char* tvm::tirx::attr::device_scope = "device_scope"

constexpr

Mark that it is in the device scope.

◆ device_type

constexpr const char* tvm::tirx::attr::device_type = "device_type"

constexpr

The device type.

◆ extern_scope

constexpr const char* tvm::tirx::attr::extern_scope = "extern_scope"

constexpr

Mark the scope as generated by extern primitive. Such scope can contain arbitrary ir program and we need to be careful when making certain assumptions about the structure of the program.

◆ kDeviceEntry

constexpr const char* tvm::tirx::attr::kDeviceEntry = "tirx.device_entry"

constexpr

Mark the device-region entry within a PrimFunc body. The AttrStmt so-keyed has a body that is the device-side region; anything before the marker (within the PrimFunc body) is host code. Value is IntImm("bool", 1) – a boolean marker, similar to kPersistentKernel.

◆ kIsEntryFunc

constexpr const char* tvm::tirx::attr::kIsEntryFunc = "tirx.is_entry_func"

constexpr

Mark the function as the entry function of the final generated runtime module.

Type: IntImm

Note: There can only be one entry function per module.

◆ kIsGlobalFunc

constexpr const char* tvm::tirx::attr::kIsGlobalFunc = "tirx.is_global_func"

constexpr

Mark the function as the global function called from the host.

Type: IntImm

◆ kIsHostFunc

constexpr const char* tvm::tirx::attr::kIsHostFunc = "tirx.is_host_func"

constexpr

Mark the function as run on the host, mutually exclusive with kTarget.

Type: IntImm

◆ kIsScheduled

constexpr const char* tvm::tirx::attr::kIsScheduled = "tirx.is_scheduled"

constexpr

Mark the function as scheduled, so the default schedule will pass will skip it.

Type: IntImm

◆ kKernelLaunchParams

constexpr const char* tvm::tirx::attr::kKernelLaunchParams = "tirx.kernel_launch_params"

constexpr

List of thread IterVar that a DeviceLaunch function corresponds to.

Type: ffi::Array<ffi::String>

We call a device kernel launch function f using the following convention:

Call(f, [arg1, arg2, ..., arg_n, work_size_1, work_size_2, ... work_size_m, dyn_shmem_size])

Flag-only launch tags do not add packed operands. The dynamic shared-memory operand is present only when its value-bearing tag is listed.

The list of kernel launch params indicates which additional parameters will be provided to the ffi::Function by the calling scope.

"threadIdx.x", "threadIdx.y", "threadIdx.z"

The extent of the thread count in x/y/z, to be used when launching the compute kernel on the device. For example, the gridDimX/Y/Z parameters passed to cuLaunchKernel when launching a CUDA kernel, or the groupCountX/Y/Z parameters passed to vkCmdDispatch when dispatching a compute pipeline to Vulkan.
"blockIdx.x", "blockIdx.y", "blockIdx.z"

The extent of the block iterators, to be used when launching the compute kernel on the device. For example, the blockDimX/Y/Z parameters passed to cuLaunchKernel when launching a CUDA kernel. For runtimes that do not require the block to be provided externally, this parameter is ignored. For example, the spv::ExecutionModeLocalSize for SPIR-V shaders on Vulkan, where this parameter is defined in the shader.
tvm::runtime::launch_param::kUseDynamicSharedMemoryTag

The dynamic shared-memory byte count passed for this launch.

Defined as "tirx.use_dyn_shared_memory".
tvm::runtime::launch_param::kUseProgramaticDependentLaunch
tvm::runtime::launch_param::kUseCooperativeLaunch
tvm::runtime::launch_param::kUseRequiredBlockDimension

Flag-only launch attributes. These tags add no packed operand.

See also: tvm::CallingConv::kDeviceKernelLaunch

◆ kLaunchBoundsMaxBlocksPerCluster

constexpr const char* tvm::tirx::attr::kLaunchBoundsMaxBlocksPerCluster

constexpr

Initial value:

=

"tirx.launch_bounds_max_blocks_per_cluster"

CUDA launch bound maximum CTAs per cluster.

Type: IntImm

◆ kLaunchBoundsMinBlocksPerSM

constexpr const char* tvm::tirx::attr::kLaunchBoundsMinBlocksPerSM = "tirx.launch_bounds_min_blocks_per_sm"

constexpr

CUDA launch bound minimum CTAs per SM.

Type: IntImm

◆ kMaxRegisters

constexpr const char* tvm::tirx::attr::kMaxRegisters = "tirx.max_registers"

constexpr

CUDA maximum registers per thread.

Emits the CUDA 13 __maxnreg__ kernel qualifier. This attribute is mutually exclusive with the launch-bounds attributes.

Type: IntImm

◆ kNoAlias

constexpr const char* tvm::tirx::attr::kNoAlias = "tirx.noalias"

constexpr

Whether to set noalias rule on the function arguments.

Type: IntImm

◆ kPersistentKernel

constexpr const char* tvm::tirx::attr::kPersistentKernel = "tirx.persistent_kernel"

constexpr

Mark the kernel as persistent.

◆ kRequiredBlockSize

constexpr const char* tvm::tirx::attr::kRequiredBlockSize = "tirx.required_block_size"

constexpr

Require CUDA to use the statically-declared block and cluster dimensions.

Emits the CUDA 13 __block_size__ kernel qualifier. Unlike __launch_bounds__, this is an exact launch contract: CUDA derives the PTX .reqntid directive from the thread extents, and interprets the launch grid in clusters using the cluster-CTA extents.

Type: IntImm (must be 1)

◆ kVolatile

constexpr const char* tvm::tirx::attr::kVolatile = "tirx.volatile"

constexpr

Annotation key on AllocBuffer marking the allocation as volatile.

◆ pragma_auto_unroll_max_step

constexpr const char* tvm::tirx::attr::pragma_auto_unroll_max_step = "pragma_auto_unroll_max_step"

constexpr

Pragma: auto-unroll, max_step.

◆ pragma_import_c

constexpr const char* tvm::tirx::attr::pragma_import_c = "pragma_import_c"

constexpr

Import C source or file into the final code gen module.

◆ pragma_import_llvm

constexpr const char* tvm::tirx::attr::pragma_import_llvm = "pragma_import_llvm"

constexpr

Import llvm source or file into the final code gen module.

◆ pragma_unroll_explicit

constexpr const char* tvm::tirx::attr::pragma_unroll_explicit = "pragma_unroll_explicit"

constexpr

Pragma: unroll explicit.

◆ storage_alignment

constexpr const char* tvm::tirx::attr::storage_alignment = "storage_alignment"

constexpr

Mark storage alignment requirement of buffers.

◆ tensorized_nki_instruction

constexpr const char* tvm::tirx::attr::tensorized_nki_instruction = "tensorized_nki_instruction"

constexpr

◆ thread_extent

constexpr const char* tvm::tirx::attr::thread_extent = "thread_extent"

constexpr

Mark launching extent of thread, used by device API.

Functions

Variables

Detailed Description

Function Documentation

◆ IsPragmaKey()

Variable Documentation

◆ buffer_allocated_addr

◆ buffer_data_alignment

◆ compute_scope

◆ device_id

◆ device_scope

◆ device_type

◆ extern_scope

◆ kDeviceEntry

◆ kIsEntryFunc

◆ kIsGlobalFunc

◆ kIsHostFunc

◆ kIsScheduled

◆ kKernelLaunchParams

◆ kLaunchBoundsMaxBlocksPerCluster

◆ kLaunchBoundsMinBlocksPerSM

◆ kMaxRegisters

◆ kNoAlias

◆ kPersistentKernel

◆ kRequiredBlockSize

◆ kVolatile

◆ pragma_auto_unroll_max_step

◆ pragma_import_c

◆ pragma_import_llvm

◆ pragma_unroll_explicit

◆ storage_alignment

◆ tensorized_nki_instruction

◆ thread_extent