Utiliser OpenCL avec Qt

Le site Qt Labs permet aux développeurs de Qt de présenter les projets, plus ou moins avancés, sur lesquels ils travaillent.

Nokia, Qt, Qt Quarterly et leurs logos sont des marques déposées de Nokia Corporation en Finlande et/ou dans les autres pays. Les autres marques déposées sont détenues par leurs propriétaires respectifs.

Cet article est la traduction de l'article Using OpenCL with Qt de Rhys Weatherley paru dans Qt Labs.

Récemment, nous avons mené des expériences avec OpenCL et Qt, pour voir comment Qt pourrait utiliser OpenCl en interne et en faciliter l'usage. Dans ce post, nous allons présenter une introduction à OpenCL, la bibliothèque QtOpenCL, montrer comment écrire votre premier programme QtOpenCL et vous indiquez comment obtenir de plus amples renseignements sur le projet.

Pour ceux qui ne connaissent cette bibliothèque, OpenCL est un standard ouvert et gratuit pour la programmation parallèle dans un environnement informatique hétérogène. L'utilisation la plus commune dont vous avez probablement entendu parler est de pouvoir exécuter n'importe quel code C sur les processeurs graphiques (GPU) de votre système. Les GPU actuels sont plus puissants que les processeurs centraux (CPU), étant conçus pour traiter des centaines de milliers de triangles texturés par seconde dans votre jeu vidéo préféré. Pour faire cela, le GPU dispose d'un accès au traitement vectoriel parallèle qui dépasse de loin la capacité des instructions x86/SSE ou ARM/NEON de votre CPU.

Pendant des années, les langages de programmation de shaders tels que GLSL ont utilisé les capacités vectorielles des GPU pour créer des effets de rendu personnalisés dans OpenGL, mais on est classiquement limité aux paramètres représentant un triangle visible. Il est également typique pour les implémentations d'OpenGL de simplifier les calculs en utilisant les calculs à virgule fixe et à faible précision. Le code source d'un shader peut utiliser un float, mais il pourrait tout aussi bien être codé sur 8 bits en vue de la précision réelle. Autant vous ne verrez pas les erreurs d'arrondis sur des grands traitements de triangles, autant cela n'est pas utilisable pour les supercalculateurs, les algorithmes mathématiques habituels, et le mélange de pixels ultraprécis.

Entrons maintenant dans OpenCL. Il définit un nouveau langage proche du C qui augmente la précision mathématique et qui permet de passer n'importe quel argument à une fonction OpenCL – appelée kernel – offrant toutes les possibilités du C. Des types de vecteurs spéciaux, tels que float4, sont fournis ainsi qu'une vaste bibliothèque mathématique. Mais sa caractéristique la plus intéressante est la découpe du travail : il est très facile de découper votre traitement en petits morceaux que le GPU pourra distribuer sur l'ensemble de ses unités de calcul (les unités de calcul englobent tous les processeurs, CPU et GPU que OpenCL trouve). Contrairement au code C habituel où vous pouvez passer beaucoup de temps à écrire des boucles externes et à lancer des threads de travail pour les sous-parties de votre problème, OpenCL le fait pour vous. Nous allons voir comment cela fonctionne sous peu.

La bibliothèque QtOpenCL encapsule l'interface (API) d'OpenCL 1.0 dans une interface de style Qt. Elle simplifie l'initialisation de OpenCL, la compilation du programme et l'exécution du kernel. Elle fournit également des fonctions pratiques pour interfacer OpenCL avec les objets Qt existants tels que QImage et QtOpenGL.

Les liens suivants devraient vous aider à démarrer le téléchargement et l'utilisation de QtOpenCL avec Qt 4.6 ou 4.7 :

• le code source : http://qt.gitorious.org/qt-labs/opencl/ ;
• les instructions de compilation : http://qt.nokia.com/doc/opencl-snapshot/building.html ;
• une documentation complète : http://qt.nokia.com/doc/opencl-snapshot/.

