CUDA 中每个 Warp 关系的数据大小到指令

Posted 2023-02-16

【中文标题】CUDA 中每个 Warp 关系的数据大小到指令

【英文标题】：Data Size to Instructions per Warp relationship in CUDA

【发布时间】：2016-06-16 14:06:47

【问题描述】：

我试图查看当数据类型大小发生变化时内核中执行的指令数

为了获得自定义大小的数据结构，我创建了如下结构，

#define DATABYTES 40

__host__ __device__
struct floatArray

    float a[DATABYTES/4];
;

然后创建一个内核只是为了将上述数据类型数组从一个数组复制到另一个数组

__global__
void copy_large_data(floatArray * d_in, floatArray * d_out)

    d_out[threadIdx.x] = d_in[threadIdx.x];

然后用一个块只调用上面的内核 32 个线程

floatArray * d_in;
floatArray * d_out;

cudaMalloc(&d_in, 32 * sizeof(floatArray));
cudaMalloc(&d_out, 32 * sizeof(floatArray));

copy_large_data<<<1, 32>>>(d_in, d_out);

当我使用nvprof 分析程序并检查instructions per warp 时，我可以看到参数值随着DATABYTES 值的变化而变化。

我的问题是，这个指令计数增加的原因是否是由于floatArray 结构内的数组。因为当我们在内核中调用copy的时候，实际上是在a结构体内部展开并复制了数组a的每个元素，创建了更多的指令。

有没有办法使用一条指令在内核中复制自定义结构变量？

【参考方案1】：

您的假设是正确的，即当您更改数组的大小时，复制指令的数量会增加。您可以在 PTX 代码和汇编中检查它，如下所示。

加载/存储指令的最大大小为 128 位，参见例如here。这意味着对于您的情况，您仍然可以通过使用 float4 而不是 float 来提高 4 倍。

或者，您可以明确指定数据结构的对齐方式，如programming guide 中所述：

#define DATABYTES 32
struct __align__(16) floatArray

    float a[DATABYTES/4];
;

---

要查看 PTX 代码，请生成目标文件 nvcc -c ... 并使用 cubobjdump --dump-ptx objfile.o。 对于您的示例，相关部分如下所示：

ld.global.f32 %f1, [%rd7];
ld.global.f32 %f2, [%rd7+4];
ld.global.f32 %f3, [%rd7+8];
ld.global.f32 %f4, [%rd7+12];
ld.global.f32 %f5, [%rd7+16];
ld.global.f32 %f6, [%rd7+20];
ld.global.f32 %f7, [%rd7+24];
ld.global.f32 %f8, [%rd7+28];
ld.global.f32 %f9, [%rd7+32];
ld.global.f32 %f10, [%rd7+36];
st.global.f32 [%rd6+36], %f10;
st.global.f32 [%rd6+32], %f9;
st.global.f32 [%rd6+28], %f8;
st.global.f32 [%rd6+24], %f7;
st.global.f32 [%rd6+20], %f6;
st.global.f32 [%rd6+16], %f5;
st.global.f32 [%rd6+12], %f4;
st.global.f32 [%rd6+8], %f3;
st.global.f32 [%rd6+4], %f2;
st.global.f32 [%rd6], %f1;

如果你进一步增加数组，你会发现编译器会选择循环而不是为每个加载/存储发出指令。

因此，您可以使用cubobjdump --dump-sass objfile.o 来检查程序集

【讨论】：

那你说的是我们单条指令可以传输的数据最大数据大小是128位？

是的，但请记住，warp 中的事务可以合并为 128 字节事务。

你的意思是在使用 L1 缓存时对吗？因为我认为如果我们只使用 L2 缓存，事务大小会减少到 32 个字节

感谢您提供有关cuobjdump 的信息。我在使用float4 时检查了说明，并且能够观察到您提到的内容，其中只给出了一条指令来加载所有 x,y,z,w 变量ld.global.v4.f32 %f1, %f2, %f3, %f4, [%rd7];

如果我创建一个包含 4 个浮点变量的自定义结构，ptx 代码显示发出 4 个单独的指令来加载完整的结构。