使用 OpenCL 对大型数据集进行增量并行求和

Question

我有一个 3-D 向量 F 当前按如下顺序填充：

// for each particle
for(int p = 0; p < pmax; p++) {
    // and each dimension (n x m x o)
    for(int i = 0; i < n; i++) {
        for(int j = 0; j < m; j++) {
            for(int k = 0; k < o; k++) {
                // do some calulations depending on p, i, j and k and add onto F
                F[i][j][k] += p * i * j * k;
            }   
        }
    }
}

F 的值是针对每个粒子递增计算的。

现在我想使用 OpenCL 加快速度。我的第一次尝试是用 3D-Range-Kernel 并行化内部循环，每个粒子都会调用它。这确实工作正常，但在 GPU 上比在 CPU 上更慢（没有应用任何优化 - 例如 F 在每次迭代中从主机复制到 GPU 并再次复制回来）。

在第二次尝试中，我尝试并行化外循环（我认为）这是比第一次尝试更好的主意，因为预计 pmax 比 m * n * o 大得多。根据我的阅读，我认为可以使用并行和减少来解决这个问题，其中有很多例子。但是，对于这种方法，我需要为每个不适合内存的线程（或 work-item）提供 F 的副本（在此处获取 CL_OUT_OF_RESOURCES）。

我现在的问题是：是否可以并行化这样一个增量总和，而无需将多个粒子的 F 多次存储在内存中？如果是，好的方法是什么样的？我应该坚持我的第一次尝试并尝试优化它吗？

请注意，我是 OpenCL 的新手，对一般的并行化技术了解不多。如果有任何提示或参考有用的讲座或示例，我将不胜感激，谢谢！

顺便说一句，我关于这个主题的谷歌搜索会议只让我计算了前缀和。

Answer 1

您可以像这样从简单的 3 维内核开始：

import numpy as np
import pyopencl as cl
import time

m=32
n=32
o=32
pmax=16

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

t1=time.time()
F=np.zeros((m,n,o)).astype(np.int32).flatten(order='F')

mf = cl.mem_flags
F_buf = cl.Buffer(ctx, mf.WRITE_ONLY, F.nbytes)

prg = cl.Program(ctx, """
    __kernel void calc(int pmax, __global int *F) {

    int i = get_global_id(0);
    int j = get_global_id(1);
    int k = get_global_id(2);
    int m = get_global_size(0);
    int n = get_global_size(1);
    int o = get_global_size(2);

    int tmp = i * j * k;
    int sum = 0;
    for(int p = 0; p < pmax; p++) 
        sum += p * tmp;
    F[i*n*o+j*o+k] = sum;
    }
    """).build()

prg.calc(queue, (m,n,o), None, np.int32(pmax), F_buf)
cl.enqueue_copy(queue, F, F_buf)

F=F.reshape((m,n,o), order='F')
F=np.transpose(F, (2, 1, 0))
t2=time.time()

t3=time.time()    
Ftest=np.zeros((m,n,o)).astype(np.int32)
for p in range(pmax):
    for i in range(m):
        for j in range(n):
            for k in range(o):
                Ftest[i][j][k] += p*i*j*k

t4=time.time()

print "OpenCL time:", (t2-t1)
print "CPU time:", (t4-t3)

测试结果：

$ python test.py 
Choose platform:
[0] <pyopencl.Platform 'Intel(R) OpenCL HD Graphics' at 0x557007fed680>
[1] <pyopencl.Platform 'Portable Computing Language' at 0x7fab67ff0020>
Choice [0]:
Set the environment variable PYOPENCL_CTX='' to avoid being asked again.
OpenCL time: 0.0124819278717
CPU time: 1.03352808952

$ python test.py 
Choose platform:
[0] <pyopencl.Platform 'Intel(R) OpenCL HD Graphics' at 0x55a2650505a0>
[1] <pyopencl.Platform 'Portable Computing Language' at 0x7fd80775d020>
Choice [0]:1
Set the environment variable PYOPENCL_CTX='1' to avoid being asked again.
OpenCL time: 0.0148649215698
CPU time: 1.11784911156

根据矩阵的大小，性能可能会有所不同。此外，这并不意味着它会比您当前的 CPU 实施更快。在 OpenCL 中，性能取决于许多因素，例如要传输到设备的数据量、是否需要进行足够的计算以使其有意义、在内核中如何访问数据：按工作项的顺序), 内存的类型 - 全局、局部、寄存器 - 等等。

使用 OpenCL 对大型数据集进行增量并行求和

Incremental parallel sum on a large dataset using OpenCL

c

reduction

opencl