LaunchParamConfig documentation?

lemo · January 21, 2022, 1:16am

I hope this is the right place to ask - if not, please redirect me.

I’m looking into a cuLaunchKernel() failure deep inside TVM, and I’d like to understand how the kernel launch are modeled.

Is the LaunchParamConfig explained anywhere other than the code? At very least I’d like to figure out what the following data members represent:

github.com

apache/tvm/blob/099ebaa7d5d1bd4862df36230f464cc3c92aa630/src/runtime/thread_storage_scope.h#L245-L250


/*! \brief base axis */
size_t base_;
/*! \brief The worker dimension */
size_t work_dim_;
/*! \brief The index mapping. */
std::vector<uint32_t> arg_index_map_;

Thanks!

lemo · February 10, 2022, 8:53pm

On a related note, can someone point me to the code responsible for calculating kernel launch parameters?

vinx13 · February 11, 2022, 7:12pm

The thread extents are part of the params of the PrimFunc for the device code, kernel launch parameters are set here c(https://github.com/apache/tvm/blob/main/src/tir/transforms/split_host_device.cc#L294-L296). In LaunchParamConfig, base and arg_index_map is used to map the index of params of PrimFunc to the thread extent which is an array of six integers representing (blockIdx.x/y/z, threadIdx.x/y/z). work_dim is the dimension of the thread extent, which is the highest dimension among x/y/z specified in the thread extent

masahi · February 11, 2022, 8:19pm

This is done by each op separately, most of them written in python, for example

Note that the launch param also includes dynamic shmem size. These params that are specified in python are retrieved in split_host_device.cc at https://github.com/apache/tvm/blob/main/src/tir/transforms/split_host_device.cc#L46

zzh_ontheway · August 30, 2022, 7:23am

what about void_args in cukernellaunch? is it packed with json graph info and params? Could you tell me where can I find the related codes? thx