pragma unroll

#pragma unroll 是 CUDA 编程中的一个编译器指令（Compiler Directive），用于 强制展开循环（Loop Unrolling），从而优化 GPU 内核（Kernel）的性能。

1. 作用

在 GPU 编程中，循环展开（Loop Unrolling）是一种常见的优化技术，通过 减少循环控制开销（如分支判断、计数器更新）来提高指令级并行性（ILP）。
#pragma unroll 告诉 CUDA 编译器：

• 强制展开指定循环（即使编译器默认不展开）。
• 减少动态分支（Branch Divergence），提高 Warp（线程束）的执行效率。

2. 基本语法

#pragma unroll [unroll_factor]
for (int i = 0; i < N; i++) {
    // 循环体
}

• unroll_factor（可选）：指定展开的倍数（如 #pragma unroll 4 表示每次迭代展开为 4 个重复代码块）。
  • 如果省略，编译器会 完全展开循环（前提是循环次数 N 是编译期常量）。
• 适用场景：仅适用于 for 循环，且循环边界在编译时可确定（常量或 constexpr）。

3. 实际示例

示例1：完全展开（Full Unroll）

// 循环次数 N=5 是编译期常量
#pragma unroll
for (int i = 0; i < 5; i++) {
    output[i] = input[i] * 2;
}

编译后的等效代码（手动展开）：

output[0] = input[0] * 2;
output[1] = input[1] * 2;
output[2] = input[2] * 2;
output[3] = input[3] * 2;
output[4] = input[4] * 2;

示例2：部分展开（Partial Unroll）

// 每次展开 2 次迭代
#pragma unroll 2
for (int i = 0; i < 8; i++) {
    output[i] = input[i] + input[i+1];
}

编译后的等效代码：

for (int i = 0; i < 8; i += 2) {
    output[i]   = input[i]   + input[i+1];
    output[i+1] = input[i+1] + input[i+2];
}

4. 为什么需要 #pragma unroll？

GPU 的性能优化需求

1. 减少分支开销：

   • GPU 的 SIMT（单指令多线程）架构中，循环的 if 条件或计数器更新可能导致 Warp 内线程分支发散（Divergence），降低并行效率。
   • 展开后，每个线程直接执行连续指令，避免分支判断。

2. 提高指令级并行（ILP）：

   • 展开后的循环体可能被编译器调度为更密集的指令流水线。

3. 显式控制编译器行为：

   • CUDA 编译器默认会根据优化级别（如 -O3）自动决定是否展开循环，但开发者可能希望更精确控制。

5. 使用注意事项

必须满足的条件

1. 循环边界必须是编译期常量：

   int N = 5;  // 变量 -> 编译错误！
   #pragma unroll
   for (int i = 0; i < N; i++) {}  // 错误：N 不是常量
   

   修正方法：

   const int N = 5;  // 或 constexpr
   #pragma unroll
   for (int i = 0; i < N; i++) {}
   

2. 避免过度展开：

   • 完全展开大循环会导致代码膨胀（Code Bloat），可能降低缓存命中率。
   • 建议对小型循环（如 N <= 32）使用完全展开，大型循环用部分展开（如 #pragma unroll 4）。

与其他优化指令结合

• #pragma unroll_and_jam：嵌套循环的展开优化（CUDA 11+）。
• __restrict__ 关键字：配合使用可进一步避免指针别名分析问题。

6. 性能对比

假设一个简单的向量加法内核：

__global__ void add(float *x, float *y, float *z, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    #pragma unroll 4  // 部分展开
    for (int j = 0; j < 4; j++) {
        if (i + j < N) z[i+j] = x[i+j] + y[i+j];
    }
}

• 无 #pragma unroll：编译器可能生成带分支的循环代码，增加开销。
• 有 #pragma unroll 4：直接展开为 4 次独立操作，减少分支和计数器更新。

7. 常见问题

Q1：#pragma unroll 和 #pragma nounroll 的区别？

• #pragma unroll：强制展开循环。
• #pragma nounroll：禁止展开循环（即使编译器默认会展开）。

Q2：如何检查循环是否真的被展开？

• 使用 NVCC 的 -Xptxas -v 选项查看 PTX 汇编代码：

  nvcc -Xptxas -v -O3 kernel.cu
  

  输出中会显示 unrolled 标记的循环。

总结