带 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.

我想我应该 link 合适的图书馆，但我也有这个：L/cuda/install/lib.

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

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

Answer 1

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

%% 转义序列在编译期间被转换为单个百分号，以便 ptxas 看到正确的特殊寄存器名称。双百分号版本在 nvcc 下也有效。

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

LLVM with CUDA inline assembly