Capítulo 11: Direcciones de investigación de GPU sobre escalarización y ejecución afín
Como se describe en el Capítulo 2, las API de computación de GPU, como CUDA y OpenCL, presentan un modelo de programación similar a MIMD que permite al programador lanzar una gran matriz de hilos escalares en la GPU. Si bien cada uno de estos hilos escalares puede seguir su propio camino de ejecución único y puede acceder a ubicaciones de memoria arbitrarias, en el caso común, todos siguen un pequeño conjunto de caminos de ejecución y realizan operaciones similares.
El flujo de control convergente entre los hilos de GPU se explota en la mayoría, si no en todas, las GPU modernas a través del modelo de ejecución SIMT, donde los hilos escalares se agrupan en warps que se ejecutan en hardware SIMD (ver Sección 3.1.1). Este capítulo resume una serie de investigaciones que explotan aún más la similitud de estos hilos escalares a través de la escalarización y la ejecución afín.
La idea clave de esta investigación radica en la observación de la estructura de valor [Kim et al., 2013] entre los hilos que ejecutan el mismo kernel de cómputo. Los dos tipos de estructura de valor, uniforme y afín, se ilustran en el kernel de cómputo en el Ejemplo 11.1.
Variable uniforme
Una variable que tiene el mismo valor constante para cada hilo en el kernel de cómputo. En el Algoritmo 11.1, la variable a, así como las literales THRESHOLD y Y_MAX_VALUE, tienen valores uniformes en todos los hilos del kernel de cómputo. Una variable uniforme se puede almacenar en un solo registro escalar y ser reutilizada por todos los hilos en el kernel de cómputo.
Variable afín
Una variable con valores que son una función lineal del ID del hilo para cada hilo en el kernel de cómputo. En el Algoritmo 11.1, la dirección de memoria de la variable y[idx] se puede representar como una transformación afín del ID del hilo threadIdx.x:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Esta representación afín se puede almacenar como un par de valores escalares, una base y un paso, que es mucho más compacta que la versión completamente expandida.
__global__ void vsadd( int y[], int a ) {
i
```Aquí está la traducción al español del archivo Markdown, con los comentarios traducidos al español y el código sin traducir:
nt idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
Algoritmo 11.1: Ejemplo de operaciones escalares y afines en un kernel de cálculo (de [Kim et al., 2013]).
Hay múltiples propuestas de investigación sobre cómo detectar y explotar variables uniformes o afines en GPUs. El resto de este capítulo resume estas propuestas en estos dos aspectos.
## Detección de variables uniformes o afines
Hay dos enfoques principales para detectar la existencia de variables uniformes o afines en un kernel de cálculo de GPU: Detección Impulsada por el Compilador y Detección a través del Hardware.
### Detección Impulsada por el Compilador
Una forma de detectar la existencia de variables uniformes o afines en un kernel de cálculo de GPU es hacerlo a través de un análisis especial del compilador. Esto es posible porque los modelos de programación de GPU existentes, CUDA y OpenCL, ya proporcionan medios para que el programador declare una variable como constante en todo el kernel de cálculo, así como proporcionar una variable especial para el ID de hilo. El compilador puede realizar un análisis de dependencia de control para detectar variables que dependen únicamente de constantes e IDs de hilo, y marcarlas como uniformes/afines. Las operaciones que trabajan únicamente con variables uniformes/afines son entonces candidatas para la escalarización.
AMD GCN [AMD, 2012] se basa en el compilador para detectar variables uniformes y operaciones escalares que pueden ser almacenadas y procesadas por un procesador escalar dedicado.
Asanovic et al. [2013] introducen un análisis convergente y variante combinado que permite al compilador determinar las operaciones en un kernel de cálculo arbitrario que son elegibles para la escalarización y/o la transformación afín. Las instrucciones dentro de las regiones convergentes de un kernel de cálculo se pueden convertir en instrucciones escalares/afines. En cualquier transición de regiones divergentes a convergentes de un kernel de cálculo, el compilador inserta una instrucción `syncwarp` para manejar las dependencias de registro inducidas por el flujo de control entre las dos regiones. Asanovic et al. [2013] adoptaron tAquí está la traducción al español del archivo Markdown, con los comentarios del código traducidos:
Este análisis para generar operaciones escalares para la arquitectura Temporal-SIMT [Keckler et al., 2011, Krashinsky, 2011].
Cálculo Afín Desacoplado (DAC) [Wang y Lin, 2017] se basa en un análisis de compilador similar para extraer candidatos escalares y afines que se desacoplarán en un warp separado. Wang y Lin [2017] aumentan el proceso con un análisis afín divergente, con el objetivo de extraer hilos de instrucciones