CUDA 编程：Occupancy、Block Size 与 Kernel 启动配置调优

Michael阿明

发布于 2026-06-29 13:09:04

960

CUDA 第 9 课：Occupancy、Block Size 与 Kernel 启动配置调优

前面我们已经看到：TILE、blockDim、shared memory、bank conflict 都会影响性能。

第 9 课进一步解决一个核心问题：

❝一个 kernel 到底应该配置多少 threads/block？block 太小、太大为什么都可能慢？Occupancy 到底怎么看？

一、本节课目标

本节课重点掌握：

1. 理解 occupancy 是什么
2. 理解 block size 如何影响 occupancy
3. 理解 occupancy 高不等于一定最快
4. 学会用 CUDA API 估算理论 occupancy
5. 通过实验测试不同 block size 下的 kernel 时间

二、核心原理

1. 什么是 Occupancy？

Occupancy 可以理解为：

❝一个 SM 上实际活跃 warp 数量，占这个 SM 理论最大 warp 数量的比例。

例如：

一个 SM 理论最多可以同时驻留 64 个 warp
当前 kernel 实际只能驻留 32 个 warp

Occupancy = 32 / 64 = 50%

GPU 为什么需要较高 occupancy？

因为 GPU 访问 Global Memory 有延迟。当一个 warp 在等内存时，SM 可以切换去执行另一个 warp：

warp A 等内存
↓
SM 执行 warp B
↓
warp B 等内存
↓
SM 执行 warp C

所以 occupancy 高，通常更容易隐藏延迟。

2. Block Size 为什么影响 Occupancy？

每个 block 会占用 SM 资源，包括：

threads
warps
registers
shared memory
block slot

例如：

block size = 64   -> 每个 block 2 个 warp
block size = 256  -> 每个 block 8 个 warp
block size = 1024 -> 每个 block 32 个 warp

block 太小：

每个 block 线程少
block 数量多
调度开销可能较大
单个 block 内并行度不足

block 太大：

一个 block 占用太多线程/寄存器/shared memory
一个 SM 上能同时放的 block 数减少
可能降低调度灵活性

所以 block size 不是越大越好，常见起点是：

128、256、512

3. Occupancy 高一定最快吗？

不一定。

Occupancy 只是说明：

SM 上有多少 warp 可供调度

但性能还取决于：

Global memory 访问是否 coalesced
shared memory 是否有 bank conflict
寄存器是否溢出
指令吞吐是否成为瓶颈
L1/L2 cache 命中情况
kernel 计算密度

所以本课要建立的核心认识是：

❝Occupancy 是重要指标，但不是唯一指标。最终仍然要用实验时间验证。

在这里插入图片描述

三、实验设计

写一个 compute kernel，让它对数组做多轮 fmaf 计算：

v = fmaf(v, 1.000001f, 0.000001f);

然后测试不同 block size：

观察：

1. 理论 occupancy 变化
2. kernel time 变化
3. GFLOPS 变化
4. 最快 block size 是否等于最高 occupancy

四、完整可运行 CUDA C++ 代码

保存为：

lesson09_occupancy_block_size.cu

#include <cuda_runtime.h>

#include <cmath>
#include <cstdlib>
#include <iomanip>
#include <iostream>
#include <string>
#include <vector>

#define CUDA_CHECK(call)                                                        \
    do {                                                                        \
        cudaError_t err = call;                                                 \
        if (err != cudaSuccess) {                                               \
            std::cerr << "CUDA Error: " << cudaGetErrorString(err)              \
                      << " at " << __FILE__ << ":" << __LINE__ << std::endl;    \
            std::exit(EXIT_FAILURE);                                            \
        }                                                                       \
    } while (0)

/*
 * 一个简单的计算型 kernel。
 *
 * 每个线程处理一个元素。
 * iters 控制每个元素做多少次 fmaf。
 *
 * fmaf(a, b, c) 大致表示：
 * a * b + c
 *
 * 通常可以粗略按 2 FLOPs 估算。
 */
__global__ void compute_kernel(const float* in,
                               float* out,
                               int n,
                               int iters) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx < n) {
        float v = in[idx];

        for (int i = 0; i < iters; ++i) {
            v = fmaf(v, 1.000001f, 0.000001f);
        }

        out[idx] = v;
    }
}

/*
 * 测量 kernel 平均执行时间。
 *
 * 注意：
 * 这里测的是 kernel time only。
 * 不包含 H2D、D2H、cudaMalloc、CPU 初始化。
 */
