这种并行计算能否在OpenCL中实现？

Question

这是我的第一篇文章。我会尽量简短，因为我很重视你的时间。这个社区对我来说非常棒。

我正在学习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核心并使其按预期工作？在这种情况下，障碍将是什么？

- Shoomla

---

你需要保存C和D的所有值，还是只关心最终结果？ - mfa

请问您能否定义function1和function2函数？ - Z boson

Answer 1

这可以在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。

- mfa

Answer 2

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

最好的方法是采用CPU的2线程方法。线程1将“生成”数据（C值），供线程2使用。

例如，一个非常天真的方法：

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。

- DarkZeros

---

此外，不是所有的全局工作项都会同时执行（它们可以被分成块来完成，必须全部完成后才能进行下一组），因此等待不是一个选项。 - Dithermaster

好的，那么我的内核基本上会保持不变，除了你提到的那些原子函数？（必须说，我仍在阅读有关openCL同步的内容，目前还无法理解这些函数的作用）。如果在CPU上运行，代码会像预期的那样在if语句上分支吗？ - Shoomla