为什么 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;
    /* ... */
}

我知道Integer Division和Modulo在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平均显示 15.2US 的kernel_A运行，而kernel_B则在 7.4US> 7.4US 平均运行。我推测NVRTC没有优化整数和模型。

在100个呼叫中平均的GeForce 750 Ti，Cuda 8.0获得了结果。给出nvrtcCompileProgram()的编译器选项是-arch compute_50。

这是预期的吗？

Answer 1

在代码库中进行了彻底的错误。事实证明，我的应用程序是在DEBUG模式下构建的。这导致其他标志-G和-lineinfo传递给nvrtcCompileProgram()

来自nvcc男人页面：

--device-debug (-G)

为设备代码生成调试信息。关闭所有优化。 不要用于分析；使用-lineinfo。