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

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

可能吗？

Answer 1

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

请参阅此示例。

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

Answer 2

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

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

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