CUDA es el lenguaje usado para programar en GPUs de NVIDIA, es vital para una gran cantidad de tareas de computación y, aún así, es un misterio para una gran cantidad de programadores. En este post, intentamos explicar parte del misterio de CUDA y ayudar a entender el paradigma especial que requiere.
Como lenguaje, la sintaxis de CUDA es extremadamente similar a C y sería posible pensar que si uno supiera programar C, podría programar CUDA sin ninguna dificultad. Pero si se intentara ejecutar un código programado para CPU en GPU, observamos que es varias magnitudes más lento que solo ejecutarlo en CPU. Este fenómeno es extraño, ya que esperaríamos que toda la potencia de la GPU hiciera que el programa sea más rápido, pero un solo thread de una CPU sería considerablemente más veloz.
Esta paradoja se debe al cambio de arquitectura. La gran mayoría de programas se escriben para una arquitectura que asume que existen uno o varios threads y cada uno ejecuta su código dado, muchas veces ni siquiera siendo el mismo código en cada thread. Esta arquitectura es conocida como SISD (Single Instruction Single Data) o SIMD (Single Instruction Multiple Data), dependiendo del set de instrucciones que la CPU tenga, lo que significa que en una CPU una sola instrucción suele tratar uno o más datos.
Cada thread tiene su núcleo de CPU asociado con su propio program counter (registro que guarda la siguiente instrucción a ejecutar en la CPU). Dos threads no comparten program counters, aun si ejecutarán el mismo código, lo harían con program counters diferentes.
Sin embargo, no es el caso con CUDA. Aquí tenemos que la unidad más pequeña de procesamiento no es el núcleo, es el SM (Streaming Multiprocessor). Un SM tiene un solo program counter, pero no tiene un solo núcleo. Mientras que en la CPU veríamos un program counter por núcleo, en una GPU vemos múltiples núcleos por SM, es decir, múltiples núcleos sólo pueden realizar la ejecución que dicta el program counter que comparten.
Por ejemplo, tenemos en el código un simple “if”. En este caso, el SM “apagaría” los núcleos que no pasan el “if” y solo ejecutaría aquellos que lo pasan. Una vez que acabe, ejecutaría los que no pasaron, ya que tener un solo program counter fuerza a que solo se pueda seguir un hilo de instrucciones a la vez. Esta arquitectura tiene el nombre de SIMT (Single Instruction Multiple Threads).
Este tipo de arquitectura no se usa porque sí, sino porque, nativamente, permite ejecutar código de forma altamente paralela. Es más, la única manera de ejecutar el código es esa. Si el código no se beneficiara de esta arquitectura, la GPU daría muy pobres resultados.
En la siguiente imagen mostramos la arquitectura de una SM, específicamente de la arquitectura de CUDA de FERMI. Podemos observar cómo l warp tiene un “file register” (donde guardan las variables declaradas en el código), una cache L1 (que solemos llamar “shared memory” y está principal y directamente administrada por el propio código), una red de interconexión (que permite que varios threads que están dentro de un warp hablen entre ellos) y varios núcleos.
Cada núcleo es asignado un thread y, como podemos ver, en esta arquitectura en específico hay dos sets de 16 núcleos, pudiendo dividir el warp en algo llamado “half-warp”. Esto se debe, principalmente, a la compatibilidad con arquitecturas anteriores a FERMI, donde los warps eran de 16 núcleos. También tiene 16 LD/ST (unidades de load/store), permitiendo leer y escribir en memoria en 16 threads a la vez. Y, finalmente, 4 unidades SFU (Special Function Unit) que ofrecen la posibilidad de ejecutar funciones especiales como cos, sin o tan, entre otras.
Debido a esta arquitectura, es necesario crear una jerarquía para separar threads entre ellos. Para hacerlo, se recurre a un sistema en 3 niveles: Primero tenemos los threads, después los thread blocks y, por último, la grid.
Esta jerarquía indica qué recursos pueden compartir los threads. Por ejemplo, threads dentro de un mismo warp pueden compartir todo usando la “interconnect network”, pero threads dentro de un mismo thread block solo pueden compartir “shared memory” (un tipo de caché controlada por el usuario), mientras que threads en bloques diferentes no comparten nada.
Aunque esta jerarquía es interesante, en este post nos centraremos en el nivel más bajo de la jerarquía, los warps. Estos son la unidad mínima de ejecución, es decir, cuando se ejecutan threads, se ejecutan dentro de un warp y un warp es ejecutado, a su vez, en un SM.
¿Qué es un Warp en CUDA?
Un warp hoy en día está formado por 32 threads, aunque en arquitecturas de GPU anteriores podemos observar warps de 16 threads. Estos 32 threads comparten un solo program counter, es decir, ejecutan la misma instrucción a la vez.
__global__ void sum(float *a)
{
__shared__ float shared_mem[32];
int idx = threadIdx.x;
shared_mem[idx] = a[idx];
for(int threshold=16; idx<threshold; threshold>>=1)
{ shared_mem[idx]+=shared_mem[idx+threshold]; }
if (idx==0) {
a[0]=shared_mem[0];
}
}
Este programa es sencillo, pero ya podemos ver la complejidad ya que tenemos varios elementos que discriminan por la ID del thread, que obtenemos con threadIdx.x. Al inicio, los 32 threads trabajan, después 16… y, finalmente, 1. Pero el código es el mismo para todos, requiriendo esta discriminación por ID.
Tal y como se puede observar por el identificador de __shared__, primero se envían los datos a “shared memory”. Mientras que se podría hacer todo desde la memoria global, es preferible hacerlo desde “shared memory” debido a la ineficiencia de operar en memoria global, así que aunque en este post no expliquemos mucho qué es la “shared memory”, haremos uso de ella por buenas prácticas.
Después de enviar a shared memory los datos, simplemente hace una reducción donde cada thread suma 2 elementos y disminuye a la mitad la cantidad de threads que trabajan hasta que solo queda 1 y finaliza, escribiendo el resultado de vuelta en la memoria global.
Ahora que vemos el programa para 32 threads podemos extenderlo a 256 threads que incluyen múltiples warps.
__global__ void sum(float *a)
{
__shared__ float shared_mem[256];
int idx = threadIdx.x;
shared_mem[idx] = a[idx];
for(int threshold=128; idx<threshold; threshold>>=1) {
shared_mem[idx]+=shared_mem[idx+threshold];
}
if (idx==0) {
a[0]=shared_mem[0];
}
}El código es esencialmente el mismo, simplemente la shared memory ahora es de 256 floats, y el threshold es 128, pero este código ya no es correcto. Aunque es verdad que normalmente dará el resultado esperado, a veces comete errores.
En CUDA normalmente no existe sincronización automática. Ahora que tenemos 256 threads, hay 8 warps trabajando simultáneamente y tienen program counters diferentes, así que puede que uno se adelante a los demás y no espere. Por ejemplo, si el warp de los primeros 32 threads se adelantara, el resultado sería solo la suma de los primeros 32 elementos.
Para esto, existen funciones primitivas de sincronización, tal y como podemos ver en el siguiente código.
__global__ void sum_synchronized(float *a)
{
__shared__ float shared_mem[256];
int idx = threadIdx.x;
shared_mem[idx] = a[idx];
__syncthreads();
for(int threshold=128; idx<threshold; threshold>>=1) {
shared_mem[idx]+=shared_mem[idx+threshold];
__syncthreads();
}
if (idx==0) {
a[0]=shared_mem[0];
}
}El script es, básicamente, el mismo, excepto que llamamos a __syncthreads() cada vez que queremos asegurarnos de que ningún warp diverge. Esta función nos asegura que todo warp que pertenece al mismo thread block estará en la misma instrucción. Esta sincronización es solo a nivel de thread block, pero en este artículo no entraremos en detalle a describirlo, solo es necesario saber que este código está contenido en un solo thread block.
Ahora tenemos un código funcional pero bastante ineficiente, ya que estamos pasando datos a una memoria compartida y después operando sobre ella, cuando hemos visto en la arquitectura de una SM, que existe algo llamado “interconnect network”. Para implementarla, podemos hacer uso de ciertas primitivas a nivel de warp como se muestra en el siguiente código.
__global__ void sum_synchronized(float *a)
{
int idx = threadIdx.x;
float val = a[idx];
for (int offset = 16; offset > 0; offset /= 2) {
val += __shfl_down_sync(0xffffffff, val, offset);
}
if(idx==0) {
a[0]=val;
}
}La función usada es __shfl_down_sync(synchronization_mask, val, offset), que suma la variable “val” de dos diferentes threads situados a un “offset” entre ellos y fuerza la sincronización de los threads dada una máscara. Por ejemplo, en el código usamos 0xffffffff, que indica que se sincronice todo thread dentro del warp. Los casos de disincronía serían raros. Esta primitiva nos permite saltarnos el usar la “shared memory” y, directamente, podemos sumar los valores a base de sumar sus variables entre ellos.
Por último, podemos aplicar estas primitivas de warp para hacer el ejemplo de 256 elementos de forma eficiente:
__global__ void sum_synchronized(float *a)
{
int idx = threadIdx.x;
float val = a[idx];
__shared__ float shared_mem[8];
for (int offset = 16; offset > 0; offset /= 2) {
val += __shfl_down_sync(0xffffffff, val, offset);
}
if(idx%32==0) {
shared_mem[idx>>5]=val;
}
__syncthreads();
if (idx<=8) {
val=shared_mem[idx];
val += __shfl_down_sync(0xffffffff, val, 4);
val += __shfl_down_sync(0xffffffff, val, 2);
val += __shfl_down_sync(0xffffffff, val, 1);
if (idx==0) {
a[0]=val;
}
}
}Hacemos la suma por warp igual que antes, pero aquí ocurre algo interesante. Una vez hemos sumado por warp, tenemos el problema de que hay 8 warps concurrentes y hay que compartir información entre ellos. En particular, un solo valor por cada uno. Para hacerlo, reservamos 8 espacios en “shared memory” y, de cada warp, el primer thread escribe la suma a “shared memory”. Después ejecutamos __syncthreads(), ya que podemos tener una disincronía. Finalmente, usamos el primer warp para sumar esos 8 elementos en 1 y escribir el resultado.
Conclusión
La arquitectura que las GPU usan y, especialmente CUDA, es muy diferente de la que nos imaginamos cuando programamos en CPU. Por este motivo, hemos de pensar en alternativas a la hora de llevar a cabo esta tarea en GPU.
Especialmente, es necesario que nos centremos en el paradigma de los threads. El código se divide en “threads”, pero todo thread ejecuta exactamente el mismo código, lo que nos fuerza a identificar cada thread con un ID y asignar una parte del trabajo siguiendo dicho ID.
Además, hay que tener en cuenta que no disponemos de mecanismos de sincronización y de que toda comunicación por encima del nivel de un warp requerirá de mecanismos para sincronizar, como es el caso de __syncthreads(), que sincroniza a nivel de thread block.
Aun así, este paradigma es importante, ya que nos permite ejecutar código altamente paralelo permitiendo que tareas que en CPU pueden tardar días, se ejecuten en GPU en minutos.
