使用 Opencl 有效地找到最小的大数组答案

【问题标题】：Efficiently find minimum of large array using Opencl使用 Opencl 有效地找到最小的大数组
【发布时间】：2014-08-07 16:05:36
【问题描述】：

我正致力于在 opencl 中实现层次聚类算法。对于每一步，我都在一个非常大的数组（大约 10^8 个条目）中找到了最小值，这样我就知道哪些元素必须组合成一个新的集群。最小值的识别必须进行 9999 次。使用我当前的内核，大约需要 200 秒才能找到最小值（在所有迭代中累积）。我解决这个问题的方法是将数组分成 2560 个大小相等的片段（我的 Radeon 7970 上有 2560 个流处理器），然后分别找到每个片段的最小值。我运行第二个内核，将这些最小值组合成全局最小值。

有没有更有效的方法来解决这个问题？最初的想法是通过使用 OpenCL 来加速 HCA，但是识别最小值所花费的时间比 CPU 上的 matlab HCA 长得多。我做错了什么？

__kernel void findMinValue(__global float * myArray, __global double * mins, __global int * elementsToWorkOn, __global int * arraysize){
int gid = get_global_id(0);
int minloc = 0;
float mymin = INFINITY;
int eltoWorkOn = *elementsToWorkOn;
int offset = gid*eltoWorkOn;
int target = offset + eltoWorkOn;

if (offset<*arraysize){
    //make sure the array size is not exceeded
    if (target > *arraysize){
        target = *arraysize;
    }

    //find minimum for the kernel
    for (int i = offset; i < target; i++){
        if (*(myArray + i) < mymin){
            mymin = *(myArray + i);
            minloc = i;
        }
    }
}
*(mins + gid * 2) = minloc;
*(mins + gid * 2 + 1) = mymin;
}


__kernel void getGlobalMin(__global double * mins, __global double * gmin, __global int * pixelsInImage){
    int nWorkitems = 2560;
    float globalMin = INFINITY;
    double globalMinLoc;
    float tempMin;
    for (int i = 0; i < nWorkitems; i++){
        tempMin = *(mins + 2 * i + 1);
        if (tempMin < globalMin){
            globalMin = tempMin;
            globalMinLoc = *(mins + 2 * i);
        }
    }
    *(gmin + 0) = globalMinLoc;
    *(gmin + 1) = globalMin;
}

更新

我根据您的建议重新设计了 findMinValue 内核。内存访问现在是合并的，我将工作分成工作组，这样我就可以减少全局内存访问的数量。以前，每个内核都将其最小值写入全局 mins 缓冲区。现在每个 worg 组只有一个内核写入一个值（即组最小值）。此外，我增加了全局工作大小以隐藏内存延迟。

这些更改可以将识别最小值所需的时间从 >200 秒减少到仅 59 秒！非常感谢您的帮助！

在优化内核时还有什么我可能遗漏的吗？你有什么进一步的建议吗？我不知道如何使用setArg()。我是否必须将指向 int 值的指针传递给它（例如：err = clSetKernelArg(kernel[2], 3, sizeof(int), &variable);）。在这种情况下，内核声明看起来如何？

这是我的新内核：

__kernel void findMinValue(__global float * myArray, __global double * mins, __global int * arraysize,__global int * elToWorkOn,__global int * dummy){
int gid = get_global_id(0);
int lid = get_local_id(0);
int groupID = get_group_id(0);
int lsize = get_local_size(0);
int gsize = get_global_id(0);
int minloc = 0;
int arrSize = *arraysize;
int elPerGroup = *elToWorkOn;
float mymin = INFINITY;


__local float lmins[128];
//initialize local memory
*(lmins + lid) = INFINITY;
__local int lminlocs[128];

//this private value will reduce global memory access in the for loop (temp = *(myArray + i);)
float temp;

//ofset and target of the for loop
int offset = elPerGroup*groupID + lid;
int target = elPerGroup*(groupID + 1);

//prevent that target<arrsize (may happen due to rounding errors or arrSize not a multiple of elPerGroup
target = min(arrSize, target);

//find minimum for the kernel
//offset is different for each lid, leading to sequential memory access
if (offset < arrSize){
    for (int i = offset; i < target; i += lsize){
        temp = *(myArray + i);
        if (temp < mymin){
            mymin = temp;
            minloc = i;
        }
    }

    //store kernel minimum in local memory
    *(lminlocs + lid) = minloc;
    *(lmins + lid) = mymin;

    //find work group minimum (reduce global memory accesses)
    lsize = lsize >> 1;
    while (lsize > 0){
        if (lid < lsize){
            if (*(lmins + lid)> *(lmins + lid + lsize)){
                *(lmins + lid) = *(lmins + lid + lsize);
                *(lminlocs + lid) = *(lminlocs + lid + lsize);
            }
        }
        lsize = lsize >> 1;
    }
}
//write group minimum to global buffer
if (lid == 0){
    *(mins + groupID * 2 + 0) = *(lminlocs + 0);
    *(mins + groupID * 2 + 1) = *(lmins + 0);
}
}

【问题讨论】：

每次搜索最小值之间的数组会经历什么样的变化？解决方案可能涉及记录这些更改，而不是重新开始每次搜索。
我已经考虑过了，但没有想出一个聪明的解决方案。事实上，每次迭代只更改两行和两列。矩阵本身是 10000x10000 项大。如果您知道任何避免每次迭代都遍历每个项目的好方法，我会非常高兴。这也将非常有帮助，因为距离矩阵将包含越来越多的 INFINITY 值，迭代越多，并且根本不需要考虑这些值。我想不出一种聪明的方法来识别全局最小值及其位置：-/

标签： c++ opencl hierarchical-clustering

【解决方案1】：

合并内存访问将计算速度提高了大约 4 倍。但是，对于我们的目的而言，这仍然会减慢速度。通过重新计算所有条目的最小值的蛮力方法是不合适的。

因此，我更改了算法，使其仅保留每行的最小值（+其位置）。在每次迭代中更改 2 行和列后，如果需要，会更新行最小值，然后通过找到行最小值的最小值来获得全局最小值。因此，如果我们有一个22500*22500 矩阵，我只需要得到22500 条目的最小值，而不是506250000。当然，这个实现需要额外的计算，但最终我们可以减少从200s（非coalescent）到59s（coalescent）一直到8s 搜索mimima的时间。

我希望这对将来的某人有所帮助:-)

【讨论】：

【解决方案2】：

WI 访问连续内存比访问分散内存效率高得多。此外，您应该先在工作组中求和，然后将其传递到全局内存。并使用单个 setArg() 整数，而不是为此目的的缓冲区。至少，你应该这样做：

__kernel void findMinValue(__global float * myArray, __global double * mins, __global int arraysize){
    int gid = get_global_id(0);
    int minloc = 0;
    float mymin = INFINITY;

    //find minimum for the kernel
    for (int i = gid ; i < arraysize; i+= get_global_size(0)){
        if (*(myArray + i) < mymin){
            mymin = *(myArray + i);
            minloc = i;
        }
    }

    *(mins + gid * 2) = minloc;
    *(mins + gid * 2 + 1) = mymin;
}

【讨论】：

我将您的大部分建议应用于内核。非常感谢你帮了很多忙。你介意看看现在的内核吗？只是看看它是否有任何主要的性能杀手。另外，如果您能帮助我使用 setArg() 命令，我将非常高兴（请参阅上面的更新）。

【解决方案3】：

如果每个工作项都遍历全局数组，则读取的合并为零。如果你改变它，让每个工作项跨过扭曲或波前大小，那么你会获得巨大的速度增益。

【讨论】：