在 AMD Radeon RX580 上解决的这个 N 体问题中，内存访问的最佳实践是什么？

Question

我计算了在它们的引力场中移动的 N 个粒子的轨迹。我编写了以下 OpenCL 内核：

#define G 100.0f
#define EPS 1.0f

float2 f (float2 r_me, __constant float *m, __global float2 *r, size_t s, size_t n)
{
    size_t i;
    float2 res = (0.0f, 0.0f);

    for (i=1; i<n; i++) {
        size_t idx = i;
//        size_t idx = (i + s) % n;
        float2 dir = r[idx] - r_me;
        float dist = length (dir);
        res += G*m[idx]/pown(dist + EPS, 3) * dir;
    }

    return res;
}

__kernel void take_step_rk2 (__constant float *m,
                             __global float2 *r,
                             __global float2 *v,
                             float delta)
{
    size_t n = get_global_size(0);
    size_t s = get_global_id(0);


    float2 mv = f(r[s], m, r, s, n);
    float2 mr = v[s];

    float2 vpred1 = v[s] + mv * delta;
    float2 rpred1 = r[s] + mr * delta;

    float2 nv = f(rpred1, m, r, s, n);
    float2 nr = vpred1;

    barrier (CLK_GLOBAL_MEM_FENCE);

    r[s] += (mr + nr) * delta / 2;
    v[s] += (mv + nv) * delta / 2;
}

然后我运行这个内核多次作为全局工作大小的一维问题=[物体数量]：

void take_step (struct cl_state *state)
{
    size_t n = state->nbodies;
    clEnqueueNDRangeKernel (state->queue, state->step, 1, NULL, &n, NULL, 0, NULL, NULL);
    clFinish (state->queue);
}

引用自AMD OpenCL Optimization Guide（2015 年）：

Under certain conditions, one unexpected case of a channel conflict is that reading from the same address is a conflict, even on the FastPath. This does not happen on the read-only memories, such as constant buffers, textures, or shader resource view (SRV); but it is possible on the read/write UAV memory or OpenCL global memory.

我队列中的工作项都试图在这个循环中访问相同的内存，所以肯定存在通道冲突：

for (i=1; i<n; i++) {
        size_t idx = i;
//        size_t idx = (i + s) % n;
        float2 dir = r[idx] - r_me;
        float dist = length (dir);
        res += G*m[idx]/pown(dist + EPS, 3) * dir;
    }

我替换了

        size_t idx = i;
//        size_t idx = (i + s) % n;

和

//        size_t idx = i;
        size_t idx = (i + s) % n;

所以第一个工作项（具有全局 ID 0）首先访问数组 r 中的第一个元素，第二个工作项访问第二个元素，依此类推。

我预计此更改一定会带来性能提升，但恰恰相反，它会导致性能显着下降（大约下降 2 倍）。我错过了什么？为什么在这种情况下所有相同的内存可以更好地访问它？

如果您有其他提高性能的技巧，请与我分享。 OpenCL 优化指南非常混乱。

Answer 1

f 函数的循环没有针对合并访问的重新收敛障碍。一旦一些项目得到他们的 r 数据，他们就开始计算，但那些不能等待他们的数据因此，失去了合并的完整性。要重新分组，至少每 10 次迭代或 2 次迭代或什至每次迭代添加 1 个障碍。但是访问global有很高的延迟。障碍 + 延迟不利于性能。您在这里需要本地内存，因为它具有低延迟和广播能力，这让它仅在大于本地线程数（64？）的粒度上失去合并性，这对全局内存访问也不错（您需要在每个线程中从全局填充本地内存第 K 次迭代，其中 N 被分成 K 个大小的组）。

2013 年的消息来源（ http://developer.amd.com/wordpress/media/2013/07/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide-rev-2.7.pdf):

Thus, the key to effectively using the LDS is to control the access pattern, so that accesses generated on the same cycle map to different banks in the LDS. One notable exception is that accesses to the same address (even though they have the same bits 6:2) can be broadcast to all requestors and do not generate a bank conflict.

为此使用 LDS(__local) 将提供良好的性能。因为 LDS 很小，所以你应该像一次 256 个粒子那样小块地做。

此外，使用 i 作为 idx 对缓存非常友好，但模数版本对缓存非常不利。一旦数据可以存在于缓存中，是否完成 N 个请求都没有关系。它们现在来自缓存。但是使用模数，您可以在缓存成分被重新使用之前销毁它们，具体取决于 N。对于小 N，它应该像您预见的那样更快。对于大 N 和小 GPU 缓存，情况会更糟。就像每个周期只有 1 个全局请求与每个周期 N-cache_size 个全局请求一样。

我猜想使用如此强大的 GPU，你有一个高 N 值，例如 64k 主体，每个主体需要 2 个变量，每个变量需要 4 个字节，总计 512kB，这不适合 L1。也许只有 L2 比 idx=i 慢到 L1。

答案：

所有到相同的一级缓存 adr 比所有到全局和二级缓存 adr 更快
在"blocking/patching"算法中使用本地内存实现高速

在 AMD Radeon RX580 上解决的这个 N 体问题中，内存访问的最佳实践是什么？

What is the best practice for memory access in this N-body problem solved on AMD Radeon RX580?

opencl

amd-gcn