HIP-Clang 在线组装

Question

此 CUDA 函数的 Hip-Clang 等价物是什么？

__device__ __forceinline__ uint32_t add_cc(uint32_t a, uint32_t b)
{
   uint32_t r;
   asm volatile ("add.cc.u32 %0, %1, %2;" : "=r"(r) : "r"(a), "r"(b));
   return r;
}

我正在将一个 CUDA 项目移植到包含内联 PTX 程序集的 HIP-Clang。该函数用于在NVIDIA GPU中实现多精度加法。我试过了：

asm volatile ("add.cc.u32 %0, %1, %2;" : "=r"(r) : "r"(a), "r"(b)); //invalid instruction
asm volatile ("V_ADD_CO_U32 %0, %1, %2;" : "=r"(r) : "r"(a), "r"(b)); //invalid operand for instruction
asm volatile ("V_ADD_CO_U32 %0, %1, %2;" : "=v"(r) : "v"(a), "v"(b)); //operands are not valid for this GPU or mode

目标硬件是 RX 6800。AMD clang 版本 14.0.0。
RDNA2 是否正确 instruction set reference?
这是LLVM user guide to AMDGPU backend适用的参考吗？

Answer 1

事实证明答案取决于硬件。对于我的编译器定义了 __gfx1030__ 的硬件，正确的语法是

asm volatile ("v_add_co_u32 %0, vcc_lo, %1, %2;" : "=v"(r) : "v"(a), "v"(b));

对于早期的架构，例如 __gfx900__ 将 vcc_lo 替换为 vcc
请参阅讨论 on the Rocm Hip Github and this AMD gpu assembly reference。

HIP-Clang 在线组装

HIP-Clang Inline Assembly

c++

clang

inline-assembly

hip