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 структуры или что-то в этом роде: данные могут быть организованы так, как мне нужно. Дело в том, что для каждого массива необходимо рассчитать среднее значение, стандартное отклонение и, в конечном итоге, гистограмму, а для одного конкретного объекта необходимо рассчитать отношения и различия между массивами.

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

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

Answer 1

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

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

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

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

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

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

#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];
....