减少最终金额的方法

Question

我继续我在上解释的第一期。

我提醒你，我想应用一种能够使用 OpenCL 进行多次和减少的方法（我的 GPU 设备仅支持 OpenCL 1.2）。我需要计算数组的总和减少以检查主循环每次迭代的收敛标准，

目前，我做了一个版本只减少了一次总和（即一次迭代 ).在这个版本中，为了简单起见，我使用了顺序 CPU 循环来计算每个部分和的总和并获得总和的最终值。

根据您在我的先例中的建议，我的问题是我不知道如何通过第二次调用 NDRangeKernel 函数（即第二次执行内核代码）来执行最后的总和。

事实上，第二次调用时，我将始终面临获取部分和的总和的相同问题（本身是根据 NDRangeKernel 的第一次调用计算的）：这似乎是一个递归问题。

我们以上图为例：如果输入数组大小为10240000且WorkGroup size为16，我们得到10000*2^10/2^4 = 10000*2^6 = 640000 WorkGroups。

所以在第一次调用之后，我得到 640000 partial sums：如何处理所有这些部分和的最终总和？如果我再次调用内核代码，例如 WorkGroup size = 16 和全局 size = 640000，我将得到 nWorkGroups = 640000/16 = 40000 partial sums，所以我必须再次调用内核代码并重复此过程直到nWorkGroups < WorkGroup size.

可能是我对第二阶段不是很了解，主要是"two-stage reduction"(on this link, I think this is the case of searching for minimum into input array)

这部分内核代码

__kernel
void reduce(__global float* buffer,
            __local float* scratch,
            __const int length,
            __global float* result) {

  int global_index = get_global_id(0);
  float accumulator = INFINITY;
  // Loop sequentially over chunks of input vector
  while (global_index < length) {
    float element = buffer[global_index];
    accumulator = (accumulator < element) ? accumulator : element;
    global_index += get_global_size(0);
  }

  // Perform parallel reduction
  ...

如果有人能解释上面这段内核代码的作用。

和归约的第二阶段即最终求和有关系吗？

如果您不明白我的问题，请随时问我更多细节。

谢谢

Answer 1

如评论中所述：声明

if input array size is 10240000 and WorkGroup size is 16, we get 10000*2^10/2^4 = 10000*2^6 = 640000 WorkGroups.

不正确。您可以选择 "arbitrary" 工作组大小和 "arbitrary" 工作组数量。此处选择的数字可能是为目标设备定制的。例如，设备可能具有特定的本地内存大小。这可以用 clDeviceGetInfo:

查询

cl_ulong localMemSize = 0;
clDeviceGetInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, 
    sizeof(cl_ulong), &localMemSize, nullptr);

这可用于计算本地工作组的大小，考虑到每个工作组都需要

sizeof(cl_float) * workGroupSize

字节的本地内存。

类似地，工作组的数量可以从其他设备特定参数中导出。

关于减少本身的关键点是工作组大小不限制可以处理的数组的大小。我对整个算法的理解也有一些困难，所以我试着在这里解释一下，希望几张图片抵得上千言万语：

如您所见，工作组的数量和工作组的大小是固定的，并且与输入数组长度无关：尽管我在示例中使用了 3 个大小为 8 的工作组（给出了全局大小为 24)，可以处理长度为 64 的数组。这主要是由于第一个循环，它只是遍历输入数组，"step size" 等于全局工作大小（此处为 24）。结果将是 24 个线程中每个线程的 one 累加值。这些然后并行减少。

减少最终金额的方法

Method to do final sum with reduction

gpu

reduction

opencl