The Gimp et Cuda

Menu principal

Exemple d’utilisation de Cuda depuis Gimp

Je pense qu’il n’est plus besoin de présenter The Gimp : LE logiciel de dessin bitmap du monde libre.

Mais connaissez-vous Cuda ?

C’est quoi Cuda ?

Cuda est, au moment où j’écris ces lignes, encore un inconnu pour bon nombre de personnes. Cuda est une technologie développée par nVidia et qui permet au développeur d’utiliser le processeur de la carte vidéo (le GPU) pour effectuer toute sorte de calcul ou de traitement.

Il faut savoir que nos chères cartes graphiques sont plus que dopées par rapport à nos Intel et autres Amd. Pour exemple, un Quad Core contient 4 cœurs (pour simplifier, cela correspond à 4 processeurs), tandis que le GPU d’une GTX280 de nVidia en contient 240… C’est ce que l’on appelle une architecture massivement parallèle.

Cuda met à disposition du développeur un compilateur (nvcc) qui lui permet de compiler du code "cu". En gros : cela rajoute quelques fonctionnalités au compilateur C/C++ traditionnel qui permet d’écrire dans le même fichier du code exécutable par le processeur central et du code exécutable par le GPU. Les fonctions GPU ainsi écrites peuvent être appelées par des fonctions classiques de manière quasi transparente.

Que faire de toute cette puissance

Et bien on peut par exemple soulager le processeur central de certains traitements graphiques. Pourquoi ne pas écrire un plug-in pour Gimp dans ce cas ?

Le plug-in fourni plus bas inverse la valeur de chaque pixel d’une image RGB. Il est évident qu’il ne s’agit pas d’une opération nécessitant une grande puissance de calcul.

Le plug-in est divisé en 2 fichiers sources :

gimpcuda.cu
gère l’interfaçage avec The Gimp. Son rôle est de récupérer le calque courant dans un tampon brut, de le transmettre à la fonction de rendu et de restituer le résultat à The Gimp.
gimpcuda_kernel.cu
gère l’interfaçage avec Cuda. Son rôle est de récupérer le tampon fourni par The Gimp, de l’envoyer sur le GPU, de lancer le traitement, de récupérer le résultat du GPU et de le redonner à The Gimp.

Fonctions de gimpcuda.cu

query
fonction appelée par The Gimp et qui lui permet d’obtenir toutes les informations concernant notre plug-in
render
crée un tampon pouvant accueillir l’image, recopie les données brutes de l’image (gimp_pixel_rgn_get_rect) dans ce tampon, appelle kernel_render pour effectuer le traitement et met à jour l’image (gimp_pixel_rgn_set_rect)
run
fonction appelée par The Gimp lorsque l’utilisateur veut exécuter notre plug-in

Fonctions de gimpcuda_kernel.cu

doProcessPicture
c’est la fonction GPU qui va effectuer le traitement à proprement parler. Ici, on ne travaille plus d’un bloc, il faut penser à morceler le travail à réaliser. Le lancement de tous les threads est réalisé automatiquement par Cuda. Chaque thread peut consulter son numéro. Comme tous les threads travaillent sur les mêmes données, cela leur permet de savoir quelle zone ils doivent traiter. Par exemple, pour une image de 100 lignes et 10 thread, le thread 0 traiterait les lignes de 0 à 9, le thread 1 traiterait les lignes de 10 à 19 etc.
kernel_render
récupère le tampon contenant les données brutes, alloue tampons sur le GPU (un qui contiendra la source et un autre qui contiendra le résultat), envoie les données brutes dans le premier tampon GPU, exécute doProcessPicture, récupère les données du deuxième tampon GPU (le résultat du traitement) dans le tampon initiale.

Pré-requis

Ce plug-in a été testé et développé dans les conditions suivantes :

Compilation et installation

La ligne ci-dessous permet de compiler le plug-in cudatest pour Gimp :

shell
nvcc -o cudatest --host-compilation C `pkg-config --cflags --libs gimp-2.0` gimpcuda.cu gimpcuda_kernel.cu

Descriptions des options :

-o cudatest
nom du fichier exécutable à créer
--host-compilation C
indique que l’on travaille en C. nvcc travaille en C++ par défaut
pkg-config --cflags --libs gimp-2.0
récupère les options de compilations nécessaires au développement de plug-in pour Gimp

Une fois le fichier cudatest généré, il ne vous reste plus qu’à l’installer dans votre répertoire de plug-ins (généralement ~/.gimp-2.4/plug-ins) et à lancer The Gimp.

Lorsque vous développez des plug-ins, il est nécessaire d’arrêter et de relancer The Gimp à chaque modification pour qu’il prenne en compte la nouvelle version du plug-in.

Liens

Voici quelques liens utiles pour le développement de plug-in sous Gimp et/ou avec Cuda :

Le code source

