Calculs sur GPU : gain ou menace pour la sécurité informatique ?

MISC

n°

mars 2012

Par

Entrés dans le domaine de la sécurité SI par le « cassage » des hashs de mots de passe, quels sont aujourd'hui les apports du GPGPU en SSI ? Tour d'horizon des applications possibles et exemple pratique par l'implémentation d'une fonction de hachage cryptographique sur GPU.

1. Les concepts du calcul sur carte graphique

Le calcul sur carte graphique (abrégé GPGPU : General-Purpose computation on Graphics Processing Units) est apparu vers 2007 avec l'idée de réaliser des calculs généraux sur les composants matériels précédemment dédiés aux opérations graphiques 2D et 3D. De nombreux champs d'applications ont bénéficié de cette nouvelle approche, des calculs scientifiques aux simulations financières [1], en passant par la classification d'images [2].

Avec le GPGPU, la carte graphique devient un « coprocesseur » qui seconde le processeur principal d'un système en réalisant une partie ou la totalité de la charge de calcul d'un problème donné. Les algorithmes utilisés doivent être adaptés pour tirer profit de ce nouveau support, qui a pour spécificité d'être une architecture massivement parallèle.

Quels sont les gains attendus ? Plus que la vitesse pure (temps mis pour réaliser une tâche), c'est une augmentation des débits (quantité d'information traitée par unité de temps) qui est espérée sur cette plate-forme.De plus, des gains énergétiques ou financiers sont aussi possibles, à puissance de calcul équivalente.

1.1. Architecture et modèle de programmation

D'un point de vue simplifié, un processeur de carte graphique (GPU) regroupe un grand nombre d'unités de calculs simples (unités arithmétiques et logiques, appelées cores) qui sont capables d'effectuer les opérations logiques et arithmétiques classiques. Un GPU peut contenir plusieurs centaines de cores, alors que les processeurs classiques contiennent actuellement moins de 10 cores, cependant l'architecture et les capacités sont différentes. Chaque cored'un processeur classique est capable d'effectuer son propre flot d'instructions, sur ses propres données, indépendamment des autres : il s'agit du modèle MultiThread (MT). Les coresdes GPU doivent par contre effectuer chacun un flot d'instructions identiques, sur des données différentes : il s'agit du modèle Single Instruction Multiple Thread (SIMT), un même programme est instancié nfois sur n jeux de données différents, en parallèle.

Différents modèles de programmation entre un CPU et un GPU : MT vs SIMT (schématique)

Ce modèle de programmation n'est pas totalement nouveau, les processeurs possédant depuis la fin des années 90 des extensions vectorielles (MMX, SSE, Altivec), instructions additionnelles permettant d'effectuer des opérations en parallèle sur des vecteurs de données (par exemple à 4 ou 8 éléments). Ces instructions vectorielles sont déjà utilisées pour accélérer, comme pour le GPGPU, des calculs multimédias (musique, vidéo) ou généraux.

Le déroulement type de l’exécution d'un programme sur carte graphique est le suivant :

- transfert des données depuis la mémoire du PC hôte vers la mémoire du GPU (via le bus PCI-Express) ;

- réalisation des calculs sur le GPU ;

- transfert des résultats du GPU vers le PC hôte et éventuellement calculs additionnels sur le CPU.

Suivant le type de calcul effectué, la bande passante du bus PCI-E et les temps de latence entre les différents niveaux de mémoire dans le GPU peuvent être limitants. Les algorithmes permettant d'exploiter au maximum les capacités des GPU sont ceux réalisant un nombre d'opérations simultanées très grand par rapport aux opérations mémoires (chargement/déchargement de données).

1.2. Quels langages et frameworks pour programmer les GPU ?

Le premier framework qui a permis le GPGPU est Cudadu fabricant Nvidia. Ce framework définit des extensions sur langage C afin d’adresser les spécificités et le parallélisme du GPGPU.

Le programme exemple suivant exécute une simple addition des composantes des vecteurs A et B. La fonction d'addition de deux composants est définie dans un kernel. Ce kernel est ensuite exécuté n fois (dans n threads), l'addition des n composantes des vecteurs est effectuée en parallèle :

