Сокращение CUDA многих небольших массивов неравномерного размера

Question

Мне интересно, может ли кто -нибудь предложить лучший подход к вычислению среднего / стандартного отклонения большого числа относительно небольших, но Массивы по -разному в Cuda?

Пример параллельного сокращения в SDK работает на одном очень большом массиве, и кажется, что размер удобно кратно количество потоков на блок, но мой случай довольно другой:

Концептуально, я, однако, имею большое количество объектов, которые содержат два компонента, upper а также lower и каждый из этих компонентов имеет x и y координировать. т.е.

upper.x, lower.x, upper.y, lower.y

Каждый из этих массивов приблизительно 800 по длине, но это варьируется между объектами (не внутри объекта), например,

Object1.lower.x = 1.1, 2.2, 3.3
Object1.lower.y = 4.4, 5.5, 6.6
Object1.upper.x = 7.7, 8.8, 9.9
Object1.upper.y = 1.1, 2.2, 3.3

Object2.lower.x = 1.0,  2.0,  3.0,  4.0, 5.0 
Object2.lower.y = 6.0,  7.0,  8.0,  9.0, 10.0
Object2.upper.x = 11.0, 12.0, 13.0, 14.0, 15.0 
Object2.upper.y = 16.0, 17.0, 18.0, 19.0, 20.0

Обратите внимание, что приведенное выше - это только мой способ представления массива, и мои данные не хранятся в C Структы или что -то в этом роде: данные могут быть организованы любым способом, который мне нужен. Дело в том, что для каждого массива среднее значение, стандартное отклонение и в конечном итоге необходимо вычислять гистограмму, и в пределах одного конкретного объекта необходимо рассчитать отношения и различия между массивами.

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

заранее спасибо

Answer 1

Если я правильно понял, вы хотите уменьшить object1.lower.x к одному результату, Object1.lower.y к другому результату и так далее. Для любого данного объекта существует четыре массива, которые должны быть уменьшены, все равные длины (для объекта).

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

Для лучшей производительности вы хотите оптимальный шаблон доступа к памяти, и вы хотите избежать дивергенции. Поскольку количество конгруэнтных массивов составляет четыре, если вы придерживаетесь наивного подхода к выполнению одного массива на поток, ниже, вы не только страдаете от плохого доступа к памяти, но и H/W должны будут проверять каждую итерацию, которая в потоке в WARP необходимо выполнить цикл - те, которые не будут отключены, что может быть неэффективным (особенно, если один массив намного длиннее других, например).

for (int i = 0 ; i < myarraylength ; i++)
    sum += myarray[i];

Вместо этого, если вы получите каждую форму, чтобы суммировать один массив, то он не только будет более эффективным, но и ваш шаблон доступа к памяти будет намного лучше, поскольку соседние потоки будут читать соседние элементы [1].

for (int i = tidwithinwarp ; i < warparraylength ; i += warpsize)
{
    mysum += warparray[i];
}
mysum = warpreduce(mysum);

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

В этом примере вы запустите четыре варианта на блок, то есть поток 128 и столько блоков, сколько у вас есть объекты.

1] Вы говорите, что вы можете выбрать любые условия памяти, которая вам нравится, часто может быть полезно для переплета массивов, чтобы массив [0] [0] рядом с массивом [1] [0], поскольку это будет означать, что соседние потоки может работать на соседних массивах и получить объединенные доступа. Однако, поскольку длина массивов не является постоянной, это, вероятно, сложная, что требует прокладки более коротких массивов.

Answer 2

Как продолжение ответа Тома, я хотел бы упомянуть, что Сокращение деформации может быть легко реализован Нога.

Вот проработанный пример:

#include <cub/cub.cuh>
#include <cuda.h>

#include "Utilities.cuh"

#include <iostream>

#define WARPSIZE    32
#define BLOCKSIZE   256

const int N = 1024;

/*************************/
/* WARP REDUCTION KERNEL */
/*************************/
__global__ void sum(const float * __restrict__ indata, float * __restrict__ outdata) {

    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x;

    unsigned int warp_id = threadIdx.x / WARPSIZE;

    // --- Specialize WarpReduce for type float. 
    typedef cub::WarpReduce<float, WARPSIZE> WarpReduce;

    // --- Allocate WarpReduce shared memory for (N / WARPSIZE) warps
    __shared__ typename WarpReduce::TempStorage temp_storage[BLOCKSIZE / WARPSIZE];

    float result;
    if(tid < N) result = WarpReduce(temp_storage[warp_id]).Sum(indata[tid]);

    if(tid % WARPSIZE == 0) outdata[tid / WARPSIZE] = result;
}

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

    // --- Allocate host side space for 
    float *h_data       = (float *)malloc(N * sizeof(float));
    float *h_result     = (float *)malloc((N / WARPSIZE) * sizeof(float));

    float *d_data;      gpuErrchk(cudaMalloc(&d_data, N * sizeof(float)));
    float *d_result;    gpuErrchk(cudaMalloc(&d_result, (N / WARPSIZE) * sizeof(float)));

    for (int i = 0; i < N; i++) h_data[i] = (float)i;

    gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));

    sum<<<iDivUp(N, BLOCKSIZE), BLOCKSIZE>>>(d_data, d_result);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    gpuErrchk(cudaMemcpy(h_result, d_result, (N / WARPSIZE) * sizeof(float), cudaMemcpyDeviceToHost));

    std::cout << "output: ";
    for(int i = 0; i < (N / WARPSIZE); i++) std::cout << h_result[i] << " ";
    std::cout << std::endl;

    gpuErrchk(cudaFree(d_data));
    gpuErrchk(cudaFree(d_result));

    return 0;
}

В этом примере массив длины N создан и результатом является сумма 32 последовательные элементы. Так

result[0] = data[0] + ... + data[31];
result[1] = data[32] + ... + data[63];
....