cuda - down - nvidia sdk




Spinta all'interno dei kernel scritti dall'utente (3)

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?


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.


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.