Empuje dentro de los núcleos escritos por el usuario

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?

Answer 1

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

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

si desea asignar vectores de empuje en el núcleo Nunca he intentado, pero no creo que va a funcionar y también si funciona no creo que proporcionará ningún beneficio.

Answer 2

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

A partir de Thrust 1.8, las primitivas CUDA Thrust se pueden combinar con la política de ejecución thrust::seq para ejecutarlas secuencialmente dentro de una sola hebra CUDA (o secuencialmente dentro de una única hebra de CPU). A continuación, se informa un ejemplo.

Si desea la ejecución en paralelo 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 de empuje con

#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

Esta es una actualización de mi respuesta anterior.

partir de empuje 1.8.1, primitivas de empuje CUDA puede combinarse con la política thrust::device ejecución para funcionar en paralelo dentro de un solo hilo CUDA explotando CUDA paralelismo dinámico. A continuación, se informa un ejemplo.

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

} 

El ejemplo anterior realiza reducciones de las filas de una matriz en el mismo sentido que Reduce matrix rows with CUDA, pero se realiza de manera diferente del poste anteriormente, es decir, llamando primitivas de empuje CUDA directamente a partir de granos escrito por el usuario. Además, el ejemplo anterior sirve para comparar el rendimiento de las mismas operaciones cuando se realiza con dos políticas de ejecución, a saber, thrust::seq y thrust::device.A continuación, algunos gráficos que muestran la diferencia en el rendimiento.

El rendimiento se ha evaluado en un K20c Kepler y Maxwell en una GeForce GTX 850M.