为什么 PTX 显示 128 位结构分配的 32 位加载操作？

Question

我像这样定义了 128 位的自定义结构-

typedef struct dtype{
int val;
int temp2;
int temp3;
int temp4;
}dtype;

然后我执行了一项任务：-

dtype temp= h_a[i]; //where h_a is dtype *

我期待 128 位加载，但 PTX 显示的内容看起来像是 32 位加载操作-

mul.wide.s32    %rd4, %r18, 16;
add.s64         %rd5, %rd1, %rd4;
ld.global.u32   %r17, [%rd5];

不应该看起来像ld.global.v4.u32 %r17, [%rd5];

我哪里错了？

Answer 1

如果内存保证与类型的大小对齐，并且使用了该类型的所有元素，编译器只会发出向量化加载或存储指令（否则向量指令将被优化为标量指令节省带宽）。

如果你这样做：

struct dtype{
int val;
int temp2;
int temp3;
int temp4;
};

struct __align__ (16) adtype{
int val;
int temp2;
int temp3;
int temp4;
};

__global__
void kernel(adtype* x, dtype* y)
{
    adtype lx = x[threadIdx.x];
    dtype ly;
    ly.val = lx.temp4;
    ly.temp2 = lx.temp3;
    ly.temp3 = lx.val;
    ly.temp4 = lx.temp2;

    y[threadIdx.x] = ly;
}

你应该得到这样的东西：

visible .entry _Z6kernelP6adtypeP5dtype(
        .param .u64 _Z6kernelP6adtypeP5dtype_param_0,
        .param .u64 _Z6kernelP6adtypeP5dtype_param_1
)
{

        ld.param.u64    %rd1, [_Z6kernelP6adtypeP5dtype_param_0];
        ld.param.u64    %rd2, [_Z6kernelP6adtypeP5dtype_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];
        add.s64         %rd7, %rd3, %rd5;
        st.global.u32   [%rd7], %r5;
        st.global.u32   [%rd7+4], %r4;
        st.global.u32   [%rd7+8], %r2;
        st.global.u32   [%rd7+12], %r3;
        ret;
}

在这里您可以清楚地看到对齐类型的矢量化加载，以及非对齐类型的非矢量化存储。如果内核被更改，以便存储到对齐的版本：

__global__
void kernel(adtype* x, dtype* y)
{
    dtype ly = y[threadIdx.x];
    adtype lx;
    lx.val = ly.temp4;
    lx.temp2 = ly.temp3;
    lx.temp3 = ly.val;
    lx.temp4 = ly.temp2;

    x[threadIdx.x] = lx;
}

你会得到这个：

.visible .entry _Z6kernelP6adtypeP5dtype(
        .param .u64 _Z6kernelP6adtypeP5dtype_param_0,
        .param .u64 _Z6kernelP6adtypeP5dtype_param_1
)
{

        ld.param.u64    %rd1, [_Z6kernelP6adtypeP5dtype_param_0];
        ld.param.u64    %rd2, [_Z6kernelP6adtypeP5dtype_param_1];
        cvta.to.global.u64      %rd3, %rd1;
        cvta.to.global.u64      %rd4, %rd2;
        mov.u32         %r1, %tid.x;
        mul.wide.u32    %rd5, %r1, 16;
        add.s64         %rd6, %rd4, %rd5;
        add.s64         %rd7, %rd3, %rd5;
        ld.global.u32   %r2, [%rd6+12];
        ld.global.u32   %r3, [%rd6+8];
        ld.global.u32   %r4, [%rd6+4];
        ld.global.u32   %r5, [%rd6];
        st.global.v4.u32        [%rd7], {%r2, %r3, %r5, %r4};
        ret;
}

现在对齐类型与向量化指令一起存储。

[使用默认 Godbolt 工具链 (10.2) 为 sm_53 编译的所有代码]

Answer 2

我要补充一点，以防有人碰巧遇到同样的问题。

{
        dtype temp = h_a[i];                  //Loading data  exactly needed

        sum.val += temp.val;
}

我按照上述^^答案中给出的步骤进行操作，尽管上述方法绝对正确，但我没有得到 128 位负载。

问题是编译器发现在结构的 4 个字段中，我在一些加法运算中只使用了 1 个字段。所以它非常聪明地只加载了我需要的块。所以无论我尝试什么，我总是得到一个 32 位的负载。

{
        dtype temp = h_a[i];                  //Loading data  exactly needed

        sum.val += temp.val;
        sum.temp2 += temp.temp2;
        sum.temp3 += temp.temp3;
        sum.temp4 += temp.temp4;
}

有点变化。 现在我正在使用所有字段。所以编译器加载了所有字段！ 是的，现在使用上面 ^^ 答案中给出的方法，使用 __align __(16) 我得到了正确的 128 位负载。 虽然这对很多人来说可能很明显，但我不是一个资深的程序员。我只在某些地方使用编码来完成我的项目。这对我来说非常有见地，我希望有人也能从中受益！