// Kernel definition

__global__ void VecAdd(float* A, float* B, float* C)

{

int i = threadIdx.x;

C[i] = A[i] + B[i];

}

int main()

{

...

// Kernel invocation with N threads

VecAdd<<<1, N>>>(A, B, C);

}

Ce langage est spécifique aux GPU de ce fabricant, ce qui permet d'optimiser au maximum pour une architecture donnée, mais limite la portabilité. Les détails sur l'architecture et le framework se trouvent dans le très riche Cuda programming guide[3], et un dossier sur la programmation GPU avec Cuda est paru dans Linux Magazine [4].

Le framework OpenCL semble aujourd'hui le choix à faire pour profiter d'une grande portabilité et de l'indépendance vis-à-vis du matériel. Un programme OpenCL peut s’exécuter à la fois sur les cartes graphiques des deux principaux constructeurs Nvidia et AMD, mais aussi sur les processeurs classiques (Intel, AMD, avec utilisation des extensions vectorielles) et sur des plateformes exotiques ou embarquées : processeur Cell de la PS3, bientôt les processeurs ARM ? Le troisième framework possible est DirectCompute de Microsoft, intégré à l'API DirectX. Le lecteur intéressé par ces différents frameworks pourra consulter cet autre article de Linux Magazine [5].

Le choix du positionnement du parallélisme

Pour exploiter au mieux le GPGPU, il est nécessaire de travailler sur des calculs massivement parallèles. Ce parallélisme peut se situer à plusieurs niveaux dans le programme, et ce choix sera structurant sur les performances et le fonctionnement de la solution.

- Parallélisme interne à l'algorithme : dans cette configuration, l'algorithme de bas niveau utilisé pour traiter le problème est lui-même parallèle. Par exemple, la multiplication de matrices de grande taille se décompose en un ensemble de calculs de composantes de la matrice résultat, indépendants les uns des autres.

- Parallélisme externe : dans cette configuration, le parallélisme est obtenu par la subdivision du problème en différentes tâches indépendantes, ou par la résolution de plusieurs problèmes simultanément. Par exemple, pour réaliser le chiffrement de données, il est possible d'utiliser le mode de chiffrement CTR, qui permet de chiffrer simultanément les blocs de données d'une même transmission, ou de choisir de chiffrer plusieurs communications en parallèle.

2. Applications du GPGPU en sécurité informatique

Un premier état des lieux de l'utilisation du GPU en SSI avait fait l'objet d'une présentation au SSTIC 2009 par A. Joux [6], les paragraphes suivants reprennent des thèmes similaires en mentionnant les derniers développements.

2.1. Autour de la cryptographie

Une des premières applications en sécurité informatique a été le cassage des hashs de mots de passe par force brute. Ce problème s'adapte simplement au modèle parallèle du GPGPU : chaque thread va hacher un mot de passe candidat et le comparer au(x) hash(s) ciblé(s). Les threads sont donc totalement indépendants, aucune communication entre eux n'est nécessaire, et la quantité de mémoire de travail par thread est faible. Le problème étant trivialement parallèle, les résultats sur GPU sont proches des accélérations maximales théoriquement possibles.

Plusieurs implémentations sont disponibles librement, comme ighashgpu [7] et oclHashcat [8]. Pour SHA-1, OclHshcat est capable de tester 400 000 000 mdp/s sur un GPU haut de gamme, et pour les mots de passe Windows NTLM, ighashgpu peut parcourir les mots de passede 8 caractères (alphanumériques etspéciaux) en une vingtaine de jours sur une carte graphique modeste [9].

Dans le domaine de la sécurité du Wi-Fi, la recherche exhaustive de preshared key WPA/WPA2 convient très bien au GPGPU, comme démontré par le logiciel Pyrit. Cependant, lors de l'utilisation d'une PSK dérivée d'un mot de passe, WPA la renforce avec la fonction PBKDF2 (avec 4096 itérations d'HMAC-SHA1) et les performances en sont diminuées d'autant : 10 000 mdp/s, cela laisse les mots de passe complexes de plus de 8 caractères hors d'atteinte [11] (du moins sur une seule carte).

