OpenCL中LDS内存的性能问题

我在使用AMD Radeon HD 6850使用LDS内存时遇到性能问题。OpenCL中LDS内存的性能问题

我有两个内核作为N粒子模拟的一部分。每个工作单元必须根据与其他粒子的相对位置计算作用于相应粒子的力。该问题的核心是：

#define UNROLL_FACTOR 8 
//Vernet velocity part kernel 
__kernel void kernel_velocity(const float deltaTime, 
          __global const float4 *pos, 
          __global float4 *vel, 
          __global float4 *accel, 
          __local float4 *pblock, 
          const float bound) 
{ 
    const int gid = get_global_id(0); //global id of work item 
    const int id = get_local_id(0); //local id of work item within work group 

    const int s_wg = get_local_size(0); //work group size 
    const int n_wg = get_num_groups(0); //number of work groups 

    const float4 myPos = pos[gid]; 
    const float4 myVel = vel[gid]; 
    const float4 dt = (float4)(deltaTime, deltaTime, 0.0f, 0.0f); 
    float4 acc = (float4)0.0f; 

    for (int jw = 0; jw < n_wg; ++jw) 
    { 
     pblock[id] = pos[jw * s_wg + id]; //cache a particle position; position in array: workgroup no. * size of workgroup + local id 
     barrier (CLK_LOCAL_MEM_FENCE); //wait for others in the work group 

     for (int i = 0; i < s_wg;) 
     { 
      #pragma unroll UNROLL_FACTOR 
      for (int j = 0; j < UNROLL_FACTOR; ++j, ++i) 
      { 
       float4 r = myPos - pblock[i]; 

       float rSizeSquareInv = native_recip (r.x*r.x + r.y*r.y + 0.0001f); 
       float rSizeSquareInvDouble = rSizeSquareInv * rSizeSquareInv; 
       float rSizeSquareInvQuadr = rSizeSquareInvDouble * rSizeSquareInvDouble; 
       float rSizeSquareInvHept = rSizeSquareInvQuadr * rSizeSquareInvDouble * rSizeSquareInv; 

       acc += r * (2.0f * rSizeSquareInvHept - rSizeSquareInvQuadr); 
      } 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
    } 
    acc *= 24.0f/myPos.w; 

    //update velocity only 
    float4 newVel = myVel + 0.5f * dt * (accel[gid] + acc); 

    //write to global memory 
    vel[gid] = newVel; 
    accel[gid] = acc; 
}

模拟运行在成果方面很好，但是使用本地内存来缓存颗粒的位置，以减轻从全局内存读取大量时的问题是性能。其实如果线路

float4 r = myPos - pblock[i];

由

float4 r = myPos - pos[jw * s_wg + i];

更换内核的运行速度更快。我真的不明白，因为从全球读书应该比从当地读书慢得多。

此外，当线

float4 r = myPos - pblock[i];

被完全去除和所有以下的r OCCURENCES由myPos - pblock[i]替换，速度和以前一样，就好像线是不存在的。这并不是因为访问r中的私有内存应该是最快的，但编译器以某种方式“优化”了这条线。

全局工作大小为4608，本地工作大小为192.它在Ubuntu 12.04中使用AMD APP SDK v2.9和Catalyst驱动程序13.12运行。

任何人都可以帮助我吗？这是我的错，还是GPU /驱动程序/ ...的问题？或者它是一个功能？ :-)

来源

2014-01-20 Vladimir Kozlowski

经过一番挖掘后发现代码导致一些LDS银行冲突。原因在于，AMD有32个长度为4个字节的存储体，但float4覆盖16个字节，因此半波阵面在同一个存储体中存取不同的地址。解决方法是制作__local float*为x和y分开坐标并单独读取他们也适当移动数组索引(id + i) % s_wg。尽管如此，性能的整体收益很小，很可能是由于@CaptainObvious提供的链接中描述的总体延迟（然后必须增加全局工作量来隐藏它们）。

来源

2014-01-22 19:44:22

我要让胡乱猜测：

当使用float4 r = myPos - pos[jw * s_wg + i];编译器是足够聪明，会注意到pblock[id]初始化之后把屏障不再是必需的，将其取下。很可能所有这些障碍（在for循环中）都会影响您的表现，因此删除它们非常明显。

是的，但全球接入成本也很高...所以我猜测背后的场景高速缓存很好地利用。还有一个事实是，你使用矢量，事实上，AMD Radeon HD 6850的架构使用VLIW处理器......也许这也有助于更好地使用高速缓存......也许。

编辑： 我刚刚发现了一个article基准GPU/APU高速缓存和内存延迟。你的GPU在列表中。你可能会得到更多答案（对不起并没有真正阅读它 - 太累了）。

来源

2014-01-20 21:30:18 CaptainObvious

障碍：这也是我的想法，但消除障碍并没有改善性能，并导致结果无意义。 –

至于基准：感谢链接。如果结果是正确的，那么，omg，这个GPU在内存延迟方面真的很糟糕！ –

OpenCL中LDS内存的性能问题

回答

相关问题