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

Posted

技术标签:

【中文标题】将结构传递给内核时是不是有任何性能下降?【英文标题】: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 将所有参数放入“常量内存”并通过经过“常量缓存”的专用内存加载函数访问它们。

【讨论】:

以上是关于将结构传递给内核时是不是有任何性能下降?的主要内容,如果未能解决你的问题,请参考以下文章

将结构传递给cupy中的原始内核

如何将结构传递到 CUDA 设备?

如何使用 PyOpenCL 将带有数组和变量的 C 结构传递给 OpenCL 内核

C:为啥你可以通过值传递(给函数)一个结构,而不是一个数组?

CLR 调试体系结构

性能下降返回光标与联合所有