Association de modèles de programmation pour
l'exploitation de clusters de GPUs dans le calcul intensif

C. Roelandt

► To cite this version:

C. Roelandt. Association de modèles de programmation pour l'exploitation de clusters de GPUs dans le calcul intensif. Calcul parallèle, distribué et partagé [cs.DC]. 2011. hal-00803304

HAL Id: hal-00803304

https://inria.hal.science/hal-00803304v1

Submitted on 31 Jul 2025

HAL is a multi-disciplinary open access
archive for the deposit and dissemination of sci-
entific research documents, whether they are pub-
lished or not. The documents may come from
teaching and research institutions in France or
abroad, or from public or private research centers.

L'archive ouverte pluridisciplinaire HAL, est
destinée au dépôt et à la diffusion de documents
scientifiques de niveau recherche, publiés ou non,
émanant des établissements d'enseignement et de
recherche français ou étrangers, des laboratoires
publics ou privés.

<figure><img src="image_1.png" /></figure>

Distributed under a Creative Commons Attribution 4.0 International License
Mémoire de Master Recherche

présenté par

Cyril ROELANDT

Association de modèles de programmation
pour l'exploitation de clusters de GPUs dans
le calcul intensif

Encadrement : Olivier AUMAGE et Samuel THIBAULT

Février – Juin 2011

Laboratoire Bordelais de Recherche en Informatique (LaBRI)
INRIA Bordeaux
Université Bordeaux 1
ENSEIRB-MATMECA
<figure><img src="image_2.png" /></figure>
# Table des matières

<table><tr><td>1 Introduction</td><td>8</td></tr><tr><td>1.1 De l'évolution des machines</td><td>8</td></tr><tr><td>1.2 Nouvelles difficultés de programmation</td><td>8</td></tr><tr><td>1.3 Un modèle hybride</td><td>9</td></tr><tr><td>1.3.1 XcalableMP</td><td>9</td></tr><tr><td>1.3.2 Le support exécutif StarPU</td><td>9</td></tr><tr><td>2 Contexte</td><td>11</td></tr><tr><td>2.1 Passage de messages et processus légers</td><td>11</td></tr><tr><td>2.1.1 Message Passing Interface</td><td>11</td></tr><tr><td>2.1.2 Threads POSIX</td><td>11</td></tr><tr><td>2.2 Deux approches de plus haut niveau</td><td>12</td></tr><tr><td>2.2.1 High Performance Fortran</td><td>12</td></tr><tr><td>2.2.2 OpenMP</td><td>12</td></tr><tr><td>2.3 Parallélisme inter-noeuds et intra-noeuds</td><td>13</td></tr><tr><td>2.3.1 XcalableMP</td><td>13</td></tr><tr><td>2.3.2 StarPU</td><td>14</td></tr><tr><td>2.4 Propositions de modèles pour les GPUs</td><td>15</td></tr><tr><td>2.4.1 HMPP</td><td>15</td></tr><tr><td>2.4.2 StarSs</td><td>15</td></tr><tr><td>3 Un modèle hybride pour l'exploitation de clusters</td><td>16</td></tr><tr><td>3.1 Clusters de machines hétérogènes</td><td>16</td></tr><tr><td>3.1.1 Architectures distribuées</td><td>16</td></tr><tr><td>3.1.2 Machines hétérogènes</td><td>16</td></tr><tr><td>3.1.3 Combinaison</td><td>16</td></tr><tr><td>3.2 Les limites de XcalableMP</td><td>17</td></tr><tr><td>3.2.1 Découpage des données</td><td>17</td></tr><tr><td>3.2.2 Sous-exploitation des noeuds</td><td>17</td></tr><tr><td>3.3 Utilisation d'un support exécutif</td><td>18</td></tr><tr><td>3.3.1 Programmer les noeuds avec StarPU</td><td>18</td></tr><tr><td>3.3.2 Vers un modèle hybride</td><td>18</td></tr></table>
<table><thead><tr><th></th><th></th><th></th></tr></thead><tbody><tr><td>4</td><td>Implémentation</td><td>20</td></tr><tr><td>4.1</td><td>Présentation du compilateur</td><td>20</td></tr><tr><td>4.1.1</td><td>Fonctionnement</td><td>20</td></tr><tr><td>4.1.2</td><td>Modifications</td><td>20</td></tr><tr><td>4.1.3</td><td>Un exemple : le pragma task</td><td>21</td></tr><tr><td>4.2</td><td>Application au cas des matrices</td><td>23</td></tr><tr><td>4.2.1</td><td>Idée générale</td><td>23</td></tr><tr><td>4.2.2</td><td>Découpage des données</td><td>24</td></tr><tr><td>4.2.3</td><td>Gestion des GPUs</td><td>24</td></tr><tr><td>5</td><td>Evaluation des performances</td><td>25</td></tr><tr><td>5.1</td><td>Description des tests</td><td>25</td></tr><tr><td>5.1.1</td><td>Environnement</td><td>25</td></tr><tr><td>5.1.2</td><td>Programme de test</td><td>25</td></tr><tr><td>5.2</td><td>Résultats</td><td>25</td></tr><tr><td>6</td><td>Discussion</td><td>28</td></tr><tr><td>6.1</td><td>Génération de code pour GPU</td><td>28</td></tr><tr><td>6.1.1</td><td>Restrictions</td><td>28</td></tr><tr><td>6.1.2</td><td>Qualificateurs spécifiques à OpenCL</td><td>28</td></tr><tr><td>6.1.3</td><td>Fonctions spécifiques à OpenCL</td><td>29</td></tr><tr><td>6.2</td><td>Protéger l'utilisateur contre lui-même</td><td>29</td></tr><tr><td>6.3</td><td>BLAS</td><td>31</td></tr><tr><td>7</td><td>Conclusion</td><td>32</td></tr><tr><td>7.1</td><td>Contributions</td><td>32</td></tr><tr><td>7.2</td><td>Perspectives</td><td>32</td></tr><tr><td>A</td><td>XcodeML</td><td>34</td></tr></tbody></table>
Table des figures

<table><tr><td>2.1 Des tâches peuvent être soumises par le programmeur et demandées par les unités de calcul. Elles sont alors ajoutées à une file. La définition de ces opérations (push et pop) permet de définir une politique d'ordonnancement.</td><td>15</td></tr><tr><td>3.1 Découpage des données</td><td>18</td></tr><tr><td>4.1 Chaîne de compilation</td><td>21</td></tr><tr><td>5.1 Test de scalabilité, effectués sur un cluster de Nehalem E5504</td><td>26</td></tr><tr><td>5.2 Certains noeuds sont équipés de GPUs plus puissants que d'autres, et terminent donc leur travail plus tôt. En vert, les temps d'exécution.</td><td>27</td></tr><tr><td>6.1 Exécution de la boucle sur deux processeurs. La dépendance implicite propage une erreur dans tout le tableau</td><td>30</td></tr></table>
# Listings

