使用 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 大得多。从我读到的内容，我认为可以使用并行求和减少来解决这个问题，其中有很多例子。但是对于这种方法，我需要为每个不适合内存的线程（或工作项）准备一份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 中，性能取决于许多因素，例如要向设备传输多少数据、是否需要进行足够的计算以使其有意义、如何在内核中访问数据：工作项的顺序)，内存的类型 - 全局、本地、寄存器 - 等等。