PyOpenCL - 多维减少内核

如何解决PyOpenCL - 多维减少内核

我是 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 意味着我们将数组分成两半，但这完全是个谜。不理解它，我无法真正实现另一个版本以获取数组的最大值，例如，对于多维数组更是如此。

此外，据我所知，如果“N 大于单个单元中的内核数”，则reduce 内核需要重新运行一次。

你能对整段代码做进一步的解释吗？或者甚至指导如何实现它以获取数组的最大值？

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

解决方法

关于 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 的多维数组功能，例如max。

3. 如果减少的数学强度更高，请尝试使用 pyopencl 的并行算法，例如reduction

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