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

Posted 2023-03-16

【中文标题】将结构传递给内核时是不是有任何性能下降？

【英文标题】：Is there any performance downside when passing a struct to a kernel?将结构传递给内核时是否有任何性能下降？

【发布时间】：2021-12-17 01:00:28

【问题描述】：

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

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

谁能推荐我一些关于这个主题的文档（在编程指南上找不到）

【问题讨论】：

我认为将内核参数作为结构传递是很常见的，我个人认为没有任何缺点......

this 可能感兴趣

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