#1361·warp

Optimize tiled CUDA launches by reading blockIdx.x directly

Author: adenzler-nvidiaCreated Apr 13, 2026Updated Sep 19, 2026
Labelstile

For tiled CUDA kernel launches, blockIdx.x already encodes the tile index and threadIdx.x the intra-tile position. The current codegen computes the linear thread index and then divides by block_dim() to recover the tile coordinate, which is redundant on the GPU side.

A dedicated codegen path for tiled CUDA kernels could read blockIdx.x directly instead of performing the linear /= block_dim() division and the multiply+add to reconstruct the linear index. The CPU path would continue using the current approach since it does not have an equivalent of blockIdx.

This came up during review of GH-1270 (templated launch_bounds_t), where the tiled launch coordinate computation was preserved as-is to limit scope.