带有 CUDA 内联汇编的 LLVM

Posted 2023-02-16

【中文标题】带有 CUDA 内联汇编的 LLVM

【英文标题】：LLVM with CUDA inline assembly

【发布时间】：2016-10-16 19:01:43

【问题描述】：

我正在尝试使用以下内联汇编编译 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）。现在，如果必须单独（而不是内联）编写和编译所提到的代码，我会迷失方向。

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

【问题讨论】：

我没有设置 atm 的工具链来验证，但我怀疑您需要使用双百分号引用特殊寄存器：%%smid。该转义序列在编译期间被转换为单个百分号，以便 ptxas 看到正确的特殊寄存器名称。双百分号版本也适用于 nvcc。这似乎是一个 nvcc 扩展，它使未知的转义序列保持不变而不是发出错误。

哦，是的，我很抱歉我没有更新这个！你说的对！！我几天前解决了它，完全忘记了我之前发布的这个问题。无论如何，谢谢你的回复。请将此添加为答案，我会接受。

【参考方案1】：

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

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

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