Modelo de Procesamiento

Los kernels CUDA son ejecutados en la GPU siguiendo un modelo de procesamiento especial que permite ocultar la latencia de las diversas operaciones que se llevan a cabo para realmente conseguir aceleración mediante paralelismo. Este modelo de procesamiento conlleva una serie de pormenores en cuanto al lanzamiento de los kernels, la planificación de su ejecución con división en malla, bloque e hilos, sus limitaciones y consideraciones de rendimiento que implica.

Como ya enunciamos en la sesión anterior, un kernel es una función invocada por la CPU pero cuya ejecución se produce en la GPU en forma de una gran cantidad de hilos siguiendo un patrón SIMT de forma paralela. A esta definición es necesario sumarle el hecho de que, por naturaleza, el lanzamiento de un kernel es asíncrono. La CPU simplemente emite las órdenes a la GPU depositándolas en un buffer en forma de cola circular. Esto quiere decir que una vez la CPU ordene el lanzamiento, el hilo en el host puede seguir su flujo de ejecución normal (a no ser que alcance una función de sincronización con la GPU o una operación bloqueante en el buffer de comandos a la GPU). Esta particularidad da pie al uso de streams para solapar cómputo y otras operaciones bloqueantes y de larga duración como las copias en memoria.

Sintaxis y Limitaciones del Lanzamiento de Kernels

Los kernels pueden ser invocados o lanzados utilizando tanto las funciones propias del driver de CUDA como el runtime simplificado. Utilizando las funciones de CUDART deberemos utilizar una sintaxis especial tal y como se muestra a continuación:

kernel <<< malla, bloque, memoria_compartida, stream >>> (parámetro1, parámetro2, ..., parámetroN);

En dicha llamada, kernel es la función __global__ llamada desde la CPU pero ejecutada en la GPU. Las dimensiones de malla y de bloque se especifican con variables del tipo dim3. Los parámetros de lanzamiento memoria_compartida y stream indican la cantidad de memoria compartida a reservar para cada bloque del kernel y el stream en el que el mismo va a ser lanzado. Estos dos últimos parámetros de lanzamiento son opcionales y se pueden omitir (en sesiones posteriores describiremos su funcionamiento). Por último, como haríamos con una función típica de C/C++, los parámetros de dicha función se especifican entre paréntesis.

Como comentamos anteriormente, esta invocación puede ser realizada utilizando la API del driver en lugar de CUDART, en conreto utilizando la función cuLaunchKernel que posee la siguiente declaración:

CUResult cuLaunchKernel (
  CUfunction kernel,
  unsigned int mallaDimX,
  unsigned int mallaDimY,
  unsigned int mallaDimZ,
  unsigned int bloqueDimX,
  unsigned int bloqueDimY,
  unsigned int bloqueDimZ,
  unsigned int memoriaCompartidaBytes,
  CUstream stream,
  void ** parametros,
  void ** extra
);

Limitaciones de Memoria

Como veremos en sesiones posteriores sobre el modelo de memoria de CUDA. Cada hilo dentro de una malla necesita una cierta cantidad de memoria para ejecutarse, además de la memoria necesaria para los datos que pueden ser accedidos por todos los hilos. Ciertas variables propias de cada hilo pueden alojarse en la asignación correspondiente a cada hilo del banco de registros del SM. No obstante, la cantidad de registros es relativamente escasa para este propósito por lo que parte de la memoria necesaria acaba siendo utilizada de una abstracción de la memoria de la GPU denominada memoria local.

Dado que generalmente se utiliza una cantidad elevada de hilos en ejecución, es presumible que también se necesite una porción considerable de la memoria de la GPU para asignar la cantidad necesaria a cada uno de los hilos. Desafortunadamente, la memoria en la GPU no es ilimitada, y por motivos de eficencia no se lleva a cabo paginación de memoria virtual en la GPU (existen formas de utilizar la memoria del host de forma transparente empleando Unified Virtual Addressing, pero se trata de u tema fuera del ámbito de esta explicación). Para evitar la posibilidad de que un kernel se quede sin memoria a mitad de ejecución, CUDA hace una reserva anticipada de la memoria necesaria por hilo. De esta forma, si la reserva de memoria falla por falta de espacio, el lanzamiento del kernel falla con el error correspondiente.

Limitaciones de Tiempo

