How to find the minimum value in an array using OpenCL

user_13097835 0
Viewed 64

I am learning opencl for the first time, and I am currently modifying the shortest path finding algorithm. I know that opencl usually uses the idea of parallel computing to solve problems. So I wonder if I can also use this parallel idea when I am dealing with finding the minimum value and its position in the array? This is my previous attempt. I think that as long as the variable is the smallest, the result can be obtained regardless of whether the operation is locked or not. Unfortunately, when I use printf to view variables, although valid nodes have been judged, I can't get the correct results.

__kernel void findWay(__global int* A, __global int* B, __global int* minNode, __global int* minDis, __global int* isFinish)
{
    //A: weightMatrix , B: usedNode
    //dijkstra algorithm , src node is 0
    size_t dst = get_global_id(1);
    size_t src = get_global_id(0);
    size_t vCount = get_global_size(0);
    int index = dst * vCount + src;

    while(isFinish[0] != vCount){
        if((src == minNode[0])&&(B[dst] == 0)&&(A[index] != INT_MAX)){
            A[dst*vCount] = min(A[dst*vCount + 0],A[minNode[0]*vCount + 0] + A[index]);
        }
        minDis[0] = INT_MAX;
        barrier(CLK_GLOBAL_MEM_FENCE);
        
        //here is the bug
        if((src == 0) &&(B[dst] == 0)){
            if(minDis[0] > A[index]){
                minDis[0] = A[index];
                minNode[0] = dst;
            }
        }
        //=========
        
        barrier(CLK_GLOBAL_MEM_FENCE);
        B[minNode[0]] = 1;
        if(index == 0){
            isFinish[0]++;
        }

    }

}

In the end, I can only use a normal way to achieve this operation.

if((src == 0) &&(dst == 0)){
    for(int i = 0 ; i < vCount ;i++){
        if(B[i] == 0 && minDis[0] > A[i *vCount]){
        minDis[0] = A[i*vCount];
        minNode[0] = i;
    }
}

I would like to ask about this search process, can the looping step be omitted?

opencl
1 Answers

Horizontal operations on the parallelized array are difficult. The general approach to them is binary-tree-like kernel passes. Start with the original array, make each GPU thread load 2 neighboring elements and choose the smaller one, write that in the same array to position of the first of the two elements. Next kernel loads two elements from the list of every second element, compares the two, writes the smaller one in the first position of the two. Repeat until there is only one element left.

I will illustrate it beloe. I mark values that are not touched by the kernel anymore with *.

original array: 5|2|1|6|9|3|4|8

after 1st kernel pass: 2 *|1 *|3 *|4 *

after 2nd kernel pass: 1 * * *|3 * * *

after 3nd kernel pass: 1 * * * * * * *

smallest element is 1.

user_91789920
Related