float time_kernel(const float* d_in,
                  float* d_out,
                  int n,
                  int iters,
                  int block_size,
                  int repeat) {
    int grid_size = (n + block_size - 1) / block_size;

    /*
     * warmup。
     */
    compute_kernel<<<grid_size, block_size>>>(d_in, d_out, n, iters);
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());

    cudaEvent_t start, stop;
    CUDA_CHECK(cudaEventCreate(&start));
    CUDA_CHECK(cudaEventCreate(&stop));

    float total_ms = 0.0f;

    for (int r = 0; r < repeat; ++r) {
        CUDA_CHECK(cudaEventRecord(start));

        compute_kernel<<<grid_size, block_size>>>(d_in, d_out, n, iters);
        CUDA_CHECK(cudaGetLastError());

        CUDA_CHECK(cudaEventRecord(stop));
        CUDA_CHECK(cudaEventSynchronize(stop));

        float ms = 0.0f;
        CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
        total_ms += ms;
    }

    CUDA_CHECK(cudaEventDestroy(start));
    CUDA_CHECK(cudaEventDestroy(stop));

    return total_ms / repeat;
}

/*
 * 使用 CUDA Runtime API 估算理论 occupancy。
 */
double estimate_occupancy_percent(int block_size,
                                  const cudaDeviceProp& prop) {
    int active_blocks_per_sm = 0;

    CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &active_blocks_per_sm,
        compute_kernel,
        block_size,
        0
    ));

    int active_warps = active_blocks_per_sm * block_size / prop.warpSize;
    int max_warps = prop.maxThreadsPerMultiProcessor / prop.warpSize;

    return 100.0 * static_cast<double>(active_warps) /
           static_cast<double>(max_warps);
}

int main(int argc, char** argv) {
    int n = 1 << 26;       // 67,108,864 floats，约 256 MB
    int iters = 256;       // 每个元素计算 256 次
    int repeat = 10;

    if (argc >= 2) {
        n = std::atoi(argv[1]);
    }
    if (argc >= 3) {
        iters = std::atoi(argv[2]);
    }
    if (argc >= 4) {
        repeat = std::atoi(argv[3]);
    }

    cudaDeviceProp prop;
    CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));

    std::cout << "CUDA Lesson 9: Occupancy and Block Size\n";
    std::cout << "GPU name                  : " << prop.name << "\n";
    std::cout << "SM count                  : " << prop.multiProcessorCount << "\n";
    std::cout << "Warp size                 : " << prop.warpSize << "\n";
    std::cout << "Max threads per block     : " << prop.maxThreadsPerBlock << "\n";
    std::cout << "Max threads per SM        : " << prop.maxThreadsPerMultiProcessor << "\n";
    std::cout << "Shared memory per block   : "
              << prop.sharedMemPerBlock / 1024 << " KB\n";

    size_t bytes = static_cast<size_t>(n) * sizeof(float);

    std::cout << "\nProblem size\n";
    std::cout << "Elements                  : " << n << "\n";
    std::cout << "Array size                : "
              << bytes / 1024.0 / 1024.0 << " MB\n";
    std::cout << "Iters per element         : " << iters << "\n";
    std::cout << "Repeat                    : " << repeat << "\n";

    std::vector<float> h_in(n);
    std::vector<float> h_out(n);

    for (int i = 0; i < n; ++i) {
        h_in[i] = static_cast<float>((i % 100) + 1) * 0.001f;
    }

    float* d_in = nullptr;
    float* d_out = nullptr;

    CUDA_CHECK(cudaMalloc(&d_in, bytes));
    CUDA_CHECK(cudaMalloc(&d_out, bytes));

    CUDA_CHECK(cudaMemcpy(d_in, h_in.data(), bytes, cudaMemcpyHostToDevice));

    std::vector<int> block_sizes = {64, 128, 256, 512, 1024};

    std::cout << "\n"
              << std::left
              << std::setw(12) << "block"
              << std::setw(14) << "grid"
              << std::setw(18) << "occupancy(%)"
              << std::setw(14) << "time(ms)"
              << std::setw(14) << "GFLOPS"
              << "status\n";

    for (int block_size : block_sizes) {
        if (block_size > prop.maxThreadsPerBlock) {
            std::cout << std::left
                      << std::setw(12) << block_size
                      << std::setw(14) << "-"
                      << std::setw(18) << "-"
                      << std::setw(14) << "-"
                      << std::setw(14) << "-"
                      << "SKIP: block_size > maxThreadsPerBlock\n";
            continue;
        }

        int grid_size = (n + block_size - 1) / block_size;

        double occupancy = estimate_occupancy_percent(block_size, prop);

        float ms = time_kernel(d_in, d_out, n, iters, block_size, repeat);

        /*
         * 每次 fmaf 粗略按 2 FLOPs 估算。
         */
        double flops = 2.0 * static_cast<double>(n) * static_cast<double>(iters);
        double gflops = flops / (ms / 1000.0) / 1e9;

        std::cout << std::fixed << std::setprecision(3)
                  << std::left
                  << std::setw(12) << block_size
                  << std::setw(14) << grid_size
                  << std::setw(18) << occupancy
                  << std::setw(14) << ms
                  << std::setw(14) << gflops
                  << "OK\n";
    }

    CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));

    /*
     * 简单检查输出是否是有限数。
     */
    bool ok = true;
    for (int i = 0; i < 10; ++i) {
        if (!std::isfinite(h_out[i])) {
            ok = false;
            break;
        }
    }

    std::cout << "\nCheck output finite: " << (ok ? "PASS" : "FAIL") << "\n";

    CUDA_CHECK(cudaFree(d_in));
    CUDA_CHECK(cudaFree(d_out));

    return ok ? 0 : 1;
}

