Одновременная сортировка множества массивов с помощью CUDA Thrust - PullRequest
Новая
       12

Одновременная сортировка множества массивов с помощью CUDA Thrust

2 голосов
/ 28 ноября 2011

Мне нужно отсортировать 20+ массивы, уже на GPU, каждый одинаковой длины, по одним и тем же ключам. Я не могу использовать sort_by_key() напрямую, поскольку он также сортирует ключи (делая их бесполезными для сортировки следующего массива). Вот что я попробовал вместо этого: 

thrust::device_vector<int>  indices(N); 
thrust::sequence(indices.begin(),indices.end());
thrust::sort_by_key(keys.begin(),keys.end(),indices.begin());

thrust::gather(indices.begin(),indices.end(),a_01,a_01);
thrust::gather(indices.begin(),indices.end(),a_02,a_02);
...
thrust::gather(indices.begin(),indices.end(),a_20,a_20);

Это, похоже, не работает, поскольку gather() ожидает другой массив для вывода, чем для ввода, т.е. это работает: 

thrust::gather(indices.begin(),indices.end(),a_01,o_01);
...

Однако я бы предпочел не выделять 20+ дополнительных массивов для этой задачи. Я знаю, что есть решение, использующее thrust :: tuple, thrust :: zip_iterator и thrust :: sort_by_keys (), аналогично здесь . Тем не менее, я могу объединять только до 10 массивов в кортеже, s.t. Мне нужно будет повторить ключевой вектор снова. Как бы вы справились с этой задачей?

Ответы [ 2 ]

3 голосов
/ 30 апреля 2015

Я думаю, что классический способ сортировки нескольких массивов - это так называемый последовательный подход, который использует thrust::stable_sort_by_key два раза.Вам необходимо создать вектор ключей , чтобы элементы в одном и том же массиве имели один и тот же ключ.Например: 

Elements: 10.5 4.3 -2.3 0. 55. 24. 66.
Keys:      0    0    0  1   1   1   1

В этом случае у нас есть два массива: первый с 3 элементами, а второй с 4 элементами.

Сначала необходимо вызвать thrust::stable_sort_by_keyимея значения матрицы в качестве ключей типа 

thrust::stable_sort_by_key(d_matrix.begin(),
                           d_matrix.end(),
                           d_keys.begin(),
                           thrust::less<float>());

После этого у вас будет 

Elements: -2.3 0 4.3 10.5 24. 55. 66.
Keys:       0  1  0    0   1   1   1

, что означает, что элементы массива упорядочены, а ключи - нет.Затем вам потребуется секунда, чтобы позвонить thrust::stable_sort_by_key 

thrust::stable_sort_by_key(d_keys.begin(),
                           d_keys.end(),
                           d_matrix.begin(),
                           thrust::less<int>());

, чтобы выполнить сортировку по ключам.После этого шага у вас есть 

Elements: -2.3 4.3 10.5 0 24. 55. 66.
Keys:       0   0   0   1  1   1   1

, который является конечным желаемым результатом.

Ниже приведен полный рабочий пример, который рассматривает следующую проблему: отдельно упорядочивает каждую строку матрицы.Это частный случай, когда все массивы имеют одинаковую длину, но подход работает с массивами, возможно, различной длины. 

#include <cublas_v2.h>

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/functional.h>
#include <thrust/random.h>
#include <thrust/sequence.h>

#include <stdio.h>
#include <iostream>

#include "Utilities.cuh"

/**************************************************************/
/* CONVERT LINEAR INDEX TO ROW INDEX - NEEDED FOR APPROACH #1 */
/**************************************************************/
template <typename T>
struct linear_index_to_row_index : public thrust::unary_function<T,T> {

    T Ncols; // --- Number of columns

    __host__ __device__ linear_index_to_row_index(T Ncols) : Ncols(Ncols) {}

    __host__ __device__ T operator()(T i) { return i / Ncols; }
};

/********/
/* MAIN */
/********/
int main()
{
    const int Nrows = 5;     // --- Number of rows
    const int Ncols = 8;     // --- Number of columns

    // --- Random uniform integer distribution between 10 and 99
    thrust::default_random_engine rng;
    thrust::uniform_int_distribution<int> dist(10, 99);

    // --- Matrix allocation and initialization
    thrust::device_vector<float> d_matrix(Nrows * Ncols);
    for (size_t i = 0; i < d_matrix.size(); i++) d_matrix[i] = (float)dist(rng);

    // --- Print result
    printf("Original matrix\n");
    for(int i = 0; i < Nrows; i++) {
        std::cout << "[ ";
        for(int j = 0; j < Ncols; j++)
            std::cout << d_matrix[i * Ncols + j] << " ";
        std::cout << "]\n";
    }

    /*************************/
    /* BACK-TO-BACK APPROACH */
    /*************************/
    thrust::device_vector<float> d_keys(Nrows * Ncols);

    // --- Generate row indices
    thrust::transform(thrust::make_counting_iterator(0),
                      thrust::make_counting_iterator(Nrows*Ncols),
                      thrust::make_constant_iterator(Ncols),
                      d_keys.begin(),
                      thrust::divides<int>());

    // --- Back-to-back approach
    thrust::stable_sort_by_key(d_matrix.begin(),
                               d_matrix.end(),
                               d_keys.begin(),
                               thrust::less<float>());

    thrust::stable_sort_by_key(d_keys.begin(),
                               d_keys.end(),
                               d_matrix.begin(),
                               thrust::less<int>());

    // --- Print result
    printf("\n\nSorted matrix\n");
    for(int i = 0; i < Nrows; i++) {
        std::cout << "[ ";
        for(int j = 0; j < Ncols; j++)
            std::cout << d_matrix[i * Ncols + j] << " ";
        std::cout << "]\n";
    }

    return 0;
}
1 голос
/ 29 ноября 2011

Ну, вам действительно нужно выделить только один дополнительный массив, если вы в порядке с манипулированием указателями на device_vector вместо: 

thrust::device_vector<int>  indices(N); 
thrust::sequence(indices.begin(),indices.end());
thrust::sort_by_key(keys.begin(),keys.end(),indices.begin());

thrust::device_vector<int> temp(N);
thrust::device_vector<int> *sorted = &temp;
thrust::device_vector<int> *pa_01 = &a_01;
thrust::device_vector<int> *pa_02 = &a_02;
...
thrust::device_vector<int> *pa_20 = &a_20;

thrust::gather(indices.begin(), indices.end(), *pa_01, *sorted);
pa_01 = sorted; sorted = &a_01;
thrust::gather(indices.begin(), indices.end(), *pa_02, *sorted);
pa_02 = sorted; sorted = &a_02;
...
thrust::gather(indices.begin(), indices.end(), *pa_20, *sorted);
pa_20 = sorted; sorted = &a_20;

Или что-то подобное должно работать в любом случае.Вам нужно исправить это, чтобы вектор временного устройства автоматически не освобождался, когда он выходит из области видимости - я предлагаю выделить указатели устройства CUDA с помощью cudaMalloc и затем обернуть их с device_ptr вместо использования автоматического device_vectors.

...