3

我在 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.

这是预期的吗？

4

1 回答 1

2

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

从nvcc手册页：

--device-debug (-G)

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

于 2017-06-10T13:37:20.137 回答