<table><tr><td>2.1</td><td>Exemple de code XcalableMP</td><td>13</td></tr><tr><td>3.1</td><td>Découpage des données par XcalableMP</td><td>17</td></tr><tr><td>3.2</td><td>L'ajout de nouveaux pragmas pourra être nécessaire</td><td>19</td></tr><tr><td>4.1</td><td>La fonction do-stuff</td><td>21</td></tr><tr><td>4.2</td><td>Fonction do-stuff compilée par xmpcc</td><td>22</td></tr><tr><td>4.3</td><td>Fonction CPU générée</td><td>22</td></tr><tr><td>4.4</td><td>Exemple d'utilisation du pragma &lt;loop&gt;</td><td>23</td></tr><tr><td>4.5</td><td>Prototype d'un noyau</td><td>23</td></tr><tr><td>6.1</td><td>Parcours d'une matrice sur un CPU</td><td>29</td></tr><tr><td>6.2</td><td>Parcours d'une matrice sur un GPU</td><td>29</td></tr><tr><td>6.3</td><td>Boucle dont les itérations sont dépendantes les unes des autres.</td><td>29</td></tr><tr><td>6.4</td><td>Boucle dont les itérations sont dépendantes les unes des autres et qui est parallélisée à l'aide d'une directive XcalableMP.</td><td>30</td></tr><tr><td>6.5</td><td>Les itérations de cette boucle ne sont pas indépendantes. On peut cependant facilement les distribuer sur deux noeuds</td><td>30</td></tr><tr><td>A.1</td><td>Programme C</td><td>34</td></tr><tr><td>A.2</td><td>Programme XML</td><td>34</td></tr></table>
Remerciements

Je tiens tout d'abord à remercier toute l'équipe Runtime, qui a su m'accueillir
pour ce stage, et notamment Raymond NAMYST, qui la dirige. Je remercie
également Olivier AUMAGE et Samuel THIBAULT, mes maîtres de stage, qui
m'ont fait confiance et ont su me guider dans mon travail. Merci à mes tuteurs,
Xavier GRANIER et François PELLEGRINI. Mention spéciale à ce dernier,
qui, au travers de ses activité d'enseignement à l'ENSEIRB-MATMECA, n'aura
cessé d'essayer de faire rentrer un peu de savoir dans ma caboche, qui semble
toujours bien étroite. Merci à Ludovic STORDEUR, toujours prêt à troller
tout en lançant vi dans emacs dans emacs... Merci également à mes camarades
stagiaires, Charles, Laurent et Lucas, qui m'ont fait découvrir le SED.

Enfin, parce qu'il existe une vie en dehors du laboratoire<fnref n="1" />, je remercie
Amélie, qui tente tant bien que mal de me faire garder un semblant de vie
sociale, Christelle, pour son exposé au sujet d'Aliénor d'Aquitaine, et Valentin
pour ses conseils concernant Haskell et les diverses discussions de nerds dans
lesquelles nous nous engageons (trop) souvent.

<footnote n="1">1. Certains vont même jusqu'à avancer qu'il existe une vie en dehors de l'informatique, sans toutefois pouvoir présenter de preuves convaincantes.</footnote>
# Chapitre 1

# Introduction

## 1.1 De l'évolution des machines

La fréquence des processeurs a atteint un seuil qui ne sera sans doute jamais dépassé, notamment à cause de la barrière thermique. Les besoins en puissance de calcul, eux, ne cessent d'augmenter. La combinaison de ces deux phénomènes a débouché sur une modification importante de l'architecture des ordinateurs.

Une première solution permettant d'améliorer la puissance de calcul sans augmenter la fréquence des processeurs consiste à utiliser des fermes de calcul. L'association de plusieurs machines permet en effet d'avoir accès à un nombre plus important d'unités de calcul, et les clusters sont désormais présents dans toutes les machines du top500[1].

Les machines multi-coeurs, encore réservées aux super-calculateurs il y a quelques années, sont aujourd'hui devenues la norme dans les machines destinées au grand public. Les processeurs graphiques étant particulièrement bons dans le domaine du calcul vectoriel, on a ensuite pensé à les utiliser pour effectuer des calculs génériques, donnant ainsi naissance aux GPGPUs (General-Purpose Processing on Graphics Processing Units).

Les architectures hétérogènes (machines composées d'unité de calcul différentes, comme des CPUs multi-coeurs, des GPUs...) actuelles posent elles aussi de nouveaux problèmes. Il n'est en effet pas trivial d'exploiter pleinement toutes les unités de calcul disponibles sur la machine. De plus, le code n'est pas assuré d'être portable : un code CUDA ne fonctionnera pas sur une machine équipée de GPUs AMD, par exemple. Quant aux performances, elles sont parfois très différentes d'une machine à une autre.

## 1.2 Nouvelles difficultés de programmation

L'émergence de machines massivement parallèles a eu pour conséquence d'amener les programmeurs à faire face à de nouveaux défis techniques.
MPI (Message Passing Interface), que nous présenterons dans le chapitre 2, est un standard actuel dans la programmation de machines à mémoire distribuée. Il utilise la technique du SPMD (Single Program, Multiple Data) : le même programme s'exécute simultanément sur plusieurs noeuds.

Toutefois, il est assez difficile d'écrire un programme MPI : il faut maîtriser le découpage des données, réécrire une grande partie d'un programme séquentiel que l'on souhaite paralléliser, et de plus le déboguage est complexe. Même s'il est parfois techniquement possible de réécrire tout un code séquentiel de façon à utiliser MPI, ce n'est pas souhaitable. En effet, cela reviendrait à accepter de développer une version séquentielle et une version parallèle. On pourrait également développer une version adaptée aux GPUs, etc. Au final, on devrait maintenir un nombre très important de branches différentes, chacune destinée à l'exploitation d'une architecture particulière, ce qui est une aberration.

Afin de diminuer l'effort de programmation nécessaire à la parallélisation d'un code séquentiel, de nouveaux modèles sont apparus et continuent à émerger. Ils proposent de paralléliser un programme simplement en l'annotant à l'aide de directives adressées au compilateur. OpenMP, un langage permettant de multi-threader facilement un programme, connaît par exemple un énorme succès.

## 1.3 Un modèle hybride

### 1.3.1 XcalableMP

XcalableMP, proposé par la communauté académique japonaise, est l'un de ces nouveaux langages à base de directives. Il offre au programmeur une couche d'abstraction placée au-dessus de MPI, lui permettant de découper facilement ses données et de les distribuer sur les noeuds composant son cluster.

Il donne également la possibilité au programmeur de paralléliser une application séquentielle de manière incrémentale, et ne l'oblige pas à réécrire l'intégralité du code.

Ce modèle de programmation ne permet toutefois pas d'exploiter au mieux les architectures rencontrées dans le monde du calcul haute performance (HPC). Il ne fournit par exemple pas de support pour les GPUs, ni pour les processeurs multi-coeurs.

Nous souhaitons donc proposer un nouveau modèle, basé sur XcalableMP, mais dont le support exécutif saurait exploiter l'hétérogénéité des machines actuelles. Ce modèle devra conserver la simplicité de XcalableMP. Un programmeur devra toujours avoir la possibilité de paralléliser un code séquentiel simplement en l'annotant. De plus, nous voulons que le code écrit à l'aide de ce modèle soit portable ; de même pour les performances d'un tel code.

### 1.3.2 Le support exécutif StarPU

StarPU est un support exécutif développé à l'INRIA. Il propose à l'utilisateur de penser son programme sous forme de tâches, potentiellement dépendantes
les unes des autres, qui peuvent être exécutées par les différentes unités de calcul disponibles sur la machine. StarPU s'occupe alors de l'ordonnancement des tâches, du transfert de données sur les GPUs et de la gestion des caches.

Il serait donc particulièrement intéressant d'utiliser StarPU en combinai-
son avec XcalableMP : le premier permettrait de découper les données sur les
différents noeuds composant un cluster, et le deuxième de maximiser l'efficacité
de chacun des noeuds.

Nous commencerons par présenter les différentes technologies permettant
d'exploiter des architectures parallèles. Ensuite, nous présenterons notre modèle
hybride ainsi que les éléments d'implémentation principaux. Nous finirons en
évaluant notre modèle et en présentant les problématiques auxquels nous serons
confrontés dans la suite de son développement.
# Chapitre 2

## Contexte

### 2.1 Passage de messages et processus légers

La programmation *multi-threadée* ainsi que la programmation de *clusters* de machines peuvent se faire à l'aide de langages et d'interfaces de bas niveau. Dans cette première section, nous présentons deux de ces interfaces : MPI et les threads POSIX.

#### 2.1.1 Message Passing Interface

MPI est une interface permettant de programmer des machines parallèles qui utilise le paradigme de passage de messages. Elle se base sur le principe SPMD (Single Process Multiple Data) : les noeuds qui constituent un *cluster* exécutent simultanément le même programme. Ils peuvent communiquer en s'échangeant des messages. Principalement pensée pour les langages C et Fortran 77 (bien que d'autres implémentations existent), elle se veut facilement portable sur différentes plateformes. Elle définit plusieurs types de communications (bloquantes ou non, collectives ou non...), ainsi que plusieurs modes de communication (bufferisée, synchrone, ...). Diverses implémentations sont disponibles : citons notamment OpenMPI et MPICH.

