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

Sho*_*mla 7 opencl barrier

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

我正在学习OpenCL,并希望从下面的算法中提取一些并行性.我只会告诉你我正在做的部分,我也尽可能简化了.

1)输入:两个长度为(n)的1D数组:A,B和n的值.还值C [0],D [0].

2)输出:两个长度为(n)的1D阵列:C,D.

C[i] = function1(C[i-1])
D[i] = function2(C[i-1],D[i-1])
Run Code Online (Sandbox Code Playgroud)

所以这些是递归定义,但是给定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;
    }
}
Run Code Online (Sandbox Code Playgroud)

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

Dar*_*ros 0

您的情况下的依赖关系是完全线性/递归的(我需要 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);
}
Run Code Online (Sandbox Code Playgroud)

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