OpenCL的代价get_local_id（）

Question

我有一个简单的扫描内核，它可以计算循环中几个块的扫描。我注意到，当get_local_id（）存储在局部变量中而不是在循环内调用时，性能有所提高。因此，为了与代码总结，这样的：

__kernel void LocalScan_v0(__global const int *p_array, int n_array_size, __global int *p_scan) 
{ 
    const int n_group_offset = get_group_id(0) * SCAN_BLOCK_SIZE; 
    p_array += n_group_offset; 
    p_scan += n_group_offset; 
    // calculate group offset 

    const int li = get_local_id(0); // *** local id cached *** 
    const int gn = get_num_groups(0); 
    __local int p_workspace[SCAN_BLOCK_SIZE]; 
    for(int i = n_group_offset; i < n_array_size; i += SCAN_BLOCK_SIZE * gn) { 
     LocalScan_SingleBlock(p_array, p_scan, p_workspace, li); 

     p_array += SCAN_BLOCK_SIZE * gn; 
     p_scan += SCAN_BLOCK_SIZE * gn; 
    } 
    // process all the blocks in the array (each block size SCAN_BLOCK_SIZE) 
} 

有吞吐量74 GB/s的GTX-780，而这一点：

__kernel void LocalScan_v0(__global const int *p_array, int n_array_size, __global int *p_scan) 
{ 
    const int n_group_offset = get_group_id(0) * SCAN_BLOCK_SIZE; 
    p_array += n_group_offset; 
    p_scan += n_group_offset; 
    // calculate group offset 

    const int gn = get_num_groups(0); 
    __local int p_workspace[SCAN_BLOCK_SIZE]; 
    for(int i = n_group_offset; i < n_array_size; i += SCAN_BLOCK_SIZE * gn) { 
     LocalScan_SingleBlock(p_array, p_scan, p_workspace, get_local_id(0)); 
     // *** local id polled inside the loop *** 

     p_array += SCAN_BLOCK_SIZE * gn; 
     p_scan += SCAN_BLOCK_SIZE * gn; 
    } 
    // process all the blocks in the array (each block size SCAN_BLOCK_SIZE) 
} 

在相同的硬件上只有70 GB /秒。唯一的区别是对get_local_id（）的调用是在循环内部还是外部。 LocalScan_SingleBlock（）中的代码在this GPU Gems article中有很多描述。

现在，这带来了一些问题。我一直认为线程标识存储在某个寄存器中，并且对任何线程局部变量的访问速度都很快。这似乎并非如此。我总是习惯于把一个本地ID缓存在一个变量中，而这个变量不愿意让一个老的“C”程序员不愿意在一个循环中调用一个函数，如果他希望每次都返回相同的值，不要认为这会有所作为。

任何想法，为什么这可能是？我没有做任何检查编译的二进制代码。有没有人有相同的经历？ CUDA中的threadIdx.x与此相同吗？ ATI平台如何？这种行为是在什么地方描述的？我通过CUDA最佳实践快速扫描，但没有发现任何内容。

Answer 1

这只是一种猜测，但按照Khronos的页面

http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/get_local_id.html

get_local_id（）没有定义为返回一个恒定值（仅仅为size_t）。这可能意味着，就编译器所知，与常量local_id相比，可能不允许执行某些优化，因为函数值的返回可能会在编译器的眼中发生变化（即使它不会按每个线程）