五、编译与运行

Tesla T4：

nvcc -O3 -arch=sm_75 lesson09_occupancy_block_size.cu -o lesson09_occupancy

运行默认实验：

./lesson09_occupancy

也可以指定参数：

./lesson09_occupancy 67108864 256 10

参数含义：

第 1 个参数：元素数量 n
第 2 个参数：每个元素 fmaf 迭代次数 iters
第 3 个参数：重复次数 repeat

六、输出现象

输出结果：

CUDA Lesson 9: Occupancy and Block Size
GPU name                  : Tesla T4
SM count                  : 40
Warp size                 : 32
Max threads per block     : 1024
Max threads per SM        : 1024
Shared memory per block   : 48 KB

Problem size
Elements                  : 67108864
Array size                : 256 MB
Iters per element         : 256
Repeat                    : 10

block       grid          occupancy(%)      time(ms)      GFLOPS        status
64          1048576       100.000           11.243        3055.998      OK
128         524288        100.000           10.850        3166.698      OK
256         262144        100.000           6.131         5604.707      OK
512         131072        100.000           6.715         5116.723      OK
1024        65536         100.000           6.729         5105.893      OK

Check output finite: PASS

七、分析实验结果

1. block=64 为什么可能不够快？

block=64 -> 每个 block 只有 2 个 warp

虽然可以启动很多 block，但每个 block 太小，调度粒度较细，单个 block 内可组织的工作较少。

可能导致：

block 数量过多
调度开销增加
SM 内资源组织效率不一定最好

2. block=256 为什么常常表现较好？

block=256 -> 每个 block 8 个 warp

这是很多 CUDA kernel 的常见默认选择。

它通常在几个方面比较均衡：

线程数足够多
warp 数适中
block 数量不会过多
SM 上通常能驻留多个 block
调度灵活性较好

之前很多实验默认使用 256，是合理的。

3. block=1024 为什么不一定最快？

block=1024 -> 每个 block 32 个 warp

它虽然线程很多，但一个 block 太大，可能导致：

一个 SM 上同时驻留 block 数减少
调度灵活性下降
寄存器资源压力增大
某些情况下 occupancy 或有效并发下降

所以：

❝最大 block size 不等于最高性能。

4. 为什么 occupancy 一样，性能还不同？

因为 occupancy 只说明：

SM 上可以驻留多少 warp

但还没有说明：

warp 执行效率
访存效率
寄存器使用
指令调度
cache 行为
内存合并访问

所以都显示 100% occupancy，运行时间仍然可能不同。

八、和前面课程的关系

你测试矩阵乘法时发现：

TILE=8、16、32 会影响性能

这不仅是 shared memory tile 大小的问题，也包括：

block = TILE × TILE

带来的线程块规模变化。

例如：

TILE=8  -> block = 64 threads
TILE=16 -> block = 256 threads
TILE=32 -> block = 1024 threads

所以第 9 课可以帮你解释：

为什么 TILE 改变时，naive kernel 也会变快或变慢

因为 block size 变了。

核心结论：

1. block size 会影响 kernel 性能
2. occupancy 表示 SM 上活跃 warp 的比例
3. occupancy 有助于隐藏延迟，但不是性能的唯一决定因素
4. block 太小可能调度效率低
5. block 太大可能降低调度灵活性
6. 256 threads/block 是常见但不是绝对最优的起点
7. 最终性能必须通过实验验证

一句话总结：