OpenCL查找数组中最小值的索引答案

【问题标题】：OpenCL finding index of minimum in arrayOpenCL查找数组中最小值的索引
【发布时间】：2017-08-13 16:48:29
【问题描述】：

我正在使用 OpenCL（通过 JOCL）在光线行进的一系列距离计算中找到最小值。伪代码看起来像这样：

Start with a point in 3d space.
There are a number of functions to calculate distances
    to that point from various other points. 
    These may be rather complex (transforms, csg etc).
Calculate all of the distances, perhaps into an array
Get the index of the minimum distance in the array..
Use that index to do up other stuff (pigmentation etc).

虽然我的实现有点垃圾。我目前没有并行化距离计算，但我想这样做。这就是我不这样做的原因：

获得最小距离很容易，但检索这个索引并不明显。我最终迭代了距离并跟踪了当前最小值及其索引，但这在并行环境中显然是垃圾。

基本上可以使用提示将我引向正确的方向，或者告诉我我是否完全吠错了树？（例如，这是 CPU 作业吗？）

谢谢！

【问题讨论】：

thread-1 可以比较两对之间的距离。 thread-2 可以对另外两对做同样的事情。然后在两个线程之间同步之后，单个线程可以检查两个结果之间并选择最小最大一个。现在，如果大小为 N，则同步数可能是 LogN，并且第一步中的最大线程数可能是 N/4。好像是减价。如果有 N/4 个核心，则只需要 LogN 同步点-迭代-伪周期。
但如果数组未排序，则可能需要进行强力 O(n^2) 检查，因此在 N/4 核心系统上需要 N*Log(N) 伪循环
@huseyintugrulbuyukisik 谢谢，我试试看。这是一个减少，所以不需要排序。只需选择最小值及其索引。谢谢！
可以调整一个简单的“reduction”内核（如jocl.org/samples/reduction.cl 中的那个）以不执行+-reduction 而是执行min-reduction，并且不返回value 但 index。你考虑过这样的事情吗？（我会尝试一下，但现在时间很紧）

标签： graphics opencl gpu jocl

【解决方案1】：

使用低端显卡 RX550 测试。

一百万元素 min() 函数：

__kernel void test(__global float * data,__global int * index)
{
    int id=get_global_id(0);
    float d1=data[id];
    float d2=data[id+get_global_size(0)];
    float f=fmin(d1,d2);
    index[id]=select( index[id+get_global_size(0)], index[id], fabs(f-d1)<=fabs(f-d2) );
    data[id]=f;
}");

具有随机值的初始化数据元素和具有自己索引的索引元素。

通过 pci-e 2.0 8x 将数据和索引上传到 GPU 耗时：3.0 ms

计算全局范围=512k,256k,128k,...,1（logN 步）耗时：0.3 ms

下载数据[0]和索引[0]耗时：0.002 ms

这是一个简单的版本，可能不是最快的实现。为了更快，可以添加工作组级别的子缩减：

适用于 OpenCL 2.0+ 的 work_group_scan_inclusive_min(x)
__local float reductionArray[256] 用于 OpenCL 1.2-

减少内核排队命令的数量以在不到一百个时间内完成工作？微秒。

【讨论】：