时间戳

首页 > 解决方案 > 为什么 PTX 显示 128 位结构分配的 32 位加载操作？

cuda - 为什么 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];

我哪里错了？

标签： cudagpuptx

解决方案

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

如果你这样做：

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 编译的所有代码]

推荐阅读