En sistemas en los cuales la GPU utilizada para cómputo está siendo a su vez utilizada para visualización, la ejecución de un kernel costoso puede provocar un impacto negativo en la interactividad del sistema puesto que la GPU no es capaz de cambiar de contexto de ejecución de cómputos CUDA a cálculo de gráficos para el sistema. Por este motivo, el driver y los sistemas operativos fuerzan un tiempo máximo de ejecución para las GPUs empleadas a su vez para visualización. Una vez expirado ese tiempo, la GPU se reinicia para servir de nuevo su propósito principal.

En el caso de Linux, el tiempo está fijado en 2 segundos. En el caso de Windows, el driver puede configurarse en dos modos: WDDM (Windows Display Driver Model), en cuyo caso el sistema operativo decide dicho tiempo, y TCC (Tesla Compute Cluster), modo en el que no se fuerza ningún tipo de tiempo máximo de cómputo.

Mallas, Bloques, Hilos, Warps y Lanes

Los kernels son lanzados en una malla, que a su vez se divide en bloques, los cuales están compuestos por hilos. Los hilos se agrupan en la unidad mínima de ejecución, conocida como warp. Los warps son agrupaciones de 32 hilos en los cuales cada uno de los hilos recibe el nombre de lane.

Cada hilo utiliza una serie de variables especiales para identificarse y decidir sobre qué dato o posiciones de memoria tiene que trabajar. La existencia de dichas variables ya fue adelantada en la sesión anterior:

gridDim: tamaño o dimensión de la malla.
blockDim: tamaño o dimensión de los bloques.
blockIdx: índice o identificador del bloque dentro de la malla.
threadIdx: índice o identificador del hilo dentro del bloque.

Dichas variables especiales son del tipo dim3 por lo que poseen tres dimensiones o campos .x, .y y .z para poder organizar tanto la malla como los bloques en 1D, 2D ó 3D tal y como se muestra en las Figuras 1 (2D) y 2 (3D). Cabe destacar que no existe ninguna ganancia de rendimiento al utilizar bloques 1D, 2D ó 3D, es una decisión puramente de aplicación.

Figura 1. Malla 2D y bloques 2D. Figura extraída de "The CUDA Handbook: A Comprehensive Guide to GPU Programming".

Figura 2. Malla 3D y bloques 3D. Figura extraída de "The CUDA Handbook: A Comprehensive Guide to GPU Programming".

Debido a las agrupaciones en warps (32 hilos), es común elegir configuraciones de bloque cuyo número de hilos sea un múltiplo del tamaño del warp. En caso contrario, ciertos warps han de ser rellenados con hilos inactivos, perdiendo cierto rendimiento dado que no se utilizan al completo todos los recursos tal y como se muestra en la Figura 3.

Figura 3. Bloque de 28 hilos con 4 hilos inactivos para completar el tamaño de warp. Figura extraída de "The CUDA Handbook: A Comprehensive Guide to GPU Programming".

Escalabilidad Transparente y Planificación de Hilos

Durante la ejecución, el hardware se encarga de asignar los bloques de la malla a los SM. De esta forma, no todos los bloques son ejecutados de forma concurrente y no se da ninguna garantía respecto al orden de ejecución de los mismos, más allá de que los hilos de un mismo bloque sean ejecutados por el mismo SM. Cada SM tiene suficientes recursos para mantener el contexto de varios bloques en ejecución (8 a la vez antes de los SM de tercera generación y 16 a partir de ella). Este modelo de ejecución es necesario para lograr que un mismo programa pueda ejecutarse y escalar su eficiencia de forma transparente dependiendo de los recursos disponibles de la GPU en la que se ejecute.

Para cubrir los costes del paralelismo (latencia tanto de la memoria como de las instrucciones) es necesario que cada SM ejecute más hilos de los que un bloque por sí solo puede contener. Así pues, como comentábamos anteriormente, varios bloques son asignados a cada SM. Una vez un bloque ha sido asignado a un SM, sus hilos son ejecutados juntos, siguiendo un modelo SIMD, en agrupaciones de 32 hilos consecutivos denominadas warps. Los hilos o lanes de un warp se ejecutan físicamente en paralelo pues el warp completo se envía los núcleos del SM.

Los warps son por lo tanto la unidad de ejecución de CUDA y el motivo de su existencia y su tamaño no son decisiones triviales. Este modelo de procesamiento basado en los warps como unidades de planificación permite a la GPU ocultar la latencia del resto de operaciones (por ejemplo, realizando accesos de memoria coalescentes para servir información a varios hilos de un warp como veremos el próximas sesiones). Los SM implementan una planificación con zero-overhead, en la cual se mantiene un conjunto de warps elegibles para ejecución (aquellos cuya siguiente instrucción tiene sus operandos preparados para ejecutarse) los cuales son elegidos para ejecución mediante colas de prioridad. Es importante destacar de nuevo que todos los hilos de un warp ejecutan la misma instrucción cuando son seleccionados.

