La vectorisation sur FPGA

Après notre première expérience avec le développement et l'optimisation d'un design single-workitem avec le cryptage AES, nous poursuivons l'aventure avec un kernel multi-workitem que nous souhaitons vectoriser.

Nous allons voir quelles ont été les difficultés rencontrées lors de cet exercice, et puis comment la mémoire constante pourrait nous aider à scaler un peu plus.

L’application AES est adaptée aux instructions SIMT (Single Instruction Multiple Threads) car tous les caractères du state peuvent être mis à jour simultanément en suivant des instructions identiques. Ceci est notamment très efficace sur GPU, dont le parallélisme mis en avant par l'architecture n'est pas le même que celui sur FPGA.

Sur FPGA il est également souhaitable d’arriver à obtenir un design SIMT / vectorisé, pour deux raisons :

- il est plus proche du modèle d’exécution GPU, qui est celui avec lequel nous souhaitons faire l’analogie. Avoir une idée du rapport coût de portage VS performance de ce genre de design (qui est notre motivation principale) est le point central de notre étude.

- il est moins coûteux en utilisation de la board de fabriquer un pipeline avec des instructions vectorielles de taille n que de répliquer n fois un pipeline identique, et cela devrait réduire les effets de la contention I/O des pipelines par un facteur n. En supposant une vitesse d’exécution semblable on a donc une meilleure efficacité du design, qu’on peut tenter de passer à l’échelle derrière.

On pourrait donc espérer pouvoir produire un design qui ressemble plus à un code type GPU - et donc nécessitant a priori moins d’efforts de portage, et reflétera mieux les performances possibles à atteindre avec le FPGA.

Contraintes de mise en oeuvre

La mise en œuvre des opérations vectorielles se fait par ajout de lignes __attribute__ précédant la définition du kernel, où sont également spécifiés la taille de workgroup et le nombre de compute units.

La taille des opérations vectorielles doit diviser le nombre de workitems présents dans le workgroup; l’exemple suivant demande au compilateur de générer un kernel dont les workgroups sont en une dimension et de taille 32 workitems, avec des opérations vectorielles opérant sur 4 éléments à la fois :

__attribute__((num_simd_work_items(4)))

__attribute__((reqd_work_group_size(32,1,1)))

__kernel void vectoradd(…)

La documentation Altera informe que la taille d’une instruction SIMD est limitée à 16, ce qui permet une implémentation naturelle de l’AES en manipulant des states entiers d’un coup.

La ligne d’attribut __attribute__((num_simd_work_items(16))) indique au compilateur aoc que nous souhaitons générer un pipeline de largeur 16; cependant cela ne veut pas dire qu’il va y parvenir. La problématique est connue depuis de nombreuses années : la vectorisation automatique du code par un compilateur est difficile et il faut bien souvent les aider.

Dans notre cas, les premières tentatives de compilation sont infructueuses ; le compilateur nous informe qu’il n’est pas capable de vectoriser le kernel - du tout - , avec le message suivant :

Compiler Warning: Kernel Vectorization: branching is thread ID dependent ... cannot vectorize.

On réalise alors que certains acquis de la programmation GPU vont poser problème ici. Sur GPU, le branching dû à une instruction conditionnelle implique seulement une séquentialisation des instructions : c’est-à-dire que dans un warp, on exécute à tour de role les différents chemins empruntés à coup d’instructions SIMT sur les threads concernés. Sur FPGA, c'est apparemment rédhibitoire d’un point de vue vectorisation.

Sous cet aspect, plusieurs instructions de control flow peuvent être considérées comme ambigües dès lors qu’elles réferrent à l’indice du workitem (ci-dessous idx). Avec un peu de recherche, on devine les coupables :

if (idx < 16)  (…)

for (int i=idx; i<256; i+=16)

.       sbox_loc[i]=sbox_d[i];

