Me gustaría limpiar los parámetros de los núcleos CUDA en mi proyecto.
Ahora, un kernel necesita 3 matrices uint32_t , lo que conduce a un código bastante feo: (id significa que la identificación del hilo global y valX es un valor arbitrario)
__global__ void some_kernel(uint32_t * arr1, uint32_t * arr2, uint32_t * arr3){arr1[id] = val1; arr2[id] = val2; arr3[id] = val3;}Me gustaría rodear todas esas matrices con una estructura:
typedef struct S{uint_32_t arr1, uint_32_t arr2, uint_32_t arr3, uint32_t size} S;donde size denota la longitud de cada arrX dentro de la estructura.
Lo que me gustaría tener, es algo como:
__global__ void some_kernel(S * s){s->arr1[id] = val1; s->arr2[id] = val2; s->arr3[id] = val3;}¿Cómo se verían cudaMalloc y cudaMemcpy correspondientes para una estructura como esta? ¿Hay algún inconveniente de rendimiento de esto, que no estoy viendo todavía?
¡Gracias por adelantado!
Tienes al menos dos opciones. Talonmies ya ofreció una excelente opción, pero le presentaré el enfoque de "aprender de la manera difícil".
Primero, su definición de estructura:
typedef struct S { uint32_t *arr1; uint32_t *arr2; uint32_t *arr3; uint32_t size; } S;... y la definición del núcleo (con alguna variable global, pero no es necesario seguir ese patrón):
const int size = 10000; __global__ void some_kernel(S *s) { int id = blockIdx.x * blockDim.x + threadIdx.x; if (id < size) { s->arr1[id] = 1; // val1 s->arr2[id] = 2; // val2 s->arr3[id] = 3; // val3 } } Tenga en cuenta que if lo protege de salirse de los límites.
Luego, venimos con alguna función que prepara datos, ejecuta kernel e imprime algún resultado. La primera parte es la asignación de datos:
uint32_t *host_arr1, *host_arr2, *host_arr3; uint32_t *dev_arr1, *dev_arr2, *dev_arr3; // Allocate and fill host data host_arr1 = new uint32_t[size](); host_arr2 = new uint32_t[size](); host_arr3 = new uint32_t[size](); // Allocate device data cudaMalloc((void **) &dev_arr1, size * sizeof(*dev_arr1)); cudaMalloc((void **) &dev_arr2, size * sizeof(*dev_arr2)); cudaMalloc((void **) &dev_arr3, size * sizeof(*dev_arr3)); // Allocate helper struct on the device S *dev_s; cudaMalloc((void **) &dev_s, sizeof(*dev_s));No es nada especial, solo asigna tres matrices y estructura. Lo que parece más interesante es cómo manejar la copia de dichos datos en el dispositivo:
// Copy data from host to device cudaMemcpy(dev_arr1, host_arr1, size * sizeof(*dev_arr1), cudaMemcpyHostToDevice); cudaMemcpy(dev_arr2, host_arr2, size * sizeof(*dev_arr2), cudaMemcpyHostToDevice); cudaMemcpy(dev_arr3, host_arr3, size * sizeof(*dev_arr3), cudaMemcpyHostToDevice); // NOTE: Binding pointers with dev_s cudaMemcpy(&(dev_s->arr1), &dev_arr1, sizeof(dev_s->arr1), cudaMemcpyHostToDevice); cudaMemcpy(&(dev_s->arr2), &dev_arr2, sizeof(dev_s->arr2), cudaMemcpyHostToDevice); cudaMemcpy(&(dev_s->arr3), &dev_arr3, sizeof(dev_s->arr3), cudaMemcpyHostToDevice);Además de la copia ordinaria de la matriz que notó, también es necesario "vincularlos" con la estructura. Para eso necesitas pasar una dirección de puntero. Como resultado, solo se copian estos punteros.
Siguiente llamada al kernel, copie los datos nuevamente al host e imprima los resultados:
// Call kernel some_kernel<<<10000/256 + 1, 256>>>(dev_s); // block size need to be a multiply of 256 // Copy result to host: cudaMemcpy(host_arr1, dev_arr1, size * sizeof(*host_arr1), cudaMemcpyDeviceToHost); cudaMemcpy(host_arr2, dev_arr2, size * sizeof(*host_arr2), cudaMemcpyDeviceToHost); cudaMemcpy(host_arr3, dev_arr3, size * sizeof(*host_arr3), cudaMemcpyDeviceToHost); // Print some result std::cout << host_arr1[size-1] << std::endl; std::cout << host_arr2[size-1] << std::endl; std::cout << host_arr3[size-1] << std::endl;Tenga en cuenta que en cualquier código serio siempre debe verificar si hay errores en las llamadas a la API de CUDA.