如何在AMD HIP中用内联GCN汇编加载多个float4到寄存器？

解决方案 #

要解决这两个核心问题（编译器复用寄存器、输入指针寄存器被覆盖），需要从约束修饰符和防止变量优化两个角度入手：

1. 使用早期clobber修饰符避免输入寄存器被覆盖

GCC风格的内联汇编中，&修饰符表示早期clobber，它会告诉编译器：对应的输出寄存器会在所有输入操作数被使用之前就被修改，因此不能将输入寄存器分配给该输出寄存器。

在尝试2的代码中，给每个输出约束添加&，即可避免输入指针的寄存器被第一个加载指令覆盖：

asm volatile(
    "global_load_dwordx4 %0,  %4, off\n\t"
    "global_load_dwordx4 %1,  %4, off offset:16\n\t"
    "global_load_dwordx4 %2,  %4, off offset:32\n\t"
    "global_load_dwordx4 %3,  %4, off offset:48\n\t"
    "s_waitcnt vmcnt(0)"
    : "=&v" (tmp11), "=&v" (tmp12), "=&v" (tmp13), "=&v" (tmp14) // 添加&标记早期clobber
    : "v" (a_ptr)
);

2. 防止编译器优化未使用的变量

编译器会自动优化掉未被使用的变量，导致寄存器复用。解决方式有两种：

• 实际使用变量：将加载后的变量写入输出参数，比如在汇编块后添加：

  out[0] = tmp11;
  out[1] = tmp12;
  out[2] = tmp13;
  out[3] = tmp14;
  

• 标记变量为已使用：用__attribute__((used))标记tmp变量，强制编译器保留它们：

  float4 tmp11 __attribute__((used)), tmp12 __attribute__((used)), tmp13 __attribute__((used)), tmp14 __attribute__((used));
  

完整修正代码

#include <hip/hip_runtime.h>
#include <cstddef>

__global__ void kernel(
    float* __restrict array,
    float4* out,
    uint32_t idx
)
{
    float* a_ptr = &array[idx];
    float4 tmp11, tmp12, tmp13, tmp14;

#ifdef __HIP_PLATFORM_AMD__
    asm volatile(
        "global_load_dwordx4 %0,  %4, off\n\t"
        "global_load_dwordx4 %1,  %4, off offset:16\n\t"
        "global_load_dwordx4 %2,  %4, off offset:32\n\t"
        "global_load_dwordx4 %3,  %4, off offset:48\n\t"
        "s_waitcnt vmcnt(0)"
        : "=&v" (tmp11), "=&v" (tmp12), "=&v" (tmp13), "=&v" (tmp14)
        : "v" (a_ptr)
    );
#endif
    // 实际使用变量，防止优化
    out[0] = tmp11;
    out[1] = tmp12;
    out[2] = tmp13;
    out[3] = tmp14;
}

int main(void)
{
}

验证效果

编译后查看汇编代码，会发现输入指针的寄存器（比如v[4:5]）和所有输出寄存器（比如v[0:3]、v[8:11]等）完全不重叠，加载指令不会覆盖指针寄存器，且每个float4都会被加载到独立的寄存器组中。

内容的提问来源于stack exchange，提问作者比尔盖子