如何使用 AMD HIP 中的内联 GCN 汇编将多个 Float4 从内存加载到寄存器？

Question

动机

我正在 AMD GPU 上进行一些微基准测试，以了解其性能特征，从而提高内核性能。我现在怀疑不同的寄存器分配和指令调度结果可能会影响有效内存带宽。我注意到编译器尝试交错内存指令和计算指令，它还尝试通过在前一个算术指令完成后立即加载新值来保存寄存器。在某些情况下，我发现可能存在显着的性能差异。如果我故意插入一些不可优化的操作，例如向 LDS 写入虚拟值，以阻止编译器进行此类交错，有时会提高性能。我怀疑原因是它改变了同时发出的内存请求的数量，导致使用的内存带宽减少。

因此，我决定在针对 AMD HIP 时使用内联汇编，以更好地控制微基准。

尝试1

以下 HIP 程序尝试使用内联汇编将 4 个浮点数加载到

float4 tmp11, tmp12, tmp13, tmp14

#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, %1, off\n\t"
        : "=v" (tmp11)
        : "v" (a_ptr)
    );
    asm volatile(
        "global_load_dwordx4 %0, %1, off, offset:16\n\t"
        : "=v" (tmp12)
        : "v" (a_ptr)
    );
    asm volatile(
        "global_load_dwordx4 %0, %1, off, offset:32\n\t"
        : "=v" (tmp13)
        : "v" (a_ptr)
    );
    asm volatile(
        "global_load_dwordx4 %0, %1, off, offset:48\n\t"
        "s_waitcnt vmcnt(0)"
        : "=v" (tmp14)
        : "v" (a_ptr)
    );
#endif
}

int main(void)
{
}

用

hipcc -S main.cpp -o main.S -O3

编译源代码并检查

main.S

，我发现生成的程序集不正确。这些值都加载到相同的寄存器中。

    ;;#ASMSTART
    global_load_dwordx4 v[0:3], v[4:5], off

    ;;#ASMEND
    ;;#ASMSTART
    global_load_dwordx4 v[0:3], v[4:5], off, offset:16

    ;;#ASMEND
    ;;#ASMSTART
    global_load_dwordx4 v[0:3], v[4:5], off, offset:32

    ;;#ASMEND
    ;;#ASMSTART
    global_load_dwordx4 v[0:3], v[4:5], off, offset:48
    s_waitcnt vmcnt(0)
    ;;#ASMEND

尝试2

编译器似乎确定加载没有影响，并且可以自由地为所有汇编指令重用相同的寄存器，这是有道理的。如果我对变量进行一些算术运算，这些值确实会被加载到不同的寄存器中，这可能是由于不同的寄存器分配所致。

因此，我的下一个尝试是使用多个指令和一条内联汇编语句将多个值加载到多个寄存器中。编译器现在知道结果应该进入不同的寄存器。

#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
}

int main(void)
{
}

不幸的是，生成的程序集仍然不正确。

    ;;#ASMSTART
    global_load_dwordx4 v[0:3],  v[0:1], off
    global_load_dwordx4 v[4:7],  v[0:1], off offset:16
    global_load_dwordx4 v[8:11],  v[0:1], off offset:32
    global_load_dwordx4 v[12:15],  v[0:1], off offset:48
    s_waitcnt vmcnt(0)
    ;;#ASMEND

第一个加载指令会破坏寄存器

v[0:1]

，因此所有后续加载将无法按预期工作。

问题

如何使用 AMD HIP 中的内联 GCN 汇编，使用正确的内联汇编语法将多个 float4 从内存加载到寄存器？

Answer 1

以相反顺序加载以防止寄存器损坏

#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 %3,  %4, off offset:48\n\t"
        "global_load_dwordx4 %2,  %4, off offset:32\n\t"
        "global_load_dwordx4 %1,  %4, off offset:16\n\t"
        "global_load_dwordx4 %0,  %4, off\n\t"
        "s_waitcnt vmcnt(0)"
        : "=v" (tmp11), "=v" (tmp12), "=v" (tmp13), "=v" (tmp14)
        : "v" (a_ptr)
    );
#endif
}

int main(void)
{
}