Cool Stack

Tutoriel n°1 : PyOpencl: premier calcul sur GPU

Pour ce premier code OpenCL sous Python, nous allons simplement calculer la somme de deux tableaux. Bien évidemment il faut tout d'abord installer un driver OpenCL et la librairie de développement ainsi que PyOpenCL.

Tout d'abord, nous avons besoin d'importer pyopencl et numpy. Pourquoi numpy ? Parce que pyopencl utilise les structures numpy.ndarray pour échanger les données entre Python et l'API OpenCL.

import pyopencl as cl
import numpy

Context, Kernel, CommandQueue et Program

Dans OpenCL, comme dans beaucoup d'API, il faut d'abord connaitre un minimum de vocabulaire avant de commencer. Pour OpenCL:

  • Device : CPU ou GPU qui peuvent être utilisés avec l'api OpenCL
  • Context : Ensemble d'un ou plusieurs devices. C'est au niveau du context que sont alloués la mémoire global (__global dans le C d'OpenCL)
  • Platform : Contient les informations concernant la machine et la version d'implémentation d'OpenCL
  • Program : Contient le code source (compiler ou non) des fonctions à exécuter en OpenCL.
  • Kernel : est un sous-ensemble d'un programme. Pour simplifier il s'agit d'une fonction principale.
  • CommandQueue :file d'attente des commandes à exécuter. Une commande peut être par exemple : Lire ou écrire un Buffer, exécuter une fonction

Dans PyOpenCL il suffit de créer un context, il n'est pas nécessaire de se préoccuper du reste. Ensuite il faut créer une queue de commande. Dans OpenCL, tout ce qui est à faire, doit être mit dans la queue, c'est ensuite l'API qui s'occupe de synchroniser les opérations pour que tout ce déroule correctement.

ctx = cl.Context()
queue = cl.CommandQueue(ctx)

Allocation de la mémoire

Et oui en utilisant PyOpenCL il va falloir s'y faire, l'allocation de mémoire est à faire à la main. On utilise créer donc les tableaux numpy dont on a besoin. Ici deux tableaux de même taille contenant des valeurs aléatoire. Et un tableau pour contenir le résultat. Il est important de bien définir aussi le type de données de ce tableau. Ici on choisit numpy.float32 qui correspond au float d'OpenCL

A = numpy.random.rand(taille).astype(numpy.float32)
B = numpy.random.rand(taille).astype(numpy.float32)
C = numpy.empty_like(A)

Cette allocation avec numpy nous permet de créer un tableau au niveau du host. On doit maintenant déclarer ce même tableau au niveau de la mémoire __global d'OpenCL. Pour faire cela il faut créer des Buffer dans le context.

A_buf = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=A)
B_buf = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=B)
C_buf = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, A.nbytes)

Le Programme en OpenCL

Un Program OpenCL est s'écrit en langage OpenCL qui est une extension du langage C. Voici pour commencer un exemple:

__kernel
void sum(__global const float* a, __global const float* b, __global float* c){
  int i = get_global_id(0);
  c[i] = a[i] + b[i];
}

Le mot clé __kernel est là pour dire qu'il s'agit d'une fonction principalle du programme. C'est une fonction que l'on peut "appeler" à partir de notre programme en python et qui s'exécute sur la carte graphique par exemple. Le mot clé __global est là pour dire qu'il s'agit d'un paramètre stocké dans la mémoire global, la mémoire du context. La fonction get_global_id(0) nous permet ici de connaitre l'index de tableau à traiter. Pour le reste c'est du C classique, vous devriez être capable de la comprendre sans difficulté. Enfin en python, pour déclarer et compiler ce program on fait:

prg = cl.Program(ctx, """
 __kernel void sum(__global const float* a, __global const float* b, __global float* c){
   int i = get_global_id(0);
   c[i] = a[i] + b[i];
 }
 """).build()

Lancer une fonction OpenCL (Kernel)

L'implémentation de pyopencl nous permet encore une fois de simplifier l'utilisation d'opencl. Il n'est plus utile de créer un objet Kernel pour l'utiliser, ni de passer les arguments un par un. Ici pour lancer sum(A,B,C) sur l'ensemble des éléments, on écrit simplement

prg.sum(queue, A.shape, A_buf, B_buf, C_buf)

Ensuite il faut transférer le résultat, stocké dans C_buf, du context vers le host.

cl.enqueue_read_buffer(queue, C_buf, C).wait()

Le .wait() est ici utilisé pour ne rendre la main qu'une fois terminé.

Conclusion

Ce premier tutoriel est maintenant terminé. Nous pouvons déjà constater que l'utilisation d'OpenCL avec Python est grandement simplifié comparer à l'utilisation direct de l'api opencl en langage C. Voici le programme python complet :

import pyopencl as cl
import numpy

A = numpy.random.rand(1000).astype(numpy.float32)
B = numpy.random.rand(1000).astype(numpy.float32)
C = numpy.empty_like(A)

ctx = cl.Context()
queue = cl.CommandQueue(ctx)

A_buf = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=A)
B_buf = cl.Buffer(ctx, cl.mem_flags.READ_ONLY | cl.mem_flags.COPY_HOST_PTR, hostbuf=B)
C_buf = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, A.nbytes)
prg = cl.Program(ctx, """
  __kernel
  void sum(__global const float* a, __global const float* b, __global float* c){
    int i = get_global_id(0);
    c[i] = a[i] + b[i];
  }
  """).build()
prg.sum(queue, A.shape, A_buf, B_buf, C_buf)
cl.enqueue_read_buffer(queue, C_buf, C).wait()

Posted Dim 03 janvier 2010 by Stéphane Planquart in programmation graphique