riduzione CUDA di tanti piccoli array, dimensionate in modo ineguale

Question

Mi chiedevo se qualcuno potrebbe suggerire il miglior approccio al calcolo del / deviazione standard media di un gran numero di relativamente piccola, ma array di diverse dimensioni in CUDA?

L'esempio parallela riduzione SDK funziona su un singolo allineamento molto grande e sembra che il formato è convenientemente un multiplo del numero di fili per blocco, ma il mio caso è piuttosto diversa:

Concettualmente, ho comunque avere un gran numero di oggetti che contengono ciascuna due componenti, upper e lower e ciascuno di questi componenti ha un x e una coordinata y. cioè.

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

Ciascuna di queste matrici è approssimativamente 800 lunghezza ma varia tra oggetti (non all'interno di un oggetto) per esempio.

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

Si prega di notare quanto sopra è solo il mio modo di rappresentare la matrice e il mio dati non sono memorizzati in struct C o qualcosa di simile: i dati possono essere organizzati in qualsiasi modo ho bisogno. Il punto è che, per ogni matrice, la media, la deviazione standard e alla fine un istogramma deve essere calcolato e entro uno scopo particolare, i rapporti e le differenze tra le matrici devono essere calcolate.

Come devo fare per inviare questi dati al dispositivo di GPU e di organizzare la mia gerarchia filo-block? Un'idea che avevo era al pad a zero tutte le mie matrici in modo che siano della stessa lunghezza, e un gruppo di blocchi che lavorano su ciascun oggetto, ma sembra che ci sono tutti i tipi di problemi con tale metodo se funzionasse affatto.

Grazie in anticipo

Answer 1

Se ho capito bene, si vuole ridurre Object1.lower.x a un risultato, Object1.lower.y ad un altro risultato e così via. Per ogni dato oggetto ci sono quattro matrici essere ridotte, tutte di uguale lunghezza (per l'oggetto).

Ci sono molte possibili approcci a questo, un fattore che influenza sarebbe il numero totale di oggetti nel vostro sistema. Darò per scontato che il numero è grande.

Per ottimizzare le prestazioni si desidera un modello ottimale di accesso alla memoria e si vuole evitare di divergenza. Poiché il numero di matrici congruenti è quattro, se si prende l'approccio naive di fare un array per thread, sotto, non solo si soffre di avere scarso accesso alla memoria ma anche il h / w necessario controllare ogni iterazione quali fili in l'ordito necessità di eseguire il ciclo -. quelli che non verrà disattivata che può essere inefficiente (soprattutto se un array è molto più lungo rispetto agli altri, per esempio)

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

Se invece si ottiene ogni ordito sommare un array, allora non solo sarà più efficiente, ma anche il vostro modello di accesso alla memoria sarà molto meglio da quando le discussioni adiacenti leggeranno elementi adiacenti [1].

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

Si dovrebbe anche prendere l'allineamento degli array in considerazione, preferibilmente allineare su un limite di 64 byte anche se si sta sviluppando per la capacità di elaborazione 1.2 o superiore, allora questo non è così importante come sulle GPU più vecchie.

In questo esempio si avvierà quattro orditi per blocco, cioè 128 thread, ed altrettanti blocchi sono presenti oggetti.

[1] Tu dici che si può scegliere qualsiasi arrangiamento di memoria che ti piace, spesso può essere utile interleave array in modo che array [0] [0] si trova accanto al array [1] [0] dal momento che questo significherà che filetti adiacenti possono operare su matrici adiacenti e ottenere accessi coalescenza. Tuttavia, poiché la lunghezza delle matrici non è costante questo è probabilmente complessa, che richiede imbottiture degli array più brevi.

Answer 2

Come follow-up di risposta di Tom, vorrei ricordare che riduzione ordito può essere facilmente implementata da CUB .

Ecco un esempio pratico:

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

In questo esempio, viene creato un array di lunghezza N e il risultato è la somma di 32 elementi consecutivi. Quindi

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