处理 OpenCL 中一小部分 GPU 线程产生的额外数据的最佳方法是什么？

Question

我是 OpenCL 的新手，遇到以下问题：

我有一个大数组（6 * 1,000,000 个浮点数）。对于数组的每个元素，我需要进行计算。基本算法在多达 16 个 GPU (Tesla K80) 上运行良好：

1.) 我为每个 GPU 设备的结果创建数组的缓冲区对象和缓冲区对象，并将其写入每个 GPU 内存。

2.) 然后，为每个数组元素生成一个线程，并在 GPU 上的内核中执行计算。

3.) 将结果写入到全局线程id对应的result数组元素中。

4.) 主机读取结果缓冲区。

我现在必须扩展这个算法。一些数组元素（10-100）实际上需要额外的计算来产生额外的结果（另外 12 个浮点数）。

这是一些伪代码。

__kernel void calculation(__global float4 *input_array,
                          __global float4 *result_array){

    int id = get_global_id(0);

    //do calculation
    float4 result = some_func(input_array[id]);
    result_array[id] = result;


    if(some_rare_condition){
        //do another, much longer calculation
        float4 result2 = another_func(input_array[id]);
    }
}

问题是我只有一些额外的结果，我不知道存储它们并让主机读取它们的最佳方法是什么。

在我计算出第一个结果之前，我不知道哪些数组元素需要额外计算。

如果这是 C++，我只需为附加结果创建一个向量，为索引创建一个向量。但是，据我所知，OpenCL 内核中没有动态内存容器。

如果我创建一个包含 1,000,000 个元素的第二个结果数组，并且只写入所需的少数几个位置，那么当我将它传递回主机时就会产生瓶颈。

如果我创建一个肯定大于所需的较小数组（例如 1000 个元素），我不确定如何让线程安全地写入它。

Answer 1

最简单的解决方案可能是 use an atomic counter 在较小的数组中分配索引。

项目在数组中出现的顺序是不可预测的，因此您还需要存储识别信息（例如原始的id），然后根据您需要处理的内容对其进行排序接下来是这个输出。

但是，这是否有效取决于许多因素；更糟糕的是，这些因素相互影响。

首先，听起来在第二个输出数组中需要一个项目的概率小于 0.1%。像这样的小数字对原子有好处 - 想要增加该计数器的工作项越多，它们尝试同时这样做并相互阻止并序列化的机会就越高。此外，分配很重要。任何集群都会对您不利。如果您的 0.1% 大部分是相邻的，那么它们在工作组中也会相邻，并且工作组中的工作项通常在 GPU 上同步运行。所以他们都需要等待对方完成增量，然后才能继续。

另一方面，如果 another_func() 的计算成本相当高（您的问题表明确实如此），那么均匀分布是不好的，因为组中的大多数工作项在单个项目正在运行 another_func() 并占用整个计算单元。

消除各种缺点的简单方法的可能变化：

1. 采用分支的项目集群，或采用分支的概率很高。 在这种情况下，您可以一次为整个工作组分配输出数组的范围，因此只有一个工作项在组增加计数器。您将需要分配整个工作组的输出槽，即使有些是不需要的，或者在整个组中运行缩减算法以计算出需要多少 ，并分配索引偏移量.
2. 具有昂贵条件计算的非集群工作项。在这种情况下，您可能希望仅使用原子方法将需要进一步计算的项的索引输出到数组中，然后调用another_func()在第二个内核中，每个工作项都在需要进一步计算的一项上工作，因此所有工作项都运行another_func()，这意味着内核已完全占用。

如果需要，当然可以将这两种变体组合起来，并且您可能还可以进行进一步的改进，特别是如果您可以使用平台的分析器来查明瓶颈。