【发布时间】：2013-09-26 12:12:15
【问题描述】：
我有一个简单的扫描内核，它在一个循环中计算几个块的扫描。我注意到当 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 最佳实践，但没有找到任何东西。
【问题讨论】：
-
请不要删除 CUDA 标签。虽然代码本身不在 CUDA 中，但问题在 NVIDIA 硬件上有所体现，并且与 CUDA 的 threadIdx 的实现方式以及它如何影响程序的 runitme 密切相关。