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

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

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

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

Что я здесь делаю неправильно?


c c++ reduce cuda minimum
person pymd    schedule 31.03.2014    source источник
comment
Пожалуйста, предоставьте полный компилируемый код, а не только ядро. Кроме того, пожалуйста, не используйте номера строк таким образом. Кому-то другому будет сложно скомпилировать ваш код.   -  person Robert Crovella    schedule 31.03.2014
comment
Если вы не выполняете эту работу в образовательных целях, возможно, рекомендуется использовать Thrust.   -  person Vitality    schedule 31.03.2014
comment
Я действительно не знал об этой библиотеке, спасибо, что познакомил меня с ней. Но я чувствую, что все же должен реализовать это самостоятельно, так как позже мне понадобится много настроек.   -  person pymd    schedule 31.03.2014
comment
Некоторые потоки пропускают вызов __syncthreads() в цикле for из-за операторов continue. Это намеренно? Это, вероятно, не будет делать то, что вы хотите.   -  person Heatsink    schedule 31.03.2014
comment
Большие блоки кода с меньшей вероятностью получат ответы. Вы можете упростить ядро ​​и объяснить предполагаемый шаблон обращений к памяти, выполняемых ядром.   -  person Heatsink    schedule 31.03.2014
comment
@Heatsink Я никогда не думал о операторах continue. Я проверю это и вернусь. Спасибо за предложение.   -  person pymd    schedule 01.04.2014

Ответы (1)


arrow_upward
2
arrow_downward

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

Ниже я привожу простой код для оценки минимума в массиве вместе с его индексом. Он основан на классическом примере, содержащемся в Введение в 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();
}
person Vitality    schedule 02.05.2014