使用CUDA中的减少来查找数组中的最小值(但跳过某些元素)

我有一大堆浮点数,我想找出数组的最小值(忽略-1 s,无论在哪里)及其索引,使用CUDA中的减少。 我已经编写了以下代码来执行此操作,在我看来应该可以工作:

  __global__ void get_min_cost(float *d_Cost,int n,int *last_block_number,int *number_in_last_block,int *d_index){ int tid = threadIdx.x; int myid = blockDim.x * blockIdx.x + threadIdx.x; int s; if(result == (*last_block_number)-1){ s = (*number_in_last_block)/2; }else{ s = 1024/2; } for(;s>0;s/=2){ if(myid+s>=n) continue; if(tid<s){ if(d_Cost[myid+s] == -1){ continue; }else if(d_Cost[myid] == -1 && d_Cost[myid+s] != -1){ d_Cost[myid] = d_Cost[myid+s]; d_index[myid] = d_index[myid+s]; }else{ // both not -1 if(d_Cost[myid]<=d_Cost[myid+s]) continue; else{ d_Cost[myid] = d_Cost[myid+s]; d_index[myid] = d_index[myid+s]; } } } else continue; __syncthreads(); } if(tid==0){ d_Cost[blockIdx.x] = d_Cost[myid]; d_index[blockIdx.x] = d_index[myid]; } return; } 

last_block_number参数是最后一个块的id, number_in_last_block是最后一个块中的元素数( 2的幂)。 因此,所有块每次都会启动1024线程,最后一个块只使用number_in_last_block线程,而其他块将使用1024线程。

运行此函数后,我希望每个块的最小值在d_Cost[blockIdx.x]并且它们的索引在d_index[blockIdx.x]

我多次调用此函数,每次更新线程和块的数量。 第二次调用此函数时,线程数现在等于剩余的块数等。

但是,上述function并没有给我所需的输出。 实际上,每次运行程序时它都会给出不同的输出,即在某些中间迭代期间它返回一个不正确的值作为最小值(尽管每次不正确的值非常接近最小值)。

我在这做错了什么?

正如我在上面的评论中提到的,我建议尽可能避免编写自己的减少并使用CUDA Thrust。 即使在您需要自定义这些操作的情况下也是如此,通过适当的重载(例如,关系操作)可以实现自定义。

下面我提供一个简单的代码来评估数组中的最小值及其索引。 它基于“ 推力简介”中的经典示例。 唯一的补充就是按照你的要求跳过计数的-1 。 这可以通过用INT_MAX替换数组中的所有-1来合理地完成,即根据IEEE浮点标准的最大可表示整数。

 #include  #include  #include  #include  #include  #include  // --- Struct returning the smallest of two tuples struct smaller_tuple { __host__ __device__ thrust::tuple operator()(thrust::tuple a, thrust::tuple b) { if (a < b) return a; else return b; } }; void main() { const int N = 20; const int large_value = INT_MAX; // --- Setting the data vector thrust::device_vector d_vec(N,10); d_vec[3] = -1; d_vec[5] = -2; // --- Copying the data vector to a new vector where the -1's are changed to FLT_MAX thrust::device_vector d_vec_temp(d_vec); thrust::replace(d_vec_temp.begin(), d_vec_temp.end(), -1, large_value); // --- Creating the index sequence [0, 1, 2, ... ) thrust::device_vector indices(d_vec_temp.size()); thrust::sequence(indices.begin(), indices.end()); // --- Setting the initial value of the search thrust::tuple init(d_vec_temp[0],0); thrust::tuple smallest; smallest = thrust::reduce(thrust::make_zip_iterator(thrust::make_tuple(d_vec_temp.begin(), indices.begin())), thrust::make_zip_iterator(thrust::make_tuple(d_vec_temp.end(), indices.end())), init, smaller_tuple()); printf("Smallest %i %i\n",thrust::get<0>(smallest),thrust::get<1>(smallest)); getchar(); }