Introduction
L’utilisation des GPGPU (General Purpose Graphical Processing Unit) n’est pas un sujet nouveau dans le monde de la sécurité. De nombreux projets de récupération de mot-de-passe utilisant d’hors et déjà cette technologie dans le but d’améliorer leurs performances (pyrit, CUDA multihash bruteforcer, …). Lors de la conférence RUXCON 2008, Daniel Reynaud a fait une présentation sur l’utilisation des GPGPU par des codes malveillants. Ces derniers délocaliseraient l’exécution de certains algorithmes (par exemple celui générant les noms de domaines permettant de se connecter à un canal de contrôle) sur la carte graphique afin d’empêcher une analyse standard de code ia32.
Deux ans après cette présentation de tels codes malveillants se font rares notament du fait que les auteurs cherchent à infecter le plus de machines possibles, ce qui est plus compliqué quand le code est compilé pour un type de chipset graphique précis. Néanmoins, si l’on regarde les caractéristiques de la gamme Apple, plateforme devenant une cible grandissante pour les auteurs de malwares, on s’aperçoit que la majorité des machines sont équipées d’une carte graphique NVIDIA et que le framework OpenCL est inclu dans l’installation de base. On pourrait imaginer que ces deux propriétés favoriseront l’utilisation des GPGPU lors d’attaques ciblant un système OSX.
La possible utilisation des GPGPU par des codes malveillants ou des solutions de protection logicielle sont des motifs poussant à se pencher sur l’analyse du code compilé pour ces plateformes (et également pour la beauté de l’Art comme diraient certains).
Les bases
Différents frameworks existent pour le développement de code pour GPGPU, comme par exemple CUDA (NVIDIA) et OpenCL (sur-couche permettant d’obtenir du code compilable sur différentes architectures). Le développement se fait dans une extension du C et le passage d’arguments est réalisé via des appels à des fonctions prédéfinies depuis le code exécuté sur le CPU (ex: CUDA!cudaMemcpy, openCL!clEnqueueReadBuffer, …). Dans le cas d’un code utilisant le GPGPU pour décoder des informations, la récupération du texte-clair lors d’une analyse serait simplifiée via la mise en place de breakpoints sur ces API. Ci-dessous, un exemple de kernel CUDA réalisant un parcours de chaine de caractères.
[sourcecode language=”cpp”]
// not so usefull function
__global__ void mystrlen(char* str, int* length)
{
int i = 0;
while (str[i])
++i;
*length = i;
}
[/sourcecode]
L’appel à cette fonction peut se faire de la manière suivante depuis le code C:
[sourcecode language=”cpp”]
cudaMalloc((void**)&devLen, sizeof (int));
cudaMalloc((void**)&devStr, MAX_SIZE * sizeof (char));
cudaMemcpy(devStr, argv[1], strlen(argv[1]) + 1, cudaMemcpyHostToDevice);
mystrlen <<< 1, 1 >>> (devStr, devLen);
cudaMemcpy(&len, devLen, sizeof (int), cudaMemcpyDeviceToHost);
[/sourcecode]
Dans le cas de CUDA, le code est compilé dans une forme intermédiaire, PTX, qui sera ensuite inclu dans le binaire. En spécifiant une architecture donnée, le PTX est compilé en CUBIN et optimisé pour cette dernière.
Extraction du code PTX
Dans le cas où l’application a été compilée sans spécifier d’architecture, le code PTX sera directement présent dans le binaire et compilé pour l’architecture virtuelle sm_10 par défaut. Une recherche sur des mots clefs comme .target permet de le localiser. Travis Goodspeed a développé un outil réalisant l’extraction du code PTX depuis ce type de binaire.
[sourcecode language=”bash”]
$ nvcc mystrlen.cu
$ cudaDump < a.out
.version 1.4
.target sm_10, map_f64_to_f32
// compiled with /usr/local/cuda/bin//../open64/lib//be
// nvopencc 3.2 built on 2010-11-11
…
.entry _Z8mystrlenPcPi (
.param .u32 __cudaparm__Z8mystrlenPcPi_str,
.param .u32 __cudaparm__Z8mystrlenPcPi_length)
{
.reg .u32 %r<10>;
.reg .pred %p<4>;
.loc 28 16 0
$LDWbegin__Z8mystrlenPcPi:
ld.param.u32 %r1, [__cudaparm__Z8mystrlenPcPi_str];
ld.global.s8 %r2, [%r1+0];
mov.u32 %r3, 0;
setp.eq.s32 %p1, %r2, %r3;
@%p1 bra $Lt_0_2306;
ld.param.u32 %r1, [__cudaparm__Z8mystrlenPcPi_str];
mov.s32 %r4, %r1;
mov.s32 %r5, 0;
$Lt_0_1794:
//<loop> Loop body line 16, nesting depth: 1, estimated iterations: unknown
.loc 28 21 0
add.s32 %r5, %r5, 1;
add.u32 %r4, %r4, 1;
ld.global.s8 %r6, [%r4+0];
mov.u32 %r7, 0;
setp.ne.s32 %p2, %r6, %r7;
@%p2 bra $Lt_0_1794;
bra.uni $Lt_0_1282;
$Lt_0_2306:
mov.s32 %r5, 0;
$Lt_0_1282:
.loc 28 22 0
ld.param.u32 %r8, [__cudaparm__Z8mystrlenPcPi_length];
st.global.s32 [%r8+0], %r5;
.loc 28 23 0 exit;
$LDWend__Z8mystrlenPcPi:
} // _Z8mystrlenPcPi
[/sourcecode]
Dans le cas d’un binaire compilé pour une architecture spécifique, comme dans l’exemple ci-dessous, l’outil decuda permet à partir du CUBIN d’obtenir le code PTX, excepté pour du code compilé à partir de CUDA 3.x du à un changement de format (CUBIN > ELF). La solution consiste à utiliser le script elfToCubin.py pour réaliser la conversion.
[sourcecode language=”bash”]
$ nvcc mystrlen.cu -arch=compute_10 -code=sm_10,sm_13
[/sourcecode]
Le binaire compilé à l’aide de la ligne de commande précédente encapsule les 2 binaires ELF correspondants à sm_10 et sm_13 dans le segment nommé __const.
Analyse du code
Une des premières choses que l’on peut remarquer en lisant le code PTX ci-dessus est le fait qu’une partie des symboles a été conservée. On retrouve mystrlen et ses deux arguments str et length.
Les instructions se décomposent de la manière suivante: opération.<espace>.type. La première instruction de la fonction mystrlen lit (ld) la valeur de type entier non-signé (u32) contenue dans le paramètre (param) str et la stocke dans le registre r1.
A la ligne 19, l’opération de comparaison setp est utilisée afin de vérifier si les registres r2 et r3 sont égaux (eq) et stocke le résultat dans le registre de prédiction p1. La ligne suivante correspond à une opération de modification du flux d’exécution. Ici, l’opération à réalisée (bra) est précédée d’un “predicate guard” (@[!]<predicate register>). Dans ce cas, le saut au label $Lt_0_2386 sera exécuté si le registre p1 contient la valeur vrai (r2 == r3). On retrouve le corps de la boucle while entre les lignes 27 et 33 avec les registres r4 et r5 qui représentent respectivement un pointeur sur str et i.
La documentation fournie par NVIDIA permet d’en apprendre plus sur ce langage, mais le jeu d’instructions reste limité et assez proche de langages assembleurs courants pour qu’il soit facilement compréhensible.
En plus d’une analyse statique, il est possible de réaliser une analyse dynamique de kernel CUDA à l’aide de gdb-cuda.