当多个线程同时读取全局内存中的某个位置时,有多少个事务?
How many transactions are there when many threads read the some location in the global memory at the same time?
我正在编写一个内核,其中将启动所有线程并同时读取全局内存中的相同位置。
我在想这个时候会发生什么?
全局内存中的值是否广播到所有线程或是否存在
任何类型的序列化?
我知道当一半的 warp 线程访问适合全局内存的同一段时,访问将合并为更少的事务。
但是当读取完全相同的位置时会发生什么?
如果是广播,是不是就不需要再优化了?
我这里用的是AMD 7470。我写了一个迭代,其中所有线程在每次迭代中都会在全局内存中读取相同的 4 个字节。
理想情况下,您希望每个工作组尽可能少地读取全局内存位置。这意味着您应该将其复制到本地内存中,进一步的读取应该针对本地缓存的数据。
我发现这更常用于大于 4 字节的数据块,但只要您避免多次全局读取,应该可以节省时间。
local int sharedInt;
int id = get_local_id(0);
if(id == 0){
sharedInt = globalVar;
}
barrier(); //or barrier(CLK_LOCAL_MEM_FENCE) if the compiler complains
现在您可以保证每个工作组仅读取一次全局数据。
此方法适用于所有设备类型。我相信它会在 CPU 上强制使用最低缓存级别,但这将取决于实现。
大多数 GPU 将相同位置的读取广播到所有工作项。优化指南中专门提到了它。
我正在编写一个内核,其中将启动所有线程并同时读取全局内存中的相同位置。 我在想这个时候会发生什么? 全局内存中的值是否广播到所有线程或是否存在 任何类型的序列化?
我知道当一半的 warp 线程访问适合全局内存的同一段时,访问将合并为更少的事务。 但是当读取完全相同的位置时会发生什么? 如果是广播,是不是就不需要再优化了?
我这里用的是AMD 7470。我写了一个迭代,其中所有线程在每次迭代中都会在全局内存中读取相同的 4 个字节。
理想情况下,您希望每个工作组尽可能少地读取全局内存位置。这意味着您应该将其复制到本地内存中,进一步的读取应该针对本地缓存的数据。
我发现这更常用于大于 4 字节的数据块,但只要您避免多次全局读取,应该可以节省时间。
local int sharedInt;
int id = get_local_id(0);
if(id == 0){
sharedInt = globalVar;
}
barrier(); //or barrier(CLK_LOCAL_MEM_FENCE) if the compiler complains
现在您可以保证每个工作组仅读取一次全局数据。
此方法适用于所有设备类型。我相信它会在 CPU 上强制使用最低缓存级别,但这将取决于实现。
大多数 GPU 将相同位置的读取广播到所有工作项。优化指南中专门提到了它。