The Gimp et Cuda
D’un côté : The Gimp
De l’autre : Cuda
Je pense qu’il n’est plus besoin de présenter The Gimp : LE logiciel de dessin bitmap du monde libre.
C’est quoi Cuda ?
Par contre, 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 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 :
- Ubuntu Hardy 8.04 (Linux),
- Cuda compilation tools, release 2.0, V0.2.1221 (nvcc -V),
- GCC 4.2.3,
- paquet libgimp2.0-dev installé.
Compilation et installation
La ligne ci-dessous permet de compiler le plug-in cudatest pour Gimp
nvcc -o cudatest --host-compilation C `pkg-config --cflags --libs gimp-2.0`
gimpcuda.cu gimpcuda_kernel.cu
Description 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 :
- Cuda learn : la base d’informations sur Cuda sur le site de nVidia. Contient tout ce qu’il faut : la documentation de référence, le SDK, des programmes d’exemple…
- Developer Gimp : le site officiel pour tous les développeurs sous Gimp.
- Plug-in tutorial : le tutoriel que j’ai suivi pour écrire la base du plug-in. Il existe également un template mais je ne sais pas comment intégrer la compilation par nvcc ou encore la prise en compte de fichier à l’extension .cu.
- PyCuda : PyCuda ou Cuda pour Python. Permet d’exécuter du code Cuda directement depuis Python sans se soucier de la compilation ou des problèmes d’allocation mémoire. Très utile pour se faire la main sur Cuda !
- Forums nVidia : le forum Cuda sur le site de nVidia. On y trouve pas mal d’astuce et de morceaux de code réutilisables.
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
#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
#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);
}