Toutefois, la réalisation d'un programme MPI n'est pas aisée. L'obtention de bonnes performances est subordonnée non seulement à une maîtrise de l'interface de programmation, mais aussi de ses concepts : communications bloquantes ou non, permanentes ou non, recouvrement calcul/communications... La parallélisation d'un code séquentiel existant demande à réécrire, si ce n'est l'intégralité du programme, tout du moins la majeure partie de celui-ci.

#### 2.1.2 Threads POSIX

L'utilisation de processus légers permet de répondre à différents besoins, que l'on souhaite améliorer la performance d'une application, comme c'est le cas en calcul haute performance, ou non. On peut par exemple citer le cas courant
des programmes où l'interface graphique est codée dans un *thread* particulier, permettant au reste du programme de s'exécuter dans un ou plusieurs autres *threads*.

Bien que fonctionnelle sur des machines dotées d'un seul processeur, la programmation multi-threadée présente beaucoup plus d'intérêt sur des machines multi-coeurs ou multi-processeurs, qui permettent d'exécuter simultanément plusieurs fils d'exécution.

Malheureusement, la programmation multi-threadée peut s'avérer relativement difficile à appréhender : au-delà de l'habituelle difficulté à maîtriser parfaitement une bibliothèque, un programmeur peut rencontrer des bogues jusque-là inconnus : deadlock, livelock, race condition.<fnref n="1" />

## 2.2 Deux approches de plus haut niveau

Les modèles présentés dans la section précédente sont d'assez bas niveau. Afin de faciliter la programmation, plusieurs approches de plus haut niveau ont été proposées dans les années 90. Nous verrons ici deux des plus connues :

- HPF, qui applique un parallélisme par distribution de données;

- OpenMP, qui utilise un parallélisme de contrôle en définissant des régions parallèles dans le code.

### 2.2.1 High Performance Fortran

Afin de programmer efficacement des machines parallèles, une extension au langage Fortran 90 a été développée. Il s'agit de High Performance Fortran, qui apporte à Fortran de nouvelles instructions, de nouvelles routines ainsi que des directives particulières.

HPF permet ainsi de distribuer facilement et efficacement les données d'un programme sur plusieurs processeurs. Les procédures peuvent désormais être annotées afin d'indiquer qu'elles sont des fonctions pures, ce qui est particulièrement intéressant dans le cadre de la programmation parallèle : une fonction pure ne dépendant pas de paramètres extérieurs (variables globales, entrées-sorties...) et n'ayant aucun effet de bord peut en effet être utilisée sans risque.

### 2.2.2 OpenMP

OpenMP est un langage de haut niveau permettant d'écrire des programmes multi-threadés à l'aide d'un ensemble de directives adressées au compilateur, de fonctions de bibliothèques et de variables d'environnement. Il est destiné aux machines à mémoire partagée.

OpenMP utilise le modèle «fork-join». Au début d'un programme OpenMP, il n'existe qu'un seul thread. Lorsque ce dernier rencontre, au cours de son

<footnote n="1">1. Le lecteur nous pardonnera l'utilisation de ces termes anglais couramment rencontrés dans la littérature. Notons que nous pourrions les traduire, en bon français, par «interblocage», «interblocage actif» et «situation de compétition».</footnote>
exécution, une région parallèle il crée un groupe de threads (fork) qui exécutent
le code de cette région avant de terminer (join).

Il a la caractéristique fort intéressante de permettre au programmeur de
paralléliser son application de manière incrémentale et sans apporter de mo-
difications très importantes au code séquentiel existant. Ainsi, l'ajout d'une
directive suffit à paralléliser une région du code et à obtenir une amélioration
des performances plus ou moin plus ou moins importante.

## 2.3 Parallélisme inter-noeuds et intra-noeuds

### 2.3.1 XcalableMP

Actuellement, MPI est un standard dans la programmation des clusters.
Toutefois, il présente quelques désavantages que nous avons détaillés dans la
section 2.1. Un nouveau langage de haut niveau, XcalableMP, espère remédier
à ceux-ci. A base de directives, il doit permettre de paralléliser simplement
un code séquentiel en l'annotant. Il serait ainsi possible de s'attaquer à la pa-
rallélisation d'une application donnée étape par étape, sans avoir à réécrire
quasi-intégralement le programme. On peut ainsi voir dans le listing 2.1 qu'il
est facile de paralléliser une boucle : l'ajout du pragma permet de répartir les
itérations de la boucle sur les différents noeuds.

1 \#pragma xmp loop on t (i)

2 for (i = 0; i < N; i++)

3 tab [i] = 2*i;

<caption>Listing 2.1 – Exemple de code XcalableMP</caption>

