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




thrust down (5)

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?


Answers

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.


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.


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);

}

Sono lo sviluppatore di VexCL , ma mi piace molto quello che Kyle Lutz , l'autore di Boost.Compute , ha Boost.Compute sullo stesso argomento sulla mailing list di Boost . In breve, dal punto di vista dell'utente Thrust , Boost.Compute, AMD's Bolt e probabilmente l' AMP C ++ di Microsoft implementano tutte un'API simile a STL, mentre VexCL è una libreria basata su modelli di espressione che è più vicina a Eigen in natura. Credo che la principale differenza tra le librerie simili a STL sia la loro portabilità:

  1. Thrust supporta solo GPU NVIDIA, ma può funzionare anche su CPU tramite i suoi backend OpenMP e TBB.
  2. Bolt utilizza le estensioni AMD per OpenCL che sono disponibili solo su GPU AMD. Fornisce anche AMP Microsoft C ++ e backbb Intel TBB.
  3. L'unico compilatore che supporta Microsoft C ++ AMP è Microsoft Visual C ++ (anche se è in corso il lavoro su come portare AMP Beyond di Windows C ++ ).
  4. Boost.Compute sembra essere la soluzione più portabile di questi, poiché è basata su OpenCL standard.

Ancora una volta, tutte queste librerie stanno cercando di implementare un'interfaccia simile a STL, quindi hanno un'applicabilità molto ampia. VexCL è stato sviluppato pensando al computing scientifico. Se Boost.Compute è stato sviluppato un po 'prima, potrei probabilmente basare VexCL su di esso :). Un'altra libreria per il calcolo scientifico che vale la pena guardare è ViennaCL , una libreria algebra lineare open source gratuita per calcoli su architetture a molti core (GPU, MIC) e CPU multi-core. Dai un'occhiata a [1] per il confronto di VexCL, ViennaCL, CMTL4 e Thrust per quel campo.

Per quanto riguarda la citata incapacità degli sviluppatori di Thrust di aggiungere un backend OpenCL: Thrust, VexCL e Boost.Compute (non ho familiarità con gli interni di altre librerie) usano tutti tecniche di metaprogrammazione per fare ciò che fanno. Ma poiché CUDA supporta i modelli C ++, il lavoro degli sviluppatori di Thrust è probabilmente un po 'più semplice: devono scrivere metaprogrammi che generano programmi CUDA con l'aiuto del compilatore C ++. Gli autori di VexCL e Boost.Compute scrivono metaprogrammi che generano programmi che generano codice sorgente OpenCL. Dai un'occhiata alle slides cui ho cercato di spiegare come viene implementato VexCL. Quindi sono d'accordo che il design attuale di Thrust proibisce loro di aggiungere un back-end OpenCL.

[1] Denis Demidov, Karsten Ahnert, Karl Rupp, Peter Gottschling, Programmazione CUDA e OpenCL: un case study che usa moderne librerie C ++ , SIAM J. Sci. Comput., 35 (5), C453-C472. (è disponibile anche una versione di arXiv ).

Aggiornamento: @gnzlbg ha commentato che non vi è alcun supporto per i funtori e i lambda C ++ nelle librerie basate su OpenCL. Infatti, OpenCL è basato su C99 ed è compilato da sorgenti memorizzate in stringhe in fase di esecuzione, quindi non esiste un modo semplice per interagire pienamente con le classi C ++. Ma per essere onesti, le librerie basate su OpenCL supportano in qualche misura le funzioni basate sull'utente e persino lambda.

Detto questo, le librerie basate su CUDA (e potrebbero essere C ++ AMP) hanno un ovvio vantaggio del compilatore compilato in tempo reale (puoi dirlo anche tu?), Quindi l'integrazione con il codice utente può essere molto più stretta.