<<<CUDA C++ 学习路线grid · block · warp · lane
返回首页
占用率

Occupancy 计算器

选择架构并设置 block 尺寸、每线程寄存器数、每 block 共享内存,算出理论占用率并指出瓶颈资源。

配置

架构

sm_80 · Ampere A100

256
32

用 nvcc -Xptxas -v 查看真实值

8 KB
理论占用率
100.0%
驻留 block
8
64 / 64 warp
瓶颈资源:线程数 / block 数上限 与 寄存器. 已经是当前资源下的最优 blockDim。

三个限制因素各自允许多少 block

线程数 / block 数上限 ← 瓶颈8 block
min(32, floor(64 / 8))
寄存器 ← 瓶颈8 block
floor(65536 / 8192) 每 block 用 8192 个
共享内存20 block
floor(164KB / 8.00KB)

占用率随 blockDim 的变化

寄存器 32 · 共享内存 8KB

锯齿状说明资源分配存在粒度:多几个线程可能就跨过一个台阶。

322565127681024
对应的 kernel 约束
1// Target: blockDim = 256, at least 8 resident blocks per SM2__global__ void __launch_bounds__(256, 8)3myKernel(const float* __restrict__ in, float* out, int n) {4    __shared__ float tile[2048];   // 8 KB5    ...6}7 8// Confirm real register count at compile time:9//   nvcc -O3 -arch=sm_80 -Xptxas -v -c kernel.cu10// Let CUDA recommend blockDim at runtime:11//   cudaOccupancyMaxPotentialBlockSize(&minGrid, &blockSize, myKernel, 0, 0);
这里算的是理论占用率,模型与 CUDA Occupancy Calculator 一致(含寄存器按 warp 粒度、共享内存按分配单元对齐)。实际运行时的 Achieved Occupancy 通常更低,因为存在尾部效应和负载不均。用 Nsight Compute 测量才是最终依据。