Etant conceptuellement similaire à OpenMP, il devrait être relativement
rapide à maîtriser. De plus, sa simplicité devrait limiter les risques d'introduire
des bogues. Il reprend également des notions apparues avec HPF, notamment
en ce qui concerne la distribution des données (on y retrouve une syntaxe quasi-
similaire pour le découpage par blocs, par blocs cycliques, etc.).

Il est enfin possible de coupler plusieurs modèles de programmation. Un
développeur pourra par exemple utiliser conjointement XcalableMP et OpenMP,
ou encore écrire une partie de son application directement à l'aide de l'API
fournie par MPI.

#### Spécification

Notons enfin que, bien qu'il ne soit pas encore largement utilisé dans la
communauté scientifique, ce langage dispose déjà d'un compilateur fonctionnel
toujours en cours de développement (xmpcc) et est capable de compiler des
programmes parallèles. Une spécification[2] ainsi qu'une page de manuel[3] ont
été publiées.
XcalableMP permet au programmeur de décrire la distribution des données
à l'aide de quelques directives. Il suffit d'aligner un tableau sur un *template* pour
que qu'il soit distribué sur les noeuds utilisés par l'application.

XcalableMP propose des directives plus ou moins basiques permettant aux
différents noeuds de se synchroniser et d'échanger des informations. L'exemple
le plus simple est le pragma «barrier», qui place dans le code compilé une
barrière de synchronisation. On peut également citer la construction «gmove»,
qui permet d'effectuer des communications «all-to-all». Un autre mécanisme
puissant est celui du «shadow/reflect» qui permet de partager certains éléments
nécessaires à un noeud lors d'un calcul mais situés sur un autre noeud.

Toutefois, XcalableMP n'est capable de gérer ni les processeurs multi-coeurs
ni les GPUs. Des supports éxécutifs peuvent cependant tirer parti des noeuds
hétérogènes.

2.3.2 StarPU

La programmation des machines hétérogènes modernes est relativement com-
plexe. Les développeurs doivent en effet exploiter toute la puissance fournie par
la machine à l'aide de nouvelles méthodes de programmation (OpenCL, Cuda
pour les GPUs Nvidia, Stream pour les GPUs AMD...) tout en conservant la
portabilité du code et des performances. Estimer correctement les coûts des
transferts de données vers et depuis les GPUs est crucial pour l'obtention de
bonnes performances.

StarPU, un support exécutif développé à l'INRIA dans le cadre de la thèse
de Cédric Augonnet, tente d'aborder ces problèmes en utilisant le paradigme
de parallélisme de tâches. Un programmeur peut utiliser la bibliothèque mise à
sa disposition afin de découper son programme en différentes tâches, qui seront
ensuite exécutées sur les différentes unités de calcul disponible sur la machine.

Tâche

Les applications développées à l'aide de StarPU sont pensées sous forme de
tâches (potentiellement non indépendantes)

Codelet

Une codelet modélise une tâche, notamment en précisant quelles unités de
calculs peuvent être utilisées pour l'effectuer (CPU, GPU avec OpenCL, GPU
avec un autre modèle de programmation...).

Ordonnancement

Distribuer les tâches sur les unités de calcul de manière statique et portable
est sans doute impossible. Il convient donc d'ordonnancer dynamiquement les
tâches en définissant des politiques d'ordonnancement, ce que permet de faire
StarPU (cf. figure 2.1). Le choix d'une politique d'ordonnancement adéquate
peut avoir un impact très important sur les performances de l'application [4].
<figure><img src="image_3.png" /><figcaption>FIGURE 2.1 – Des tâches peuvent être soumises par le programmeur et demandées par les unités de calcul. Elles sont alors ajoutées à une file. La définition de ces opérations (push et pop) permet de définir une politique d'ordonnancement.</figcaption></figure>

## Gestion de la mémoire

StarPU dispose d'un mécanisme de mémoire virtuelle partagée (Distributed Shared memory). Le programmeur doit enregistrer ses données auprès de cette DSM, qui se charge ensuite de gérer les transferts de données vers les GPUs. Le programmeur n'a donc pas à se soucier des mouvements de données[5].

## 2.4 Propositions de modèles pour les GPUs

### 2.4.1 HMPP

HMPP (Hybrid Multicore Parallel Programming) vise à faciliter la programmation des GPU. C'est un langage à base de directives, qui permet de définir des codelets (des fonctions pures qui seront exécutées sur un accélérateur) qui peuvent ensuite être appelés dans le code.[6]

HMPP gère également les transferts de données vers les accélérateurs et propose des générateurs de code permettant de compiler du code C vers du code spécifique aux GPUs (OpenCL, par exemple).

### 2.4.2 StarSs

StarSs est un autre modèle de programmation proposant d'annoter du code séquentiel afin d'indiquer quelles parties de ce dernier peuvent être exécutées en parallèle. Destiné à être utilisé dans des codes C et Fortran, il laisse la possibilité au programmeur d'utiliser explicitement MPI.
# Chapitre 3

# Un modèle hybride pour
l'exploitation de clusters

## 3.1 Clusters de machines hétérogènes

Bien qu'encore peu répandus dans les machines du top500, les GPUs semblent être l'avenir des machines dédiées au calcul haute performance. Cela ne remet toutefois pas en cause la notion de *cluster* de machines. Les deux approches sont complémentaires : il est tout à fait envisageable d'associer un grand nombre de machines constituées de GPUs et de CPU multi-coeurs. Encore faut-il trouver un modèle permettant de les programmer.

### 3.1.1 Architectures distribuées

Plusieurs langages permettent d'obtenir de bonnes performances en utilisant des grappes de machines. Ainsi, HPF et MPI sont des solutions envisageables pour programmer ce type de machines.

### 3.1.2 Machines hétérogènes

Il est plus compliqué d'utiliser des machines très hétérogènes de façon portable. Des propositions ont été faites dans ce sens (StarPU, HMPP...) et donnent des résultats parfois plus qu'encourageants.

### 3.1.3 Combinaison

Il n'existe toutefois pas de modèle permettant de programmer des grappes de machines hétérogènes. Certes, un programmeur peut écrire un code hybride. Il lui suffirait de distribuer les données sur les différents noeuds à sa disposition à l'aide de MPI, par exemple, et de gérer le parallélisme fourni par chacun des noeuds à l'aide d'OpenMP et de langages dédiés aux accélérateurs. Cette solution a deux inconvénients majeurs :
- Elle n'est pas portable : la gestion des threads et des accélérateurs se fera différemment sur une autre grappe;

- Elle n'est pas simple : le développement d'un tel code sera à la fois extrêmement complexe et chronophage.

Il convient donc de comprendre en quoi les modèles permettant d'exploiter des grappes de machines échouent à gérer l'hétérogénéité des noeuds, puis d'imaginer une solution permettant de remédier à ce problème.

## 3.2 Les limites de XcalableMP

### 3.2.1 Découpage des données

XcalableMP permet au programmeur de répartir ses données sur plusieurs noeuds de différentes manières (distribution par blocs, cyclique, cyclique par blocs) grâce à un code similaire à celui présenté au listing 3.1, illustré par la figure 3.1.

