Atras Menú

Black-byte

Conceptos de CUDA

Threading

CUDA utiliza un concepto similar al de los lenguajes de programación convencionales, el “hilo”, con la diferencia que CUDA posee una organización de éstos en diferentes elementos agrupadores de acuerdo a la posibilidad de compartir pequeños espacios de memoria y de ser sincronizados. Los elementos a tener en cuenta son:

  • Grid: Es un conjunto de blocks que se están ejecutando en un llamado a la GPU.
  • Block: Son formas de agrupar un conjunto de thread. Los blocks se pueden organizar en arreglos de una, dos o tres dimensiones, donde cada bloque tiene un índice único que lo permite identificar al momento de realizar las operaciones en paralelo. Dentro de los Block los hilos pueden sincronizarse y acceder a memoria compartida.
  • Thread: Son segmentos de un programa que se ejecutan de manera paralela en la GPU. Al igual que con los Block, los hilos se pueden instanciar en una, dos o tres dimensiones.

La cantidad de Block y el tamaño de cada uno se definen de manera dinámica al crear un Grid. Así por ejemplo al lanzar un kernel de ejecución puedo crear 9 blocks en una matriz bidimensional <3,3> y dentro de cada Block instanciar 16 hilos en una matriz <4,4>, para un total de 144 hilos paralelos procesándose en el Grid. Aunque en principio parece que esta distribución de procesamiento es algo exótica, en realidad obedece a la facilidad de realizar ciertas operaciones en paralelo, especialmente las operaciones matriciales y también a que el hardware está optimizado para la ejecución de cientos de hilos simultáneamente.
Cuando se está ejecutando el código dentro de los segmentos paralelos del kernel se pueden invocar variables globales para identificar en que bloque e hilo se encuentra ejecutando el código. Las contantes son:

  • threadIdx : Identifica el thread donde se está ejecutando el código.
  • blockidx; Identifica el block donde se está ejecutando el código.}
  • blockDim: Indica la cantidad de threads asociadas al bloque.
  • gridDim: Indica la cantidad de blocks instanciados para la ejecución del kernel.

Estas variables globales se acceden con operadores punto, por ejemplo threadIdx.x indica la posición del hilo en “fila x” (la denominación de fila, columna es arbitraria) y threadIdx.y indica la “columna y”. Si se eligen crear blocks con threads organizados en matrices bidimensionales el indicador z siempre será cero.

Los hilos se instancian en el dispositivo y son manejados internamente por éste.
Memoria

En CUDA existen varios tipos de memoria independientes de la memoria principal del sistema. Debido a que la transferencia de memoria entre la GPU y la memoria del sistema es costosa , aún a través del puerto PCI Express, uno de los de mayor ancho de banda para placas comerciales, el hardware de las GPU ha ido incorporando memoria en el chip y memoria especializada de video, generalmente con características iguales o superiores a las memorias de sistema.

CUDA hace explícita la localización de la memoria, así que la creación y liberación de la memoria se hace directamente sobre el device y los kernel no pueden acceder directamente a punteros de memoria en el host. Esto implica que si se quieren procesar datos en el GPU que residen en memoria principal, es necesario copiarlos hacia el device, procesarlos y copiar el resultado de vuelta al host. Dada la latencia que existe entre ambas memorias, es importante desde el principio planear para minimizar la transferencia de datos entre éstas.
Los tipos de memoria que CUDA utiliza son:

  • Local Storange: Es la memoria utilizada por los threads, generalmente manejada por el compilador. Solo tiene visibilidad dentro del thread y es eliminada una vez el thread se termina de ejecutar.
  • Shared Memory: Cada block tiene conjunto de memoria compartida que es visible entre todos los threads del block. Es eliminada una vez el bloque termina su ciclo de ejecución (todos los hilos del bloque se ejecutan).
  • Global device memory: Es la memoria de video de la tarjeta. Accesible por todos los threads de todos los block y también por los hilos corriendo en la CPU. Se puede separar y eliminar dinámicamente de la misma manera que se hace con la memoria convencional, únicamente indicando que es memoria de tipo dispositivo.
  • Constant Memory. Es un pequeño espacio de memoria cacheado, accesible a través de todos los threads. Como su nombre lo indica sirve para colocar constantes, debido a su muy rápido acceso.
  • Texture Memory: Este es un tipo de memoria Read-only, a partir del momento que se lanza el kernel no se puede modificar, que normalmente está asociado a un cierto espacio físico de procesamiento de la GPU. Tiene un ancho de banda mayor que la Global Memory. Esta memoria puede ser accedida por todos los threads, pero es de tipo binding.

Adicional a estos tipos de memorias en CUDA se puede declarar la Page-Locked memory, que es un tipo de memoria de host que puede ser directamente manipulada por el device. Existen varios tipos de Page-Locked memory a saber Portable, Write-Combining y Mapped memory. Por ahora no se entrará en el detalle de éstas.
En principio las variedades de memoria disponibles en CUDA parecen dificultar el uso de esta arquitectura, aún así los programas pueden ejecutarse utilizando solo Local, Shared y Global, y cuando se requiere mejorar el desempeño se desarrollan.

A partir de la versión 4 de CUDA la memoria tanto de device como de host se referencian con UVA (Unified Virtual Address Space), en donde a ambas memorias se les asigna un único espacio de referencias, conservando las reglas de copia y acceso desde los threads.

Siguientes Tutoriales
En los siguientes tutoriales se presentará el Hola Mundo de CUDA y aplicación práctica de los principios vistos en éste tutorial. Para éstos se requerirá de una tarjeta de video que soporte CUDA, de los developer drivers, sdk y toolkit de CUDA disponibles en http://developer.nvidia.com/cuda-toolkit-40 .

Bibliografía:

  • CUDA by Example: An Introduction to General-Purpose GPU Programming, Jason Sanders and Edward Kandrot, Addison-Wesley , 2010.
  • Programming Massively Parallel Processors: A Hands-on Approach, David Kirk and Wen-mei W Hwu, Morgan Kaufmann, 2010
  • Cuda_C_Programming Guide, Nvidia Corporation 2011.
Páginas: 1 2 3