Les GPU ne serviraient-ils donc qu'à bruteforcerdes hashs de mots de passe ? Non, des projets plus complexes ont aussi été réalisés, comme les rainbows tables probabilistes [12] (cf. MISC n°58), ou encore la recherche de collisions pour la fonction de hachage SHA-1 [13] (implémentation qui a permis de calculer des collisions pour SHA-1 réduit à 75 tours, nombre de tours proche des 80 de la fonction nominale).

Il existe aussi des approches élaborées sur l'utilisation de la cryptographie sur GPGPU plutôt que l’accélération des attaques. La voie a été ouverte par différentes implémentations de l'algorithme de chiffrement AES sur GPU [14] [15], apportant des performances supérieures aux CPU. Cependant, de telles performances s'obtiennent en chiffrant plusieurs messages différents en parallèle avec le mode CBC (le chiffrement en mode CBC n’étant pas parallélisable pour un seul message, contrairement au déchiffrement), ou en utilisant le mode CTR qui permet générer plusieurs blocs de keystream simultanément, pour chiffrer un même message.

Concernant l'algorithme RSA, son implémentation est aussi possible sur GPU, mais nécessite des adaptations fines [16]. De façon similaire au chiffrement symétrique, les meilleures performances sont obtenues lorsque des centaines d'opérations RSA sont effectuées en parallèle.

Armé des trois fonctions cryptographiques AES, SHA-1 et RSA adaptées aux GPU, il est possible de réaliser un accélérateur SSL complet [17], qui permet de décharger le déchiffrement SSL d'un serveur web sur le GPU, avec de meilleures performances (principalement un meilleur débit) sous une charge importante (grand nombre de sessions simultanées). Le système décrit dans ce papier (reverse proxy en coupure SSL, avec accélération opportuniste sur GPU) est prometteur par ses applications pratiques immédiates.

2.2. Recherche de signatures et forensic

La recherche de signatures (chaînes de caractères dans un flux de données, expressions régulières) a été implémentée sur GPU, comme l'adaptation de l'IDS Snort [18], qui apporte un léger gain de performance.