<figure><img src="image_4.png" /><figcaption>Listing 3.1 – Découpage des données par XcalableMP</figcaption></figure>

### 3.2.2 Sous-exploitation des noeuds

Ainsi, lors de la parallélisation d'une boucle for, les itérations de la boucle sont réparties sur chacun des noeuds. Si la répartition des itérations est satisfaisante (dans le cas d'un cluster de noeuds homogènes), on peut légitimement penser que les noeuds sont souvent sous-exploités. En effet, un code séquentiel y sera exécuté. Peu importe la topologie de la machine (multi-coeurs, dotée de GPU...), elle n'exécutera pas d'instructions en parallèle. Certes, il est possible de lancer plusieurs processus sur un même noeud, mais cela ne donnera pas les meilleures performances et n'utilisera de toute manière pas les GPUs.

XcalableMP pourrait profiter des avantages d'un support exécutif : gestion du parallélisme (CPUs, GPUs) de chacun des noeuds, simplicité de programmation et niveau d'abstraction plus élevé. Nous nous refusons toutefois à complexifier inutilement le modèle existant. L'utilisation d'un support exécutif doit rester transparente pour l'utilisateur.
<figure><img src="image_5.png" /><figcaption>FIGURE 3.1 – Découpage des données</figcaption></figure>

## 3.3 Utilisation d'un support exécutif

### 3.3.1 Programmer les noeuds avec StarPU

StarPU, que nous avons présenté en 2.3.2, permet d’exploiter au mieux un noeud constitué de différents types d’unités de calculs : CPUs multi-coeurs, GPUs, processeurs Cell... Il garantit de plus la portabilité du code et des performances.

On comprend alors tout l’intérêt qu’un programmeur pourrait avoir à découvrir ses données en utilisant XcalableMP (parallélisme inter-noeuds), et à les traiter sur chaque noeud à l’aide de StarPU (parallélisme intra-noeuds). Toutefois, dans le cas où il souhaite paralléliser un code séquentiel existant, cela l’obligerait à réécrire une grande partie de son application afin d’utiliser StarPU. Il perdrait alors une des fonctionnalités les plus intéressantes de XcalableMP : la possibilité de paralléliser un code de manière incrémentale, simplement en l’annotant.

### 3.3.2 Vers un modèle hybride

Nous nous devons donc de proposer un modèle de programmation permettant d’abstraire à la fois le découpage des données et leur traitement. Nous présentons dans le présent mémoire un modèle hybride reprenant les principes et la syntaxe de XcalableMP, et tirant parti de StarPU dans le back-end afin d’utiliser au mieux chacun des noeuds sur lesquels s’exécuteront les applications l’utilisant.

L’ajout de StarPU au modèle proposé par XcalableMP nous obligera sans doute à effectuer quelques changements dans le langage. Ainsi, certaines in-
formations seront nécessaires à StarPU, mais ne pourront sans doute pas être déduites du code source. Citons l'exemple de l'enregistrement des données auprès de StarPU : son API permet d'enregistrer des matrices, à condition de connaître leur taille. Leur découpage n'est pas non plus trivial. On peut penser à demander à l'utilisateur de nous guider en lui demandant des renseignements très simples concernant son application, à l'aide de nouveaux pragmas, avec une syntaxe similaire à celle présentée au listing 3.2.

```
#pragma xmp register_matrix m width=1024 height=512
```

1

```
#pragma xmp register_matrix n width=512 height=1024
```

2

3

```
/* Code using m and n */
```

4

5

```
/* m and n are not needed any more */
```

6

```
#pragma xmp unregister_matrix m
```

7

```
#pragma xmp unregister_matrix n
```

8

<caption>Listing 3.2 – L’ajout de nouveaux pragmas pourra être nécessaire</caption>

## Simplicité

Nous voulons faire en sorte que ce modèle hybride garde la simplicité de XcalableMP. La plupart des changements seront donc effectués dans le compilateur, afin de conserver le niveau actuel d’abstraction fourni par XcalableMP.

## Performances

Notre objectif principal est bien entendu de faire en sorte que notre modèle hybride obtienne de meilleures performances que XcalableMP dans sa version non modifiée. Il conviendra également de vérifier que notre modèle passe à l’échelle. Pour ce faire, nous présenterons une évaluation de ce modèle dans diverses configurations matérielles.

## Générer le code des accélérateurs

Nous avons choisi de programmer les GPUs à l'aide d'OpenCL, un standard ouvert proposé par le Khronos Group, et supporté largement par de nombreux constructeurs. Dans une version plus évoluée du prototype, on pourra éventuellement proposer un support pour d'autres technologies similaires. Comme on le verra par la suite, la compilation de code OpenCL depuis un code C est un problème particulièrement complexe. Nous considérerons la plupart du temps que le programmeur fournit lui-même les noyaux OpenCL qu'il souhaite utiliser.
# Chapitre 4

# Implémentation

## 4.1 Présentation du compilateur

### 4.1.1 Fonctionnement

Attardons-nous tout d'abord sur la figure 4.1, qui décrit le fonctionnement du compilateur fourni par l'équipe de XcalableMP (xmpcc). Après avoir été traité par le préprocesseur, le code source est analysé par le module L2X, qui génère une sortie XcodeML, qui est la représentation interne de xmpcc (cf. annexe A). Des modifications sont ensuite apportées par le module LX2X afin de remplacer les directives XcalableMP par du code. Enfin, du code C est généré depuis le code XcodeML par le module X2L. Ce code est enfin compilé par le compilateur C installé sur la machine de l'utilisateur.

### 4.1.2 Modifications

Afin de pouvoir insérer du code StarPU dans le backend, nous avons modifié le module LX2X en suivant la même démarche que les développeurs de xmpcc : nous avons ainsi défini nos propres fonctions dans une bibliothèque C (la libxmp) et inséré des appels à ces fonctions dans le code source généré par le compilateur. Cette architecture rappelle celle adoptée par GCC pour l'implémentation d'OpenMP[7], et nous permet de minimiser notre empreinte sur le code, et donc de faciliter l'intégration des changements apportés par ses développeurs. Les modifications les plus lourdes ont été effectuées dans le traitement du pragma «loop» : création du kernel CPU, collecte des arguments, création de noyau OpenCL lorsque c'est possible... Ces modifications restent localisées dans un seul fichier. D'autres modifications sont plus légères mais touchent un plus grand nombre de fichiers : c'est le cas de l'intégration de la syntaxe OpenCL dans le backend.
<figure><img src="image_6.png" /><figcaption>FIGURE 4.1 – Chaîne de compilation</figcaption></figure>

### 4.1.3 Un exemple : le pragma task

Donnons ici un exemple simple de modification permettant d'amener le compilateur à insérer des appels à StarPU dans le code source qu'il génère. XcalableMP fournit un pragma «task» qui permet de définir une tâche à exécuter sur un ou plusieurs noeuds. Il semble naturel de modifier la manière dont ce pragma est traité par le compilateur afin que ce dernier génère une tâche StarPU.