QtOpenCL est toujours en cours de développement, distribué comme module autonome et hébergé dans des dépôts différents de Qt. Suggestions et commentaires sont les bienvenus pour l'améliorer.

On va écrire un programme simple qui modifie une image en récupérant sa version en niveau de gris et qui la multiplie par une couleur. Le code est dans les dépôts de QtOpenCL dans le répertoire « examples/opencl/colorize/ ». Nous allons utiliser les variables membres suivantes dans la classe ColorizeWidget :

QCLContext context;
QCLProgram program;
QCLKernel colorize;
QImage dstImage;
QCLImage2D srcImageBuffer;
QCLImage2D dstImageBuffer;
QColor color;

La première chose que nous devons faire est de créer QCLContext, qui détermine quel périphérique de calcul (CPU ou GPU) peut être utilisé et de l'ouvrir pour notre usage :

if (!context.create())
    qFatal("Could not create OpenCL context");

Dans cet exemple, nous ne faisons pas vraiment attention si le périphérique est un CPU ou GPU, mais si nous voulons vraiment utiliser le GPU comme dans l'implémentation d'OpenGL, nous pouvons faire cela :

if (!context.create(QCLDevice::GPU))
    qFatal("Could not create OpenCL context");

La chose suivante à faire, c'est de compiler notre programme OpenCL depuis le fichier source colorize.cl :

program = context.buildProgramFromSourceFile(QLatin1String(":/colorize.cl"));

C'est maintenant le bon moment pour regarder le code OpenCL lui-même à l'intérieur colorize.cl :

// code
const sampler_t samp = CLK_ADDRESS_CLAMP_TO_EDGE |
                       CLK_FILTER_LINEAR;

__kernel void colorize(__read_only image2d_t srcImage,
                       __write_only image2d_t dstImage,
                       float4 color)
{
    int2 pos = (int2)(get_global_id(0), get_global_id(1));
    float4 srcColor = read_imagef(srcImage, samp, pos);
    float gray = srcColor.x * 11.0f / 32.0f +
                 srcColor.y * 16.0f / 32.0f +
                 srcColor.z * 5.0f / 32.0f;
    float4 pixel = (float4)(color.xyz * gray, srcColor.w);
    write_imagef(dstImage, pos, clamp(pixel, 0.0f, 1.0f));
}

Nous allons le décomposer étape par étape :

• le mot-clé __kernel introduit un point d'entrée pour la fonction spéciale appelée « colorize » que nous utiliserons plus tard dans notre code C++ ;
• le point d'entrée de la fonction « colorize » prend trois paramètres correspondant à l'image source, l'image de destination, et la couleur à combiner avec l'image ;
• la variable « pos » est définie comme vecteur à deux dimensions d'entiers qui contient les deux premiers identifiants globaux. Qu'est-ce que c'est ? Eh bien, dans OpenCL, pour chaque exécution du kernel est donné un argument implicite qui indique quelle partie de l'ensemble du traitement est en cours. Dans notre cas, nous utilisons les coordonnées (x, y) du pixel de l'image que nous voulons traiter. Cette ligne sert simplement à récupérer le pixel courant et la fonction « colorize » ne travaille donc que sur un seul pixel à la fois ;
• la valeur « srcColor » est lue à partir de l'image source à la position « pos ». Vous pouvez ignorer l'échantillonneur « samp » pour l'instant (c'est une technique d'OpenCL pour ajuster la façon dont les valeurs sont extraites à partir de l'objet image), nous utilisons un échantillonneur linéaire simple ;
• nous convertissons la « srcColor » en niveaux de gris puis la combinons avec « color » ;
• enfin, nous écrivons le pixel sur l'image de destination.

Cela est assez simple. La principale différence avec une fonction C habituelle qui ferait la même chose que ce qui précède, c'est que nous n'avons pas inclus les boucles for pour itérer sur X et Y : OpenCL le fera pour nous. Retournons au code C + + maintenant. Les fonctions kernel d'OpenCL en cours d'exécution dans un processeur de calcul ne peuvent pas accéder directement à la mémoire globale, nous devons donc prendre des dispositions pour copier notre image source dans un tampon d'image :

