为什么 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 的块启动，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.

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

Why isn't NVRTC optimizing out my integer division and modulo operations?

optimization

cuda

compiler-flags

nvrtc