Ces lignes de code sont tout à fait banales et très présentes en GPU mais vont empêcher la vectorisation du kernel FPGA ici. Si la première ne pose aucun problème à contourner ici, la seconde est cependant beaucoup plus ennuyante : aoc n’est pas capable de déterminer, qu’avec une taille de workgroup forcée de 16, la boucle ci-dessus donne un travail parfaitement identique à tous les workitems et peut donc être vectorisée.

C’est-à-dire qu’il va nous falloir réécrire la boucle pour le lui montrer:

for (int i=0; i<16; i++)

.       sbox_loc[idx_l+16*i]=sbox_d[idx_g+16*i];

Si ici la résolution est simple (bien qu’ennuyeuse), le problème est plus profond : si aoc n’arrive pas à gérer ce genre de cas, alors on peut douter de ses capacités à vectoriser des codes dont on ne peut ré-écrire les boucles. Par exemple celles dépendant d’un paramètre non connu à la compilation du kernel OpenCL, qui se fait à part. C’est-à-dire que :

- le cas où le pas de la boucle dépend du nombre de workgroups (stride de la taille de la grille par exemple) ne sera pas vectorisable

- le cas où le critère d’arrêt porte sur la taille d’une variable non connue à la compilation (comme la taille d’un tableau lu en entrée) non plus

Après quelques retouches rapides (et plusieurs heures de compilation infructueuses) le compilateur accepte de vectoriser notre kernel. Il nous fait cependant remarquer que certaines opérations de read and write n’ont pu être vectorisées :

Compiler Warning: Vectorized kernel contains loads/stores that cannot be vectorized. This might reduce performance.

Nous aurions beaucoup aimé savoir lesquelles! L’impact d’une ou deux instructions non-vectorisées pourraient ralentir tout le pipeline et réduire plus ou moins à néant le reste des efforts de vectorisation.

Sur CPU, tout ce que vous vectorisez sans détruire le code initial c'est un gain immédiat: votre temps d'exécution total c'est la somme des temps de chaque portion du code et vous l'avez effectivement réduite. Sur FPGA, toutes les données avancent de manière cadencée dans un pipeline d'instructions, dont la latence est effectivement régie par quelque chose qui se rapproche de la somme des latences de chaque portion de code mais dont le débit est donné par la portion la plus lente... Et plus vous avez de données à traiter, plus c'est le débit qui vous intéresse. Aïe.

L’efficacité de ce design est bien en-dessous du kernel single-workitem avec 16 cu : son temps d’exécution est à peu près 6 fois plus long. On émet alors immédiatement l’hypothèse que les opérations de read/write non-vectorisées ralentissent considérablement le reste de l’application.

Les deux boucles lisant les look-up tables sbox et roundkeys depuis la mémoire globale présentent une burst-rate de 2, tandis que celle lisant les states est à 16 (ce qui est un comportement identique à celui de notre kernel single-workitem). Sachant qu’il n’y a aucune différence visible entre ces 3 boucles à la compilation du kernel et que nous les avons aussi ré-écrites pour qu’elles ne posent pas de problème au compilateur, le problème vient d’ailleurs. Et il ne semble y avoir qu’un seul bon candidat : l’étape de substitution par la sbox avec un accès mémoire indirect (state[i] = sbox[state[i]]). Bien que les accès non-contigüs à la mémoire partagée aient un impact limité sur GPU, nous n’avons pas vraiment idée des conséquences de cette situation vis-à-vis de notre design FPGA.

Utiliser la mémoire constante pour ces look-up tables permettrait tout de même de se débarrasser de ces deux boucles à la performance ennuyeuse. Sur GPU, la mémoire constante à plusieurs spécificités qui doivent être connues afin d’assurer une meilleure performance que la mémoire globale. Sur FPGA, nous ne lui connaissons pas de caractéristiques particulières.

