Нахождение минимума в массиве (но пропуская некоторые элементы), используя сокращение в CUDA

Question

Нахождение минимума в массиве (но пропуская некоторые элементы), используя сокращение в CUDA

У меня есть большой массив чисел с плавающей запятой, и я хочу узнать минимальное значение массива (игнорируя -1s везде, где присутствует), а также его индекс, используя сокращение в 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 Аргумент является идентификатором последнего блока, и number_in_last_block это количество элементов в последнем блоке (который является степенью 2). Таким образом, все блоки будут запущены 1024 потоки каждый раз, и последний блок будет использовать только number_in_last_block темы, в то время как другие будут использовать 1024 потоки.

После запуска этой функции, я ожидаю, что минимальные значения для каждого блока будут в d_Cost[blockIdx.x] и их показатели в d_index[blockIdx.x],

Я вызываю эту функцию несколько раз, каждый раз обновляя количество потоков и блоков. Во второй раз, когда я вызываю эту функцию, количество потоков теперь становится равным количеству оставшихся блоков и т. Д.

Однако вышеприведенная функция не дает мне желаемого результата. Фактически, каждый раз, когда я запускаю программу, он выдаёт разные выходные данные, то есть возвращает неверное значение как минимум во время некоторой промежуточной итерации (хотя это неверное значение довольно близко к минимуму каждый раз).

Что я здесь не так делаю?

1

c++cuda minimum reduce

Решение

Другие решения

Других решений пока нет …

Источник

Accepted Answer

Как я упоминал в своем комментарии выше, я бы рекомендовал избегать написания собственных сокращений и по возможности использовать CUDA Thrust. Это справедливо даже в том случае, когда вам необходимо настроить эти операции, причем настройка возможна путем правильной перегрузки, например, реляционных операций.

Ниже я предоставляю простой код для оценки минимума в массиве вместе с его индексом. Он основан на классическом примере, содержащемся в Введение в тягу презентация. Единственным дополнением является пропуск, как вы и просили, -1это от подсчета. Это может быть разумно сделано путем замены всех -1в массиве INT_MAXто есть максимальное представимое целое число согласно стандартам IEEE с плавающей запятой.

#include <thrust\device_vector.h>
#include <thrust\replace.h>
#include <thrust\sequence.h>
#include <thrust\reduce.h>
#include <thrust\iterator\zip_iterator.h>
#include <thrust\tuple.h>

// --- Struct returning the smallest of two tuples
struct smaller_tuple
{
__host__ __device__ thrust::tuple<int,int> operator()(thrust::tuple<int,int> a, thrust::tuple<int,int> 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<int> 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<int> 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<int> indices(d_vec_temp.size());
thrust::sequence(indices.begin(), indices.end());

// --- Setting the initial value of the search
thrust::tuple<int,int> init(d_vec_temp[0],0);

thrust::tuple<int,int> 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();
}

2