为什么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图形处理器。然而，我认为这种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个大小的块启动时，nvprof显示kernel_A平均运行15.2us，而kernel_B平均运行7.4us。我推测NVRTC没有优化出整数除法和模数。

结果是在一个GeForce 750 Ti，CUDA8.0，平均从100个电话。提供给nvrtcCompileProgram()的编译器选项是-arch compute_50。

这是意料之中吗？

Answer 1

在代码库中做了一个彻底的错误。原来我的应用程序是在DEBUG模式下构建的。这将导致传递给-G和-lineinfo的附加标志nvrtcCompileProgram()。

来自nvcc手册页：

--device-debug (-G) 为设备代码生成调试信息。关闭所有优化。不要使用分析，而是使用-lineinfo。