多次调用 get_global_id() vs 将结果保存在局部变量中？答案

【问题标题】：Call multiple times get_global_id() vs save the result in the local variable?多次调用 get_global_id() vs 将结果保存在局部变量中？
【发布时间】：2011-03-16 10:28:18
【问题描述】：

这可能是一个愚蠢的问题，但是：在 OpenCL 内核中调用一些 get_* 函数有多贵？将结果保存在某些局部变量中以供将来使用，还是在需要时调用所需的函数更好？

还是依赖于平台？

PS 我认为，cuda 使用各种 threadIdx 变量可以更好地解决它。

【问题讨论】：

标签： cuda opencl

【解决方案1】：

我认为这应该对所有 GPU 架构都是免费的。它应该被相应的硬件寄存器或缓存组中的常量替换。

编译器也可以对其进行持续传播。您可以使用 AMD Stream Analyser 检查自己：

OpenCL：

__kernel 
void testKernel(__global uint * uintArray)
{
    uint threadId = get_global_id(0);

    uintArray[threadId] = 0xbaadf00d;
}

Radeon HD 5870 (Cypress) 组件：

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。所以，为了回答你的问题，我会使用最易读的形式。

【讨论】：