有没有办法将 128 位从内存直接加载到寄存器？

Posted 2023-03-16

【中文标题】有没有办法将 128 位从内存直接加载到寄存器？

【英文标题】：Is there a way to load 128 bits from memory directly to registers?

【发布时间】：2021-11-26 13:19:36

【问题描述】：

我知道 CUDA 中存在 32、64 和 128 位的全局内存负载，但大多数示例使用宽负载（即 LD.E.128）从内存复制到内存。

我有多个类似于这个的结构数组：

typedef struct 
    uint32_t a,
    uint32_t b,
    uint32_t c,
    uint32_t d
 mystruct_t;

每个 CUDA 线程都必须加载 a、b、c 和 d。使用它们进行操作，并将结果存储在另一个 mystruct_t 类型的数组中。

独立加载结构的四个元素会导致性能不佳，因为内存访问没有合并。这可以通过使用数组结构而不是结构数组模式来解决，但是当数组很大时（如我的情况），L2 缓存性能很低，因为内存局部性很小。

是否可以将 128 位从全局内存直接加载到寄存器以对其进行操作？我尝试过类似的方法：

int4 vector_data = *(int4*)(&myarray[threadIdx.x]);
uint32_t a = vector_data.x;
uint32_t b = vector_data.y;
uint32_t c = vector_data.z;
uint32_t d = vector_data.w;
// Do computations with a, b, c, and d

但我仍然得到单独的 32 位未合并负载。我不知道是否可以将 128 位存入寄存器，因为大多数 GPU 中的寄存器都是 32 位的。

【问题讨论】：

inform the compiler 您的结构已正确对齐以进行 128 位加载。

【参考方案1】：

如 cmets 中所述，GPU 编译器只会在保证与加载和存储大小对齐的类型上生成向量加载和存储指令。

在您的问题中考虑以下对 sn-ps 的粗略改编：

#include <cinttypes>

struct mystruct

    uint32_t a, b, c, d;
;

struct alignas(16) mystructa

    uint32_t a, b, c, d;
;

template<typename T>
__global__ void kernel(T* in, uint32_t* out)

   T ival = in[threadIdx.x];
   out[threadIdx.x] = (ival.d - ival.c) * (ival.a + ival.b);


template __global__ void kernel<mystruct>(mystruct*, uint32_t*);
template __global__ void kernel<mystructa>(mystructa*, uint32_t*);

这里我们有一个简单的内核，它使用对齐或未对齐的 128 位类型进行实例化。编译器（本例中为 nvcc 11.3.1）发出以下 PTX：

.visible .entry _Z6kernelI8mystructEvPT_Pj(
        .param .u64 _Z6kernelI8mystructEvPT_Pj_param_0,
        .param .u64 _Z6kernelI8mystructEvPT_Pj_param_1
)


        ld.param.u64    %rd1, [_Z6kernelI8mystructEvPT_Pj_param_0];
        ld.param.u64    %rd2, [_Z6kernelI8mystructEvPT_Pj_param_1];
        cvta.to.global.u64      %rd3, %rd2;
        cvta.to.global.u64      %rd4, %rd1;
        mov.u32         %r1, %tid.x;
        mul.wide.u32    %rd5, %r1, 16;
        add.s64         %rd6, %rd4, %rd5;
        ld.global.u32   %r2, [%rd6];
        ld.global.u32   %r3, [%rd6+4];
        ld.global.u32   %r4, [%rd6+8];
        ld.global.u32   %r5, [%rd6+12];
        sub.s32         %r6, %r5, %r4;
        add.s32         %r7, %r3, %r2;
        mul.lo.s32      %r8, %r6, %r7;
        mul.wide.u32    %rd7, %r1, 4;
        add.s64         %rd8, %rd3, %rd7;
        st.global.u32   [%rd8], %r8;
        ret;


.visible .entry _Z6kernelI9mystructaEvPT_Pj(
        .param .u64 _Z6kernelI9mystructaEvPT_Pj_param_0,
        .param .u64 _Z6kernelI9mystructaEvPT_Pj_param_1
)


        ld.param.u64    %rd1, [_Z6kernelI9mystructaEvPT_Pj_param_0];
        ld.param.u64    %rd2, [_Z6kernelI9mystructaEvPT_Pj_param_1];
        cvta.to.global.u64      %rd3, %rd2;
        cvta.to.global.u64      %rd4, %rd1;
        mov.u32         %r1, %tid.x;
        mul.wide.u32    %rd5, %r1, 16;
        add.s64         %rd6, %rd4, %rd5;
        ld.global.v4.u32        %r2, %r3, %r4, %r5, [%rd6];
        sub.s32         %r10, %r5, %r4;
        add.s32         %r11, %r2, %r3;
        mul.lo.s32      %r12, %r10, %r11;
        mul.wide.u32    %rd7, %r1, 4;
        add.s64         %rd8, %rd3, %rd7;
        st.global.u32   [%rd8], %r12;
        ret;

您可以看到，在第一个内核实例中，提供了您问题中的类型，四个单独的ld.global.u32 用于将结构带入寄存器，而在第二个内核实例中，提供对齐类型，单个改为使用ld.global.v4.u32。