OpenCL中LDS内存的性能问题

Question

我有一个性能问题，当使用LDS内存与AMD Radeon HD 6850。

我有两个核作为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都被myPos - pblock[i]所取代，速度和以前一样，就好像根本没有这条线一样。这一点我没有得到更多，因为访问私有内存在r应该是最快的，但编译器在某种程度上“优化”了这一行。

全球工作规模为4608，本地工作规模为192。它在Ubuntu12.04中与AMD应用SDKV2.9和催化剂驱动程序13.12一起运行。

有人能帮我吗？这是我的错还是GPU /驱动程序/.的问题？或者是一种特征？:-)

Answer 1

经过进一步的挖掘，结果发现代码导致了一些LDS银行冲突。原因是对于AMD来说，有32个4字节长的银行，但是float4覆盖了16个字节，因此半波前访问相同银行中的不同地址。解决方案是将x坐标和y坐标分别生成__local float*，并分别读取它们，并以数组索引的适当偏移作为(id + i) % s_wg。尽管如此，性能的总体收益是很小的，这很可能是@CaptainObvious提供的链接中所描述的总体延迟(那么我们必须增加全局工作规模来隐藏它们)。

Answer 2

我要做个疯狂的猜测：

当使用float4 r = myPos - pos[jw * s_wg + i];时，编译器足够聪明地注意到，在pblock[id]初始化之后设置的屏障不再必要，并删除它。很可能所有这些障碍(在for循环中)都会影响您的性能，因此消除它们是非常明显的。

是的，但是全局访问花费了很多too...So，我猜场景缓存内存的背后被很好地利用了。还有一个事实是，您使用向量，事实上，and 6850的架构使用VLIW processors...maybe，这也有助于更好地利用缓存memories...maybe。

编辑：我刚刚发现了一个文章基准测试GPU/APU缓存和内存延迟。你的GPU在名单上。你可能会得到更多的答案(很抱歉没有读到--太累了)。