Gouyon

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

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 - 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 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
EXPORTS
trt_init
trt_imgtobmp

Le cœur de la fonction

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	
}

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 - Visualiser dans une fenêtre à part

__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 - Visualiser dans une fenêtre à part

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 - Visualiser dans une fenêtre à part

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

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

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