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

问题描述

我定义了这样的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];

我要去哪里错了？

解决方法

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

如果您这样做：

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,%rd1;
        cvta.to.global.u64      %rd4,%rd2;
        mov.u32         %r1,%rd5;
        add.s64         %rd7,%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,%r5,%r4};
        ret;
}

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

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

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

{
        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位负载。虽然这对于许多人来说可能非常明显，但是我不是资深编码员。我只在某些地方使用编码来制定项目。这对我来说是非常有见地的见解，我希望有人也能从中受益！

cuda cuda gpu gpu ptx