-->

Tutoriel n°1 : PyOpencl: premier calcul sur GPU


By Stéphane Planquart - Posted on 03 janvier 2010

Pour ce premier code OpenCL sous Python, nous allons simplement calculer la somme de 2 tableaux. Bien évidement 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 structure numpy.ndarray pour échanger les données entre Python et l'API OpenCL.

  import pyopencl as cl
  import numpy

1 - Context, Kernel, CommandQueue et Program

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

  • Device : CPU ou GPU qui peuvent être utiliser avec l'api OpenCL
  • Context : Ensemble d'une ou plusieurs devices. C'est au niveau du context que sont alloué 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 à executer. Une commande peut être par exemple: Lire ou ecrire un Buffer, exécuter une fonction

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

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

2 - 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 creer donc les tableaux numpy dont on as besoin. Ici 2 tableau 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)

3 - 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 principal 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és. 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()

4 - 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 argument 1 par 1. 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 resultat, stocké dans C_buf, du context vers le host.

  cl.enqueue_read_buffer(queue, C_buf, C).wait()
Le .wait() est ici utiliser pour ne rendre la main qu'une fois terminer.

Conclusion

Ce premier tutoriel est maintenant terminer. Nous pouvons déjà constater que l'utilisation d'OpenCL avec Python est grandement simplifier 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()