将结构传递给内核时是否有任何性能下降？

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 将所有参数放入“常量内存”并通过经过“常量缓存”的专用内存加载函数访问它们。

将结构传递给内核时是否有任何性能下降？

Is there any performance downside when passing a struct to a kernel?

cuda