Dites-nous en plus sur votre projet

Nous vous proposerons une formule adaptée à vos besoins, ainsi qu’une estimation de devis.
Laissez-vous guider

Pourquoi choisir Aneo pour votre projet ?

Aneo est une agence conseil en transformation qui vous accompagne tout au long de la vie de vos projets sur des problématiques de :

Le +  d’Aneo : le conseil sans frontière !

Notre richesse, c’est de vous proposer des équipes pluridisciplinaires, chaque membre avec une approche, une expérience et une sensibilité différente : coach agile, formateur soft skills, Data Scientist designer, Architecte etc. On mise sur la complémentarité des disciplines pour vous apporter la meilleure réponse !

Pourquoi choisir Aneo  pour votre projet ? - Aneo

Aneo, une organisation à part entière !

Depuis 2015, Aneo est une organisation plate à gouvernance plate. Aplatir la hiérarchie, supprimer les silos, instaurer des bonus collectifs, ce nouveau modèle d’organisation avec comme objectif: apporter plus de valeur  à l’entreprise, aux collaborateurs et aux clients.

Le + d’Aneo : l’inattendu !

La sérendipité, c’est « le don de faire par hasard des découvertes fructueuses ». Chez Aneo, nous croyons fermement que les meilleures solutions sont parfois les plus inattendues. Ne rien s’interdire, expérimenter, oser s’entourer de profils atypiques et avoir une obsession : apporter la juste valeur.

Aneo, une organisation à part entière !  - Aneo

Qui êtes-vous ?

Vous êtes pour

Votre secteur

1 seul choix possible
Assurance & Protection sociale
Banque et Finance
Industrie & Services
Santé

Vos besoins

Plusieurs choix possibles
IT & Digital
Transformation des Organisations
Stratégie Business
Pilotage de projets

Détails

Des précisions à ajouter sur votre projet ? (facultatif)
C'est noté !
Nous avons pris en compte les spécificités de votre projet.
Nos équipes vous contacteront sous 48h pour en discuter plus amplement.
Votre prénom *
Votre nom *
Votre adresse email pro *
Votre numéro de téléphone *
Bien reçu !
Nos équipes vous contacteront sous peu
pour discuter de votre projet.

Description et portage du cryptage AES 128-bit sur GPU – Billet #4

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 (4×4), contenant initialement les 16 caractères à crypter.

new_state : autre tableau de travail (4×4), 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.

La suite de la série FPGA est disponible ici.

Crédit : Damien DUBUC

Ça peut aussi vous intéresser
Répartition du profit : satisfaire les actionnaires et la planète, c’est possible !

Répartition du profit : satisfaire les actionnaires et la planète, c’est possible !

Cette intervention s’inscrit dans le cadre du sommet NextGen sur la transformation des entreprises. J’aurai le plaisir d’y participer les 26-27 novembre prochains en animant deux tables rondes avec des décideurs du monde économique, politique et social. Pourquoi parler aujourd’hui de répartition des bénéfices ? Dans un article…
2 décembre 2020
Et si on créait un capitalisme plus vertueux grâce à la Blockchain et à l’Intelligence Artificielle ?

Et si on créait un capitalisme plus vertueux grâce à la Blockchain et à l’Intelligence Artificielle ?

Cette intervention s’inscrit dans le cadre du sommet NextGen sur la transformation des entreprises, auquel j’aurai le plaisir de participer le 26-27 novembre prochains en animant deux tables rondes avec des décideurs du monde économique, politique et social.     Le capitalisme est un sport collectif Chemins de fer, médicaments contre…
18 novembre 2020
Aneo vous présente Soter-Vision, outil de reconnaissance faciale !

Aneo vous présente Soter-Vision, outil de reconnaissance faciale !

Aldwin, la marque Aneo dédiée aux expertises HPC, IA, Cloud et Big Data a développé Soter-Vision. Il s’agit d’un outil de reconnaissance faciale. En effet, Soter-Vision est capable de reconnaître les visages en temps réel dans une foule. Il compare des photos de référence à des vidéos captées en direct.
2 novembre 2020
Vous avez un projet de transformation
digitale pour votre entreprise ?