OpenACC 中是否有更快的 argmin/argmax 实现？答案

【问题标题】：Is there a faster argmin/argmax implementation in OpenACC?OpenACC 中是否有更快的 argmin/argmax 实现？
【发布时间】：2021-08-26 22:36:12
【问题描述】：

在 OpenACC 中计算 argmin 是否有比在一个最小缩减循环和另一个循环中拆分工作以实际找到最小值索引更快的替代方法？

这看起来很浪费：

    float minVal = std::numeric_limits<float>::max();
    #pragma acc parallel loop reduction(min: minVal)
    for(int i = 0; i < arraySize; ++i) {
        minVal = fmin(minVal, array[i]);
    }
    #pragma acc parallel loop
    for(int i = 0; i < arraySize; ++i) {
        if(array[i] == minVal){
            minIndex = i;
        }
    }

事实上，这成了我当前项目的瓶颈。

【问题讨论】：

请注意，如果minval 在数组中出现多次，则您的代码包含竞争条件。
@JérômeRichard 没错，但这在应用程序中重要吗？换言之，minIndex = i上方应该有#pragma acc atomic write吗？或者你的意思是，多个相等的minVals 的结果是不确定的？在大多数应用程序中，后者应该无关紧要，afaik。
首先，结果可能不是确定性的，而我在您的应用程序中可能不是问题。话虽如此，是的，我认为至少进行一次原子写入对于避免由于竞争条件而导致的与硬件相关的奇怪影响很重要。事实上，虽然我不应该成为大多数 GPU 的关键问题，但没有什么能阻止某些 GPU 以非原子方式写入 minIndex 导致错误结果。请注意，我认为几乎所有主流现代 GPU 都以原子方式写入 4 字节值，因此在实践中不应出现这种效果。您可以使用原子最小值/最大值获得确定性结果。

标签： c++ optimization openacc argmax

【解决方案1】：

我们收到了对 minloc/maxloc 的请求，但它很困难，而且很可能性能不佳，所以没有添加。您正在使用的方法是为此推荐的解决方案。

【讨论】：

这里的计算在大多数 GPU 上可能会受到内存限制。只要实现可以扩展，即使是非常糟糕的 minloc/maxloc 实现在许多 GPU 上也会更快。可以搜索每个块的 minloc/maxloc，然后使用原子 CAS 来执行块之间的缩减。或者，可以对旧的/无功能的 GPU 执行两次缩减。由此产生的实现可以在这里快 2 倍（并且可能会在许多现代 GPU 上）。
@JérômeRichard 您能否在答案中用一些代码概述您的建议？这对我来说将是一个巨大的帮助。
@Dunkelkoon 请注意，我说的是可能的后端（部分）实现。因此，不是 OpenACC 代码，而是类似于 Cuda 代码的东西（我想同样的事情可以用 OpenCL 实现，但我非常熟悉它）。你可以吗？
@JérômeRichard 啊，我的错。感谢您清除它。