是的,你的内存访问模式几乎是最优的。每个半warp正在访问16个连续的32位字。此外,访问是64字节对齐的,因为缓冲区本身是对齐的,每个半warp的起始索引是16的倍数。因此,每个半warp将生成一个64字节的事务。因此,您不应通过未合并的访问浪费内存带宽。
既然你在上一个问题中询问了示例,让我们修改此代码以进行其他(不太优化的)访问模式(由于循环实际上没有做任何事情,我将忽略它):
kernel void vecAdd(global int* a, global int* b, global int* c)
{
int gid = get_global_id(0);
a[gid+1] = b[gid * 2] + c[gid * 32];
}
首先让我们看看这在1.3(GT200)硬件上是如何工作的。
对于写入a,这将生成一个稍微不太优化的模式(按其ID范围识别半线程束和相应的访问模式):
gid | addr. offset | accesses | reasoning
0- 15 | 4- 67 | 1x128B | in aligned 128byte block
16- 31 | 68-131 | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
32- 47 | 132-195 | 1x128B | in aligned 128byte block
48- 63 | 196-256 | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
基本上,我们浪费了大约一半的带宽(奇数半warp的访问宽度不到两倍并没有帮助太多,因为它会生成更多访问,这不如浪费更多字节更快)。
对于从b读取的数据,线程仅访问数组的偶数元素,因此每个半warp的所有访问都位于128字节对齐块中(第一个元素位于128B边界处,因为对于该元素,gid是16的倍数 => 索引是32的倍数,对于4字节的元素,这意味着地址偏移量是128B的倍数)。访问模式延伸到整个128B块,因此每个半warp将执行128B传输,再次浪费一半的带宽。
从c中读取则产生最坏的情况之一,其中每个线程在自己的128B块中索引,因此每个线程需要其自己的传输,这一方面是一种序列化场景(虽然不像正常情况那么糟糕,因为硬件应该能够重叠传输)。更糟糕的是,这将为每个线程传输32B块,浪费7/8的带宽(我们每个线程访问4B,32B / 4B = 8,因此只有1/8的带宽被利用)。由于这是朴素矩阵转置的访问模式,强烈建议使用本地内存来执行转置(根据经验说的)。
Compute 1.0(G80)
在这里,唯一会创建良好访问的模式是原始模式,在示例中的所有模式都会创建完全未集合的访问,浪费7/8的带宽(每个线程的32B传输,详见上文)。对于G80硬件,每个半warp中第n个线程不访问第n个元素的每个访问都会创建此类未集合的访问。
Compute 2.0(Fermi)
在这里,对内存的每次访问都创建128B事务(尽可能多地收集所有数据,因此在最坏情况下为16x128B),但这些缓存,使得数据将要传输到何处不太明显。暂时假设缓存足够大以容纳所有数据且没有冲突,因此最多每个128B缓存行将传输一次。进一步假设半warp的执行是串行化的,因此我们有一个确定性的缓存占用。
对于访问b,仍将始终传输128B块(没有其他线程索引相应的内存区域)。访问c将为每个线程生成128B传输(最坏的访问模式)。
对于a的访问是这样的(目前将其视为读取):
gid | offset | accesses | reasoning
0- 15 | 4- 67 | 1x128B | bringing 128B block to cache
16- 31 | 68-131 | 1x128B | offsets 68-127 already in cache, bring 128B for 128-131 to cache
32- 47 | 132-195 | - | block already in cache from last halfwarp
48- 63 | 196-259 | 1x128B | offsets 196-255 already in cache, bringing in 256-383
因此,对于大型数组,对a的访问在理论上几乎不会浪费带宽。
对于这个例子,现实当然不是那么好,因为对c的访问会很好地破坏缓存。
对于分析器,我会假设超过1.0的效率仅仅是浮点不准确性的结果。
希望这可以帮到你。