CUDA -Reduktion vieler kleiner, ungleich großer Arrays

Question

Ich frage mich, ob jemand den besten Ansatz zur Berechnung der mittleren / Standardabweichung einer großen Anzahl relativ kleiner aber vorschlagen könnte, aber unterschiedlich große Arrays in Cuda?

Das Beispiel für parallele Reduktion im SDK funktioniert auf einem einzigen sehr großen Array und es scheint, dass die Größe bequem ein Vielfaches der Anzahl der Threads pro Block ist, aber mein Fall ist ziemlich unterschiedlich:

Konzeptionell habe ich jedoch eine große Anzahl von Objekten, die jeweils zwei Komponenten enthalten. upper und lower und jede dieser Komponenten hat eine x und ein y Koordinate. dh

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

Jeder dieser Arrays ist ungefähr 800 in Länge, aber es variiert zwischen Objekten (nicht innerhalb eines Objekts), z.

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

Bitte beachten Sie, dass die oben genannte Art nur meine Art, das Array darzustellen, und meine Daten nicht gespeichert sind C Strukturen oder ähnliches: Die Daten können auf jede Art und Weise organisiert werden, die ich brauche. Der Punkt ist, dass für jedes Array der Mittelwert, die Standardabweichung und schließlich ein Histogramm berechnet werden müssen und innerhalb eines bestimmten Objekts Verhältnisse und Unterschiede zwischen Arrays berechnet werden müssen.

Wie soll ich diese Daten an das GPU-Gerät senden und meine Thread-Block-Hierarchie organisieren? Eine Idee, die ich hatte, war, alle meine Arrays zu null, so dass sie die gleiche Länge haben und eine Gruppe von Blöcken an jedem Objekt arbeiten, aber es scheint, dass es alle möglichen Probleme mit dieser Methode gibt, wenn sie überhaupt funktionieren würde.

Danke im Voraus

Answer 1

Wenn ich richtig verstanden habe, möchten Sie Object1.lower.x auf ein Ergebnis reduzieren, Object1.lower.y auf ein anderes Ergebnis und so weiter. Für ein bestimmtes Objekt müssen vier Arrays reduziert werden, alle gleiche Länge (für das Objekt).

Es gibt viele mögliche Ansätze dafür, ein Einflussfaktor wäre die Gesamtzahl der Objekte in Ihrem System. Ich gehe davon aus, dass diese Zahl groß ist.

Für die beste Leistung möchten Sie ein optimales Speicherzugriffsmuster und Sie möchten eine Divergenz vermeiden. Da die Anzahl der kongruenten Arrays vier beträgt, leiden Sie nicht nur unter schlechten Speicherzugriff, sondern auch die H/W müssen nicht nur einen schlechten Speicherzugriff haben, sondern auch die H/W muss nach jeder Iteration, in der sich die Threads in den Threads befassen Der Verrückte muss die Schleife ausführen - diejenigen, die nicht deaktiviert sind, was ineffizient sein kann (insbesondere wenn ein Array viel länger ist als die anderen).

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

Wenn Sie jedes Kett auf ein Array summieren, ist es nicht nur effizienter, sondern auch Ihr Speicherzugriffsmuster viel besser, da benachbarte Threads benachbarte Elemente lesen [1].

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

Sie sollten auch die Ausrichtung der Arrays berücksichtigen, vorzugsweise an einer 64 -Byte -Grenze ausgerichtet sind. Wenn Sie sich jedoch für die Berechnungsfähigkeit 1.2 oder höher entwickeln, ist dies nicht ganz so wichtig wie beim älteren GPUs.

In diesem Beispiel würden Sie vier Warps pro Block, dh 128 Threads und so viele Blöcke wie Sie Objekte haben.

1] Sie sagen, dass Sie die Speicheranordnung auswählen können, die Sie mögen. Oft kann es nützlich sein, Arrays zu verschärfen, so dass Array [0] [0] neben Array [1] [0] liegt, da dies diese benachbarten Threads bedeutet Kann mit benachbarten Arrays arbeiten und zusammenhängender Zugriffe erhalten. Da jedoch die Länge der Arrays nicht konstant ist, ist dies wahrscheinlich komplex, was die Polsterung der kürzeren Arrays erfordert.

Answer 2

Als Follow-up von Toms Antwort möchte ich das erwähnen Kettenreduzierung kann leicht implementieren durch Jungtier.

Hier ist ein bearbeitetes Beispiel:

#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 diesem Beispiel eine Reihe von Länge N wird erstellt und das Ergebnis ist die Summe von 32 aufeinanderfolgende Elemente. So

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