OpenCL全局内存提取

我在考虑重新调整我的GPU OpenCL内核以加快速度。问题是有很多全球内存不合并，并且提取真的会降低性能。所以我打算将尽可能多的全球内存复制到本地，但我必须选择要复制的内容。OpenCL全局内存提取

现在我的问题是：做很多小块内存伤害更多更少的大块更多？

2012-03-28 Nigel

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

小读取可能会很麻烦，但如果您从同一个缓存行读取数据，则应该没问题。简短的回答：你需要保持你的“小块”紧密联系在一起，以保持速度。

我有两个函数来演示两种访问内存的方法 - vectorAddFoo和vectorAddBar。第三个函数copySomeMemory（...）具体适用于您的问题。两个矢量函数的工作项都添加了一部分要添加的矢量，但使用不同的内存访问模式。 vectorAddFoo获取每个工作项以处理一组矢量元素，从其在阵列中的计算位置开始，并向前移动其工作负载。 vectorAddBar的工作项目在他们的gid处开始，并且在获取和添加下一个元素之前跳过gSize（= global size）元素。

vectorAddBar将执行得更快，因为读写操作落入内存中的同一缓存行。每4个浮点读取将落在同一缓存行中，并且只从内存控制器执行一个动作。在阅读本文的[]和b []后，所有四个工作项都可以进行添加，并将写入队列排列为c []。

vectorAddFoo将保证读取和写入不在同一缓存行中（除非是很短的向量〜totalElements < 5）。每次从工作项目读取都需要内存控制器的操作。除非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... 
}

来源

2012-03-29 14:20:27 mfa

一般而言，较大尺寸的情感会更少效率更高。如果没有看到您的代码，我不能给您具体的建议，但要确保从工作项目访问连续的块，以启用“流式传输”。在将数据带入本地存储器之后，请执行任何转置或随机存储器访问。

来源

2012-03-29 10:28:42

我无法理解你的问题正确，但如果你有大的全球访问如果这些比使用重复使用使用本地内存。

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

来源

2012-03-30 08:44:07 Megharaj

OpenCL全局内存提取

回答

相关问题