cuda - resueltos - Empuje dentro de los núcleos escritos por el usuario
principio de arquimedes ejercicios julioprofe (4)
Esta es una actualización de mi respuesta anterior.
A partir de Thrust 1.8.1, las primitivas de empuje CUDA pueden combinarse con la política de ejecución thrust::device para ejecutarse en paralelo dentro de un único hilo CUDA que explote el paralelismo dinámico de CUDA. 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 filas de matriz con CUDA , pero se hace de manera diferente a la publicación anterior, es decir, llamando primitivas de empuje CUDA directamente desde los núcleos escritos por el usuario. Además, el ejemplo anterior sirve para comparar el rendimiento de las mismas operaciones cuando se hace 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 ha sido evaluado en un Kepler K20c y en un Maxwell GeForce GTX 850M.
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?
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);
}
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.
Como fue escrito originalmente, Thrust es puramente una abstracción del lado del host. No puede usarse dentro de los kernels. Puede pasar la memoria del dispositivo encapsulada dentro de un thrust::device_vector a su propio núcleo de esta manera:
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() );
y también puede usar la memoria del dispositivo no asignada por empuje dentro de los algoritmos de empuje instanciando un empuje :: device_ptr con el puntero de memoria del dispositivo bare cuda.
Editado cuatro años y medio después para agregar que según la respuesta de @ JackOLantern, el empuje 1.8 agrega una política de ejecución secuencial, lo que significa que puede ejecutar versiones de subprocesos únicos de los algoritmos de empuje en el dispositivo. Tenga en cuenta que aún no es posible pasar directamente un vector de dispositivo de empuje a un kernel y que los vectores de dispositivo no se pueden usar directamente en el código del dispositivo.
Tenga en cuenta que también es posible utilizar la política de ejecución thrust::device en algunos casos para tener una ejecución de empuje paralela iniciada por un núcleo como una cuadrícula secundaria. Esto requiere una compilación / enlace de dispositivos y hardware por separado que sea compatible con el paralelismo dinámico. No estoy seguro de si esto es realmente compatible con todos los algoritmos de empuje o no, pero ciertamente funciona con algunos.