Intéressons-nous à la fonction `do'emulate` du listing 4.1 et à sa version compilée présentée dans le listing 4.2.

```c
1  static int answer = 0;
2  void do'emulate(int param) {
3   #pragma xmp task on p(1)
4   {
5    answer = param;
```
6
7
}
}

<caption>Listing 4.1 – La fonction do_stuff</caption>

1
2
3
4
5
6
7
8
9
10
11
{
    if (_XMP_exec_task_NODEs_PART(0, _XMP_DESC_p, (int
    )(0), (int)(1), (int)(1), (int)(1))) {
        void *xmp_starpu_argv[2];
        (*xmp_starpu_argv) = ((void *)(&*(param)));
        (*xmp_starpu_argv + (1))) = ((void *)(&(
        answer)));
        xmp_starpu_do_task(xmp_starpu_cpu_func_0,
                           xmp_starpu_argv, 2);
        XMP_pop_n_free_nodes();
    }
}
}

<caption>Listing 4.2 – Fonction do_stuff compilée par xmpcc</caption>

On remarque qu’un branchement conditionnel a été ajouté lors de la compi-
lation (ligne 3). Il permet de s’assurer que le noeud qui est en train d’exécuter
cette partie du code est bien le noeud 1 (comme demandé par le programmeur
à la ligne 3).

Le corps du «if» est intéressant. Là où la version originelle de xmpcc aurait
simplement laissé le code existant (ligne 5 du programme non compilé), notre
modèle hybride a ajouté un appel à la fonction xmp_starpu_do_task que nous
avons nous-même définie. Elle prend en paramètre un pointeur vers le noyau
StarPU pouvant être exécuté sur un CPU (xmp_starpu_cpu_func_0), un tableau
de pointeurs vers les variables utilisées dans le bloc mais définies à l'extérieur
de celui-ci, et la taille de ce tableau.

1
2
3
4
5
6
7
8
9
10
11
12
{
    void * * arg;
    int * pp_param;
    int * pp_answer;
    arg = argv;
    pp_param = ((int *)(*(arg)));
    pp_answer = ((int *)(*(arg + (1))));
    {
        (*(pp_answer)) = (*(pp_param));
    }
}

static void xmp_starpu_cpu_func_0(void * * buffers, void * argv)
<caption>Listing 4.3 – Fonction CPU générée</caption>

La fonction `xmp_starpu_do_task` crée alors une tâche StarPU qui exécute `xmp_starpu_cpu_func_0` en lui passant les arguments nécessaires. Le paramètre «buffers» n'est pas utilisé ici, puisqu'aucune donnée distribuée n'est utilisé : «answer» et «param» sont ici de simples paramètres.

## 4.2 Application au cas des matrices

Notre prototype est bien sûr imparfait, et ne pourra donc être parfaitement fonctionnel avec tous les codes possibles et imaginables. Toutefois, il devrait être opérationnel avec des codes simples. Nous nous intéressons ici au cas du traitement des matrices.

### 4.2.1 Idée générale

Un développeur utilisant XcalableMP peut utiliser la directive «loop» pour paralléliser le traitement d'une boucle for. Il lui est par exemple possible d'écrire le code présenté au listing 4.4.

```c
1 #define MAX 16
2 #pragma xmp nodes p(*)
3 #pragma xmp template t(0:15)
4 #pragma xmp distribute t(block) onto p
5
6 ...
7
8 #pragma xmp loop on t(i)
9 for (i = 0; i < MAX; i++)
10 /* Instructions */
```

<caption>Listing 4.4 – Exemple d'utilisation du pragma «loop»</caption>

Les itérations de la boucle seront alors réparties sur les noeuds utilisés par l'application et y seront traitées.

Afin d'utiliser le support exécutif StarPU sur chacun des noeuds, nous commençons par générer les noyaux dont il aura besoin. Le noyau CPU est composé des instructions de la boucle parallélisée. Les variables non locales à cette boucle sont passées en paramètres, ainsi que les «buffers» (cf. listing 4.5). Nous aborderons le cas des noyaux OpenCL en 4.2.3.

```c
1 void kernel (void **buffers, void *args);
```

<caption>Listing 4.5 – Prototype d'un noyau</caption>
Nous introduisons ensuite lors de la compilation un appel à la fonction `xmp_starpu_do_matrix_loop`, développée par nos soins. Elle commence par enregistrer les données auprès de StarPU, qui gérera les transferts vers les GPUs, et éventuellement les découper. Elle crée ensuite des tâches, qui seront finalement ordonanncées par StarPU.

### 4.2.2 Découpage des données

Dans un premier temps, notre prototype pouvait se permettre de ne pas découper les données traitées par StarPU. Toutefois, il est vite devenu indispensable de le faire. Prenons l'exemple de la multiplication de matrices. Soient A et B deux matrices, et C leur produit. Découper C est obligatoire : c'est en effet la seule matrice accédée en écriture, ce qui implique que plusieurs tâches ne peuvent y accéder de manière concurrente. Afin de pouvoir exploiter du parallélisme, il est donc nécessaire de fournir à chaque unité de calcul une sous-matrice qui pourra être modifée sans crainte.

Les matrices A et B ne sont accédées qu'en lecture, et peuvent donc être utilisées simultanément par plusieurs tâches. Toutefois, on gagnerait à les diviser en sous-matrices, afin de ne transférer vers les GPUs que des données utiles.

### 4.2.3 Gestion des GPUs

Les GPUs posent des problèmes à la fois syntaxiques et sémantiques. Voyons pourquoi les seconds sont bien plus complexes à résoudre que les premiers.

#### Support au sein du compilateur

Le compilateur xmpcc est prévu pour générer du code C, et ne sait donc compiler vers un langage spécialement développé pour les accélérateurs. OpenCL étant tout de même assez proche syntaxiquement du C99, il nous a été possible de modifier le backend afin de supporter de nouveaux qualificateurs de types («_global») et de fonctions («_kernel»).

#### Génération de noyaux OpenCL

Le support des éléments syntaxiques propres à OpenCL ne suffit pourtant pas à générer des noyaux corrects. Reprenons l'exemple de la multiplication de matrices. Le code séquentiel calculera chaque $C(i, j)$ à l'aide de l'habituelle triple boucle «for». Ce code pourra être repris tel quel dans un noyau CPU, mais pas dans un noyau OpenCL. En effet, afin d'exploiter correctement le parallélisme offert par les GPUs, on veut pouvoir calculer les $C(i, j)$ en parallèle.

La compilation de C vers OpenCL est donc loin d'être un problème trivial. Toutefois, il est raisonnable de penser qu'il est possible d'écrire un compilateur générant des noyaux OpenCL depuis un code C destiné à un CPU. Nous considérons donc pour l'instant que l'utilisateur fournit ses noyaux, et ce tant que leur génération ne sera pas réalisable automatiquement.
# Chapitre 5

## Evaluation des performances

### 5.1 Description des tests