La recherche de données (en-tête d'un format de fichier, ou base de fichiers recherchés) par fonction de checksumsimple (CRC32 ou similaire) est utilisée en informatique légale, et peut être accélérée par GPU [19]. Les performances obtenues dépassent largement les débits des disques et permettent de faire face aux capacités croissantes de ceux-ci.

À côté de ces adaptations réussies, le monde de l'antivirus semble encore échapper à la mode du GPGPU, probablement à cause de la complexité des algorithmes utilisés (heuristiques, émulation), par rapport à la recherche simple de signature ou d'expressions régulières. Quelques essais [20] [21] ont été effectués, mais peu d'applications concrètes (antivirus complet) ou de publications récentes existent sur le sujet.

3. Exemple d’application : accélération d’une fonction de hachage en mode « arbre »

Ce paragraphe présente l'implémentation de la fonction de hachage cryptographique Keccak [22] sur GPU. Keccak est une fonction de hachage proposée pour la compétition SHA-3 [23] qui vise à établir un nouveau standard de primitive de hachage. Il a pour objectif le remplacement de la famille de fonctions SHA-2 (SHA256, SHA512) étant données les avancées en cryptanalyse dans ce domaine (par exemple sur SHA-1). Le NIST, organisateur, a défini un cahier des charges et une interface que les algorithmes candidats doivent respecter.

L'implémentation GPU de Keccak [24] présentée ici repose sur le framework Cuda, qui était le plus mature à la date de réalisation du projet en 2010. Elle a été réalisée pour répondre à un concours organisé par l’équipe responsable de la proposition Keccak afin d'évaluer la portabilité et les performances de l'algorithme sur des architectures exotiques (hors x86 et x86-64).

3.1. Choix du mode d'opération

Le choix d'un parallélisme externe à la fonction a été fait pour ce projet : le mode d'opération en arbre de hachage (ou arbre de Merkle), qui permet de paralléliser le hachage d'un seul message. Dans ce mode, défini récursivement, le hash d'un nœud de l'arbre est égal au hash de l'assemblage des hashs des éléments fils du nœud. Les données à traiter sont réparties dans les feuilles de l'arbre, et le résultat final est le hash du nœud racine.

Un arbre de hachage a plusieurs avantages : les calculs de chaque feuille et des nœuds intermédiaires sont parallélisables. De plus, si la structure de l'arbre et les résultats intermédiaires des nœuds sont conservés, il est possible de réaliser du hash incrémental : si une partie du fichier haché est modifiée, il suffit de recalculer la branche de l'arbre de cette partie, et non l'arbre en entier. Les arbres de hachage sont utilisés sur les réseaux P2P, où les fichiers sont divisés en morceaux et pour lesquels l’intégrité doit pouvoir être vérifiée sans attendre le téléchargement complet du fichier.

Le déroulement de l'algorithme implémenté sur GPU est le suivant :

- Le message à hacher est découpé pour former les feuilles de l'arbre, les données sont transférées dans la mémoire du GPU.

- Chaque feuille de l'arbre est hachée par la fonction Keccak (par un thread sur le GPU).

- Les résultats du hash des feuilles sont rassemblés et sont hachés par un nœud supérieur (ou par le nœud racine qui produit le résultat final).

Pour simplifier l'implémentation, un arbre à deux niveaux a été choisi : les nœuds de niveau 1 sont calculés sur le GPU, et l'unique nœud (racine) de niveau 2 est calculé sur le CPU. La taille des données de chaque feuille a été fixée : le nombre total de feuilles varie donc en fonction de la taille du message à traiter. Le programme itère les opérations pour chaque nouveau groupe de feuilles, tant qu'il reste des données à hacher.

Arbre de hachage implémenté sur le GPU, avec un nœud racine calculé par le CPU

Le surplus de calcul (par rapport à un hachage séquentiel) dû aux calculs des nœuds supérieurs fait que le mode de hachage par arbre n'est performant que si la taille des données des feuilles est importante. Ce mode est donc adapté pour le hachage de fichiers de taille importante (plusieurs Go).

3.2. Portage de la fonction Keccak sur GPU et optimisation des performances

3.2.1. Choix du type d’implémentation

La fonction de hachage Keccak est entièrement définie par des opérations logiques (principalement XOR, AND et des rotations) sur des mots de 64 bits, et vise ainsi des performances optimales sur les processeurs 64 bits. Cependant, les GPU sont aujourd'hui des architectures avec des opérations scalaires sur 32 bits, donc une implémentation 32 bits a été utilisée pour ce portage.

Écrire du code efficace pour GPU nécessite de connaître certaines spécificités du matériel. Les cores des cartes graphiques Nvidia sont regroupés en Streaming Multiprocessors (SM), chaque SM possède N cores (quantité variable selon les modèles, de 8 à 48) et une mémoire partagée rapide (shared memory). Les threads sont répartis en blocs, et chaque bloc s’exécute sur un SM. Les threads d'un même bloc peuvent utiliser la mémoire partagée pour communiquer ou comme espace de travail, en plus des registres qui sont privés pour chaque thread.

Architecture simplifiée d'un GPU Nvidia dans le contexte Cuda

Le Cuda programming guide recommande d'utiliser un nombre important de threads par bloc afin de masquer les latences des opérations de calculs et de mémoire (lorsqu'un thread attend le résultat d'une opération, il est remplacé par un autre thread en attente d’exécution). Cependant, un nombre important de threads entraîne une grande quantité de mémoire utilisée par l'ensemble des threads d'un bloc, et si cette mémoire ne peut loger dans la mémoire partagée, des appels à la mémoire globale très lente vont pénaliser les performances.

Chaque thread qui exécute une instance de Keccak doit donc occuper une mémoire de travail la plus faible possible pour stocker les données nécessaires aux calculs du hash : l'implémentation en langage C, qui al'empreinte mémoire la plus compacte, a donc été choisie et directement utilisée avec Cuda.

3.2.2. Superposition des calculs GPU/CPU et des transferts mémoire

L’implémentation la plus simple réalisée effectue les tâches suivantes dans l'ordre :

pour chaque groupe de feuilles :

- transfert des données à traiter au GPU ;

- calculs sur le GPU ;

- transfert des résultats vers le CPU ;

- calculs sur le CPU.

Cette exécution séquentielle des transferts mémoire et des calculs CPU et GPU peut être optimisée. Premièrement, l’exécution des commandes Cuda (transfert mémoire et calculs GPU) peut être rendue asynchrone et réalisée en même temps qu'un calcul sur CPU. Dans le cas général de groupes de calculs s'effectuant d'abord sur GPU puis finalisés sur CPU, il est possible de lancer les transferts et le calcul du groupe suivant sur GPU pendant que le CPU travaille sur le résultat du groupe précédent.

Une seconde optimisation possible est la superposition des transferts mémoire et des calculs GPU. Cuda définit le concept de streams, séquence d'opérations dont le déroulement chronologique est assuré (par contre, la chronologie des opérations entres les différents streams n'est pas définie). Pour le calcul du hash d'une feuille, il est possible de commencer le calcul même si la totalité des données ne sont pas encore transférées sur le GPU. En découpant les données en entrée des fonctions de hachage, il est donc possible de superposer les calculs GPU et les transferts mémoire entre l'hôte et le GPU.

