这种并行性可以在 OpenCL 中实现吗答案

【问题标题】：Can this parallelism be implemented in OpenCL这种并行性可以在 OpenCL 中实现吗
【发布时间】：2016-05-15 03:49:09
【问题描述】：

这是我的第一篇文章。我会尽量保持简短，因为我珍惜你的时间。这个社区对我来说太不可思议了。

我正在学习 OpenCL，想从下面的算法中提取一点并行性。我只会向您展示我正在处理的部分，我也尽可能地简化了它。

1) 输入：两个长度为 (n) 的一维数组：A、B 和 n 的值。还有值 C[0], D[0]。

2) 输出：两个长度为 (n) 的一维数组：C、D。

C[i] = function1(C[i-1])
D[i] = function2(C[i-1],D[i-1])

所以这些是递归定义，但是对于给定的 i 值的 C 和 D 的计算可以并行完成（它们显然更复杂，以便有意义）。一个天真的想法会为以下内核创建两个工作项：

__kernel void test (__global float* A, __global float* B, __global float* C,
                    __global float* D, int n, float C0, float D0) {
    int i, j=get_global_id(0);

    if (j==0) {
       C[0] = C0;
       for (i=1;i<=n-1;i++) {
          C[i] = function1(C[i-1]);
          [WAIT FOR W.I. 1 TO FINISH CALCULATING D[i]];
       }
       return;
    }
    else {
       D[0] = D0;
       for (i=1;i<=n-1;i++) {
          D[i] = function2(C[i-1],D[i-1]);
          [WAIT FOR W.I. 0 TO FINISH CALCULATING C[i]];
       }
       return;
    }
}

理想情况下，两个工作项（数字 0,1）中的每一个都会进行一次初始比较，然后进入各自的循环，为每次迭代同步。现在考虑到 GPU 的 SIMD 实现，我假设这将不起作用（工作项将等待所有内核代码），但是是否可以将这种类型的工作分配给两个 CPU 内核并使其按预期工作？在这种情况下，障碍是什么？

【问题讨论】：

是要求保存C和D的所有值，还是只关心最终结果？
你能定义function1和function2吗？

标签： opencl barrier

【解决方案1】：

在您的情况下，依赖性是完全线性/递归的（我需要 i-1）。甚至不像其他问题（归约、求和、排序等）那样是对数的。因此，这个问题不适用于 SIMD 设备。

你能做的最好的就是在 CPU 中采用 2 线程方法。线程 1 将为线程 2“生成”数据（C 值）。

一个非常幼稚的方法例如：

Thread 1:
for(){
    ProcessC(i);
    atomic_inc(counter); //This function should unlock
}

Thread 2:
for(){
    atomic_dec(counter); //This function should lock
    ProcessD(i);
}

其中atomic_inc 和atomic_dec 可以通过例如计数信号量来实现。

【讨论】：

此外，并非所有全局工作项都一起执行（它们可能会在必须在下一组完成之前完全完成的块中完成），因此不能选择等待。
好的，所以除了你提到的那些原子函数之外，我的内核基本上保持不变？（不得不说，我仍在阅读有关 openCL 同步的内容，因此无法掌握这些函数在 atm 中的作用）。如果在 CPU 上运行，代码会像预期的那样在 if 语句上分支？

【解决方案2】：

这可以在 opencl 中实现，但就像其他答案所说，您最多只能使用 2 个线程。

你的函数的我的版本应该用一个有两个工作项的工作组来调用。

__kernel void test (__global float* A, __global float* B, __global float* C, __global float* D, int n, float C0, float D0)
{
    int i;
    int gid = get_global_id(0);

    local float prevC;
    local float prevD;

    if (gid == 0) {
        C[0] = prevC = C0;
        D[0] = prevD = D0;
    }

    barrier(CLK_LOCAL_MEM_FENCE);

    for (i=1;i<=n-1;i++) {
        if(gid == 0){
            C[i] = function1(prevC);
        }else if (gid == 1){
            D[i] = function2(prevC, prevD);
        }

        barrier(CLK_LOCAL_MEM_FENCE);
        prevC = C[i];
        prevD = D[i];
    }
}

这应该可以在任何 opencl 硬件上运行。如果您不关心保存所有 C 和 D 值，则可以简单地以两个浮点数而不是整个列表返回 prevC 和 prevD。由于对中间值的所有读取和写入都坚持较低的缓存级别（即本地内存），这也将使其更快。本地内存提升也应该适用于所有 opencl 硬件。

那么在 GPU 上运行它有什么意义吗？不是为了并行性。你被 2 个线程困住了。但是，如果您不需要返回 C 和 D 的所有值，您可能会看到显着的加速，因为 GPU 的内存要快得多。

所有这些都假设 function1 和 function2 并不过分复杂。如果是，请坚持使用 CPU——可能还有其他多处理技术，例如 OpenMP。

【讨论】：