Utilisation de CUDA sous Delphi : comment profiter de la puissance de calcul des GPU

Un billet blog de Gouyon

Le 1^er mai 2021, par Gouyon

0PARTAGES

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

Code :

Sélectionner tout

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;

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.

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

Code :

Sélectionner tout

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 fichier de définition qui lui par contre permet à Delphi de retrouver les fonctions dans la DLL

Code :

Sélectionner tout

1
2
3
4
 
EXPORTS 
trt_init 
trt_imgtobmp

Le cœur de la fonction

Code :

Sélectionner tout

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	 
}

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).

Ce qu'il faut savoir c'est que la fonction de transformation est

Code :

Sélectionner tout

__global__ void valtocoul(int N, unsigned short *brt, int *bmp, float a, float b)

et qu'elle est appelée de cette manière un peu particulière dans la fonction trt_imgtobmp

Code :

Sélectionner tout

1
2
 
valtocoul <<<256, 256 >>> (640 * 512, dev_in, dev_out, aa, min);

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.

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.

Code :

Sélectionner tout

1
2
 
bmp[tid] = 0xFF000000 + (vb << 16) + (vb << 8) + vb;

Maintenant je vais expliquer comment utiliser ce type de fonction dans Delphi

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.

Code :

Sélectionner tout

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;

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

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;

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.
Je déclare aussi une Bitmap qui contiendra l'image créée à partir des données de la caméra.

Code :

Sélectionner tout

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;

Dans la procédure FormCreate je rempli mon tableau et je créé mon objet bitmap.

Code :

Sélectionner tout

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;

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.
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.

Une erreur dans cette actualité ? Signalez-nous-la !

Votre nom : Votre e-mail :

Décrivez l'erreur que vous souhaitez porter à notre connaissance :

Identifiez-vous

Créer un compte

Utilisation de CUDA sous Delphi : comment profiter de la puissance de calcul des GPU

Un billet blog de Gouyon

Identifiant
Mot de passe

Mot de passe oublié ?

Identifiez-vous

Créer un compte

Utilisation de CUDA sous Delphi : comment profiter de la puissance de calcul des GPU Un billet blog de Gouyon

Utilisation de CUDA sous Delphi : comment profiter de la puissance de calcul des GPU

Un billet blog de Gouyon