[Actualité] Utilisation de CUDA sous Delphi
par
, 21/04/2021 à 08h32 (5987 Affichages)
Pour ceux qui ne le savent pas CUDA est un langage semblable au C qui permet d’exécuter des algorithmes sur les cartes graphiques de chez NVDIA. L'avantage de ce type d’exécution c'est qu'il est massivement parallèle, car vous allez disposer selon le type de carte de plusieurs centaines à plusieurs milliers de processeurs. Du coup vous allez pouvoir distribuer vos données sur tous ces processeurs et ce que vous auriez fait en 1000 coups d'horloges va être réalisé en 1.
Ainsi si dans une application Delphi vous avez des fonctions qui nécessitent de la puissance de calcul on voit bien l'avantage d'utiliser CUDA. Dans mon cas j'utilise ce type de fonctionnalité pour faire du traitement d'image en "temps réel". J'ai une application Delphi qui récupère, traite et affiche le résultat du traitement des informations issues d'une caméra. Concrètement je dispose d'une caméra qui filme une scène, les images issues de la caméra sont récupérées sous la forme d'un tableau d'entier court non signé que je dois traiter et afficher sous la forme d'une image en niveau de gris à une cadence de 25Hz soit un temps de traitement de 40ms. Ce qui en fonction du type de traitement peut être très très court. C'est pourquoi j'ai développé mes fonctions de traitement en CUDA afin d'accélérer au maximum les traitements pour tenir dans la limite des 40ms.
Comme il n'est pas possible d'intégrer directement du code CUDA dans du code Delphi, je suis passé par une DLL en CUDA qui sera utilisée par le programme Delphi. A titre d'exemple je présente ici la conversion d'un tableau d'entier court non signé en une image en niveau de gris.
Description de la fonction
Pour chaque point du tableau on va calculer un niveau de gris qui sera égal à 0 si la valeur du tableau est inférieur ou égale à un minimum et égal à 255 si la valeur du tableau est supérieure ou égale à un maximum. Pour cela il suffit de calculer un facteur d'échelle qui sera:
facteur=255/(maximum-minimum)
et de l'appliquer à chaque point du tableau de cette manière:
niveau=facteur (valeur-minimum) ensuite on corrigera le niveau s'il est inférieur à 0 ou supérieur à 255. Voici ce que ça donne
Si la taille de l'image est de 1024x768 cette opération devra être faite 786432 fois. Maintenant si on dispose d'une carte graphique avec 1024 processeurs alors cette opération devra être répétée seulement 768 fois. Ce qui offre un gain de temps non négligeable.
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4
5
6
7
8
9
10
11 facteur:=255/(maximum-minimum); for i:=0 to TaiileImage-1 do begin vv = facteur * (tableau[i] - minimum); if (vv > 255) vv = 255; if (vv < 0) vv= 0; niveau[i]:=vv; end;
Présentation de la DLL
J'ai développé la DLL avec VisualStudio et je présent ici les différent éléments.
le fichier des entêtes qui n'est pas utilisé avec Delphi
Le fichier de définition qui lui par contre permet à Delphi de retrouver les fonctions dans la DLL
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4 #pragma once int __stdcall trt_init(int *nthb); int __stdcall trt_imgtobmp(unsigned short *dtatin, int *dtaout, int taille, unsigned short min, unsigned short max);
Le cœur de la fonction
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4 EXPORTS trt_init trt_imgtobmp
Pour comprendre comment fonctionne exactement cette fonction je vous suggère de regarder la documentation de CUDA ainsi que les différents tutoriel qui existent sur le site de developpez (exemple introduction à CUDA).
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79 #include "cuda_runtime.h" #include "device_launch_parameters.h" int __stdcall trt_init(int *nthb) { int nbacc; cudaDeviceProp prop; if (cudaGetDeviceCount(&nbacc) == cudaSuccess) { if (cudaGetDeviceProperties(&prop, 0) == cudaSuccess) { *nthb = prop.maxThreadsPerBlock; } else { *nthb = -1; } return nbacc; } else return -1; } __global__ void valtocoul(int N, unsigned short *brt, int *bmp, float a, float b) { int tid = threadIdx.x + blockIdx.x*blockDim.x; while (tid < N) { float v = a * (brt[tid] - b); if (v > 255) v = 255; if (v < 0) v = 0; unsigned char vb = (unsigned char)v; bmp[tid] = 0xFF000000 + (vb << 16) + (vb << 8) + vb; tid += blockDim.x*gridDim.x; } } int __stdcall trt_imgtobmp(unsigned short *dtatin, int *dtaout, int taille, unsigned short min, unsigned short max) { unsigned short *dev_in; int *dev_out; float aa; if (cudaMalloc((void**)&dev_in, taille * sizeof(unsigned short)) == cudaSuccess) { if (cudaMalloc((void**)&dev_out, taille * sizeof(int)) == cudaSuccess) { if (cudaMemcpy(dev_in, dtatin, taille * sizeof(unsigned short), cudaMemcpyHostToDevice) == cudaSuccess) { if (max != min) aa = 255.0 / (float)(max - min); else aa = 0; valtocoul << <256, 256 >> > (640 * 512, dev_in, dev_out, aa, min); if (cudaMemcpy(dtaout, dev_out, taille * sizeof(int), cudaMemcpyDeviceToHost) == cudaSuccess) { cudaFree(dev_in); cudaFree(dev_out); return 1; } else return -4;//erreur copie tableau sortie vers l'host } else return -3;//erreur copie tableau entree sur le device } else return -2;//erreur allocation tableau sortie } else return -1;//erreur allocation tableau entrée }
Ce qu'il faut savoir c'est que la fonction de transformation est
et qu'elle est appelée de cette manière un peu particulière dans la fonction trt_imgtobmp
Code : Sélectionner tout - Visualiser dans une fenêtre à part __global__ void valtocoul(int N, unsigned short *brt, int *bmp, float a, float b)
les valeurs entre <<< et >>> indiquent sur comment le code va être distribué entre les processeurs. Ici on a indiqué qu'on allait utiliser 256 processeurs et chaque processeur va exécuter 256 thread. Je ne vais pas m'étendre plus sur CUDA.
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2 valtocoul <<<256, 256 >>> (640 * 512, dev_in, dev_out, aa, min);
Au niveau de la sortie le niveau de gris doit être traduit en une couleur utilisable par les bitmap. J'ai choisi le format ARGB qui est le plus simple. ma couleur sera donc égale à $FFvvvvvv. C'est la ligne de code ci dessous qui réalise cette opération.
Maintenant je vais expliquer comment utiliser ce type de fonction dans Delphi
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2 bmp[tid] = 0xFF000000 + (vb << 16) + (vb << 8) + vb;
Appel de la DLL dans Delphi
Pour tester ma fonction j'ai créé une application en FMX qui contient un Timage pour l'affichage de mon image en niveaux de gris et un TTimer qui va simuler la cadence de la caméra.
Au début du programme je déclare le prototype es fonctions de ma DLL. Comme en C on utilise des pointeurs pour accéder aux tableaux, j'ai déclaré les paramètres dtain et dtaout comme étant des pointeurs du même type que les valeurs qu'ils contiennent à savoir des entiers court non signé pour les données d'entrées et des entiers non signées pour les données de sortie.
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4
5
6
7
8
9
10
11
12 unit testdll; interface uses System.SysUtils, System.Types, System.UITypes, System.Classes, System.Variants, Windows, FMX.Types, FMX.Controls, FMX.Forms, FMX.Graphics, FMX.Dialogs, FMX.Controls.Presentation, FMX.StdCtrls, FMX.Objects; //DECLARATION DES PROTOTYPES function trt_init(var a: integer): integer; cdecl; external 'traitement.dll' name 'trt_init'; function trt_imgtobmp(dtatin: PWord; dtaout: PUINT; taille: integer; min: uint16; max: uint16): integer; cdecl;
En variable privée je déclare un tableau de 640x512 entier court non signé qui contiendra les valeurs censées êtres issues de ma caméra.
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20 type TForm1 = class(TForm) vuecam: TImage; Timer1: TTimer; procedure FormCreate(Sender: TObject); procedure Timer1Timer(Sender: TObject); private { Déclarations privées } buffcam: array [0 .. 640 * 512 - 1] of uint16; bmpNB: TBitmap; public { Déclarations publiques } end; var Form1: TForm1;
Je déclare aussi une Bitmap qui contiendra l'image créée à partir des données de la caméra.
Dans la procédure FormCreate je rempli mon tableau et je créé mon objet bitmap.
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16 implementation {$R *.fmx} procedure TForm1.FormCreate(Sender: TObject); var m: integer; begin for m := 0 to 640 * 512 - 1 do begin buffcam[m] := m mod 8000 + 2000; end; bmpNB := TBitmap.Create(640, 512); end;
La procédure Timer1Timer va être appelée toutes les 40ms Dans cette procédure je modifie les données caméra et je fait appel à la fonction de transformation dans la DLL.
Code : Sélectionner tout - Visualiser dans une fenêtre à part
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22 procedure TForm1.Timer1Timer(Sender: TObject); var ret, i, m: integer; dta: TBitmapData; begin for i := 1 to 200 do begin m := random(640 * 512); buffcam[m] := 10; end; if bmpNB.Map(FMX.Graphics.TMapAccess.Write, dta) then begin ret := trt_imgtobmp(@buffcam[0], PUINT(dta.Data), 640 * 512, 1000, 12000); if ret = 1 then begin bmpNB.Unmap(dta); vuecam.Bitmap.Assign(bmpNB); end; end; end;
Afin de mettre directement dans l'image le résultat du traitement, j'utilise les procédures Map et Unmap qui permettent d'accéder directement aux pixels de la bitmap. Pour le pointeur de sortie, je force le type (PUINT(dta.Data)) afin que le compilateur ne génère pas d'erreur.
Quand je fais tourner le code voici j'obtiens bien une image qui évolue dans le temps de
vers
Conclusion
Je vous ai montré un petit aperçu de ce qu'il est possible de faire en combinant Delphi et CUDA. Il est bien entendu que ce petit exemple n'a qu'une valeur pédagogique et que le gain de temps de calcul ne vaut pas le temps passé à développer la DLL. Cependant lorsque on envisage des fonctions plus lourdes comme de la corrélation, du filtrage etc etc là le gain est beaucoup plus important. Enfin je dirais que cet exemple illustre le fait qu'on ne peut pas tout faire avec un même environnement de développement et que la combinaison de différentes technique de programmation peut s'avérer être un plus. Ce qui m'amène à dire que ce serait un plus si CBuilder tout comme VisualStudio pouvait intégrer la possibilité de faire des développement CUDA.