Denis Salem | Blog

CATEGORIES

ARCHIVES

Avant de se lancer dans ce tutoriel il convient d'être à l’aise avec la notion de shaders et les opérations basiques d'OpenGL.

Si vous ne connaissez pas OpenGL, il est sans doute trop tôt pour se lancer dans les compute shaders. Il existe une pléthore de très bons tutoriels sur l'OpenGl moderne qui vous permettrons de vous familiariser avec l'API. J'en retiens notamment trois

Pour comprendre cette introduction aux compute shaders vous n'avez pas besoin d'aller trop loin dans les tutoriels OpenGL ci-dessus. Le minimum étant de savoir rendre à l'écran un simple triangle, auquel on applique une texture, en OpenGL moderne, avec les vertex/fragment shader minimaux qu'il convient d'avoir.

Nous allons voir ici une fonctionnalité introduite depuis la version 4.3 de notre API favorite qui va nous permettre d'étendre nos possibilités.

Rappelons qu’OpenGL définit une pipeline de rendu bien précise dont certaines étapes (stages) sont programmables en GLSL (OpenGL Shading Language) ; Ces étapes programmables sont appelées shaders. On connaissait déjà les vertex shaders qui permettent de manipuler les vertices et les fragments shaders qui permettent de réaliser et d’appliquer des effets graphiques à notre image de sortie.

Avec les compute shaders il est maintenant possible de manipuler arbitrairement des données de la façon dont on le souhaite. Ce gain en flexibilité représente un grand intérêt puisque l’on peut tirer parti de la puissance de la GPU pour, par exemple, créer à la volée des textures de synthèses, générer procéduralement des objets, ou effectuer des calculs parallèles qui ne sont pas forcément en rapport avec de l’imagerie 3D sans dépendre des APIs comme CUDA ou OpenCL dont le propos n'est de faire QUE du calcul parallèle. En effet, embrayer dans un même programme d'OpenGL vers l'une des APIs mentionnées à l'instant (ou vis versa) est très chronophage, les compute shaders constituent une solution commode à ce problème.

Aperçu

Lorsque l’on manipule les vertex shaders ou les fragment shaders on se rend bien compte que chacun des appels de ces shaders sont indépendant est ignorent les autres. Ces shaders, s’inscrivant dans une pipeline de rendu bien définit, leurs entrées/sorties sont elles aussi bien définies, tout comme leur fréquence d’exécution (par exemple un vertex shader donné s’exécutera pour chaque vertex qu'on lui donne).

Les computes shaders, eux, sont indépendant de cette pipeline de rendu et lorsque ceux là sont invoqués les entrées/sorties sont définies par le programmeur et non plus par la pipeline. C’est également le programmeur qui détermine la charge de travail à effectuer en définissant des groupes de travail (workgroups).

Pour illustrer le fonctionnement des computes shaders on se donne le programme exemple suivant :

damnSimpleComputeShader

Ce code source sert de support à ce tutoriel mais relire le code d'un autre peut être laborieux et/ou peut-être n'utilisez vous pas les mêmes librairies que moi. En réalité, il est très simple d'intégrer le compute shader rudimentaire que je présente ici dans un petit programme OpenGL tout aussi rudimentaire. N'ayez donc pas peur d'écrire votre propre programme et d'intégrer au fur et à mesure les éléments manquants qui nous intéressent.

Pour se représenter un peu le tableau de loin, on se propose à l’aide d’un compute shader de générer le contenu d’une texture vide pour réaliser un damier, et l’afficher. En fait, c’est assez simple ; dans notre exemple, tout se passe comme si on allait créer une texture vierge et qu’on l’affichait normalement sur un quad aux dimensions du viewport, à la différence près qu’avant d’entrer dans la boucle de rendu, on génère le damier en appelant notre compute shader et en s’assurant que la boucle de rendu n’interfère pas avec la génération des données. En effet, sans explicitement bloquer le fil d’exécution OpenGL va écrire dans la texture et simultanément essayer de l’afficher, parallélisme oblige, ce qui est évidemment très peu commode.

Il faut insister sur un point important. L'architecture de la carte graphique est conçu pour exécuter parallèlement des instructions. Ce qui impose nécessairement de prendre en considération cet état de fait lors de l'implémentation d'un algorithme, en particulier quand celui là est à l'origine conçu pour des architectures plus traditionnelles où le multi-threading est en option.

Création de la texture

Maintenant que l’on a un aperçu de la façon dont se déroule notre programme, allons un peu plus dans le détail. La première chose à faire est de créer la texture qui nous servira de support.

glGenTextures(1, &quadTextureID);
  glActiveTexture(GL_TEXTURE0);
  glBindTexture(GL_TEXTURE_2D, quadTextureID);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
    glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR);
    glGenerateMipmap(GL_TEXTURE_2D);  
    glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA32F, 640, 480, 0, GL_RGBA, GL_FLOAT, NULL);
    glBindImageTexture (0, quadTextureID, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_RGBA32F);
    glBindImageTexture (0, 0, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_RGBA32F);
  glBindTexture(GL_TEXTURE_2D, 0);

