带有 CUDA 内联汇编的 LLVM

Question

我正在尝试使用以下内联程序集编译 CUDA 代码：

static __device__ uint get_smid(void) {
    uint ret;
    asm("mov.u32 %0, %smid;" : "=r"(ret) );
    return ret;
}

nvcc带有标志的代码编译得很好-Xptxas -v。

当我尝试使用clang++（版本 4.0）和相应的标志-Xcuda-ptxas -v（我认为这是正确的，但我可能错了）编译它时，出现以下错误：

../../include/cutil_subset.h:23:25: error: invalid % escape in inline assembly string asm("mov.u32 %0, %smid;" : "=r"(ret) );

它指向%smid.

我想我应该链接正确的库，但我也有这个：L/cuda/install/lib.

另一种可能是 NVPTX asm 不兼容。在这个页面上，解释了 LLVM 对所有 PTX 变量都有不同的定义（smid 和 warpid 也有一些）。现在，如果提到的代码必须单独（不是内联）编写和编译，我会迷失方向。

有没有人以前处理过类似的问题？欢迎提出建议。

Answer 1

您需要使用双百分号引用特殊寄存器：%%smid。

该%%转义序列被编译过程中转换成燎百分号，使ptxas认为正确的特殊寄存器名称。双百分号版本也适用于 nvcc。

nvcc内联汇编程序中的转义序列似乎比现在更宽容clang++，并且保持未知的转义序列不变，而不是像在这种情况下那样发出错误。