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

Question

我试图在数据类型的大小更改

时尝试查看内核中执行的指令数

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

#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];
}

然后仅使用一个单个块

调用上述内核

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结构内的数组。因为当我们在内核中调用副本时，它实际上会扩展并复制floatArray结构内数组a的每个元素，从而创建更多指令。

是否有一种方法可以使用单个指令在内核中复制自定义结构变量？

Answer 1

您在假设更改数组大小时的复制指令数量会增加。您可以按照我将在下面显示的PTX代码和组装中检查此内容。

负载/商店指令的最大尺寸为128位，例如这里。这意味着对于您的情况，您仍然可以使用float4而不是float来提高4倍。

另外，您可以按照编程指南中的说明明确指定数据结构的对齐方式：

#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