Rien de bien nouveau à la différence qu’ici on bind la texture courante à une image unit. Ce qui signifie dans la pratique que l’on va pouvoir travailler dans un tableau avec des indices entiers, et non plus dans un sampler2D où les coordonnées sont normalisées.

La plupart du temps, dans les tutoriels que l'on peut lire en ligne (en français ou en anglais), il n'est jamais question d'image unit et de texture unit. Pourtant c'est important dans la mesure où même s'il semble que l'on peut travailler sans, dans des cas simples, il est beaucoup plus rigoureux (mais plus fastidieux) d'expliciter le comportement de votre application. J'ouvre donc une parenthèse pour éclaircir un peu ce dont il s'agit: glBindTexture associe l'objet texture à la cible d'une texture unit active (Cette cible c'est GL_TEXTURE_2D en général). Ce qui signifie qu'il faut… L'activer. Sauf dans le cas où l'on utilise qu'une seule texture où, en effet, la texture unit par défaut est fixée à GL_TEXTURE0. Cette petite opération est donc invisible la plupart du temps.

La texture unit permet de récupérer la texture sous la forme d'un sampler2D dans le shader qui souhaite l'utiliser, et par là, accéder à l'image de la texture avec le mode d'interpolation définit avec glTexParameteri. Dans ce cas là les coordonnées utilisées pour accéder à l'image de la texture sont normalisées, d'où l'interpolation.

L'image unit associée avec glBindImageTexture permet de récupérer dans le shader l'image à proprement parler, sans coordonnées normalisées. On se déplace donc dans un tableau traditionnel avec des coordonnées entières.

Remarquons également qu'en général on passe à glTexImage2D un pointeur vers un tableau de données pour initialiser la texture avec celle-ci. Comme on alloue la mémoire côté GPU et qu'on générera ces données toujours sur la GPU ce pointeur peut-être NULL.

Comme on le sait déjà, on unbind généralement ses buffers par précaution après avoir fini de travailler avec. C’est la même chose pour glBindImageTexture. Plus loin dans le code de l’exemple on peut voir qu’à chaque appels de glBindImageTexture un unbind de circonstance est placé avant glBindTexture(GL_TEXTURE_2D, 0).

Enfin, l’image unit associée à la texture est en WRITE_ONLY. En effet pour ce que l’on se propose de faire on a pas besoin de lire l’image depuis le compute shader, mais seulement d’y écrire.

Création du compute shader

Un peu plus bas dans le programme, on arrive au bloc d’instructions où l’on crée notre compute shader.

 1 PrintWorkGroupsCapabilities();
 2 
 3   GLuint computeShaderID;
 4   GLuint csProgramID;
 5   char * computeShader = 0;
 6 
 7   GLint Result = GL_FALSE;
 8   int InfoLogLength = 1024;
 9   char ProgramErrorMessage[1024] = {0};
10 
11   computeShaderID = glCreateShader(GL_COMPUTE_SHADER);
12 
13   loadShader(&computeShader, "compute.shader");
14   compileShader(computeShaderID, computeShader);
15 
16   csProgramID = glCreateProgram();
17 
18   glAttachShader(csProgramID, computeShaderID);
19   glLinkProgram(csProgramID);
20   glDeleteShader(computeShaderID);

