对浮点数求和的最佳 OpenCL 2 内核是什么？

C++ 17引入了许多新算法来支持并行执行，特别是标准::减少 http://en.cppreference.com/w/cpp/algorithm/reduce是一个并行版本std::累积 http://en.cppreference.com/w/cpp/algorithm/accumulate这允许non-deterministic的行为non-commutative运算，例如浮点加法。我想使用 OpenCL 2 实现一个归约算法。

英特尔有一个例子here https://www.intel.com/content/dam/develop/external/us/en/documents/opencl-workgroupfunctions-531084.pdf它使用 OpenCL 2work group内核函数来实现std::exclusive_scan http://en.cppreference.com/w/cpp/algorithm/exclusive_scanOpenCL 2 内核。下面是内核对浮点数求和，基于 Intel 的exclusive_scan例子：

kernel void sum_float (global float* sum, global float* values)
{
  float sum_val = 0.0f;

  for (size_t i = 0u; i < get_num_groups(0); ++i)
  {
    size_t index = get_local_id(0) + i * get_enqueued_local_size(0);
    float value = work_group_reduce_add(values[index]);
    sum_val += work_group_broadcast(value, 0u);
  }

  sum[0] = sum_val;
}

上面的内核可以工作（或者看起来可以！）。然而，exclusive_scan需要的work_group_broadcast传递最后一个值 1 的函数work group到下一个，而这个内核只需要将work_group_reduce_add的结果添加到sum_val, so an atomic add更合适。

OpenCL 2 提供了atomic_int它支持atomic_fetch_add。上述使用atomic_int 的内核的整数版本是：

kernel void sum_int (global int* sum, global int* values)
{
  atomic_int sum_val;
  atomic_init(&sum_val, 0);

  for (size_t i = 0u; i < get_num_groups(0); ++i)
  {
    size_t index = get_local_id(0) + i * get_enqueued_local_size(0);
    int value = work_group_reduce_add(values[index]);
    atomic_fetch_add(&sum_val, value);
  }

  sum[0] = atomic_load(&sum_val);
}

OpenCL 2 还提供了atomic_float but it doesn't支持atomic_fetch_add.

实现 OpenCL2 内核对浮点数求和的最佳方法是什么？

kernel void sum_float (global float* sum, global float* values)
{
  float sum_val = 0.0f;

  for (size_t i = 0u; i < get_num_groups(0); ++i)
  {
    size_t index = get_local_id(0) + i * get_enqueued_local_size(0);
    float value = work_group_reduce_add(values[index]);
    sum_val += work_group_broadcast(value, 0u);
  }

  sum[0] = sum_val;
}

这有一个将数据写入 sum 的零索引元素的竞争条件，所有工作组都在执行相同的计算，这使得这个 O(N*N) 而不是 O(N) 并且需要超过 1100 毫秒才能完成 1M 元素数组和。

对于相同的 1-M 元素数组， this(global=1M, local=256)

kernel void sum_float2 (global float* sum, global float* values)
{
      float sum_partial = work_group_reduce_add(values[get_global_id(0)]);
      if(get_local_id(0)==0)
        sum[get_group_id(0)] = sum_partial; 
}

接下来是这个（全局=4k，本地=256）

kernel void sum_float3 (global float* sum, global float* values)
{
  float sum_partial = work_group_reduce_add(sum[get_global_id(0)]);
  if(get_local_id(0)==0)
    values[get_group_id(0)] = sum_partial; 
}

除了第三步之外，在几毫秒内完成同样的事情。第一个将每个组的总和放入其组 ID 相关项中，第二个内核将这些总和放入 16 个值中，这 16 个值可以很容易地由 CPU 求和（微秒或更短）（作为第三步）。

程序的工作原理如下：

values: 1.0 1.0 .... 1.0 1.0 
sum_float2
sum: 256.0 256.0 256.0
sum_float3
values: 65536.0 65536.0 .... 16 items total to be summed by cpu 

如果你需要使用原子，你应该尽可能稀疏地使用。最简单的示例可以是使用局部原子对每个组的许多值进行求和，然后使用每个组的单个全局原子函数来执行最后一步来添加所有值。我现在还没有为 OpenCL 准备好 C++ 设置，但我猜当您使用 OpenCL 2.0 原子时会更好多个设备具有相同的内存资源（可能是流模式或SVM）和/或CPU使用 C++17 函数。如果您没有多个设备同时在同一区域进行计算，那么我认为这些新原子只能是在已经运行的 OpenCL 1.2 原子之上的微优化。我没有使用这些新的原子，所以对所有这些都持保留态度。