réduction CUDA de nombreux petits, les tableaux de taille inégale

Question

Je me demande si quelqu'un pourrait suggérer la meilleure approche pour calculer la moyenne / écart-type d'un grand nombre de tableaux relativement faible, mais différentes tailles CUDA?

L'exemple de réduction parallèle du kit de développement fonctionne sur un seul tableau très grand et il semble que la taille est commodément un multiple du nombre de threads par bloc, mais mon cas est assez différent:

Conceptuellement, je cependant un grand nombre d'objets qui contiennent chacune deux composantes, upper et lower et chacun de ces composants a une x et un y de coordonnées. i.e..

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

Chacune de ces matrices est d'environ 800 en longueur, mais il varie entre les objets (non à l'intérieur d'un objet) par exemple.

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

S'il vous plaît noter ce qui précède est juste ma façon de représenter le tableau et mes données ne sont pas stockées dans structs C ou quelque chose comme ça: les données peuvent être organisées de quelque manière que j'ai besoin. Le fait est que, pour chaque réseau, la moyenne, l'écart type et éventuellement un histogramme doit être calculée et dans un but particulier, les rapports et les différences entre les matrices doivent être calculées.

Comment dois-je aller sur l'envoi de ces données au dispositif de GPU et d'organiser ma hiérarchie thread-bloc? Une idée que j'avais été à pad zéro tous mes tableaux de sorte qu'ils sont de la même longueur, et un groupe de blocs de travail sur chaque objet, mais il semble qu'il ya toutes sortes de problèmes avec cette méthode si cela fonctionnerait du tout.

Merci d'avance

Answer 1

Si je comprends bien, vous voulez réduire Object1.lower.x à un résultat, Object1.lower.y à un autre résultat et ainsi de suite. Pour tout objet donné, il y a quatre matrices à être réduite, toutes de même longueur (de l'objet).

Il existe de nombreuses approches possibles, un facteur d'influence serait le nombre total d'objets dans votre système. Je suppose que ce nombre est important.

Pour de meilleures performances que vous voulez un modèle d'accès mémoire optimale et que vous voulez éviter une divergence. Étant donné que le nombre de tableaux congruents est quatre, si vous prenez l'approche naïve de faire un tableau par fil, ci-dessous, vous souffrez non seulement d'avoir accès à la mémoire pauvre, mais aussi h / w devez vérifier à chaque itération qui se visse dans la chaîne a besoin pour exécuter la boucle -. ceux qui ne seront désactivées qui peut être inefficace (surtout si un réseau est beaucoup plus longue que les autres, par exemple)

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

Au lieu de cela si vous obtenez chaque chaîne pour résumer un tableau, alors non seulement il sera plus efficace, mais aussi votre modèle d'accès mémoire sera beaucoup mieux puisque fils adjacents lisent des éléments adjacents [1].

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

Vous devez également prendre l'alignement des tableaux en considération, aligner de préférence sur une limite de 64 octets, bien que si vous développez pour calculer la capacité 1.2 ou plus, alors ce n'est pas tout à fait aussi importante que sur les processeurs graphiques plus.

Dans cet exemple, vous lancer quatre chaînes par bloc, à savoir 128 threads, et autant de blocs que vous avez des objets.

[1] Vous dites que vous pouvez choisir ce arrangement de mémoire que vous voulez, souvent, il peut être utile pour entrelacer les tableaux de telle sorte que tableau [0] [0] est à côté de tableau [1] [0] car cela signifie que les fils adjacents peuvent fonctionner sur des réseaux adjacents et obtenir des accès coalescées. Toutefois, étant donné la longueur des tableaux n'est pas constante ce qui est probablement complexe, nécessitant un rembourrage des tableaux plus courts.

Answer 2

En tant que suivi de la réponse de Tom, je voudrais mentionner que réduction de chaîne peut être facilement mis en oeuvre par CUB .

Voici un exemple concret:

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

Dans cet exemple, un réseau de N de longueur est créé, et le résultat est la somme des éléments consécutifs 32. Donc,

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