La première chose à faire, bien que ce ne soit pas obligatoire ici, c’est de connaître la capacité des groupes de travail offerte par le GPU. Ça ne mange pas de pain et ça va être l’occasion de détailler comment les groupes de travail sont organisés. Les groupes de travail contiennent un certain nombre d’invocations du compute shader, définit à l’intérieur du shader lui même comme nous le verrons plus bas. Ce nombre est appelé taille locale du groupe de travail. On définit également au moment de l’exécution du compute shader le nombre de groupes de travail à effectuer.

On peut schématiser ce qui vient d’être dit ainsi :

Groupes de travail

Où chaque petit carré représente une invocation de shader contenu dans un groupe de travail.

Ceci étant dit on devine que la carte graphique a une capacité limitée de groupes de travail et de nombre d’invocations locales. OpenGL permet de déterminer la capacité de la carte graphique en la matière. En effet dans la fonction printWorkGroupsCapabilities on récupère ces informations et on les affiches dans la sortie standard :

 1 void printWorkGroupsCapabilities() {
 2   int workgroup_count[3];
 3   int workgroup_size[3];
 4   int workgroup_invocations;
 5 
 6   glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 0, &workgroup_count[0]);
 7   glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 1, &workgroup_count[1]);
 8   glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 2, &workgroup_count[2]);
 9 
10   printf ("Taille maximale des workgroups globaux:\n\tx:%u\n\ty:%u\n\tz:%u\n",
11   workgroup_size[0], workgroup_size[1], workgroup_size[2]);
12 
13   glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, 0, &workgroup_size[0]);
14   glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, 1, &workgroup_size[1]);
15   glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, 2, &workgroup_size[2]);
16 
17   printf ("Taille maximale des workgroups locaux:\n\tx:%u\n\ty:%u\n\tz:%u\n",
18   workgroup_size[0], workgroup_size[1], workgroup_size[2]);
19 
20   glGetIntegerv (GL_MAX_COMPUTE_WORK_GROUP_INVOCATIONS, &workgroup_invocations);
21   printf ("Nombre maximum d'invocation de workgroups locaux:\n\t%u\n", workgroup_invocations);
22 }

Ce qu’on peut remarquer, et c’est très important, c’est que la taille locale d’un groupe de travail ainsi que le nombre de groupes de travail est définit en trois dimensions. Autrement dit les compute shaders travaillent dans l’espace. C’est très pratique, évidement, pour faire du traitement d’image, ou travailler en volume. Nous verrons plus bas comment organiser notre shader en fonction de ces informations.

Précisons également que ces groupes de travail sont lancés en parallèles selon la capacité et la disponibilité de la carte graphique. Il n’est pas possible de connaître l’ordre d’exécution des invocations et dans certains cas il est absolument vital de composer avec cet état de fait sous peine d’obtenir des résultats pour le moins inattendus ou pire, de planter le système. Heureusement, la question de la synchronisation n’est pas le sujet ici, et le programme que l’on étudie ne nécessite pas que l’on soit attentif à l’ordre d’exécution des groupes de travail et d’invocations des compute shaders. Notre exemple n’étant pas non plus gourmand en ressource il ne nécessite pas que l’on implémente une vérification particulière de la capacité des groupes de travail. La fonction printWorkGroupsCapabilities est donc ici purement informative.

Comme on le voit dans le listing plus haut, les instructions qui servent à créer le compute shader sont en fait tout à fait classique.

Rappelons tout de même que l’on crée un programme spécifiquement POUR le compute shader. Un programme qui n’est donc PAS celui où sera lié le vertex shader et le fragment shader qui servent au rendu. On le voit bien, plus bas dans le code ; la création du programme de rendu est indépendante et ne fait intervenir à aucun moment le compute shader.

Exécution du shader

Une fois que tout est prêt on peut enfin lancer notre compute shader.

