如何使用 OpenCL 查找数组中的最小值
How to find the minimum value in an array using OpenCL
第一次学习opencl,目前正在修改最短路径查找算法。我知道opencl一般都是用并行计算的思想来解决问题。所以我想知道我在处理查找最小值及其在数组中的位置时是否也可以使用这个并行的想法?
这是我以前的尝试。我觉得只要变量最小,不管有没有加锁操作都可以得到结果。不幸的是,当我使用printf
查看变量时,虽然已经判断了有效节点,但我无法得到正确的结果。
__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]++;
}
}
}
最后只能用正常的方式来实现这个操作
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;
}
}
请问这个搜索过程,可以省略循环步骤吗?
并行化数组的水平操作很困难。对它们的一般方法是类似二叉树的内核传递。从原始数组开始,让每个 GPU 线程加载 2 个相邻元素并选择较小的一个,将其写入同一数组中两个元素中第一个元素的位置。下一个内核从每隔一个元素的列表中加载两个元素,比较两者,将较小的写入两者的第一个位置。重复直到只剩下一个元素。
我会在下面说明。我用 *.
标记内核不再触及的值
原始数组:5|2|1|6|9|3|4|8
第一次内核传递后:2 *|1 *|3 *|4 *
第二次内核传递后:1 * * *|3 * * *
第三次内核传递后:1 * * * * * * *
最小的元素是 1。
第一次学习opencl,目前正在修改最短路径查找算法。我知道opencl一般都是用并行计算的思想来解决问题。所以我想知道我在处理查找最小值及其在数组中的位置时是否也可以使用这个并行的想法?
这是我以前的尝试。我觉得只要变量最小,不管有没有加锁操作都可以得到结果。不幸的是,当我使用printf
查看变量时,虽然已经判断了有效节点,但我无法得到正确的结果。
__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]++;
}
}
}
最后只能用正常的方式来实现这个操作
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;
}
}
请问这个搜索过程,可以省略循环步骤吗?
并行化数组的水平操作很困难。对它们的一般方法是类似二叉树的内核传递。从原始数组开始,让每个 GPU 线程加载 2 个相邻元素并选择较小的一个,将其写入同一数组中两个元素中第一个元素的位置。下一个内核从每隔一个元素的列表中加载两个元素,比较两者,将较小的写入两者的第一个位置。重复直到只剩下一个元素。
我会在下面说明。我用 *.
标记内核不再触及的值原始数组:5|2|1|6|9|3|4|8
第一次内核传递后:2 *|1 *|3 *|4 *
第二次内核传递后:1 * * *|3 * * *
第三次内核传递后:1 * * * * * * *
最小的元素是 1。