PyOpenCL-多维约简内核

Question

我是OpenCL的新手。

我正在尝试为多维数组编写一个沿一个轴求和的归约内核。我偶然发现了来自这里的代码：https://tmramalho.github.io/blog/2014/06/16/parallel-programming-with-opencl-and-python-parallel-reduce/

__kernel void reduce(__global float *a, __global float *r, __local float *b) {
uint gid = get_global_id(0);
uint wid = get_group_id(0);
uint lid = get_local_id(0);
uint gs = get_local_size(0);
b[lid] = a[gid];
barrier(CLK_LOCAL_MEM_FENCE);

for(uint s = gs/2; s > 0; s >>= 1) {
if(lid < s) {
b[lid] += b[lid+s];
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if(lid == 0) r[wid] = b[lid];
}

我不理解for循环部分。我知道uint s = gs/2意味着我们将数组一分为二，但这完全是个谜。如果不理解它，我就无法真正实现另一个以数组最大值为例的版本，更不用说多维数组了。

此外，据我所知，reduce内核需要再次运行，如果"；N大于单个单元中的核的数量"；。

你能就整段代码给我进一步的解释吗？甚至指导如何实现数组的最大值？

完整的代码可以在这里找到：https://github.com/tmramalho/easy-pyopencl/blob/master/008_localreduce.py

Answer 1

关于for循环含义的第一个问题：

for(uint s = gs/2; s > 0; s >>= 1)

这意味着你将局部大小gs除以2，并保持除以2(移位部分s >>= 1相当于s = s/2(，而s>0，换句话说，直到s＝1。该算法取决于你的数组大小是2的幂，否则你必须处理超过2的幂的问题，直到你缩小了整个数组，或者你必须用中性值来填充你的数组，直到完成2的幂。

第二个问题是，当N大于GPU的容量时，你是对的：你必须对适合的部分进行缩减，然后合并结果。

最后，当你要求指导如何实现缩减以获得阵列的最大值时，我建议如下：

1. 对于像max或sum这样的简单归约，请尝试使用numpy，尤其是在处理按轴编程归约的情况下。

2. 如果你认为GPU会给你带来优势，请先尝试使用pyopencl的多维阵列功能，例如最大

3. 如果减少是更密集的数学，尝试使用pyoppencl的并行算法，例如减少

我认为使用pyopencl的全部目的是避免处理底层GPU的架构。否则，直接处理CUDA或HIP比处理OpenCL更容易。