1 glUseProgram(csProgramID);
2     glBindTexture(GL_TEXTURE_2D, quadTextureID);
3       glBindImageTexture (0, quadTextureID, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_RGBA32F);
4         glDispatchCompute(40,30,1);
5         glMemoryBarrier(GL_SHADER_IMAGE_ACCESS_BARRIER_BIT);
6       glBindImageTexture (0, 0, 0, GL_FALSE, 0, GL_WRITE_ONLY, GL_RGBA32F);
7     glBindTexture(GL_TEXTURE_2D, 0);
8   glUseProgram(0);

OpenGL étant une machine à état on bind notre texture et l’image unit qui lui est associé exactement comme on l’a fait au moment de la création de la texture.

C’est finalement avec glDispatchCompute() que l’on lance l’exécution du compute shader et plus exactement que l’on lance l’exécution d’un certains nombre de groupes de travail. En effet glDispatchCompute prend trois paramètres x, y et z spécifiant chacun les dimensions de l’espace des workgroups à lancer. En particulier, dans notre example, on travaille sur une image, donc dans un espace en deux dimensions, le troisième paramètre est naturellement fixé à 1.

Dans le programme on voit que j’ai choisi pour l’espace des groupes de travail une largeur de 40 et une hauteur de 30 (j’expliquerais le choix de ces valeurs plus bas), ce qui signifie qu’OpenGL va exécuter 40 x 30 = 1200 groupes de travail. Cela peut paraître beaucoup mais c’est en fait relativement peu.

Après l’appel de glDispatchCompute, OpenGL se met au travail et le fil d’exécution continue côté CPU. Sauf qu’en fait non ; étant vigilant, on s’assure que le fil d’exécution de notre programme s’interrompe le temps qu’OpenGL termine l’exécution de tous les groupes de travail à l’aide de la commande glMemoryBarrier qui permet comme on le devine de mettre en pause l’exécution du programme tant que la carte graphique n’a pas terminé ses transactions mémoires. En particulier, cela permet de bloquer l’exécution en fonction du type de mémoire que l’on veut. Ici, comme nous travaillons sur une image on utilise la valeur prédéfinie GL_SHADER_IMAGE_ACCESS_BARRIER_BIT.

On pourrait également, si cela était nécessaire, combiner plusieurs valeurs prédéfinies par OpenGL pour bloquer l’exécution du programme en fonction de plusieurs zones mémoire. Enfin si GL_ALL_BARRIER_BITS est utilisé alors le programme se verrouille en fonction de toutes les zones mémoire précédemment sollicitées.

Ceci étant fait, les workgroups étant tous exécutés, on peut reprendre le fil d’exécution du programme et unbind nos états OpenGL.

Après quoi, on rentre tout simplement dans la boucle de rendu qui ne fait absolument rien de nouveau.

Détail du compute shader

C’est donc le moment de voir à quoi ressemble le compute shader que l’on a exécuté.

 1 #version 430
 2 
 3 layout (local_size_x = 16, local_size_y = 16) in;
 4 
 5 layout (rgba32f, binding = 0) uniform image2D img_output;
 6 
 7 void main() {
 8   // Aucun tableau de donnée n'étant passé au moment de la création de la texture,
 9   // c'est le compute shader qui va dessiner à l'intérieur de l'image associé
10   // à la texture.
11 
12   // gl_LocalInvocationID.xy * gl_WorkGroupID.xy == gl_GlobalInvocationID
13   ivec2 coords = ivec2(gl_GlobalInvocationID);
14 
15   // Pour mettre en evidence. Les groupes de travail locaux on dessine un damier.
16   vec4 pixel;
17   if ( ((gl_WorkGroupID.x & 1u) != 1u) != ((gl_WorkGroupID.y & 1u) == 1u)) {
18     pixel = vec4(1.0,.5,.0,1.0);
19   }
20   else {
21     pixel = vec4(.0,.5,1.0,1.0);
22   }
23 
24   imageStore(img_output, coords, pixel);
25 }

Un compute shader se présente en fait de façon très similaire aux shaders classiques. La nouveauté c’est que l’on définit à l’intérieur du shader lui même la taille des groupes de travail. Comme on le voit, la taille locale d’un groupe de travail est de 16 par 16. Rappelons que les dimensions de l’espace de travail des workgroups étaient de 40 par 30. On aura donc 16² x 1200 invocations du compute shader. On peut également remarquer que 16 * 40 = 640 et 16 * 30 = 480. Et ça tombe bien parce que c’est justement les dimensions de notre image et de notre viewport.

