c++ - 带 CUDA 内联汇编的 LLVM

Question

我正在尝试使用以下内联汇编编译 CUDA 代码:

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

代码可以通过带有标志 -Xptxas -v 的 nvcc 正常编译。

当我尝试用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 不兼容。关于这个page ，解释说 LLVM 对所有 PTX 变量都有不同的定义(smid 和 warpid 也有一些定义)。现在，如果上述代码必须单独(而不是内联)编写和编译，我会迷路。

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

Answer 1

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

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

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