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

Posted 2023-02-21

【中文标题】对浮点数求和的最佳 OpenCL 2 内核是啥？

【英文标题】：What is the optimum OpenCL 2 kernel to sum floats?对浮点数求和的最佳 OpenCL 2 内核是什么？

【发布时间】：2017-10-21 08:33:59

【问题描述】：

C++ 17 引入了许多新算法来支持并行执行，特别是 std::reduce 是 std::accumulate 的并行版本，它允许 non-deterministic 行为用于 non-commutative 操作，例如浮点加法。我想使用 OpenCL 2 实现一个 reduce 算法。

英特尔有一个示例 here，它使用 OpenCL 2 work group 内核函数来实现 std::exclusive_scan OpenCL 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 函数将一个work group 的最后一个值传递给下一个，而这个内核只需要将work_group_reduce_add 的结果添加到sum_val，因此atomic add 更合适.

OpenCL 2 提供了一个支持atomic_fetch_add 的atomic_int。上面使用 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，但它不支持atomic_fetch_add。

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

【问题讨论】：

sum[0] = atomic_load(&sum_val); 不是数据竞争条件吗？您的意思是在循环中添加寄存器然后最后只添加一次原子吗？

@huseyintugrulbuyukisik 我不知道。我能找到的关于使用 OpenCL 2 atomics 的最佳描述是 here。我找不到任何使用 OpenCL 2 atomics 的示例内核，这就是我问这个问题的原因。如果您有 sum_int 内核的正确实现（理想情况下，还有 sum_float 内核），请提供它作为答案。

【参考方案1】：

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; 

除了第三步之外，在几毫秒内完成相同的操作。第一个将每个组的总和放入他们的 group-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++ 设置，但我猜当您使用具有相同内存资源（可能是流模式或在 SVM 中）和/或的 多个设备 时，OpenCL 2.0 atomics 会更好使用 C++17 函数的 CPU。如果您没有多个设备同时在同一区域进行计算，那么我认为这些新的原子只能是在已经工作的 OpenCL 1.2 原子之上的微优化。我没有使用这些新的原子，所以把所有这些当成一粒盐。

【讨论】：

哇，真快！您的回答是一个很好的解决方案，并感谢您关于原子的观点。但是，我仍然想看看如何正确使用 OpenCL 2 原子，尤其是在 atomic_compare_exchange_strong 和 atomic_compare_exchange_weak 之间进行选择。

我以前没用过。让我摆弄一些内核 :) 这可能需要一段时间，也许明天我可以完全回答。由于我是直接使用 C#，所以现在可能对我来说是不可能的。

我只实现了你的一个内核，在适当的地方交换 CPU 代码中的输入和输出缓冲区，它就像一个魅力。这比使用atomics 好多了！ ;)

不错。稍后我会寻找 opencl 2.0 atomics，但我现在只能做内核内的事情，比如 atomics 本身，但不尝试它们是否与 C++17 锁保护一起工作。