Quelques explications s’imposent. Nous pourrions effectivement générer un damier dans une image de façon itérative dans un seul thread. Mais ça ne serait pas vraiment optimal si on tient compte du fait que l’on peut faire la même chose en tirant parti de l’aptitude du GPU à faire du calcul parallèle.

Comme on l'a dit plus haut, la raison pour laquelle le nombre de groupes de travail et la taille de ceux là sont formulés en terme d’espace est que les compute shaders travaillent dans l'espace. Pour se repérer dans l’espace de la structure de donnée dans laquelle on travaille on utilise des variables introduites avec les compute shaders qui nous renseignent sur l’identifiant de l’invocation courante du compute shader. Cet identifiant étant tridimensionnel, on sait donc où l’on se trouve au moment d’une invocation donnée. Comme les invocations sont parallèles (mais également dans un ordre arbitraire décidé par OpenGL) on peut donc ici travailler simultanément à différents endroits de notre image, ce qui fait substantiellement gagner du temps.

Du coup, la seule façon de savoir où l’on se trouve est d’interroger ces variables spéciales qui identifient l’invocation courante.

Ces variables sont des vec3. À titre de rappel, OpenGL travail le plus souvent avec des vecteurs de deux, trois ou quatre composantes; respectivement vec2, vec3 et vec4. Les vecteurs peuvent contenir des entiers signés ou non signés ou, comme c'est le cas la plupart du temps, des nombres à virgules flottantes. Pour en savoir plus rendez vous sur le wiki d'OpenGL.

  • gl_LocalInvocationID est un vec3 qui nous dit où l’on se trouve relativement au groupe de travail courant. Les coordonnées xyz ainsi retournées ne peuvent donc pas être en dehors du volume décrit par la taille du groupe locale, à savoir ici : 16 par 16 par 1.
  • gl_WorkGroupID est un vec3 qui nous renseigne sur le groupe de travail courant, à ne pas confondre avec l’invocation courante donc. Les coordonnées xyz ainsi retournées ne peuvent pas être en dehors du volume décrit par la taille de l’espace des groupes de travail, à savoir ici : 40 par 30 par 1.
  • gl_GlobalInvocationID est un vec3 qui est en fait le produit des deux variables vues précédemment. Il repère l’invocation courante non plus relativement au groupe de travail courant mais dans l’espace global des invocations. Autrement dit, dans notre programme, cette variable nous dit où l’on se trouve dans l’image.

Finalement pour dessiner notre damier, on détermine la couleur de sortie en fonction de gl_WorkGroupID ; si sa composante x est paire et que sa composante y ne l’est pas alors la couleur de sortie est orange, sinon, elle est bleue. Toutes les invocations à l’intérieurs d’un groupe de travail seront donc soit oranges, soit bleues.

Pour terminer, on utilise la fonction GLSL imageStore pour écrire dans notre image. Remarquez qu’en fait on n’enregistre qu’un seul et unique pixel pour une invocation donnée du compute shader.

Résultat

En compilant et en exécutant le programme exemple (en vous assurant préalablement de satisfaire les dépendances) on obtient l’image ci-dessous.

Damier de 40x30 secteurs de 16 pixels²

Comme on pouvait s’y attendre notre damier comporte des secteurs de 16 pixels² au nombre de 40 en largeur et 30 en hauteur. On illustre ainsi sans ambiguïté l’utilisation d’un compute shader et des groupes de travail.

Pour terminer ce cours introductif vous pouvez jouer avec le compute shader présenté ici et modifier la taille du groupe de travail locale et adapter l'espace des groupes en conséquence pour changer la taille des cases du damier ou pour rendre à l'écran des formes géométriques et les positionner comme bon vous semble.

Bibliographie

Remerciement

Cet article à été rédigé dans le cadre d’un stage à l’INRIA, un grand merci à Sylvain Lefebvre et son équipe pour leur aide et leur disponibilité.

Un très grand merci également à Ge0 et Shenzyn pour leurs conseils et leurs relectures.