是否可以将汇编指令放入CUDA代码中？

Question

我想在CUDA C代码中使用汇编代码,以减少昂贵的执行,因为我们在c编程中使用asm.

可能吗？

Answer 1

自CUDA 4.0以来,CUDA工具链支持内联PTX.工具箱中有一个描述它的文档:Using_Inline_PTX_Assembly_In_CUDA.pdf

下面是一些代码,演示了如何在CUDA 4.0中使用内联PTX.请注意,此代码不应该用作CUDA内置__clz()函数的替代,我只是编写它来探索新的内联PTX功能的各个方面.

__device__ __forceinline__ int my_clz (unsigned int x)
{
    int res;

    asm ("{\n"
         "        .reg .pred iszero, gezero;\n"
         "        .reg .u32 t1, t2;\n"
         "        mov.b32         t1, %1;\n"
         "        shr.u32         %0, t1, 16;\n"
         "        setp.eq.b32     iszero, %0, 0;\n"
         "        mov.b32         %0, 0;\n"
         "@iszero shl.b32         t1, t1, 16;\n"
         "@iszero or.b32          %0, %0, 16;\n"
         "        and.b32         t2, t1, 0xff000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 8;\n"
         "@iszero or.b32          %0, %0, 8;\n"
         "        and.b32         t2, t1, 0xf0000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 4;\n"
         "@iszero or.b32          %0, %0, 4;\n"
         "        and.b32         t2, t1, 0xc0000000;\n"
         "        setp.eq.b32     iszero, t2, 0;\n"
         "@iszero shl.b32         t1, t1, 2;\n"
         "@iszero or.b32          %0, %0, 2;\n"
         "        setp.ge.s32     gezero, t1, 0;\n"
         "        setp.eq.b32     iszero, t1, 0;\n"
         "@gezero or.b32          %0, %0, 1;\n"
         "@iszero add.u32         %0, %0, 1;\n\t"
         "}"
         : "=r"(res)
         : "r"(x));
    return res;
}

Answer 2

不，你不能，没有什么比 C/C++ 的 asm 构造更好的了。您可以做的是调整生成的 PTX 程序集，然后将其与 CUDA 一起使用。

请参阅此示例。

但对于 GPU 来说，组装优化不是必需的，您应该首先进行其他优化，例如内存合并和占用。有关更多信息，请参阅CUDA 最佳实践指南。