Sincronización y Control de Flujo

Debido al modelo de ejecución descrito en esta sesión, existe una serie de limitaciones en cuanto a las operaciones de sincronización que pueden realizar los hilos, así como una forma especial de tratar las condiciones o control de flujo.

Métodos de Sincronización

Anteriormente hemos descrito cómo el modelo de procesamiento de CUDA fuerza a que a que cada bloque sea ejecutado por un SM y cómo, para conseguir transparencia en la escalabilidad, no se garantiza ningún tipo de orden en la ejecución de los mismos. Esto implica que los hilos de diferentes bloques no pueden cooperar de una forma sencilla. A pesar de todo, utilizando diferentes funciones de CUDA es posible programar mecanismos de sincronización entre bloques. No obstante, estos mecanismos artificiales suelen presentar muchas complicaciones y suponen una pérdida de eficiencia total.

Las opciones de sincronización eficientes se producen entre hilos de un mismo bloque ya que sí tenemos la garantía de que se ejecutarán en un mismo SM. Estos hilos de un mismo bloque pueden sincronizarse de tres maneras distintas:

Barreras de sincronización (ver Figura 4): la función intrínseca __syncthreads() puede ser utilizada para imponer una barrera de sincronización en la que todos los hilos del bloque deben esperar hasta que todos hayan alcanzado dicho punto en la ejecución. Una vez todos los hilos han alcanzado la barrera, la ejecución puede continuar con normalidad.

Es muy importante tener en cuenta que es posible generar una situación de bloqueo si se utiliza código condicional con barreras de sincronización, pues si no todos los hilos del bloque alcanzan la instrucción de barrera la ejecución de los hilos a la espera no podrá continuar.

Figura 4. Ilustración de barrera de sincronización de hilos.

Operaciones atómicas: las funciones atómicas son aquellas que llevan acabo operaciones de lectura-modificación-escritura en memoria global o compartida de forma que garantizan que dicha operación será realizada sin interferencia de otros hilos. En otras palabras, garantizan que esa lectura y modificación de la memoria será realizada por un único hilo a la vez, el cual impone un bloqueo sobre la posición en cuestión que solamente es liberado cuando finaliza la función. Esto permite realizar operaciones sobre una misma posición de memoria sin caer en condiciones de carrera, obviamente con cierta penalización al rendimiento dada la naturaleza secuencial del bloqueo.
Memoria compartida: la porción de memoria compartida por los hilos de un mismo bloque puede ser utilizada a su vez para realizar ciertos tipos de sincronización. Este tipo de memoria y su funcionamiento será detallado la siguiente sesión sobre el modelo de memoria de CUDA.

Control de Flujo

Las operaciones condicionales pueden ser utilizadas dentro de kernels CUDA, sin embargo, dado que todos los hilos de un warp deben ejecutar las mismas instrucciones al ser planificados, el tratamiento de las condiciones se realiza de una forma especial. Para ello, el compilador predica las instrucciones de condición de forma que todos los hilos ejecuten las mismas instrucciones. Esto significa que, por ejemplo, ante la siguiente situación:

if (x < 0.0)
  z = x - 2.0;
else
  z = sqrt(x);

El compilador predicará las instrucciones para que todos los hilos ejecuten las mismas instrucciones. Así pues, todos evaluarán la condición, posteriormente se llevará a cabo una pasada con todos los hilos ejecutando una rama y seguidamente una segunda ejecutando la otra rama.

cond: p = (x < 0.0);
p: z = x - 2.0;
!p: z = sqrt(x);

La diferencia reside en que en la primera pasada, todos los hilos que no cumplan la condición son enmascarados y marcados con un flag para ejecutar el equivalente a un NOP en el procesador. En la segunda pasada ocurre lo mismo pero con los hilos del warp que cumplan la condición.

Así pues, el caso común de un condicional implica que hilos de un mismo warp tomen caminos diferentes y por lo tanto se provoca una penalización al rendimiento al tener que ejcutar ambas ramas. Obviamente, este problema se agrava con el anidamiento de condiciones.

