HIP-Clang内嵌组件

Question

什么是与这个CUDA功能相对应的Hip？

__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。该功能用于实现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是正确的指令集参考吗？

这是一个适用的AMDGPU后端的LLVM用户指南参考吗？

Answer 1

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

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

对于早期的体系结构，如__gfx900__，将vcc_lo替换为vcc

请参阅讨论关于Rocm Hip Github和这个AMD gpu组件参考。