cuda using - Spinta all'interno dei kernel scritti dall'utente





thrust down (5)


Questo è un aggiornamento alla mia precedente risposta.

A partire da Thrust 1.8.1, i primitivi di CUDA Thust possono essere combinati con il criterio di esecuzione di thrust::device per essere eseguito in parallelo all'interno di un singolo thread CUDA che sfrutta il parallelismo dinamico CUDA. Di seguito, viene riportato un esempio.

#include <stdio.h>

#include <thrust/reduce.h>
#include <thrust/execution_policy.h>

#include "TimingGPU.cuh"
#include "Utilities.cuh"

#define BLOCKSIZE_1D    256
#define BLOCKSIZE_2D_X  32
#define BLOCKSIZE_2D_Y  32

/*************************/
/* TEST KERNEL FUNCTIONS */
/*************************/
__global__ void test1(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) {

    const unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;

    if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::seq, d_data + tid * Ncols, d_data + (tid + 1) * Ncols);

}

__global__ void test2(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) {

    const unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x;

    if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::device, d_data + tid * Ncols, d_data + (tid + 1) * Ncols);

}

/********/
/* MAIN */
/********/
int main() {

    const int Nrows = 64;
    const int Ncols = 2048;

    gpuErrchk(cudaFree(0));

//    size_t DevQueue;
//    gpuErrchk(cudaDeviceGetLimit(&DevQueue, cudaLimitDevRuntimePendingLaunchCount));
//    DevQueue *= 128;
//    gpuErrchk(cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, DevQueue));

    float *h_data       = (float *)malloc(Nrows * Ncols * sizeof(float));
    float *h_results    = (float *)malloc(Nrows *         sizeof(float));
    float *h_results1   = (float *)malloc(Nrows *         sizeof(float));
    float *h_results2   = (float *)malloc(Nrows *         sizeof(float));
    float sum = 0.f;
    for (int i=0; i<Nrows; i++) {
        h_results[i] = 0.f;
        for (int j=0; j<Ncols; j++) {
            h_data[i*Ncols+j] = i;
            h_results[i] = h_results[i] + h_data[i*Ncols+j];
        }
    }

    TimingGPU timerGPU;

    float *d_data;          gpuErrchk(cudaMalloc((void**)&d_data,     Nrows * Ncols * sizeof(float)));
    float *d_results1;      gpuErrchk(cudaMalloc((void**)&d_results1, Nrows         * sizeof(float)));
    float *d_results2;      gpuErrchk(cudaMalloc((void**)&d_results2, Nrows         * sizeof(float)));
    gpuErrchk(cudaMemcpy(d_data, h_data, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice));

    timerGPU.StartCounter();
    test1<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(d_data, d_results1, Nrows, Ncols);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    printf("Timing approach nr. 1 = %f\n", timerGPU.GetCounter());

    gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost));

    for (int i=0; i<Nrows; i++) {
        if (h_results1[i] != h_results[i]) {
            printf("Approach nr. 1; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]);
            return 0;
        }
    }

    timerGPU.StartCounter();
    test2<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(d_data, d_results1, Nrows, Ncols);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
    printf("Timing approach nr. 2 = %f\n", timerGPU.GetCounter());

    gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost));

    for (int i=0; i<Nrows; i++) {
        if (h_results1[i] != h_results[i]) {
            printf("Approach nr. 2; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]);
            return 0;
        }
    }

    printf("Test passed!\n");

}

L'esempio precedente esegue riduzioni delle righe di una matrice nello stesso senso di Ridurre le righe della matrice con CUDA , ma è fatto in modo diverso dal post precedente, vale a dire chiamando i primitivi di spinta CUDA direttamente dai kernel scritti dall'utente. Inoltre, l'esempio sopra riportato serve a confrontare le prestazioni delle stesse operazioni quando viene eseguito con due criteri di esecuzione, ovvero, thrust::seq e thrust::device . Di seguito, alcuni grafici che mostrano la differenza di prestazioni.

Le prestazioni sono state valutate su Kepler K20c e su una Maxwell GeForce GTX 850M.

Sono un principiante di Thrust. Vedo che tutte le presentazioni e gli esempi di Thrust mostrano solo il codice host.

Mi piacerebbe sapere se posso passare un device_vector al mio kernel? Come? Se sì, quali sono le operazioni consentite su di esso all'interno del codice kernel / dispositivo?




Come è stato scritto originariamente, Thrust è puramente un'astrazione dal lato host. Non può essere usato all'interno dei kernel. Puoi passare la memoria del dispositivo incapsulata all'interno di thrust::device_vector nel tuo kernel in questo modo:

thrust::device_vector< Foo > fooVector;
// Do something thrust-y with fooVector

Foo* fooArray = thrust::raw_pointer_cast( &fooVector[0] );

// Pass raw array and its size to kernel
someKernelCall<<< x, y >>>( fooArray, fooVector.size() );

e si può anche usare la memoria del dispositivo non allocata dalla spinta all'interno degli algoritmi di spinta istanziando un thrust :: device_ptr con il puntatore di memoria del dispositivo cuda nudo.

