CUDA kernel函数中有依赖关系的while循环嵌套问题

我有一个while循环嵌套问题,谁能帮我解答一下?

global void kdtree_bbf_kernel(struct kd_node* kd_root,struct kd_node* kd_node_num,struct feature* feat,int n2,int k,int num,int max_nn_chks, struct min_pq min_pq,struct kd_node* expl,struct kd_node* unexpl)
{
int i,j;
int tid=threadIdx.x+blockIdx.xblockDim.x;
if(tid<n2)
{
int t=0;
min_pq[tid].n=0;
min_pq[tid].nallocd=MINPQ_INIT_NALLOCD;
minpq_insert(&min_pq[tid],kd_root,0);
while(min_pq[tid].n>0 && t<max_nn_chks)
{
minpq_extract_min(&min_pq[tid],&expl[tid]);
if(! &expl[tid])
{
printf(“Warning: PQ unexpectedly empty.\n”);
return;
}
explore_to_leaf(&min_pq[tid],&expl[tid],&unexpl[tid],&feat[tid],kd_node_num);
if(! &expl[tid])
{
printf(“Warning: PQ unexpectedly empty.\n”);
return;
}
t++;
}
tid+=blockDim.x
gridDim.x;
}
}
device int minpq_insert(struct min_pq* min_pq,struct kd_node* data,int key)
{
int n=min_pq->n;
min_pq->pq_array[n].key=INT_MAX;
min_pq->pq_array[n].data[0]=data[0];
decrease_pq_node_key( min_pq->pq_array, min_pq->n, key );
min_pq->n++;
return 0;
}
device void minpq_extract_min( struct min_pq* min_pq,struct kd_node* expl )
{
if( min_pq->n < 1 )
{
return;
}
expl[0] = min_pq->pq_array[0].data[0];
min_pq->n–;
min_pq->pq_array[0] = min_pq->pq_array[min_pq->n];
restore_minpq_order( min_pq->pq_array, 0, min_pq->n );
}
device void explore_to_leaf(struct min_pq* min_pq,struct kd_node* expl,struct kd_node* unexpl,struct feature* feat,struct kd_node* kd_node_num)
{
double kv;
int ki;
while( expl && ! expl->leaf )
{
ki=expl[0].ki;
kv=expl[0].kv;
if(ki>=feat->d)
{
return;
}
if(feat->descr[ki]<=kv)
{
unexpl[0] = kd_node_num[expl[0].index_right];
expl[0] = kd_node_num[expl[0].index_left];
}
else
{
unexpl[0] = kd_node_num[expl[0].index_left];
expl[0] = kd_node_num[expl[0].index_right];
}
if(minpq_insert( min_pq,unexpl,ABS(kv-feat->descr[ki] ) ) )
{
return;
}
}
}
如果去掉explore_to_leaf函数中的 if(minpq_insert( min_pq,unexpl,ABS(kv-feat->descr[ki] ) ) ) { return; }kernel中的while只循环一次,运行时间为30多ms,如果加上if(minpq_insert( min_pq,unexpl,ABS(kv-feat->descr[ki] ) ) ) { return; },那么它就和kernel函数中的while循环构成循环嵌套,kernel中的while循环多次,时间就会升到900多毫秒,还没有在CPU上串行快,这是怎么回事啊?应该怎么改啊?

LZ您好:

1:任何算法的优化和改进都只能在保证原逻辑等价的前提下实现。您在修改代码的时候需要保证修改前后的等价性,而不是拿一个循环多次的部分和一个只计算单次的部分比较。

2:并不保证任何在GPU上的算法实现都比相应的CPU版本快。这个和算法是否适合GPU有关,也和具体实现的写法有关。

3:您需要在吃透您算法的基础上,考虑该算法是否适合GPU实现,如果适合,再考虑出一个具体的CUDA实现,而我无法告知您这一点。

大致如此,祝您编码顺利~

忙着考试,现在才来谢谢ice版主的回复,我会再把代码的思路整理一下。祝工作顺利~

欢迎您常来论坛,祝您考试顺利~