get_*
这可能是一个愚蠢的问题,但是:在 OpenCL 内核中调用某些函数有多昂贵?将结果保存以供将来在某些本地变量中使用还是在需要时调用所需的函数更好?
还是依赖于平台?
PS我认为,cuda使用各种threadIdx变量可以更好地解决它。
我认为这应该对所有 GPU 架构都是免费的。它应该被相应的硬件寄存器或高速缓存组中的常量替换。
编译器也可以对其进行持续传播。您可以使用 AMD Stream Analyzer检查自己:
__kernel
void testKernel(__global uint * uintArray)
{
uint threadId = get_global_id(0);
uintArray[threadId] = 0xbaadf00d;
}
0 ALU: ADDR(32) CNT(10) KCACHE0(CB0:0-15) KCACHE1(CB1:0-15)
0 x: MOV R1.x, (0xBAADF00D, -0.001327039325f).x
t: MULLO_INT ____, R1.x, KC0[1].x
1 x: ADD_INT ____, R0.x, PS0
2 w: ADD_INT ____, PV1.x, KC0[6].x
3 z: LSHL ____, PV2.w, (0x00000002, 2.802596929e-45f).x
4 y: ADD_INT ____, KC1[0].x, PV3.z
5 x: LSHR R0.x, PV4.y, (0x00000002, 2.802596929e-45f).x
01 MEM_RAT_CACHELESS_STORE_RAW: RAT(1)[R0].x___, R1, VPM
这里get_global_id(0)
映射到常量缓存组值KC0[1].x
。所以,为了回答你的问题,我会使用最易读的形式。