Ici devrait se trouver tout le bla-bla habituel sur le code fourni : je ne garantis rien, il n’est pas une référence en matière de commentaires ni en matière de programmation. Son but est uniquement de montrer le principe permettant d’utiliser Cuda dans un plug-in Gimp.

gimpcuda.cu

c
#include <libgimp/gimp.h>

void kernel_render(unsigned char *buffer,int width,int height,int channels);

static void query(void) {
  static GimpParamDef args[]={
    {GIMP_PDB_INT32   ,"run-mode","Run mode"      },
    {GIMP_PDB_IMAGE   ,"image"   ,"Input image"   },
    {GIMP_PDB_DRAWABLE,"drawable","Input drawable"}
  };

  gimp_install_procedure(
    "plug-in-cuda-test",
    "Plug-in test cuda de Fred",
    "Teste l'intégration de Cuda dans un plug-in Gimp",
    "Frédéric BISSON",
    "Copyright Frédéric BISSON",
    "2008",
    "_Cuda test...",
    "RGB",
    GIMP_PLUGIN,
    G_N_ELEMENTS (args), 0,
    args, NULL
  );

  gimp_plugin_menu_register("plug-in-cuda-test","<Image>/Filters/Misc");
}

static void render(GimpDrawable *drawable) {
  gint         channels;
  gint         x1,y1,x2,y2;
  GimpPixelRgn src,dst;
  guchar       *buffer;
  gint         width,height;

  gimp_drawable_mask_bounds(drawable->drawable_id,&x1,&y1,&x2,&y2);
  channels=gimp_drawable_bpp(drawable->drawable_id);

  width =x2-x1;
  height=y2-y1;

  gimp_pixel_rgn_init(&src,drawable,x1,y1,width,height,FALSE,FALSE);
  gimp_pixel_rgn_init(&dst,drawable,x1,y1,width,height,TRUE ,TRUE );

  buffer=g_new(guchar,channels*width*height);

  gimp_pixel_rgn_get_rect(&src,buffer,x1,y1,width,height);

  kernel_render(buffer,width,height,channels);

  gimp_pixel_rgn_set_rect(&dst,buffer,x1,y1,width,height);

  g_free(buffer);

  gimp_drawable_flush(drawable);
  gimp_drawable_merge_shadow(drawable->drawable_id,TRUE);
  gimp_drawable_update(drawable->drawable_id,x1,y1,width,height);
}

static void run(const gchar *name,gint nparams,const GimpParam *param,gint *nreturn_vals,GimpParam **return_vals) {
  static GimpParam  values[1];
  GimpPDBStatusType status = GIMP_PDB_SUCCESS;
  GimpRunMode       run_mode;
  GimpDrawable      *drawable;

  // Setting mandatory output values
  *nreturn_vals=1;
  *return_vals =values;

  values[0].type         =GIMP_PDB_STATUS;
  values[0].data.d_status=status;

  // Getting run_mode - we won't display a dialog if we are in NONINTERACTIVE mode
  run_mode=(GimpRunMode)param[0].data.d_int32;

  if(run_mode!=GIMP_RUN_NONINTERACTIVE) {
    g_message("Cuda test in progress...\n");
  }

  // Get specified drawable
  drawable=gimp_drawable_get(param[2].data.d_drawable);

  render(drawable);

  gimp_displays_flush();
  gimp_drawable_detach(drawable);
}

GimpPlugInInfo PLUG_IN_INFO={
  NULL,
  NULL,
  query,
  run
};

MAIN()

gimpcuda_kernel.cu

c
#include <stdio.h>

__global__ void doProcessPicture(unsigned char *src,unsigned char *dst,int width,int height,int channels) {
  int i,j;
  int base=threadIdx.x*width*channels*gridDim.x;
  int offset;

  for(j=0;j<gridDim.x;j++) {
    for(i=0;i<width*channels;i++) {
      offset=base+i+j*width*channels;
      dst[offset]=255-src[offset];
    }
  }
}

void kernel_render(unsigned char *buffer,int width,int height,int channels) {
  unsigned char *devSrc;
  unsigned char *devDst;
  int bufferSize;
  int nbIter;
  int nbThread;

  bufferSize=width*height*channels;
  nbThread=192;
  nbIter=height/nbThread;

  // Allocate two buffers on the GPU
  cudaMalloc((void**)&devSrc,bufferSize);
  cudaMalloc((void**)&devDst,bufferSize);

  // Copy our buffer into the GPU input buffer
  cudaMemcpy(devSrc,buffer,bufferSize,cudaMemcpyHostToDevice);

  // Run the GPU routine
  doProcessPicture<<<nbIter,nbThread,0>>>(devSrc,devDst,width,height,channels);

  // Retrieve the GPU output buffer into our buffer
  cudaMemcpy(buffer,devDst,bufferSize,cudaMemcpyDeviceToHost);

  // Free allocated GPU buffers
  cudaFree(devSrc);
  cudaFree(devDst);
}