我有一个简单的扫描内核，它计算循环中几个块的扫描。我注意到当 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() 中的代码在这篇 GPU Gems 文章中有详细描述。
现在这带来了一些问题。我一直认为线程 id 存储在某个寄存器中，并且访问它的速度与访问任何线程局部变量一样快。情况似乎并非如此。我一直习惯于将本地 id 缓存在变量中，而老“C”程序员不愿意在循环中调用函数，如果他希望它每次都返回相同的值，但我没有t认真地认为它会有所作为。
关于为什么会这样的任何想法？我没有对编译的二进制代码进行任何检查。有没有人有同样的经历？threadIdx.x在CUDA中是否相同？ATI 平台怎么样？这种行为是在某处描述的吗？我快速浏览了 CUDA 最佳实践，但没有找到任何东西。