tvm::tirx::attr Namespace Reference

Function specific attribute names. More...

Variables
constexpr const char *	kKernelLaunchParams = "tirx.kernel_launch_params"
	Ordered launch-parameter tags for a device kernel.
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 *	kPersistentKernel = "tirx.persistent_kernel"
	Mark the kernel as persistent.
constexpr const char *	kPragmaUnroll = "pragma_unroll"
constexpr const char *	kDisableUnroll = "disable_unroll"
constexpr const char *	kAutoUnrollMaxStep = "auto_unroll_max_step"
	For annotation: maximum work for automatic unrolling.
constexpr const char *	kUnrollExplicit = "unroll_explicit"
	For annotation: expand unrolled bodies instead of preserving unrolled loops.
constexpr const char *	kVolatile = "tirx.volatile"
	Annotation key on AllocTensor marking the allocation as volatile.
constexpr const char *	kTensorDataAlignment = "tensor_data_alignment"
	Mark tensor initial addr alignment in bytes.

Detailed Description

Function specific attribute names.

Statement attribute and loop annotation keys.

See also

tvm::attr

Variable Documentation

kAutoUnrollMaxStep

constexpr const char* tvm::tirx::attr::kAutoUnrollMaxStep = "auto_unroll_max_step"

constexpr

For annotation: maximum work for automatic unrolling.

Integer policy inherited by nested loops unless they override it. Consumed by UnrollLoop.

kDisableUnroll

constexpr const char* tvm::tirx::attr::kDisableUnroll = "disable_unroll"

constexpr

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

kKernelLaunchParams

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

constexpr

Ordered launch-parameter tags for a device kernel.

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

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.

kPragmaUnroll

constexpr const char* tvm::tirx::attr::kPragmaUnroll = "pragma_unroll"

constexpr

kTensorDataAlignment

constexpr const char* tvm::tirx::attr::kTensorDataAlignment = "tensor_data_alignment"

constexpr

Mark tensor initial addr alignment in bytes.

kUnrollExplicit

constexpr const char* tvm::tirx::attr::kUnrollExplicit = "unroll_explicit"

constexpr

For annotation: expand unrolled bodies instead of preserving unrolled loops.

Integer policy inherited by nested loops unless they override it. Consumed by UnrollLoop.

kVolatile

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

constexpr

Annotation key on AllocTensor marking the allocation as volatile.