c-使用OpenCL在大型数据集上进行增量并行求和

Question

我有一个三维矢量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范围内核来并行内部循环，该内核为每个粒子调用。这确实很好，但在GPU上甚至比在CPU上慢(没有应用任何优化——例如，F在每次迭代中从主机复制到GPU并再次复制回来(。

在第二次尝试中，我尝试将外循环并行化，这(我认为(比第一次尝试更好，因为pmax预计要比m * n * o大得多。从我读到的内容来看，我认为这个问题可以通过并行和约简来解决，其中有很多例子。然而，对于这种方法，我需要为每个不适合内存的线程(或工作项(提供F的副本(在这里获得CL_OUT_OF_RESOURCES(。

我现在的问题是：在不将F多次存储在内存中的情况下，是否有可能并行化这样一个增量和？如果是，一个好的方法会是什么样子？我应该坚持我的第一次尝试并尝试优化它吗？

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

顺便说一句，我在谷歌上搜索这个话题只会导致我计算前缀和。

Answer 1

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

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中，性能取决于许多因素，如有多少数据要传输到设备，是否有足够的计算要进行，以使其有意义，如何在内核中访问数据：按工作项的顺序，内存的类型-全局、本地、寄存器-等等