Enfin, l’actuel pipeline vectorisé n’est pas réplicable et consomme trop de ressources RAM. Une tentative de compilation avec 2cu renvoie l’erreur de compilation finale (puisqu’il faut attendre plusieurs heures avant de la voir) : Error : Kernel doesn’t fit

Nous pensons qu’il faut un changement notable en mémoire à ce design pour avoir l’espoir de le répliquer et espérons que la mémoire constante puisse nous venir en aide.

La mémoire constante pour scaler plus loin ?

L’utilisation de la mémoire constante n’a pas encore été abordée dans le cadre de nos expérimentations. Dans la documentation Altera, ce cache de mémoire est décrit comme idéal pour utiliser des look-up tables, avec un pourcentage de cache-hit très élevé. Contrairement à l’architecture GPU, où la taille de cette zone mémoire est figée, elle est réglable à la compilation sur FPGA. Par défaut elle vaut 16 ko, l’utilisateur peut régler sa taille à l’aide d’un flag de compilation :

-- const-cache-bytes N

où N est sa taille en bytes (et doit être une puissance de 2). Nous l’utilisons pour stocker les tables sbox et roundkeys, ce qui donne la valeur 512 pour n. En outre, les tables en questions doivent être passées en arguments au kernel précédées du mot-clef __constant.

L’utilisation de cette mémoire constante permet de s’affranchir des copies de look-up tables en mémoire locale, une approche privilégiée sur GPU.

Il est dit que sur GPU que si des threads d’un même warp souhaitent accèder à des adresses différentes en mémoire constante alors ces requêtes sont séquentialisées et que la transaction mémoire pourrait alors être moins efficace qu’une lecture groupée en mémoire globale. Sur GPU, la mémoire constante a été conçue pour idéalement diffuser une unique valeur à un warp entier en une seule requête. Aucun des documents Altera (Programming Guide & Best practices) ne mentionnent des particularités de la mémoire constante FPGA, qui autorise donc probablement les accès groupés (coalesced).

Dans notre cas, tous les threads accèdent à des valeurs différentes. Pour la roundkey, il s’agit de 16 valeurs successives en mémoire ; pour la sbox, il s’agit de 16 valeurs dont on ignore la ventilation. Quand bien même la mémoire constante permettrait de faire des accès groupés pour les éléments de la roundkey, on voit mal comment nous vectoriser les accès à la sbox. Une solution serait de ne plus utiliser la sbox en tant que look-up table, mais de retrouver ses valeurs par le calcul, puisque celle-ci est pré-déterminée.

Au-delà d’un gain en termes de temps d’exécution, nous attendons de voir aussi comment ce changement va se répercuter sur l’occupation de la board. En utilisant le cache de mémoire constante (qui fait partie de la mémoire globale) au lieu de stocker l’ensemble de nos données dans des registres de mémoire partagée, nous devrions réduire l’utilisation en blocs mémoire de notre kernel. Ceci, quitte à perdre en performance immédiate après cette modification, pourrait nous permettre de scaler à plusieurs compute units. Après compilation, nous obtenons les chiffres suivants :

Ce design utilisant la mémoire constante permet d’avoir un temps d’exécution plus petit et réduit d’environ 30% la quantité de RAM utilisée, ce qui nous permet de scaler jusqu’à 2 cu.

En revanche, le compilateur nous prévient toujours qu’il existe des opérations load / store non vectorisables et que par conséquent la performance de notre kernel n’est pas optimale (tout pointe vers la sbox). Nous remarquons qu’il n’est pas aisé de maximiser l’utilisation de la board, avec la réplication d’un pipeline de base dont l’occupation des ressources est aussi élevée : le niveau de granularité n'est pas le même que dans le cas d’un kernel single-workitem.

Ce billet conclut donc notre retour d'expérience utilisateur concernant le SDK OpenCL d'Altera (désormais Intel). Nous revenons dans un dernier billet afin de partager avec vous nos réflexions et perspectives sur le FPGA dans le monde du calcul scientifique et HPC.