choose-launch-configuration

Skill: Choose CUDA Launch Configuration

Purpose

Guide the agent through selecting the correct and efficient thread block dimensions and grid dimensions for a CUDA kernel, covering occupancy analysis, register and shared memory constraints, tail effects, persistent kernels, and when to use cudaOccupancyMaxActiveBlocksPerMultiprocessor as a decision tool.

Use this when

• Writing a new CUDA kernel and choosing the initial block and grid dimensions
• A kernel is underperforming and you suspect the launch configuration limits occupancy or causes tail effects
• Selecting block dimensions for a kernel with significant register or shared memory usage
• Implementing a persistent kernel or a cooperative grid launch
• Tuning a kernel where block size is a compile-time template parameter and you want to choose the best default

Do not use this when

• The kernel's performance is dominated by a compute or memory bottleneck unrelated to occupancy: fixing the launch config will not help if the kernel is register-bound or has severe memory access inefficiency
• The launch configuration is already determined by an API contract (e.g., cuBLAS, cuDNN, cooperative kernels with fixed grid requirements): do not override the library's choices
• The block size is constrained by the algorithm's data decomposition to a specific value (e.g., exactly 32 threads for a warp-level primitive): respect the algorithmic constraint and do not change it to improve occupancy

Inputs the agent should gather first

• Kernel register usage: from ptxas -v output, nvcc --ptxas-options=-v, or from Nsight Compute profiling. Registers per thread is the primary constraint on occupancy when register usage is high.
• Static shared memory per block: declared __shared__ arrays. Size in bytes.
• Dynamic shared memory per block: the third kernel launch argument. Size in bytes.
• Problem size: total number of work items (elements, rows, output tiles, etc.). Determines the minimum grid size needed to cover all work.
• Hardware target: SM architecture and SM count. Key occupancy limits differ by architecture. Also: the maximum threads per block (1024 on all current GPUs), the maximum blocks per SM (varies: 16–32 depending on SM gen), the register file size per SM (65536 32-bit registers on SM70+), the smem per SM.
• Whether the kernel is latency-sensitive: for latency-critical single-sample inference, occupancy may matter less than minimizing launch overhead and maximizing L2 utilization for a fixed small grid.
• Whether cooperative groups are needed: cooperative grid launches have additional constraints (entire grid must fit resident on the GPU).