Подсчет вхождений чисел в массиве CUDA

Question

У меня есть массив целых чисел без знака, хранящихся в графическом процессоре с CUDA (обычно 1000000 элементов).Я хотел бы посчитать вхождение каждого числа в массиве.Есть только несколько отдельных чисел (около 10), но эти числа могут варьироваться от 1 до 1000000.Около 9/10 числа чисел 0, мне не нужно их количество.Результат выглядит примерно так:

58458 -> 1000 occurrences
15 -> 412 occurrences

У меня есть реализация, использующая atomicAdd s, но она слишком медленная (многие потоки пишут на один и тот же адрес).Кто-нибудь знает быстрый / эффективный метод?

Answer 1

Вы можете реализовать гистограмму, сначала отсортировав числа, а затем выполнив сокращение на ключ.

Самый простой способ - использовать thrust::sort, а затем thrust::reduce_by_key. Это также часто намного быстрее, чем специальное бининг, основанное на атомарности. Вот пример .

Answer 2

Я сравниваю два подхода, предложенных в повторяющемся вопросе Количество отсчетов , а именно

1. Используя thrust::counting_iterator и thrust::upper_bound, следуя примеру гистограммы Thrust;
2. Использование thrust::unique_copy и thrust::upper_bound.

Ниже приведен полностью обработанный пример.

#include <time.h>       // --- time
#include <stdlib.h>     // --- srand, rand
#include <iostream>

#include <thrust\host_vector.h>
#include <thrust\device_vector.h>
#include <thrust\sort.h>
#include <thrust\iterator\zip_iterator.h>
#include <thrust\unique.h>
#include <thrust/binary_search.h>
#include <thrust\adjacent_difference.h>

#include "Utilities.cuh"
#include "TimingGPU.cuh"

//#define VERBOSE
#define NO_HISTOGRAM

/********/
/* MAIN */
/********/
int main() {

    const int N = 1048576;
    //const int N = 20;
    //const int N = 128;

    TimingGPU timerGPU;

    // --- Initialize random seed
    srand(time(NULL));

    thrust::host_vector<int> h_code(N);

    for (int k = 0; k < N; k++) {
        // --- Generate random numbers between 0 and 9
        h_code[k] = (rand() % 10);
    }

    thrust::device_vector<int> d_code(h_code);
    //thrust::device_vector<unsigned int> d_counting(N);

    thrust::sort(d_code.begin(), d_code.end());

    h_code = d_code;

    timerGPU.StartCounter();

#ifdef NO_HISTOGRAM
    // --- The number of d_cumsum bins is equal to the maximum value plus one
    int num_bins = d_code.back() + 1;

    thrust::device_vector<int> d_code_unique(num_bins);
    thrust::unique_copy(d_code.begin(), d_code.end(), d_code_unique.begin());
    thrust::device_vector<int> d_counting(num_bins);
    thrust::upper_bound(d_code.begin(), d_code.end(), d_code_unique.begin(), d_code_unique.end(), d_counting.begin());  
#else
    thrust::device_vector<int> d_cumsum;

    // --- The number of d_cumsum bins is equal to the maximum value plus one
    int num_bins = d_code.back() + 1;

    // --- Resize d_cumsum storage
    d_cumsum.resize(num_bins);

    // --- Find the end of each bin of values - Cumulative d_cumsum
    thrust::counting_iterator<int> search_begin(0);
    thrust::upper_bound(d_code.begin(), d_code.end(), search_begin, search_begin + num_bins, d_cumsum.begin());

    // --- Compute the histogram by taking differences of the cumulative d_cumsum
    //thrust::device_vector<int> d_counting(num_bins);
    //thrust::adjacent_difference(d_cumsum.begin(), d_cumsum.end(), d_counting.begin());
#endif

    printf("Timing GPU = %f\n", timerGPU.GetCounter());

#ifdef VERBOSE
    thrust::host_vector<int> h_counting(d_counting);
    printf("After\n");
    for (int k = 0; k < N; k++) printf("code = %i\n", h_code[k]);
#ifndef NO_HISTOGRAM
    thrust::host_vector<int> h_cumsum(d_cumsum);
    printf("\nCounting\n");
    for (int k = 0; k < num_bins; k++) printf("element = %i; counting = %i; cumsum = %i\n", k, h_counting[k], h_cumsum[k]);
#else
    thrust::host_vector<int> h_code_unique(d_code_unique);

    printf("\nCounting\n");
    for (int k = 0; k < N; k++) printf("element = %i; counting = %i\n", h_code_unique[k], h_counting[k]);
#endif
#endif
}

Первый подход показал, чтобыстрый.На карте NVIDIA GTX 960 у меня было следующее время для ряда N = 1048576 элементов массива:

First approach: 2.35ms
First approach without thrust::adjacent_difference: 1.52
Second approach: 4.67ms

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

Answer 3

Полагаю, вы можете найти помощь в примерах CUDA, в частности, в гистограммах. Они являются частью вычислений на GPU SDK. Вы можете найти его здесь http://developer.nvidia.com/cuda-cc-sdk-code-samples#histogram. У них даже есть технический документ, объясняющий алгоритмы.