CUDA error: cooperative launch of kernel exceeds device occupancy (cudaErrorCooperativeLaunchTooLarge)
ID: cuda/cooperative-launch-too-large
Version Compatibility
| Version | Status | Introduced | Deprecated | Notes |
|---|---|---|---|---|
| 12 | active | — | — | — |
Root Cause
A cooperative kernel launch (cudaLaunchCooperativeKernel) requires all thread blocks to be resident simultaneously on the GPU. The requested grid size exceeds the number of blocks the device can run concurrently given the kernel's resource usage (registers, shared memory). Unlike regular launches, cooperative kernels cannot oversubscribe the GPU.
genericWorkarounds
-
88% success Reduce block count and block size to fit within device occupancy limits
Use cudaOccupancyMaxActiveBlocksPerMultiprocessor to query the max blocks per SM for your kernel, then multiply by the number of SMs (cudaDeviceProp.multiProcessorCount). Set grid size to at most this value. Reduce shared memory or register usage to increase occupancy if needed.
-
85% success Split the algorithm into multiple non-cooperative kernel launches with explicit synchronization
Replace the single cooperative kernel with multiple regular kernel launches separated by cudaDeviceSynchronize() or stream synchronization. Each kernel handles a phase of the computation. This avoids the occupancy constraint at the cost of launch overhead and requiring intermediate buffers for inter-phase communication.
Dead Ends
Common approaches that don't work:
-
Increasing GPU memory to allow more concurrent blocks
90% fail
Cooperative launch occupancy is limited by SM count, registers per SM, and shared memory per SM — not global memory. Adding more VRAM does not increase the number of SMs or change per-SM resource limits.
-
Switching to a regular kernel launch without redesigning the algorithm
80% fail
Cooperative kernels use grid-wide synchronization (cooperative_groups::grid_group::sync). Removing cooperative launch without removing grid sync calls causes undefined behavior or hangs. The algorithm must be restructured to not require global synchronization.