Spinta all'interno dei kernel scritti dall'utente

Question

Sono un principiante di Spinta. 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?

Answer 1

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

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

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

Answer 2

Vorrei fornire una risposta aggiornata a questa domanda.

Partendo dalla spinta 1.8, CUDA primitive di spinta può essere combinata con la politica thrust::seq di esecuzione per eseguire in sequenza all'interno di un singolo filo CUDA (o sequenzialmente all'interno di un singolo filo CPU). Di seguito, viene riportato un esempio.

Se si vuole l'esecuzione parallela all'interno di un thread, allora si può considerare l'utilizzo di CUB che fornisce le routine di riduzione che può essere chiamato dall'interno di un threadblock, a condizione che la scheda permette il parallelismo dinamico.

Ecco l'esempio con spinta

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

} 

Answer 3

Questo è un aggiornamento per la mia risposta precedente.

Partendo dalla spinta 1.8.1, CUDA primitive di spinta può essere combinata con la politica thrust::device di esecuzione per l'esecuzione in parallelo all'interno di un singolo filo CUDA CUDA sfruttando parallelismo dinamico. 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 sopra esegui riduzioni delle righe di una matrice nello stesso senso come Reduce matrix rows with CUDA, ma è fatto in modo diverso dal post sopra, cioè, chiamando CUDA primitive spinta direttamente dal scritti dall'utente kernel. Inoltre, l'esempio precedente 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.

La performance è stata valutata su un Keplero K20c e su Maxwell GeForce GTX 850M.