[cuda] Spinta all'interno dei kernel scritti dall'utente


Answers

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

}
Question

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.







Tags

cuda   thrust