cuda - OpenCL get_local_id() 的成本

Question

我有一个简单的扫描内核，它计算循环中几个 block 的扫描。我注意到当 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)
}

在 GTX-780 上的吞吐量为 74 GB/s，同时:

__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/s。唯一的区别是对 get_local_id() 的调用是在循环内部还是外部。 LocalScan_SingleBlock() 中的代码在 this GPU Gems article 中有详细描述。 .

现在这带来了一些问题。我总是想象线程 id 存储在某个寄存器中，并且访问它与访问任何线程局部变量一样快。似乎并非如此。我总是习惯于将本地 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 相比，它可能不允许执行某些优化，因为函数值的返回值在编译器看来可能会发生变化(即使它不会针对每个线程)