Optimisations successives par superposition des transferts mémoire et des calculs CPU et GPU

3.2.3. Performances obtenues

Les mesures de performances ont été réalisées sur un CPU Core i5-750 (en mode monothread) et sur une carte Nvidia GTS 250 (128 Cuda cores, ancienne génération). En termes de performance pure, les résultats (surtout sur CPU) sont loin des records car l’implémentation n’est pas optimisée pour la vitesse, mais choisie pour une utilisation faible de la mémoire. L'objectif est de comparer les performances du CPU et du GPU, ainsi que de mesurer les apports des différentes optimisations.

Plateforme	Core-i5-750 2.6Ghz Nvidia GTS 250
Optimisation
CPU seul	20 Mo/s
CPU + GPU simple	700 Mo/s
CPU + GPU calculs superposés	1000 Mo/s
CPU + GPU transferts et calculs superposés	1200 Mo/s

Sur une carte récente, les performances peuvent être multipliées par 4 ou plus, et se rapprocher des vitesses maximum du bus PCI-E 2.0.

3.3. Un mode d'opération pour le chiffrement par flot (streamcipher)

Plusieurs algorithmes candidats à SHA-3, dont Keccak, proposent un mode d'utilisation de la fonction de hachage en chiffrement par flot, similaire à RC4.L'idée générale est d'appliquer la fonction de hachage sur des données d'entrée qui contiennent la clé de chiffrement, un IV (Initialisating Vector, aléatoire pour chaque communication), et un index. Le résultat du hachage constitue la clé de chiffrement par flot, ou keystream, qui est combiné par XOR aux données à chiffrer.

Cette configuration de Keccak tire aisément profit du GPGPU : chaque thread sur le GPU calcule une partie du keystream correspondant à un index donné, et les différentes partie du keystream sont calculées en parallèle. De plus, keccak possède un mode d'opération où la taille du résultat est arbitrairement longue, appelé squeeze mode : chaque thread peut ainsi calculer un keystream différent de taille arbitraire, en fonction de la quantité globale de données à chiffrer.

Keccak en mode chiffrement par flot (streamcipher)

Les performances obtenues sont similaires au mode de hachage, plus d'1Go/s ont été obtenus sur la modeste GTS 250. Ce mode est adapté au chiffrement de données à la volée, et pourrait être utilisé pour la couche confidentialité d'un VPN à haut débit.

Avec quelques adaptations, il est possible d'utiliser ce mode comme générateur de nombres pseudo aléatoires (PRNG), comme spécifié dans la publication du NIST SP 800-90 [25]. Le GPU permet alors de dériver des bits ou nombres pseudo-aléatoires, avec un débit important, à partir d'une graine (remplaçant la clé de chiffrement) qu'il faut choisir provenant d'une bonne source d'entropie (générateur matériel, entropie du système d'exploitation). Des applications possibles sont, par exemple, les simulations par la méthode de Monte-Carlo.

4. Conclusion et perspectives

4.1. Hype, Buzz & Cloud