#### 5.1.1 Environnement

Plusieurs machines sont utilisées pour ces tests :

- une machine dotée d'un Nehalem X5550 à 2.67Ghz et de 3 Nvidia Quadro FX5800;

- une machine équipée d'un Westmere X5650 à 2.67Ghz et de 3 Nvidia Tesla C2050;

- un clusters de Nehalem E5504.

#### 5.1.2 Programme de test

Nous avons effectué un produit de matrices sur des CPUs et des GPUs. Le noyau CPU utilisé par StarPU est généré par notre compilateur, tandis que le noyau OpenCL est fourni.

### 5.2 Résultats

Nous avons d'abord testé notre programme sur un seul noeud. En utilisant 8 threads, nous obtenons des résultats très encourageants. La multiplication de 2 matrices de taille 2048 * 2048 prend 36,72 secondes au lieu de 229,31 secondes lorsque le programme est compilé avec le compilateur originel. Nous observons donc un gain de 6.24.

L'utilisation de GPUs est rendue simple par notre modèle, qui donne des performances intéressantes. L'utilisation de 2 NVIDIA Quadro FX5800 nous permet d'effectuer le même calcul que précédemment en seulement 28,36 secondes, soit 8 fois plus rapidement qu'avec la première version de xmpcc.
Le programme compilé par notre compilateur hybride semble de plus passer correctement à l'échelle, comme le montre la figure 5.1

<figure><img src="image_7.png" /><figcaption>FIGURE 5.1 – Test de scalabilité, effectués sur un cluster de Nehalem E5504</figcaption></figure>

L'exploitation de GPUs répartis sur plusieurs noeuds donne de moins bons résultats. En effet, certains noeuds ont une capacité de calcul bien plus importante que d'autres, comme en témoigne la figure 5.2. Ce n'est pas une surprise : notre modèle hybride sait gérer l'hétérogénéité de chacun des noeuds d'un *cluster*, mais pas l'hétérogénéité des noeuds du *cluster* ! Il est donc adapté à des grappes de machines hétérogènes étant toutes similaires. Il conviendra cependant de réfléchir à l'ordonnancement global des tâches afin de pouvoir donner plus de travail aux noeuds les plus puissants.
<figure><img src="image_8.png" /><figcaption>FIGURE 5.2 – Certains noeuds sont équipés de GPUs plus puissants que d'autres, et terminent donc leur travail plus tôt. En vert, les temps d'exécution.</figcaption></figure>
# Chapitre 6

## Discussion

Nous discutions ici des évolutions envisagées pour notre modèle hybride.

### 6.1 Génération de code pour GPU

Comme nous l'avons vu dans ce mémoire, la compilation d'un code C vers un code OpenCL (ou vers tout autre langage permettant de programmer des accélérateurs) est particulièrement difficile et pourrait, à elle seule, faire l'objet d'un travail de thèse.

La difficulté principale vient du fait qu'OpenCL est un langage proche du C99, mais toutefois assez différent[8]. Ainsi, certaines fonctionnalités du C99 sont «illégales» en OpenCL, et vice-versa.

#### 6.1.1 Restrictions

Certaines fonctionnalités du C99 ne sont pas permises en OpenCL (pointeurs de fonctions, champs de bits, fonctions récursives...). Ces restrictions poseront inévitablement problème lors de la compilation. Certaines d'entre elles pourront sans doute être contournées; d'autres empêcheront la production d'un code correct. Il serait alors bon d'indiquer au programmeur la raison précise de l'échec de la compilation, afin qu'il puisse aisément corriger son programme.

#### 6.1.2 Qualificateurs spécifiques à OpenCL

L'utilisation d'OpenCL induit l'introduction de nouveaux qualificateurs. Le qualificateur de fonction « _kernel » est ainsi utilisé pour définir les noyaux OpenCL. D'autres qualificateurs sont également introduits afin de désigner quatres régions de l'espace d'adressage définies par OpenCL : « _global », « _local », « _constant », « _private ». Dans l'état actuel de notre prototype, seul le qualificateur « _global » est géré (les arguments des fonctions sont par défaut dans l'espace « _private »).
### 6.1.3 Fonctions spécifiques à OpenCL

Certaines fonctions sont spécifiques à OpenCL et n'auraient aucun sens dans un code C. Citons par exemple `get_global_id()`, qui retourne l'identifiant du « work-item » dans une dimension donnée, ou encore `get_work_dim` qui retourne le nombre de dimensions du problème. Leur utilisation est toutefois indispensable dans la grande majorité des noyaux.

Ainsi, pour effectuer une opération sur chacun des éléments d'une matrice, un programmeur écrira sans doute un code semblable à celui présenté dans le listing 6.1. Le code OpenCL sera sensiblement différent, puisque l'on pourra s'affranchir de cette double boucle (cf. listing 6.2).

```
1 for (i = 0; i < HEIGHT; i++)
2 for (j = 0; j < WIDTH; j++)
3 /* do something with M[i*WIDTH+j] */
```

<caption>Listing 6.1 – Parcours d'une matrice sur un CPU</caption>

```
1 int i = get_global_id(0);
2 int j = get_global_id(1);
3
4 /* Do something with M[i * WIDTH + j] */
```

<caption>Listing 6.2 – Parcours d'une matrice sur un GPU</caption>

Ces différences syntaxiques ne sont que le reflet de différences sémantiques. En effet, là où un programmeur CPU veut exprimer « le contenu de la double-boucle doit être exécuté `WIDTH * HEIGHT` fois », un programmeur spécialiste d'OpenCL souhaite « exécuter une suite d'instructions sur `WIDTH * HEIGHT` work-items ». La différence architecturale entre un CPU et un GPU se ressent dans le code.

## 6.2 Protéger l'utilisateur contre lui-même

XcalableMP ne gère pas les dépendances entre les itérations d'une boucle. Considérons le listing 6.3. A la fin de l'exécution de cette boucle, on devrait avoir `t[0] = t[1] = t[2] = t[3] = 42`

```
1 for (i = 0; i < 4; i++)
2 tab[i] = (i == 0) ? 42:tab[i-1];
```

<caption>Listing 6.3 – Boucle dont les itérations sont dépendantes les unes des autres.</caption>

Un programmeur naïf voulant paralléliser cette boucle sur 2 noeuds écrira sans doute le code présenté dans le listing 6.4. Son programme produira alors un résultat erroné. En effet, avec une distribution par blocs, le premier noeud exécutera la boucle pour $i \in [0; 1]$ et le second pour $i \in [2; 3]$ (cf. figure 6.1).
1
2
3
```c
        #pragma xmp loop on t(i)
        for (i = 0; i < 4; i++)
            tab[i] = (i == 0) ? 42:tab[i-1];
```

<figure><img src="image_9.png" /><figcaption>FIGURE 6.1 – Exécution de la boucle sur deux processeurs. La dépendance implicite propage une erreur dans tout le tableau</figcaption></figure>

