OpenCL全局内存获取

Question

我正在考虑重新编写我的GPU OpenCL内核以加快速度。问题在于有大量未对齐的全局内存，提取操作会降低性能。因此，我计划将尽可能多的全局内存复制到本地内存，但我必须选择要复制的内容。

现在我的问题是：许多小块内存的提取是否比较少的大块内存提取更具有伤害性？

- Nigel

Answer 1

您可以使用clGetDeviceInfo来查找设备的缓存行大小。(clGetDeviceInfo, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE) 在许多设备上，这个值通常是16字节。

小读取可能会有问题，但如果您从同一缓存行中读取，应该没问题。简短的答案是：您需要将“小块”紧密地放在内存中以保持速度。

我下面有两个函数来演示两种访问内存的方式- vectorAddFoo和vectorAddBar。第三个函数copySomeMemory(...)特别适用于您的问题。两个向量函数都有它们的工作项添加正在添加的向量的一部分，但使用不同的内存访问模式。vectorAddFoo让每个工作项处理一个向量元素块，从计算出的数组位置开始，并向前移动其工作负载。vectorAddBar使工作项从gid开始，在获取并添加下一个元素之前跳过gSize（=全局大小）个元素。

由于读写都在同一缓存行中，vectorAddBar 的执行速度更快。每 4 次浮点数读取将落在同一缓存行中，并且只需要从内存控制器执行一次操作。以这种方式读取 a[] 和 b[] 后，所有四个工作项都能够进行加法运算，并将其写入 c[]。

除非向量长度很短（总元素<5），否则 vectorAddFoo 将保证读写不在同一缓存行中。每个工作项的读取都需要从内存控制器执行操作。除非 GPU 在每种情况下缓存以下 3 个浮点数，否则这将导致 4 倍的内存访问。

__kernel void  
vectorAddFoo(__global const float * a,  
          __global const float * b,  
          __global       float * c,
          __global const totalElements) 
{ 
  int gid = get_global_id(0); 
  int elementsPerWorkItem = totalElements/get_global_size(0);
  int start = elementsPerWorkItem * gid;

  for(int i=0;i<elementsPerWorkItem;i++){
    c[start+i] = a[start+i] + b[start+i]; 
  }
} 
__kernel void  
vectorAddBar(__global const float * a,  
          __global const float * b,  
          __global       float * c,
          __global const totalElements) 
{ 
  int gid = get_global_id(0); 
  int gSize = get_global_size(0);

  for(int i=gid;i<totalElements;i+=gSize){
    c[i] = a[i] + b[i]; 
  }
} 
__kernel void  
copySomeMemory(__global const int * src,
          __global const count,
          __global const position) 
{ 
  //copy 16kb of integers to local memory, starting at 'position'
  int start = position + get_local_id(0); 
  int lSize = get_local_size(0);
  __local dst[4096];
  for(int i=0;i<4096;i+=lSize ){
    dst[start+i] = src[start+i]; 
  }
  barrier(CLK_GLOBAL_MEM_FENCE);
  //use dst here...
} 

- mfa

Answer 2

通常来说，较大尺寸的少量读取会更有效率。我无法在看到您的代码之前给出具体建议，但请确保从工作项中访问连续块以实现“流式处理”。在将数据带入本地内存后进行任何转置或随机内存访问操作。

- Lubo Antonov

Answer 3

我无法完全理解你的问题，但如果你有大量的全局访问并且这些被重复使用，那么请使用本地内存。

注意：小的本地工作大小共享的数据较少，因此没有用处，而大的本地工作大小则会减少并行线程。因此，你需要选择最佳方案。

- Megharaj