QImage img(QLatin1String(":/qtlogo.png"));
srcImageBuffer = context.createImage2DCopy(img,QCLMemoryObject::ReadOnly);

Nous précisons que la source de l'image est en lecture seule parce que le kernel ne fera que la lire. L'image de destination est créée de la même façon, mais en écriture seule :

dstImage = QImage(img.size(), QImage::Format_ARGB32);
dstImageBuffer = context.createImage2DDevice(
    dstImage.format(), dstImage.size(), QCLMemoryObject::WriteOnly);

La fonction createImage2DDevice() copie l'image dans la mémoire locale du périphérique utilisé par OpenCL. La dernière étape d'initialisation est de trouver le point d'entrée du kernel :

// code
colorize = program.createKernel("colorize");
colorize.setGlobalWorkSize(img.size());
colorize.setLocalWorkSize(8, 8);

Nous avons ajusté la taille du traitement global aux dimensions de l'image, ce qui permet à OpenCL de créer implicitement les boucles qui parcourent les valeurs X et Y pour nous. Nous avons également ajusté la taille du traitement local à 8x8, ce qui indique que OpenCL peut traiter les données en blocs de 8x8 et que chaque élément du bloc peut être traité en parallèle. C'est ainsi que OpenCL obtient un gain de performance : en ajustant la taille du traitement local, nous pouvons adapter le parallélisme pour un usage plus efficace des ressources. J'ai trouvé que des blocs de taille 8x8 fonctionnent très bien pour les images, c'est donc ce que nous allons utiliser dans cet exemple. Maintenant que nous avons initialisé notre environnement OpenCL et le kernel, regardons la fonction paintEvent() :

colorize(srcImageBuffer, dstImageBuffer, color);
dstImageBuffer.read(&dstImage);
QPainter painter(this);
painter.drawImage(0, 0, dstImage);

La première ligne exécute le kernel pour nous par l'intermédiaire de l'opérateur surchargé () de QCLKernel(). La seconde ligne copie alors le contenu de « dstImageBuffer » de l'interface OpenCL vers « dstImage » de la mémoire globale. Nous pouvons ensuite peindre la fenêtre de manière habituelle. Et voilà la base !

Eh bien… pas si vite ! J'ai mis de côté un petit détail : le noyau s'exécute en arrière-plan et retourne au programme C + + immédiatement. Ainsi, après la première ligne, l'exécution se poursuivra. Mais la fonction read() va se bloquer automatiquement pour attendre que le kernel termine son exécution : tout va bien dans cet exemple. Mais si nous voulions vraiment marquer une pause pour attendre que le kernel termine son exécution, nous pouvons utiliser un QCLEvent :

QCLEvent event = colorize(srcImageBuffer, dstImageBuffer, color);
event.waitForFinished();
dstImageBuffer.read(&dstImage);
QPainter painter(this);
painter.drawImage(0, 0, dstImage);

Le dépôt de QtOpenCLcontient un certain nombre d'exemples avec lesquels vous pouvez jouer :

• un exemple d'addition vectorielle - une autre introduction à QtOpenCL ;
• un programme de visualisation de Mandelbrot qui démontre la production de données QImage et textures GL via OpenCL ;
• un exemple de flou gaussien et des tests de performance qui compare ce programme avec les effets graphiques de Qt ;
• un exemple de subdivision de courbes de Bézier pour démontrer comment utiliser OpenCL comme shader géométrique pour générer un grand nombre de sommets ;
• un exemple de tracé simple et de mélange et de dessin d'image.

Et maintenant une capture d'écran est nécessaire. Le programme de visualisation de courbes de Mandelbrot permet d'agrandir la zone bien connue pour donner cette image :

