optimization - 为什么 NVRTC 不优化我的整数除法和模运算？

Question

我在 NVRTC 编译了一个内核:

__global__ void kernel_A(/* args */) {
    unsigned short idx = threadIdx.x;
    unsigned char warp_id = idx / 32;
    unsigned char lane_id = idx % 32;
    /* ... */
}

我知道整数除法和模运算在 CUDA GPU 上的成本非常高。但是我认为这种 2 的幂除法应该优化为位运算，直到我发现它不是:

__global__ void kernel_B(/* args */) {
    unsigned short idx = threadIdx.x;
    unsigned char warp_id = idx >> 5;
    unsigned char lane_id = idx & 31;
    /* ... */
}

kernel_B 似乎运行得更快。当忽略内核中的所有其他代码时，以 1024 个大小为 1024 的 block 启动时，nvprof 显示 kernel_A 平均运行 15.2us，而 kernel_B 平均运行 7.4us。我推测 NVRTC 没有优化整数除法和模数。

结果是在 GeForce 750 Ti、CUDA 8.0 上获得的，取 100 次调用的平均值。提供给 nvrtcCompileProgram() 的编译器选项是 -arch compute_50。

这是预期的吗？

Answer 1

在代码库中进行了彻底的错误扫描。原来我的应用程序是在 DEBUG 模式下构建的。这会导致额外的标志 -G 和 -lineinfo 传递给 nvrtcCompileProgram()

来自 nvcc 手册页:

--device-debug (-G)

Generate debug information for device code. Turns off all optimizations. Don't use for profiling; use -lineinfo instead.