La suite de ces billets s’articule autour d’une application de cryptage de données par AES.
Les raisons de ce choix se basent sur sa pertinence - très utilisée dans son domaine -, sur le parallélisme naturel mais non-trivial qui la caractérise et enfin sur le fait qu’elle fait uniquement appel à des opérations logiques et entières.
Ce choix, concernant le dernier point, se base sur le fait que les performances des opérations flottantes FPGA actuelles sont un cran en dessous de celles d’autres architectures. Altera continue d'améliorer ses IP de calcul flottant à chaque génération de FPGA, mais il est pour le moment préférable de se baser sur des opérations logiques et entières pour se faire une idée plus juste des différences entre ces deux architectures et envisager de comparer des performances.
C'est quoi l'AES 128-bit ?
Le cryptage de données par AES (Advanced Encryption Standard) est devenu depuis le début des années 2000 le nouveau standard de chiffrement de données ; il est actuellement le plus utilisé et le plus sécuritaire. Il se base sur la connaissance d’une clef unique, utilisée à la fois pour l’encryptage et le décryptage des données, et consiste en une suite d’opérations logiques et entières à appliquer sur l’ensemble du texte.
Dans sa version de base, l’algorithme prend en entrée un bloc de 128-bits (16 octets) et une clef de 128 bits (d’autres tailles de clef sont possibles, et altèrent légèrement l’implémentation). Pour crypter un message d’une longueur donnée, il faut donc le découper en paquets de 128 bits (appelés « states»), sur lesquels un cryptage identique va être réalisé : l’application expose donc un parallélisme naturel, dont la granularité est le « state », contenant 16 caractères.
C’est cette version de l’algorithme que nous utilisons. Le cryptage et le décryptage étant symétriques, nous nous focalisons uniquement sur la première phase, décrite ci-dessous.
Le cryptage d’un state (ensemble de 16 caractères, de taille 128 bits ici) se résume à l’application d’un ensemble de 4 opérations, répétées au cours de 10 tours de cryptage. Ces quatre opérations sont nommées ici subbytes, addroundkey, shiftrows et mixcolumns, dont vous trouverez des descriptions imagées ci-dessous, reprises sans vergogne de la jolie page wikipedia ;)
Le pseudo code se présente comme suit :
key : clef initiale ( 16 caractères)
roundkeys : clef étendue, pour les 10 tours
sbox : Rindjael substitution box requise par l’AES (look-up table, 256 caractères)
state : tableau de travail (4x4), contenant initialement les 16 caractères à crypter.
new_state : autre tableau de travail (4x4), contiendra les 16 caractères cryptés à la fin.
subbytes substitue l’ensemble des caractères du state par d’autres (in-place), par l’intermédiaire de la look-up table sbox, qui fait partie de la méthode AES : elle est pré-déterminée.
shiftrows shifte vers la gauche de n positions les éléments de la ligne n du state (la première ligne étant la ligne 0).
mixcolumns est l’équivalent un produit-scalaire de taille 4 pour chaque caractère, qui est donc remplacé par une combinaison linéaire (pré-définie, notée c(x) sur le schéma) de lui-même et des autres caractères de sa colonne.
addroundkey applique à chaque tour une partie différente de la clef étendue (générée à partir de la clef de base), simplement par le biais d’une opération XOR, à chaque caractère.
L’ensemble des dépendances des données au travers des différentes étapes, ainsi que la position dans le state des caractères lors de l’enchaînement des opérations, est donnée par la figure suivante.
Elle montre essentiellement la nécessité que le state soit entièrement à jour avant de passer d’une opération à l’autre, et donc l’incapacité à revenir à une granularité plus fine de parallélisme : mixcolumns a en particulier besoin de 4 éléments mis-à-jour par output, mais d’un tour à l’autre ces 4 éléments occupent des positions différentes et dépendent eux-mêmes d’autres éléments du mixcolumns précédant.
Implémentation sur GPU
Cette application s’implémente très naturellement sur GPU, où les threads / workitems sont exécutés de manière cadencée à coup d’instructions SIMT, par demi-warps de 16 threads. Dès lors que les instructions des 16 threads traitant chacun un caractère d'un state donné sont identiques et exécutées en simultané, chaque fonction (subbytes, etc) est embarassingly parallel.
Les deux tableaux de travail state et state_new ne contiennent que 16 caractères non-signés et chacun de leurs indices va être accédé de multiples fois par des threads différents : ils vont donc être alloués en mémoire partagée (__shared__ / __local).
La sbox est propre à l’AES (donc pré-définie), ne pèse que 256 octets, et va être lue de manière non structurée et multiple par tous les threads.
Enfin, les roundkeys sont calculées à partir de la clef de base, mais ces calculs ne sont ni lourds, ni parallélisables : cette table pesant 176 octets peut être calculée une unique fois sur CPU, puis envoyée sur GPU. Cette table sera aussi consultée en lecture par tous les threads.
Nous décidons de mettre ces tables en mémoire partagée dans un premier temps, bien que la mémoire constante puisse parfois être une alternative à envisager.
Essentiellement, pour un block / workgroup, le kernel se résume à :
- Copie d’un state (portion de texte en clair), de la sbox et des roundkeys en mémoire partagée (448 octets au total)
- Enchaînement des opérations uniquement sur des tables en mémoire partagée
- Copie du state crypté dans le tableau output, en mémoire globale (16 octets)
Les accès à la mémoire globale - une lecture et une écriture par state - sont efficaces, puisque des threads successifs accèdent à des adresses contigües.
Les warps étant constitués de 32 threads, il n’est pas efficace de seulement utiliser une taille de block / workgroup de 16 threads : le code GPU a été adapté pour une taille quelconque (de la forme 16*n), cryptant en parallèle n states, dont 2 par warp, afin de réduire le nombre de blocs ordonnancés et utiliser pleinement les warps. Ceci entraine également une diminution des copies en mémoire partagée des lookup tables sbox et roundkeys, désormais ré-utilisées n fois.
Sous ces optimisations raisonnables, dont l’implémentation et la vérification ne nécessitent que quelques heures, le kernel OpenCL donnant le cryptage d’un texte de taille 2^25, soit à peu près 33 millions de caractères, prend de l’ordre de 480 ms sur un GPU Nvidia Tesla C2075 (architecture Fermi). Nous gardons à l’esprit ce rapport qualitatif entre temps d’optimisation et performance obtenue, avant de passer au même exercice sur FPGA.D’autres optimisations plus fines, transformant sensiblement le kernel, ainsi que l’assistance du profiler d’NVidia permettent de ramener le temps d’exécution autour de 60 ms avec 2 jours-homme de travail supplémentaire. Néanmoins, nous sommes repartis de la formulation naturelle de l’AES (à 480 ms) afin de ne pas biaiser le développement FPGA, parce qu'on est des gens biens.
Dans le prochain billet, nous nous attaquerons à l'implémentation d'un premier design FPGA. Nous parlerons des problématiques rencontrées et des solutions qui nous permis d'aboutir à un premier design tout en obtenant des performances de calcul qui nous ont paru correctes.