Sur la carte graphique NVIDIA GeForce GTX 275 de mon ordinateur de bureau sous Linux, cela permet d'obtenir jusqu'à 120 images par seconde en l'exécutant sur 30 unités de calcul, sans faire couler une goutte de sueur. Pour mettre cela en perspective, le même algorithme s'exécutant sur le processeur lutte pour atteindre cinq images par seconde. Le décharger de ce travail et découper le traitement en morceaux de 8x8 fait une énorme différence (initialement, la performance n'a pas été excellente jusqu'à ce que je réalise que j'utilisais une taille pour les traitements de 1x1).

Il y a peu d'interactions entre QtOpenCL et QtConcurrent, comme décrit ici (http://qt.nokia.com/doc/opencl-snapshot/concurrent.html). Du fait que QtConcurrent fonctionne en arrière-plan sur une structure multicore homogène, un peu de travail doit être réalisé pour marier vraiment les deux mondes, mais rien n'est impossible. Pour l'instant, la caractéristique la plus utile de l'interaction est que vous pouvez obtenir un QFuture pour une exécution le kernel et le transmettre à un QFutureWatcher pour bénéficier des signaux :

QCLEvent event = kernel(arg1, arg2);
QFutureWatcher<void> *watcher = new QFutureWatcher<void>(this);
watcher->setFuture(event.toFuture());
connect(watcher, SIGNAL(finished()), this, SLOT(eventFinished()));

Ou bien, en utilisant la conversion implicite entre QCLEvent et QFuture :

QFutureWatcher<void> *watcher = new QFutureWatcher<void>(this);
watcher->setFuture(kernel(arg1, arg2));
connect(watcher, SIGNAL(finished()), this, SLOT(eventFinished()));

À l'heure actuelle, QtOpenCL fonctionne très bien avec les implémentations non embarquées d'OpenCL comme celle de NVIDIA (nous l'avons testé sous Linux, Mac et Windows). Actuellement, les systèmes embarqués sont une autre affaire : les implémentations d'OpenCL sont encore très rudimentaires dans cet espace. Le gain en performance est à peine meilleur en utilisant les processeurs embarqués qu'en utilisant les instructions ARM/NEON par exemple. Et les processeurs graphiques intégrés sont généralement configurés en dur pour GLSL/ES et manquent des nombreuses fonctionnalités qui rendent OpenCL vraiment intéressant. Mais comme tout dans le monde de l'embarqué, les choses sont susceptibles de changer très rapidement. En sortant une version stable de QtOpenCL, j'espère que nous pourrons stimuler les fournisseurs de systèmes embarqués pour accélérer le développement en leur donnant quelque chose à tester. « Soyez le premier appareil embarqué du marché à proposer un programme de démonstration des courbes de Mandelbrot tournant à 10, 20 ou 60 images par secondes ! »

Beaucoup de choses restent à faire, en particulier sur la façon dont nous pouvons utiliser OpenCL à l'intérieur de Qt lui-même. Il y a beaucoup d'endroits où OpenCL pourrait être utile :

• accélérer le mélange d'images et le tracé de courbes dans le moteur de dessin ;
• décompression rapide à la volée et dimensionnement des images JPEG ;
• effets graphiques : flou, colorié, effet bloom, etc., etc., etc. ;
• les effets de particules et autres simulations physiques ;
• subdivision de meshs et algorithmes de morphing en Qt/3D.

Les possibilités sont infinies. Nous attendons avec impatience vous commentaires ! :-)

N.D.L.T. : pour plus d'informations sur le calcul sur processeurs graphiques et l'architecture de ces processeurs, je vous conseille l'article d'introduction à CUDA de Thibaut Cuvelier et la FAQ sur le GPGPU :

• https://tcuvelier.developpez.com/gpgpu/cuda/introduction/
• https://jeux.developpez.com/faq/gpgpu/

Merci à dourouc05, à worm83 et à Pierre Fauconnier pour leurs relectures et pour leurs conseils.

Au nom de toute l'équipe Qt, j'aimerais adresser le plus grand remerciement à Nokia pour nous avoir autorisés à traduire de cet article !