Chapitre 11 : Directions de recherche sur les GPU pour la scalarisation et l'exécution affine
Comme décrit dans le Chapitre 2, les API de calcul GPU, telles que CUDA et OpenCL, présentent un modèle de programmation de type MIMD qui permet au programmeur de lancer un grand nombre de fils d'exécution scalaires sur le GPU. Bien que chacun de ces fils d'exécution scalaires puisse suivre son propre chemin d'exécution unique et puisse accéder à des emplacements mémoire arbitraires, dans le cas le plus courant, ils suivent tous un petit ensemble de chemins d'exécution et effectuent des opérations similaires.
Le flux de contrôle convergent des fils d'exécution GPU est exploité sur la plupart, sinon la totalité, des GPU modernes via le modèle d'exécution SIMT, où les fils d'exécution scalaires sont regroupés en warps qui s'exécutent sur du matériel SIMD (voir la Section 3.1.1). Ce chapitre résume une série de recherches qui exploitent davantage la similarité de ces fils d'exécution scalaires via la scalarisation et l'exécution affine.
L'idée clé de cette recherche réside dans l'observation de la structure des valeurs [Kim et al., 2013] entre les fils d'exécution exécutant le même noyau de calcul. Les deux types de structure de valeur, uniforme et affine, sont illustrés dans le noyau de calcul de l'Exemple 11.1.
Variable uniforme
Une variable qui a la même valeur constante pour chaque fil d'exécution dans le noyau de calcul. Dans l'Algorithme 11.1, la variable a, ainsi que les littéraux THRESHOLD et Y_MAX_VALUE, ont tous des valeurs uniformes pour tous les fils d'exécution dans le noyau de calcul. Une variable uniforme peut être stockée dans un seul registre scalaire et réutilisée par tous les fils d'exécution dans le noyau de calcul.
Variable affine
Une variable dont les valeurs sont une fonction linéaire de l'ID du fil d'exécution pour chaque fil d'exécution dans le noyau de calcul. Dans l'Algorithme 11.1, l'adresse mémoire de la variable y[idx] peut être représentée comme une transformation affine de l'ID du fil d'exécution threadIdx.x :
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Cette représentation affine peut être stockée sous la forme d'une paire de valeurs scalaires, une base et un pas, ce qui est beaucoup plus compact que la version entièrement développée.
__global__ void vsadd( int y[], int a ) {
i
```Voici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :
int idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
Algorithme 11.1 : Exemple d'opérations scalaires et affines dans un noyau de calcul (tiré de [Kim et al., 2013]).
Il existe plusieurs propositions de recherche sur la façon de détecter et d'exploiter les variables uniformes ou affines dans les GPU. Le reste de ce chapitre résume ces propositions sur ces deux aspects.
## Détection des variables uniformes ou affines
Il existe deux approches principales pour détecter l'existence de variables uniformes ou affines dans un noyau de calcul GPU : la détection pilotée par le compilateur et la détection via le matériel.
### Détection pilotée par le compilateur
Une façon de détecter l'existence de variables uniformes ou affines dans un noyau de calcul GPU est de le faire via une analyse de compilateur spéciale. Cela est possible car les modèles de programmation GPU existants, CUDA et OpenCL, fournissent déjà des moyens pour le programmeur de déclarer une variable comme constante tout au long du noyau de calcul, ainsi qu'une variable spéciale pour l'ID de thread. Le compilateur peut effectuer une analyse de dépendance de contrôle pour détecter les variables qui dépendent uniquement de constantes et d'ID de thread, et les marquer comme uniformes/affines. Les opérations qui fonctionnent uniquement sur des variables uniformes/affines sont alors des candidates à la scalarisation.
AMD GCN [AMD, 2012] s'appuie sur le compilateur pour détecter les variables uniformes et les opérations scalaires qui peuvent être stockées et traitées par un processeur scalaire dédié.
Asanovic et al. [2013] introduisent une analyse convergente et variante combinée qui permet au compilateur de déterminer les opérations dans un noyau de calcul arbitraire qui sont éligibles à la scalarisation et/ou à la transformation affine. Les instructions dans les régions convergentes d'un noyau de calcul peuvent être converties en instructions scalaires/affines. À toute transition des régions divergentes aux régions convergentes d'un noyau de calcul, le compilateur insère une instruction `syncwarp` pour gérer les dépendances de registre induites par le contrôle-flux entre les deux régions. Asanovic et al. [2013] ont adopté tVoici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :
Cette analyse permet de générer des opérations scalaires pour l'architecture Temporal-SIMT [Keckler et al., 2011, Krashinsky, 2011].
Le calcul affine découplé (DAC) [Wang et Lin, 2017] s'appuie sur une analyse de compilateur similaire pour extraire les candidats scalaires et affines à découpler dans un warp séparé. Wang et Lin [2017] augmentent le processus avec une analyse affine divergente, dans le but d'extraire des séquences d'instructions