将结构传递给内核时，是否存在性能方面的缺点

Question

我有一个内核，它接受几个数组作为输入。为了提高可读性，最好将它们分组到一个结构中，并(在为每个输入分配适当的内存和复制之后(将该结构传递给内核，而不是指针的长列表。

在访问内核内部的数组时，这两种情况在内存方面是否相同？

有人能给我推荐一些关于这个主题的文档吗(在编程指南上找不到(

Answer 1

不，应该没有区别。您可以读取PTX输出以确定。这里有一个简单的例子：

struct Foo
{
int* a, *b, *c;
};
__global__ void bar(Foo f)
{ f.a[0] = f.b[0] + f.c[0]; }
__global__ void baz(int* a, int* b, int* c)
{ a[0] = b[0] + c[0]; }
struct Quz
{
int* a, *b, *c;
~Quz() {}
};
__global__ void quuz(Quz f)
{ f.a[0] = f.b[0] + f.c[0]; }

这是PTX组件。请注意，函数之间基本上没有区别。

.visible .entry _Z3bar3Foo(
.param .align 8 .b8 _Z3bar3Foo_param_0[24]
)
{
.reg .b32   %r<4>;
.reg .b64   %rd<7>;
ld.param.u64    %rd1, [_Z3bar3Foo_param_0+16];
ld.param.u64    %rd2, [_Z3bar3Foo_param_0+8];
ld.param.u64    %rd3, [_Z3bar3Foo_param_0];
cvta.to.global.u64  %rd4, %rd3;
cvta.to.global.u64  %rd5, %rd2;
cvta.to.global.u64  %rd6, %rd1;
ld.global.u32   %r1, [%rd5];
ld.global.u32   %r2, [%rd6];
add.s32     %r3, %r2, %r1;
st.global.u32   [%rd4], %r3;
ret;
}
.visible .entry _Z3bazPiS_S_(
.param .u64 _Z3bazPiS_S__param_0,
.param .u64 _Z3bazPiS_S__param_1,
.param .u64 _Z3bazPiS_S__param_2
)
{
.reg .b32   %r<4>;
.reg .b64   %rd<7>;

ld.param.u64    %rd1, [_Z3bazPiS_S__param_0];
ld.param.u64    %rd2, [_Z3bazPiS_S__param_1];
ld.param.u64    %rd3, [_Z3bazPiS_S__param_2];
cvta.to.global.u64  %rd4, %rd1;
cvta.to.global.u64  %rd5, %rd3;
cvta.to.global.u64  %rd6, %rd2;
ld.global.u32   %r1, [%rd6];
ld.global.u32   %r2, [%rd5];
add.s32     %r3, %r2, %r1;
st.global.u32   [%rd4], %r3;
ret;
}
.visible .entry _Z4quuz3Quz(
.param .align 8 .b8 _Z4quuz3Quz_param_0[24]
)
{
.reg .b32   %r<4>;
.reg .b64   %rd<7>;

ld.param.u64    %rd1, [_Z4quuz3Quz_param_0+16];
ld.param.u64    %rd2, [_Z4quuz3Quz_param_0+8];
ld.param.u64    %rd3, [_Z4quuz3Quz_param_0];
cvta.to.global.u64  %rd4, %rd3;
cvta.to.global.u64  %rd5, %rd2;
cvta.to.global.u64  %rd6, %rd1;
ld.global.u32   %r1, [%rd5];
ld.global.u32   %r2, [%rd6];
add.s32     %r3, %r2, %r1;
st.global.u32   [%rd4], %r3;
ret;
}

这一切都是一样的，因为CUDA将所有参数放入"；恒定存储器"；并通过专门的存储器加载功能访问它们；常量缓存">