[cuda] Empuje dentro de los núcleos escritos por el usuario


1 Answers

Me gustaría brindar una respuesta actualizada a esta pregunta.

A partir de Thrust 1.8, las primitivas de empuje de CUDA pueden combinarse con la política de ejecución thrust::seq para ejecutarse secuencialmente dentro de un solo hilo CUDA (o secuencialmente dentro de un único hilo de CPU). A continuación, se informa un ejemplo.

Si desea la ejecución paralela dentro de un hilo, entonces puede considerar usar CUB que proporciona rutinas de reducción que se pueden llamar desde un threadblock, siempre que su tarjeta permita el paralelismo dinámico.

Aquí está el ejemplo con 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

Soy un novato en Thrust. Veo que todas las presentaciones y ejemplos de Thrust solo muestran código de host.

Me gustaría saber si puedo pasar un device_vector a mi propio kernel? ¿Cómo? En caso afirmativo, ¿cuáles son las operaciones permitidas dentro del código kernel / dispositivo?




Si quiere utilizar los datos asignados / procesados ​​por empuje sí puede, simplemente obtenga el puntero sin formato de los datos asignados.

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

si quieres asignar vectores de empuje en el kernel, nunca lo intenté, pero no creo que funcione, y si funciona no creo que brinde ningún beneficio.






Related



Tags

cuda   thrust