Sin embargo, dependiendo de la granularidad del salto, puede no haber divergencia. En el caso anterior, el más frecuente, la divergencia se produce en base al identificador del hilo, por ejemplo:

if (threadIdx.x > 2)
  dosomething;
else
  dootherthing;

En esta situación, como adelantamos anteriormente, los hilos de un mismo warp divergen y debemos ejecutar ambas ramas predicándolas. De forma técnica, se dice que la granularidad del salto es inferior al tamaño del warp. No obstante, si la granularidad del salto es múltiplo del tamaño del warp, como por ejemplo:

if (threadIdx.x/WARP_SIZE > 2)
  dosomething;
else
  dootherthing;

todos los hilos del warp siguen el mismo camino y el compilador puede optimizar el código generado evitando la necesidad de predicar las instrucciones. En otras, palabras, no existe divergencia.

Streams

Además del paralelismo de grano fino expuesto por el modelo de procesamiento a la hora de procesar los hilos de un mismo warp de forma físicamente paralela y las jerarquías superiores de forma concurrente, el modelo de procesamiento de CUDA también admite concurrencia de un grano más grueso para mejorar la eficiencia de los programas:

Concurrencia entre CPU y GPU: como hemos destacado, los kernels se ejecutan de forma asíncrona, por lo que una vez emitida la orden de lanzamiento por parte del host, éste puede reanudar su ejecución y realizar otro tipo de cómputos de forma independiente a la GPU.
Concurrencia entre cómputo de kernel y transferencia de memoria: las GPUs poseen copy engines que actúan de forma independiente a los SM, por lo que se pueden realizar transferencias de memoria de forma independiente al trabajo realizado en los SMs.
Concurrencia entre kernels: aquellas GPUs con SMs de generación 2.X o posterior pueden ejecutar hasta cuatro kernels en paralelo.
Concurrencia multi-GPU: múltiples GPUs pueden operar en paralelo (este tema está fuera del ámbito de la explicación).

La forma de expresar esta concurrencia en CUDA es utilizando streams. Los streams son abstracciones para referirse a una secuencia de operaciones que se ejecutan de forma secuencial según el orden de envío en la GPU. La posibilidad de tener múltiples streams ejecutándose en la GPU implica que las operaciones de diferentes streams pueden ser ejecutadas de forma concurrente según los recursos disponibles y que además operaciones de diferente naturaleza pueden ser solapadas para utilizar todos los recursos de forma independiente.

El ejemplo más sencillo y efectivo de streams se puede dar considerando una aplicación típica que realiza una copia de memoria host-to-device, la ejecución de un kernel y posteriormente otra copia de memoria device-to-host. Las GPUs modernas disponen de dos copy engines (uno para transferencias host-to-device y otro para las device-to-host). Podemos utilizar varios streams, tal y como se muestra en la Figura 5, para solapar las copias en ambas direcciones junto con la ejecución de un kernel para conseguir un máximo aprovechamiento de todos los recursos del chip gráfico de forma simultánea.

Figura 5. Concurrencia de copias de memoria y kernels mediante streams.

Medición de Tiempos, Sincronización Host-Device y Eventos

Analizar el rendimiento de nuestro código CUDA, como ya comentamos en sesiones anteriores, es vital para determinar la ganancia que estamos obteniendo con nuestra implementación y así establecer un punto de partida para optimizaciones posteriores. La medición de tiempos para calcular el rendimiento de la implementación se suele realizar en la parte host, bien utilizando temporizadores de la CPU o bien temporizadores específicos de CUDA. De cualquier manera, el modelo de procesamiento de CUDA impone ciertas restricciones que deben ser tenidas en cuenta a la hora de realizar las mediciones para asegurar su validez.

Sincronización Host-Device

Como ya adelantamos, la ejecución de los kernels es asíncrona, lo cual quiere decir que una vez la CPU emite la orden y ésta se almacena en el buffer de comandos de la GPU, la parte host continúa su ejecución. Este hecho condiciona en gran medida la medición de tiempos. Por ejemplo, observemos el siguiente código.

cudaMemcpy(d_x_, h_x_, vector_size_bytes_, cudaMemcpyHostToDevice);
cudaMemcpy(d_y_, h_y_, vector_size_bytes_, cudaMemcpyHostToDevice);

kernel<<<grid, block>>>(d_x_, d_y_, N);

cudaMemcpy(h_y_, d_y_, vector_size_bytes_, cudaMemcpyDeviceToHost);