Il serait donc sans doute bon d'intégrer à xmpcc des vérifications concernant
le contenu des boucles qui est exécuté sur chaque noeud, afin de détecter des bugs
liés au parallélisme, et notamment les dépendances implicites entre itérations.
On pourrait éventuellement adapter le code afin que les résultats soient corrects,
bien que cela ne soit pas toujours possible. Ainsi, si l'on s'intéresse au listing
6.5, on peut remarquer qu'il suffirait de distribuer les itérations paires sur un
noeud et les itérations impaires sur un autre afin de paralléliser correctement et
efficacement la boucle présentée.

1
2
```c
        for (i = 2; i < N; i++)
            t[i] = 2 * t[i-2];
```

<caption>Listing 6.5 – Les itérations de cette boucle ne sont pas indépendantes. On peut cependant facilement les distribuer sur deux noeuds</caption>
## 6.3 BLAS

XcalableMP et StarPU sont principalement destinés à l'écriture de codes scientifiques. Il serait donc particulièrement intéressant d'implémenter les opérations les plus couramment rencontrées dans les codes scientifiques de manière optimisée.

On pourrait ainsi implémenter les opérations BLAS en ajoutant des fonctions à la bibliothèque fournie par XcalableMP. L'effort de programmation demandé aux utilisateurs serait toujours acceptable. De plus, le code pourrait être adapté à chaque machine en automatisant des tests afin de trouver les meilleurs paramètres pour cette machine particulière (taille des blocs, méthode d'ordonnancement...).
# Chapitre 7

## Conclusion

### 7.1 Contributions

Dans ce document, nous avons tout d'abord fait un tour d'horizon des solutions techniques permettant d'exploiter le parallélisme proposé par les machines modernes, que ce soit au niveau de clusters (MPI, HPF) ou au niveau d'un noeud hétérogène (HMPP, StarPU) comme ceux que l'on rencontre de plus en plus fréquemment, à la fois dans les machines dédiées au calcul parallèle et dans celles à la disposition de tout un chacun.

Le modèle de programmation proposé par XcalableMP nous semblant permetteur, et connaissant les bonnes performances de StarPU dans l'exploitation de machines hétérogènes, nous avons ensuite décidé de développer un prototype associant ces deux approches. Nous avons donc entrepris de modifier le compilateur fourni par l'équipe de développement de XcalableMP afin de lui faire générer du code C faisant appel à StarPU.

Nous avons donc ajouté à ce compilateur le support de la syntaxe d'OpenCL. Il peut désormais générer du code OpenCL dans des fichiers «.cl», bien que, dans la plupart des cas, la compilation de C vers OpenCL soit trop compliquée à effectuer. Nous lui avons ensuite permis d'utiliser StarPU dans le *backend*. Le découpage des données est donc assuré par XcalableMP, et les calculs par StarPU. Au cours de notre travail, nous avons également signalé différents bogues aux développeurs du compilateur, et parfois proposé des solutions éventuelles.

Enfin, l'évaluation de ce modèle hybride nous semble satisfaisante et nous amène donc à proposer des améliorations possibles dans la dernière partie de ce mémoire.

### 7.2 Perspectives

Le problème le plus important est malheureusement loin d'être le plus simple : la compilation d'un code C vers un code OpenCL. Comme nous l'avons vu
précédemment, les difficultés sont à la fois syntaxiques et sémantiques. L'im-
possibilité de générer du code OpenCL n'est pas un problème bloquant : le
modèle hybride que nous proposons peut fonctionner sans cette fonctionnalité,
mail il gagnerait à la voir implémentée.

La gestion des dépendances implicites (notamment entre les itérations d'une
boucle) est un problème moins important, mais auquel il conviendra tout de
même de réfléchir. Le compilateur pourrait sans doute être enrichi afin de pro-
duire des messages d'avertissement lorsqu'il pense reconnaître un motif qui
semble contenir un bogue, notamment afin d'aider les utilisateurs débutants.

Enfin, proposer une implémentation des BLAS permettra d'obtenir des per-
formances optimales, ce qui est particulièrement souhaitable dans le cadre du
développement d'application scientifique.
# Annexe A

## XcodeML

XcodeML est la représentation intermédiaire de xmpcc. La documentation est disponible en japonais. Donnons ici un exemple simple : un programme «hello world», dans sa version C et sa version XcodeML.

1 main()

2 {

3 printf("hello_world\n");

4 }

<caption>Listing A.1 – Programme C</caption>

```xml
<?xml version="1.0" encoding="ISO-8859-1"?>
<XCodeProgram source="t1.c">
  <typeTable>
    <pointerType type="P0" is_restrict="1" ref="char"/>
    <pointerType type="P1" ref="F1"/>
    <functionType type="F0" return_type="int">
      <params/>
    </functionType>
    <functionType type="F1" return_type="int">
      <params/>
    </functionType>
  </typeTable>
  <globalSymbols>
    <id type="F0" sclass="extern_def">
      <name>main</name>
    </id>
    <id type="F1" sclass="extern_def">
      <name>printf</name>
    </id>
  </globalSymbols>
</XCodeProgram>
```
22 <functionDefinition>
23 <name>main</name>
24 <symbols/>
25 <params/>
26 <body>
27 <exprStatement>
28 <functionCall type="int">
29 <function>
30 <funcAddr type="F1">printf</funcAddr>
31 </function>
32 <arguments>
33 <stringConstant>hello world
</stringConstant>
34 </arguments>
35 </functionCall>
36 </exprStatement>
37 </body>
38 </functionDefinition>
39 </globalDeclarations>
40 </XcodeProgram>

<caption>Listing A.2 – Programme XML</caption>
Bibliographie

[1] Overview of recent supercomputers - 2007 edition. http://www.top500.org/, 2007.

[2] Specification of xcalablemp. http://www.xcalablemp.org/xmp-spec-0.7a.pdf, 2010.

[3] Xcalablemp manual page. http://www.xcalablemp.org/manual/, 2011.

[4] Cédric Augonnet, Samuel Thibault, Raymond Namyst, and Pierre-André Wacrenier. StarPU : A Unified Platform for Task Scheduling on Heterogeneous Multicore Architectures. *Concurrency and Computation: Practice and Experience, Special Issue : Euro-Par 2009*, 23 :187–198, February 2011.

[5] Cédric Augonnet. StarPU : un support exécutif unifié pour les architectures multicœurs hétérogènes. In *19èmes Rencontres Francophones du PARALLÉLISME*, Toulouse / France, September 2009. Best Paper Award.

[6] S. Bihan R. Dolbeau and F. Bodin. A hybrid multi-core parallel programming environment. In *the proceedings of the Workshop on General Purpose Processing on Graphcis Processing Units (GPGPU 2007)*, Boston, Massachusetts, October 2007.

[7] D. Novillo. Openmp and automatic parallelization in gcc. In *the Proceedings of the GCC Developer's Summit*, June 2006.

[8] AMD. Introduction to opencl programming, May 2010.

[9] Alan Leung. Automatic parallelization for graphics processing units inkesRVM. Master's thesis, University of Waterloo, May 2008.