0

我计算了 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 优化指南（2015 年）的引述：

在某些情况下，通道冲突的一个意外情况是从同一地址读取是一个冲突，即使在 FastPath 上也是如此。这不会发生在只读存储器上，例如常量缓冲区、纹理或着色器资源视图 (SRV)；但可以在读/写 UAV 内存或 OpenCL 全局内存上。

我队列中的工作项都尝试在此循环中访问相同的内存，因此肯定存在通道冲突：

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 优化指南非常混乱。

4

1 回答 1

1

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

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

因此，有效使用 LDS 的关键是控制访问模式，以便在同一周期上生成的访问映射到 LDS 中的不同存储体。一个值得注意的例外是对同一地址的访问（即使它们具有相同的 6:2 位）可以广播给所有请求者，并且不会产生银行冲突。

为此使用 LDS( __local) 将提供良好的性能。由于 LDS 很小，因此您应该一次处理 256 个粒子这样的小块。

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

我猜对于如此强大的 GPU，你有一个很高的 N 值，例如 64k 个主体，每个主体需要 2 个变量，每个变量需要 4 个字节，总计 512kB，无法容纳 L1。也许只有 L2 通过 L1 比 idx=i 慢。

回答：

  • all to same L1 cache adr 比 all to global and L2 cache adr 快

  • 在“blocking/patching”算法中使用本地内存来实现高速

于 2019-09-09T17:30:33.770 回答