Получить статистику для списка номеров с использованием графического процессора

Question

У меня есть несколько списков чисел в файле.Например,

.333, .324, .123 , .543, .00054
.2243, .333, .53343 , .4434

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

.333 = 2 times in entire file
.324 = 1 time

и т. Д.

Я ищу общее решение.Не тот, который работает только на устройствах с определенными вычислительными возможностямион обрабатывает массивы произвольного размера.Т.е. обрабатывать условие, когда общее количество потоков <количество элементов </p>

Answer 1

Вот как бы я делал код в matlab

A = [333, .324, .123 , .543, .00054 .2243, .333, .53343 , .4434];
[values, locations] = unique(A);   % Find unique values and their locations
counts = diff([0, locations]);     % Find the count based on their locations

Нет простого способа сделать это простым языком, но вы можете использовать существующие библиотеки для этого.

1) Тяга

Он также поставляется с инструментарием CUDA из CUDA 4.0.

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

float _A[] = {.333, .324, .123 , .543, .00054 .2243, .333, .53343 , .4434};
int _I[] = {0, 1, 2, 3, 4, 5, 6, 7, 8};
float *A, *I; 
// Allocate memory on device and cudaMempCpy values from _A to A and _I to I
int num = 9;
// Values vector
thrust::device_vector<float>d_A(A, A+num);
// Need to sort to get same values together    
thrust::stable_sort(d_A, d_A+num);
// Vector containing 0 to num-1
thrust::device_vector<int>d_I(I, I+num);
// Find unique values and elements
thrust::device_vector<float>d_Values(num), d_Locations(num), d_counts(num);
// Find unique elements
thrust::device_vector<float>::iterator valiter;
thrust::device_vector<int>::iterator idxiter;
thrust::pair<valiter, idxiter> new_end;
new_end = thrust::unique_by_key(d_A, d_A+num, d_I, d_Values, d_Locations);

Теперь у вас есть местоположения первого экземпляра каждого уникального значения. Теперь вы можете запустить ядро, чтобы найти различия между соседними элементами от 0 до new_end в d_Locations. Вычтите окончательное значение из числа, чтобы получить счетчик для окончательного местоположения.

РЕДАКТИРОВАТЬ (Добавление кода, который был предоставлен в чате)

Вот как нужно сделать код разницы

#define MAX_BLOCKS 65535
#define roundup(A, B) = (((A) + (B) - 1) / (B))

int uniqueEle = newend.valiter – d_A;
int* count;
cudaMalloc((void**)&count, uniqueEle * sizeof(int));

int TPB = 256;
int num_blocks = roundup(uniqueEle, TPB);
int blocks_y = roundup(num_blocks, MAX_BLOCKS);
int blocks_x = roundup(num_blocks, blocks_y);
dim3 blocks(blocks_x, blocks_y);

kernel<<<blocks,TPB>>>(d_rawI, count, uniqueEle);

__global__ void kernel(float *i, int* count, int n)
{
int tx = threadIdx.x;
int bid = blockIdx.y * gridDim.x + blockIdx.x;
int id = blockDim.x * bid + tx;
__shared__ int indexes[256];

if (id < n) indexes[tx] = i[id];
__syncthreads();

if (id < n - 1) {
if (tx < 255) count[id] = indexes[tx + 1] - indexes[tx];
else count[id] = i[id + 1] - indexes[tx];
}

if (id == n - 1) count[id] = n - indexes[tx];
return;
}

2) ArrayFire

Это простая в использовании, бесплатная библиотека на основе массива.

В ArrayFire вы можете сделать следующее.

using namespace af;
float h_A[] = {.333, .324, .123 , .543, .00054 .2243, .333, .53343 , .4434};
int num = 9;
// Transfer data to device
array A(9, 1, h_A);
array values, locations, original;
// Find the unique values and locations
setunique(values, locations, original, A);
// Locations are 0 based, add 1.
// Add *num* at the end to find count of last value. 
array counts = diff1(join(locations + 1, num));

Раскрытие информации: я работаю на AccelerEyes , которая разрабатывает это программное обеспечение.

Answer 2

Чтобы ответить на последнее дополнение к этому вопросу - ядро ​​diff, которое должно завершить тягу метод , предложенный Паваном , может выглядеть примерно так:

template<int blcksz>
__global__ void diffkernel(const int *i, int* count, const int n) { 
    int id = blockDim.x * blockIdx.x + threadIdx.x; 
    int strd = blockDim.x * gridDim.x;
    int nmax = blcksz * ((n/blcksz) + ((n%blcksz>0) ? 1 : 0));

    __shared__ int indices[blcksz+1]; 

    for(; id<nmax; id+=strd) {
        // Data load
        indices[threadIdx.x] = (id < n) ? i[id] : n; 
        if (threadIdx.x == (blcksz-1)) 
            indices[blcksz] = ((id+1) < n) ? i[id+1] : n; 

        __syncthreads(); 

        // Differencing calculation
        int diff = indices[threadIdx.x+1] - indices[threadIdx.x];

        // Store
        if (id < n) count[id] = diff;

        __syncthreads(); 
    }
} 

Answer 3

вот решение:

__global__ void counter(float* a, int* b, int N)
{
    int idx = blockIdx.x*blockDim.x+threadIdx.x;

    if(idx < N)
    {
        float my = a[idx];
        int count = 0;
        for(int i=0; i < N; i++)
        {
            if(my == a[i])
                count++;
        }

        b[idx]=count;
    }
}

int main()
{

    int threads = 9;
    int blocks = 1;
    int N = blocks*threads;
    float* h_a;
    int* h_b;
    float* d_a;
    int* d_b;

    h_a = (float*)malloc(N*sizeof(float));
    h_b = (int*)malloc(N*sizeof(int));

    cudaMalloc((void**)&d_a,N*sizeof(float));
    cudaMalloc((void**)&d_b,N*sizeof(int));

    h_a[0]= .333f; 
    h_a[1]= .324f;
    h_a[2]= .123f;
    h_a[3]= .543f;
    h_a[4]= .00054f;
    h_a[5]= .2243f;
    h_a[6]= .333f;
    h_a[7]= .53343f;
    h_a[8]= .4434f;

    cudaMemcpy(d_a,h_a,N*sizeof(float),cudaMemcpyHostToDevice);

    counter<<<blocks,threads>>>(d_a,d_b,N);

    cudaMemcpy(h_b,d_b,N*sizeof(int),cudaMemcpyDeviceToHost);

    for(int i=0; i < N; i++)
    {
        printf("%f = %d times\n",h_a[i],h_b[i]);
    }

    cudaFree(d_a);
    cudaFree(d_b);
    free(h_a);
    free(h_b);
    getchar();
    return 0;
}