将结构传递给内核时是不是有任何性能下降?
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 将所有参数放入“常量内存”并通过经过“常量缓存”的专用内存加载函数访问它们。
【讨论】:
以上是关于将结构传递给内核时是不是有任何性能下降?的主要内容,如果未能解决你的问题,请参考以下文章
如何使用 PyOpenCL 将带有数组和变量的 C 结构传递给 OpenCL 内核