LaunchParamConfig documentation?

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:


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

The thread extents are part of the params of the PrimFunc for the device code, kernel launch parameters are set here c( 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


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 at

1 Like

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