Edito quattro anni e mezzo dopo per aggiungere che, come per la risposta di @ JackOLantern, la spinta 1.8 aggiunge una politica di esecuzione sequenziale che significa che è possibile eseguire versioni a thread singolo degli alogritmi di spinta sul dispositivo. Si noti che non è ancora possibile passare direttamente un vettore del dispositivo di spinta a un kernel ei vettori del dispositivo non possono essere direttamente utilizzati nel codice del dispositivo.

Si noti che è anche possibile utilizzare la politica di esecuzione di thrust::device in alcuni casi per avere un'esecuzione parallela spinta lanciata da un kernel come griglia figlio. Ciò richiede un collegamento separato di compilazione / dispositivo e hardware che supporti il ​​parallelismo dinamico. Non sono sicuro se questo sia effettivamente supportato in tutti gli algoritmi di spinta o no, ma sicuramente funziona con alcuni.




Vorrei fornire una risposta aggiornata a questa domanda.

A partire da Thrust 1.8, le primitive di spinta di CUDA possono essere combinate con il criterio di esecuzione thrust::seq da eseguire sequenzialmente all'interno di un singolo thread CUDA (o sequenzialmente all'interno di un singolo thread della CPU). Di seguito, viene riportato un esempio.

Se si desidera l'esecuzione parallela all'interno di un thread, è possibile prendere in considerazione l'utilizzo di CUB che fornisce routine di riduzione che possono essere richiamate da un threadblock, a condizione che la scheda abiliti il ​​parallelismo dinamico.

Ecco l'esempio di Thrust

#include <stdio.h>

#include <thrust/reduce.h>
#include <thrust/execution_policy.h>

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

__global__ void test(float *d_A, int N) {

    float sum = thrust::reduce(thrust::seq, d_A, d_A + N);

    printf("Device side result = %f\n", sum);

}

int main() {

    const int N = 16;

    float *h_A = (float*)malloc(N * sizeof(float));
    float sum = 0.f;
    for (int i=0; i<N; i++) {
        h_A[i] = i;
        sum = sum + h_A[i];
    }
    printf("Host side result = %f\n", sum);

    float *d_A; gpuErrchk(cudaMalloc((void**)&d_A, N * sizeof(float)));
    gpuErrchk(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));

    test<<<1,1>>>(d_A, N);

}



Se si intende utilizzare i dati allocati / elaborati dalla spinta sì, è sufficiente ottenere il puntatore raw dei dati allocati.

int * raw_ptr = thrust::raw_pointer_cast(dev_ptr);

se si desidera allocare i vettori di spinta nel kernel non ho mai provato, ma non penso che funzionerà e anche se funziona, non penso che fornirà alcun vantaggio.




È probabile che queste informazioni siano valide solo per calcolare capabality 1.x, o cuda 2.0. Architetture più recenti e cuda 3.0 hanno un accesso alla memoria globale più sofisticato e in effetti i "carichi globali coalizzati" non sono nemmeno profilati per questi chip.

Inoltre, questa logica può essere applicata alla memoria condivisa per evitare conflitti bancari.

Una transazione di memoria a coalescenza è quella in cui tutti i thread in una memoria globale di accesso a half-warp allo stesso tempo. Questo è troppo semplice, ma il modo corretto per farlo è di avere thread consecutivi che accedono a indirizzi di memoria consecutivi.

Quindi, se i thread 0, 1, 2 e 3 leggono la memoria globale 0x0, 0x4, 0x8 e 0xc, dovrebbe essere una lettura a coalescenza.

In un esempio di matrice, tenere presente che si desidera che la matrice risieda linearmente in memoria. Puoi farlo come vuoi, e l'accesso alla tua memoria dovrebbe riflettere il modo in cui è strutturata la tua matrice. Quindi, la matrice 3x4 sotto

0 1 2 3
4 5 6 7
8 9 a b

potrebbe essere fatto una riga dopo l'altra, in questo modo, in modo che (r, c) sia mappato alla memoria (r * 4 + c)

0 1 2 3 4 5 6 7 8 9 a b

Supponi di aver bisogno di accedere all'elemento una volta e di avere quattro thread. Quali thread saranno utilizzati per quale elemento? Probabilmente entrambi

thread 0:  0, 1, 2
thread 1:  3, 4, 5
thread 2:  6, 7, 8
thread 3:  9, a, b

o

thread 0:  0, 4, 8
thread 1:  1, 5, 9
thread 2:  2, 6, a
thread 3:  3, 7, b

Che è migliore? Quale risulterà in letture coalescenti e quali no?

Ad ogni modo, ogni thread fa tre accessi. Diamo un'occhiata al primo accesso e vediamo se i thread accedono alla memoria consecutivamente. Nella prima opzione, il primo accesso è 0, 3, 6, 9. Non consecutivi, non coalizzati. La seconda opzione, è 0, 1, 2, 3. Consecutiva! Fusero! Sìì!

Il modo migliore è probabilmente scrivere il tuo kernel e poi tracciarlo per vedere se hai carichi e depositi globali non coalescati.







cuda thrust