Un código elemental: suma de dos vectores

 

Conceptos que pondremos en práctica

  • Cómo alojar y liberar memoria.
  • Cómo transferir datos entre el host y el dispositivo (y viceversa).
  • Cómo invocar un kernel. Configuración y ejecución.
  • Cómo programar un kernel sencillo.

 

Cómo acceder a los ficheros fuente

Los códigos a completar para la realización de este ejercicio se encuentran en el directorio NVIDIA_GPU_Computing_SDK/C/src/lab0-vectorAdd dentro de tu cuenta de usuario.

 

Ejercicio: suma de vectores

El objetivo del ejercicio es realizar la suma de dos vectores de números reales. El resultado se almacenará en un tercer vector.
Para ello, en primer lugar deberemos realizar la reserva e inicialización de los vectores en memoria del host. Puesto que CUDA no dispone de mecanismos para diferenciar a qué espacio de memoria apuntan los punteros, es habitual seguir la siguiente convención:
  • Los nombres de punteros que apuntan a memoria de host vendrán precedidos del prefijo "h_". Así, en nuestro caso, los punteros en memoria central se declararán como h_A, h_B y h_C.
  • Los nombres de punteros que apuntan a memoria del dispositivo vendrán precedidos del prefijo "d_". Así, en nuestro caso, los punteros en memoria del dispositivo se declararán como d_A, d_B y d_C.
Observa las declaraciones de los punteros al inicio del fichero vectorAdd.cu.
El siguiente paso (tras reservar e inicializar los vectores en memoria del host), será hacer lo propio en memoria del dispositivo. Para ello, CUDA nos proporciona la siguiente función:
cudaMalloc( void ** device_ptr, size_t size );
Busca en la documentación de CUDA el funcionamiento de estas rutinas, y añádelas al fichero fuente. Por supuesto, toda reserva conlleva una posterior liberación de memoria. En CUDA, la liberación se realiza usando la siguiente función:
cudaFree( void * device_ptr );
Ya tenemos nuestros vectores reservados en memoria del dispositivo. El siguiente paso será transferir los vectores que hemos inicializado en memoria del host al dispositivo. Para ello, utilizaremos funciones de tipo (búscalas en la documentación para más información):
cudaMemcpy( void * destino, void * origen, size_t size, enum cudaMemcpyKind kind);
Añade el código necesario para transferir los dos vectores (h_A y h_B) de entrada a la GPU, y traer de vuelta el (futuro) resultado (d_C).
Sólo nos queda un detalle: sumar los valores de d_A y d_B en la GPU.
Identifica en el fichero vectorAdd.cu la función que se ejecutará en GPU (es decir, el kernel), y averigua su funcionamiento.
Ten en cuenta que el kernel va a ser ejecutado por cada uno de los hilos que creemos sobre la GPU en el momento de la invocación. En este caso, cada hilo creado va a obtener un sólo elemento del vector resultado.
Todo kernel CUDA suele presentar la siguiente estructura:
  1. Obtención de un identificador único para cada hilo dentro del grid de ejecución. En nuestro caso, este identificador se almacena en la variable i.
  2. Trabajo sobre la estructura de datos de salida (en este caso, C). SIEMPRE deberemos decidir sobre qué elementos de C va a operar cada hilo, en función de su identificador único. En nuestro caso, cada hilo i obtendrá el valor de la i-ésima posición del vector de salida. Alternativamente, podríamos hacer que cada hilo obtuviese, por ejemplo, un bloque distinto de datos del vector.
Intenta obtener un identificador único para cada hilo de ejecución, de forma que el código se ejecute correctamente. Para ello, observa el número de bloques, y de hilos por bloque, con los que se lanza a ejecución el kernel.
Observa también la sintaxis utilizada para invocar el kernel. ¿Cómo se especifica el número de bloques de hilos, y el número de hilos por bloque en el momento de la ejecución? Intenta variar estos parámetros sin que el programa deje de funcionar correctamente.

 

Modificaciones a realizar sobre el código CUDA

Accede al fichero vectorAdd.cu, que incluye tanto el programa main que se ejecuta en la CPU como el kernel que se ejecuta en la GPU. Localiza en él las etiquetas "COMPLETAR", que indican la falta de elementos de CUDA que tú debes rellenar en el orden en que se describe a continuación.

En el host (CPU)

  1. Añadir las invocaciones necesarias para reservar y liberar memoria en GPU.
  2. Transferir los vectores h_A y h_B a la GPU.
  3. Invocar al kernel CUDA. 
  4. Transferir el vector resultado (d_C) a la CPU.

En el device (GPU)

  1. Obtener un identificador único para cada hilo de ejecución en GPU (variable i), y utilizarlo para acceder al índice del vector que le corresponde computar.

 

Edición

Para modificar los archivos y completar las sentencias que faltan, puedes utilizar el editor vi incluido en el kernel del Linux o el editor nano (que resulta más sencillo al disponer en su parte inferior de un breve listado con los comandos más útiles). Para ello, lanza el correspondiente editor desde el intérprete de comandos (shell) de Linux tecleando vi nombre-de-archivo-a-editar o nano nombre-de-archivo-a-editar.

 

Compilación

Se puede utilizar el comando nvcc desde el shell del Linux. O mejor, apoyarse en el fichero Makefile que verás en el mismo directorio de trabajo donde está el código fuente. Para ello, ejecutar el comando make. Si lo que quieres es recompilar tras haber modificado algún código, asegúrate de que has borrado los ficheros generados en la compilación previa ejecutando make clean previamente a make.

 

Ejecución

La compilación de tu código a través del fichero Makefile genera el archivo ejecutable vectorAdd en el mismo directorio en el que has compilado. Para ejecutarlo, debes teclear ./vectorAdd. Alternativamente, otros Makefile generan dicho archivo ejecutable dentro de tu cuenta de usuario en la ruta NVIDIA_GPU_Computing_SDK/C/bin/linux/release, con el nombre vectorAdd. En ese caso, puedes moverte a ese directorio y lanzar el fichero ejecutable tecleando ../vectorAdd, o bien lanzar ese ejecutable desde el mismo directorio en que estabas compilando, tecleando para ello ../../bin/linux/release/vectorAdd

 

Validación: Cotejar los resultados de salida

Verás que al final del código del host, la CPU computa secuencialmente en C la misma suma de vectores que acabamos de realizar en CUDA, con objeto de comprobar que los resultados que ha recibido de la GPU son correctos. En caso afirmativo, el programa imprimirá "RESULTADOS CORRECTOS Y VALIDADOS CON LA CPU". En los ejercicios siguientes, este código C que realiza la computación en la CPU y valida los resultados se ubicará en un fichero aparte con sufijo "gold". 

 

Ayuda: Puedes echar un vistazo a las soluciones :-)

Este primer ejercicio sirve como calentamiento previo al resto del tutorial. En el caso de que no consigas finalizarlo, puedes acceder a la versión del código que contiene las soluciones, y que encontrarás en el directorio NVIDIA_GPU_Computing_SDK/C/src/lab0-vectorAdd.sol . Pero por favor, intenta esforzarte para lograr el objetivo antes de darte por vencido. Así aprenderás mucho más. ¡Suerte!