De manière générale, un effet de mode sur le GPGPU a eu lieu, poussé par les fabricants et probablement des chercheurs en manque de sponsors. De nombreuses applications ont été portées sur GPU, et celles permettant les meilleures accélérations ont été mises en avant : 10 fois, 50x, 100x plus rapide qu'un CPU !? Peut-être, mais seulement sur certaines applications et au prix d'optimisations spécifiques et non portables !Des contraintes spécifiques (non détaillées dans cet article), comme la limitation du nombre de branchements (if, then, else) dans le programme, la gestion délicate de la hiérarchie mémoire, le débit limitant du bus PCI-E, font barrage à un portage simple des programmes existant sur GPU.

L'intégration du GPGPU profite néanmoins d'avancées très rapides : concernant les progrès matériels, l'intégration du GPU proche du CPU sur les APU (Accelerated Processing Unit) d'AMD va grandement améliorer les transferts mémoire entre CPU et GPU, les deux composants étant concentrés sur la même puce. De même, le bus PCI-Express 3.0 augmentera grandement la bande passante entre le GPU et le système hôte. Le support des frameworks s'améliore, comme OpenCL intégré aux dernières distributions Linux, et l'utilisation possible d'autres langages que le C (via des interfaces) tels que C++, Java, Python.

Dans le domaine de la sécurité et des attaques, des titres comme « A German hacker used cloud computing to crack passwords stored in an algorithm that was developed by the NSA » ou encore « How to Crack Passwords in the Cloud with Amazon's Cluster GPU Instances »ont fait un buzz énorme alors qu'il s'agissait seulement d'un benchmark de cassage de mots de passe de 6 caractères (!), hachés par SHA-1, sur l'offre Cloud GPU d'Amazon [26], avec un logiciel pré-existant [27][28].

Bien sûr, le GPU (et le Cloud) permet d'obtenir une puissance de calcul beaucoup plus importante, mais les gains restent « linéaires » : le cloud GPU permet un gain d'un facteur 100000 en rapidité de cassage ? Ce gain est effacé en ajoutant 3 caractères à votre mot de passe, ou par du renforcement de clé, comme PBKDF2 avec 100000 itérations.

Un hash mot de passe Windows est attaquable avec des rainbowtables probabilistes ? Il serait temps que le stockage local des hashs des mots de passe sous Windows soit mieux sécurisé, avec l'utilisation d'un diversifiant pour chaque mot de passe et d'une fonction de type PBKDF2 avec un grand nombre d'itérations ; cela permettrait d’éviter les titres racoleurs comme « Cheap GPUs are rendering strong passwords useless » [29].

4.2. GPGPU, gain ou menace pour la sécurité SI ?

Le GPGPU a donc placé la barre plus haut, par exemple les standards de robustesse de mots de passe sont montés d'un cran (8 → 10 caractères, ou plus...), mais c'est faire porter à l'utilisateur les conséquences de défauts de conception : les implémentations et les protocoles doivent réduire au maximum les possibilités et l’efficacité des attaques, qu'elles soient en ligne (limitation des tentatives d'authentification, pas de divulgation d'information sur les secrets) mais aussi hors ligne. Même sile GPU est un concentré de puissance de calcul, il ne permet pas de remettre en cause la sécurité des algorithmes existants, lorsqu'ils sont correctement utilisés : RSA ne sera pas cassé dans un futur proche sur carte graphique, ni AES !

Concernant les applications du GPU en tant qu’accélérateur cryptographique, que ce soient le reverse proxy de déchiffrement SSL ou l’accélération de chiffrement ou de hachage, ces projets restent encore des Proof of Concept, et leur utilité et pertinence « en production » doivent être démontrées. Les performances n'étant au rendez-vous que sur des calculs massivement parallèles, des modes d'opération spéciaux doivent être utilisés, et ces modes ne sont pas encore intégrés aux standards. Peut-être qu'un jour le chiffrement intégral du disque dur sera accéléré par GPU de façon invisible à l'utilisateur ? Mais ce n'est pas encore pour demain, et les applications qui apporteront des gains significatifs seront probablement du côté des serveurs et des infrastructures que de l'utilisateur final. À quand Google ou Facebook en HTTPs intégral, accéléré par GPU ?