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

【问题标题】：Is it possible to put assembly instructions into CUDA code?是否可以将汇编指令放入 CUDA 代码中？
【发布时间】：2011-04-10 06:50:30
【问题描述】：

我想在 CUDA C 代码中使用汇编代码为了减少昂贵的处决就像我们在 c 编程中使用 asm 一样。

有可能吗？

【问题讨论】：

与 SU 相关：superuser.com/questions/668019/… 对于 Intel 的 Xeon Phi，它看起来是可能的。
How to create or manipulate GPU assembler?的可能重复

标签： c assembly cuda inline-assembly ptx

【解决方案1】：

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

有关示例，请参阅 this。

但对于 GPU，组装优化不是必需的，您应该先进行其他优化，例如内存合并和占用。请参阅CUDA Best Practices guide 了解更多信息。

【讨论】：

第二个！根据我的经验，CUDA 程序几乎总是受内存限制，而不是计算限制。
感谢以上两位。我只是想减少除法和模运算的数量，但现在我将重点关注内存问题。
注意，如果您正在针对最新的架构进行编译（使用标志 -arch sm_20），那么最新的 API 现在完全了吗？符合 IEEE 浮点除法和平方根规范。如果您有很多部门并且还使用 -arch sm_20，那么您可能会考虑使用以下标志切换回“较少”兼容版本以获得性能提升：-prec-div=false forums.nvidia.com/lofiversion/index.php?t170749.html
建议不接受这个答案并接受 njuffa 的，因为由于新功能，时间已经使这个答案变得不那么有用了。

【解决方案2】：

自 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;
}

【讨论】：