【CUDA学习笔记】__launch_bounds__

语法

• __global__ void __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) kernel(...) {// 内核代码
  }

• maxThreadsPerBlock
  指定每个线程块（Block）可包含的最大线程数。这是硬性限制，如果启动内核时超过此值，会触发运行时错误。

• minBlocksPerMultiprocessor（可选）
  提示编译器确保每个流式多处理器（SM）至少驻留的线程块数量，以优化资源分配。默认值为 0（由编译器自动决定）。

作用

1. 控制寄存器使用

   • 通过限制每个线程块的线程数，编译器可以调整寄存器的分配，从而减少寄存器溢出（Spill）到本地内存的情况，提高性能。
   • 如果寄存器使用过多，编译器可能减少每个线程的寄存器数量，通过牺牲部分性能来支持更多的并行线程。

2. 优化线程块调度

   • 帮助编译器更好地分配线程块到 SM，提高 SM 的利用率（Occupancy）。
   • 显式指定 minBlocksPerMultiprocessor 可以确保足够的线程块驻留在 SM 上，隐藏延迟。

Demo

// 限制每个 Block 最多 256 线程，并提示每个 SM 至少驻留 4 个 Blocks
__global__ void __launch_bounds__(256, 4) myKernel(float* data) {int idx = blockIdx.x * blockDim.x + threadIdx.x;data[idx] *= 2.0f;
}int main() {float* d_data;cudaMalloc(&d_data, 1024 * sizeof(float));// 启动内核：必须满足 blockDim.x <= 256myKernel<<<1024 / 256, 256>>>(d_data);cudaFree(d_data);return 0;
}

• --ptxas-options=-v
  编译时添加此选项（如 nvcc -Xptxas -v），可输出寄存器和共享内存的使用信息。