CUDA中每翘曲关系指令的数据大小
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结构中的数组。因为当我们在内核中调用复制时，它实际上扩展并复制了a数组中的floatArray结构中的每个元素，从而创建了更多的指令。
是否有一种方法可以使用单个指令在内核中复制自定义结构变量？
回答 1
Stack Overflow用户
回答已采纳
发布于 2016-06-16 14:48:55
您正确的假设是，当您更改数组的大小时，复制指令的数量会增加。您可以在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检查程序集。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/37861601
复制相关文章
相似问题
腾讯云开发者
