Эффективно найти минимум большого массива с помощью Opencl

Я работаю над реализацией алгоритма иерархической кластеризации в opencl. Для каждого шага я нашел минимальное значение в очень большом массиве (около 10 ^ 8 записей), чтобы я знал, какие элементы должны быть объединены в новый кластер. Определение минимального значения должно быть сделано 9999 раз. С моими текущими ядрами требуется около 200 секунд, чтобы найти минимальное значение (накопленное за все итерации).
Как я подошел к проблеме, разделив массив на 2560 фрагментов одинакового размера (на моем Radeon 7970 есть 2560 потоковых процессоров), чтобы найти минимум каждого фрагмента в отдельности. Я запускаю второе ядро, которое объединяет эти минимумы в глобальный минимум.

Есть ли более эффективный способ решения этой проблемы? Первоначальная идея состояла в том, чтобы ускорить HCA с помощью OpenCL, но с учетом того, что на определение минимума уходит гораздо больше времени, чем на 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. Теперь только одно ядро ​​на группу ворга записывает одно значение (то есть минимум группы). Кроме того, я увеличил общий размер работы, чтобы скрыть задержку памяти.

Эти изменения позволили сократить время, необходимое для определения минимумов, с> 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);
}
}