为什么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只是运行得更快。当在kernel中省略所有其他代码，以1024个大小为1024的块启动时，nvprof显示平均kernel_A运行15.2us，而kernel_B平均运行7.4us。我猜测NVRTC没有优化掉整数除法和模数。

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

这符合预期吗？

- Kh40tiK

---

不需要猜测。通过 cuobjdump -sass 运行可执行文件即可找到答案。 - tera

@tera 我正在通过NVRTC进行JIT编译，因此没有可执行文件。有什么方法可以获取汇编转储吗？ - Kh40tiK

NVRTC会生成PTX代码并传递给驱动程序进行即时编译。您可以使用nvrtcGetPTX函数提取PTX代码。这样就不需要猜测了。 - talonmies

@talonmies 谢谢你的建议，我会尽力深入调查。 - Kh40tiK

不是针对你的问题的答案，但实际上你不需要费心去处理这两个问题；warp和lane id是免费提供的。 - einpoklum

我认为这两种情况之间不太可能有任何差异。当使用nvcc编译到SASS时，每种情况生成的代码（包括PTX和SASS）都是相同的。我认为通过nvcc源->PTX再通过JIT将PTX->SASS不太可能有任何不同。由于在这个问题中没有给出完整的重现案例，我的猜测是差异可能存在于其他地方。 - Robert Crovella

Answer 1

对代码库进行了彻底的漏洞扫描，结果发现我的应用程序是在 DEBUG 模式下构建的。这会导致向 nvrtcCompileProgram() 传递附加标志 -G 和 -lineinfo

从 nvcc 手册中获取：

--device-debug (-G)

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

- Kh40tiK