OpenCL 全局内存获取答案

【问题标题】：OpenCL global memory fetchesOpenCL 全局内存获取
【发布时间】：2012-03-30 16:44:07
【问题描述】：

我正在考虑重新设计我的 GPU OpenCL 内核以加快速度。问题是有很多未合并的全局内存，并且获取确实降低了性能。所以我打算将尽可能多的全局内存复制到本地，但我必须选择要复制的内容。

现在我的问题是：多次提取小块内存是否比提取更少大块更有害？

【问题讨论】：

【解决方案1】：

您可以使用 clGetDeviceInfo 找出设备的缓存线大小。 (clGetDeviceInfo, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE) 在当今的许多设备上，此值通常为 16 个字节。

小读取可能很麻烦，但如果您从同一个缓存行读取，应该没问题。简短的回答：你需要在内存中将你的“小块”放在一起以保持速度。

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

vectorAddBar 将执行得更快，因为读取和写入落在内存中的同一缓存行中。每 4 次浮点读取将落在同一个高速缓存行上，并且只从内存控制器执行一项操作。读完这件事的 a[] 和 b[] 后，四个工作项都可以做它们的加法，并将它们的写入排队到 c[]。

vectorAddFoo 将保证读取和写入不在同一缓存行中（非常短的向量 ~totalElements

__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...
}

【讨论】：

【解决方案2】：

一般来说，大尺寸的更少的 fecthe 会更有效。如果没有看到您的代码，我无法给您具体的建议，但请确保从工作项访问顺序块以启用“流式传输”。将数据带入本地内存后进行任何转置或随机内存访问。

【讨论】：

【解决方案3】：

我无法正确理解您的问题，但是如果您有大量的全局访问权限并且如果这些访问权限被重用，则使用本地内存。

注意：本地工作量小，共享数据少，所以没用，本地工作量大，并行线程少。所以你需要选择最好的。

【讨论】：