Redução de CUDA de muitas pequenas matrizes de tamanho desigual

Question

Gostaria de saber se alguém poderia sugerir a melhor abordagem para calcular a média / desvio padrão de um grande número de relativamente pequeno, mas matrizes de tamanho diferente Em Cuda?

O exemplo de redução paralela no SDK funciona em uma única matriz muito grande e parece que o tamanho é convenientemente um múltiplo do número de fios por bloco, mas meu caso é bastante diferente:

Conceitualmente, no entanto, tenho um grande número de objetos que cada um contém dois componentes, upper e lower e cada um desses componentes tem um x e a y coordenada. ou seja

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

Cada uma dessas matrizes é aproximadamente 800 em comprimento, mas varia entre objetos (não dentro de um objeto), por exemplo,

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

Observe que o exposto acima é apenas a minha maneira de representar a matriz e meus dados não são armazenados em C estruturas ou qualquer coisa assim: os dados podem ser organizados da maneira que eu precisar. O ponto é que, para cada matriz, a média, o desvio padrão e, eventualmente, um histograma precisa ser calculado e, dentro de um objeto específico, proporções e diferenças entre as matrizes precisam ser calculadas.

Como devo enviar esses dados para o dispositivo GPU e organizar minha hierarquia de threadblock? Uma idéia que eu tive era zero Pad todas as minhas matrizes para que elas tenham o mesmo comprimento e tenham um grupo de blocos trabalhando em cada objeto, mas parece que existem todos os tipos de problemas com esse método, se funcionasse.

desde já, obrigado

Answer 1

Se eu entendi corretamente, você deseja reduzir o object1.lower.x para um resultado, object1.lower.y para outro resultado e assim por diante. Para qualquer objeto, existem quatro matrizes a serem reduzidas, todo o mesmo comprimento (para o objeto).

Existem muitas abordagens possíveis para isso, um fator de influência seria o número total de objetos em seu sistema. Vou assumir que esse número é grande.

Para o melhor desempenho, você deseja um padrão de acesso à memória ideal e deseja evitar a divergência. Como o número de matrizes congruentes é quatro, se você adotar a abordagem ingênua de fazer uma matriz por tópico, abaixo, não apenas sofrerá por ter acesso de memória de mau A urdidura precisa executar o loop - aqueles que não serão desativados, o que pode ser ineficiente (especialmente se uma matriz for muito mais longa que as outras, por exemplo).

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

Em vez disso, se você conseguir cada urdidura para somar uma matriz, não apenas será mais eficiente, mas também o seu padrão de acesso à memória será muito melhor, pois os threads adjacentes lerão elementos adjacentes [1].

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

Você também deve levar em consideração o alinhamento das matrizes, de preferência alinhado em um limite de 64 bytes, embora se você estiver desenvolvendo a capacidade de computação 1.2 ou superior, isso não é tão importante quanto nas GPUs mais antigas.

Neste exemplo, você lançaria quatro distorções por bloco, ou seja, 128 threads e tantos blocos quanto objetos.

1] Você diz que pode escolher o arranjo de memória que desejar, geralmente pode ser útil intercalar as matrizes para que a matriz [0] [0] seja a próxima da matriz [1] [0], pois isso significa que threads adjacentes pode operar em matrizes adjacentes e obter acessos coalescedentes. No entanto, como a duração das matrizes não é constante, isso provavelmente é complexo, exigindo preenchimento das matrizes mais curtas.

Answer 2

Como acompanhamento da resposta de Tom, gostaria de mencionar que redução de urdidura pode ser facilmente implementado por FILHOTE.

Aqui está um exemplo trabalhado:

#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;
}

Neste exemplo, uma variedade de comprimento N é criado e o resultado é a soma de 32 elementos consecutivos. Então

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