En este código de ejemplo, las instrucciones cudaMemcpy son síncronas o bloqueantes, lo cual significa que no se ejecutarán hasta que todos los comandos CUDA anteriores hayan sido completados y, de igual manera, los comandos posteriores no comenzarán hasta que las copias hayan terminado. Esto contrasta con el hecho de que la llamada al kernel sea asíncrona, la CPU emitirá el comando apropiado para que la GPU ejecute el kernel (dicho comando se almacenará en el buffer de la GPU para ser ejecutado en el momento que corresponda) y continuará su ejecución hasta alcanzar la instrucción de copia de memoria. No obstante, la naturaleza síncrona de las copias de memoria asegura que no se produzca ninguna condición de carrera.

Si tuviéramos que calcular el tiempo de ejecución empleado por el kernel, midiendo desde el host, podríamos hacerlo de la siguiente manera:

cudaMemcpy(d_x_, h_x_, vector_size_bytes_, cudaMemcpyHostToDevice);
cudaMemcpy(d_y_, h_y_, vector_size_bytes_, cudaMemcpyHostToDevice);

unsigned long t_start_ = cpu_timer();
kernel<<<grid, block>>>(d_x_, d_y_, N);
unsigned long t_end_ = cpu_timer();

cudaMemcpy(h_y_, d_y_, vector_size_bytes_, cudaMemcpyDeviceToHost);

Sin embargo, dada la naturaleza asíncrona del kernel, únicamente estaríamos midiendo el tiempo que se tarda en emitir la orden de lanzamiento y no el tiempo de ejecución del kernel en sí. Para poder medir ese tiempo necesitamos imponer una barrera de sincronización explícita entre la CPU y la GPU para bloquear la ejecución en el host hasta que los comandos en el buffer del dispositivo hayan sido completados. Esto lo podemos conseguir empleando la instrucción cudaDeviceSynchronize.

cudaMemcpy(d_x_, h_x_, vector_size_bytes_, cudaMemcpyHostToDevice);
cudaMemcpy(d_y_, h_y_, vector_size_bytes_, cudaMemcpyHostToDevice);

unsigned long t_start_ = cpu_timer();
kernel<<<grid, block>>>(d_x_, d_y_, N);
cudaDeviceSynchronize();
unsigned long t_end_ = cpu_timer();

cudaMemcpy(h_y_, d_y_, vector_size_bytes_, cudaMemcpyDeviceToHost);

El problema con esta aproximación es que bloquean tanto el flujo de ejecución en la CPU como en la GPU. Existe una forma alternativa para, entre otras utilidades, medir el tiempo de ejecución de partes del programa en la GPU. Para ello podemos utilizar los denominados eventos con la API de CUDA. Los eventos son esencialmente marcas temporales (del inglés timestamps) en la GPU que son establecidas en puntos especificados por el usuario. Estas operaciones generan una sobrecarga prácticamente nula en la ejecución pues es la propia GPU la que genera el timestamp.

Su utilización con la API de CUDA RT permite medir tiempos de forma sencilla. Con la función cudaEventCreate podemos crear eventos del tipo cudaEvent_t, la función cudaEventRecord nos permite especificar esos puntos en los que se va a almacenar el timestamp determinado en el evento. Por último, es importante sincronizar y bloquear la ejecución en el host para evitar leer el temporizador antes de que la GPU haya llegado a almacenar el timestamp. Para ello utilizaremos la función cudaEventSynchronize. Hecho esto, la función cudaEventElapsedTime nos permite calcular el número de milisegundos entre dos timestamps con una resolución de medio microsegundo aproximadamente.

A continuación se muestra un ejemplo de código completo para medir el tiempo de ejecución correctamente de un kernel en la GPU.

cudaEvent_t t_start_, t_stop_;
cudaEventCreate(&t_start_);
cudaEventCreate(&t_stop_);

cudaMemcpy(d_x_, h_x_, vector_size_bytes_, cudaMemcpyHostToDevice);
cudaMemcpy(d_y_, h_y_, vector_size_bytes_, cudaMemcpyHostToDevice);

cudaEventRecord(t_start_);
kernel<<<grid, block>>>(d_x_, d_y_, N);
cudaEventRecord(t_stop_);

cudaMemcpy(h_y_, d_y_, vector_size_bytes_, cudaMemcpyDeviceToHost);

cudaEventSyncrhonize(t_stop_);
float milliseconds_ = 0.0f;
cudaEventElapsedTime(&milliseconds_, t_start_, t_stop_);

Sesión 3: Modelo de Procesamiento