Variable tvm::tirx::attr::kKernelLaunchParams#
Defined in File function.h
Variable Documentation#
-
constexpr ffi::reflection::AttrKey<ffi::Array<ffi::String>> tvm::tirx::attr::kKernelLaunchParams{"tirx.kernel_launch_params"}#
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.