Bonjour à tous

Actuellement élève ingénieur à Grenoble je suis en stage dans un laboratoire où je dois "développer un environnement d'applications dataflow(systemC) pour cluster de GPU".

Sans vouloir trop rentrer dans les détails, je bloque actuellement sur un code que j'ai codé en systemC et qui 'simule' l'exécution en pipeline de 3 kernels "de type CUDA".

En entrée je considère une matrice d'entier et je vais venir effectuer 3 séries d'opérations successivement sur ses éléments.

Voilà la structure de mon main :

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
#include "systemc.h"
#include "common.h"  
#include "kernel1.h"
#include "kernel2.h"
#include "kernel3.h"
 
 
 
int sc_main (int argc , char *argv[]) 
{
  const sc_time       HLG_PERIOD(10,SC_NS) ;
  sc_clock            HLG("HLG",HLG_PERIOD);
  sc_signal<int**>      F1;
  sc_signal<int**>      F2;
  sc_signal<int**>      F3;
  sc_signal<int**>      F4;
 
 
  sc_report_handler::set_actions("/IEEE_Std_1666/deprecated", SC_DO_NOTHING);
 
//Variables de boucle
  static int i,j;
 
//Définition des 3 kernels
  kernel1 kern1( "Kern1" );
  kernel2 kern2( "Kern2" );
  kernel3 kern3( "Kern3" );
 
  //Matrice de travail
  static int** A;
  A=init_mat(0,N);
  ecrire_mat("pos_init.txt",A,N);  
 
 
 
  F1.write(A);	// Connexion signal 1 
 
  //Premier kernel
  kern1.pos_in1(F1);
  kern1.pos_mod(F2);
  kern1.clk(HLG);
 
  //Deuxième kernel
  kern2.pos_in2(F2);
  kern2.pos_diag(F3);
  kern2.clk(HLG);
 
  //Troisième kernel 
  kern3.pos_in3(F3);
  kern3.pos_nul(F4);
  kern3.clk(HLG);
 
 
  sc_start(HLG_PERIOD); 
 
 
//Debug et écriture dans un fichier
 ecrire_mat("pos_mod.txt",F2,N);
 ecrire_mat("pos_diag.txt",F3,N);
 /*ecrire_mat("pos_nul.txt",F4,N);*/
 
 DEBUG_DATA("Fin de simulation appuyer sur ENTREE :\n");
 getchar();
 return 0;
 
}
Je travaille donc sur des int** et je désire stocker les différents résultats des différents signaux dans des fichiers .txt pour venir les charger/visualiser avec matlab. Mais je n'arrive pas à faire tourner le code correctement. J'ai passé beaucoup de temps à essayer de le déboguer.

Je pense sincèrement que le problème doit être tout bête mais, débutant avec systemC, je n'arrive pas à trouver le pourquoi du comment.
J'utilise 4 signaux pour décrire toute la chaine de traitement de la matrice. F1 en entrée du kernel1 puis ensuite je branche F2 entre la sortie kernel1 et l'entrée kernel2 etc.

Voilà le code de un des kernels (kernel2.h) sachant que les 3 sont assez similaires mis à part l'opération effectuée sur les éléments de la matrice.

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
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
//macro kernel2 : <<< >>> sont des caractères externes  
#define LAUNCHKERNEL2(numBlocks,threadsPerBlock,blockIdx,blockDim,A) for(i=0;i<numBlocks;i++)\
	 {\
		 threadIdx=init_dim3(0,0,0);\
		for(j=0;j<threadsPerBlock.x;j++)\
		 {\
			 for(k=0;k<threadsPerBlock.y;k++)\
			 {\
				m = blockIdx.x*blockDim.x + threadIdx.x;\
				n = blockIdx.y*blockDim.y + threadIdx.y;\
				if(m==n)\
				{\
				A[m][n]+=100;\
				}\
				else\
				{\
				A[m][n]=0;\
				}\
				threadIdx.x++;\
			}\
			 threadIdx.x=0;\
			 threadIdx.y++;\
		 }\
		blockIdx=init_dim3(blockIdx.x++,0,0);\
	 }\
	 pos_diag.write(A)\
 
 
 
 
class kernel2 : public sc_module {
public:
 
    sc_in <int**> pos_in2;
	sc_out <int**> pos_diag;
	sc_in <bool> clk;
 
  SC_HAS_PROCESS(kernel2);
  kernel2(sc_module_name name) : sc_module(name) 
  {
    SC_THREAD(main2);
    sensitive_pos << clk;
  }
 
 
   void  main2()
   { 
	 // Variable de boucle
	 static int i,j,k;
 
	//Variable concernant la grille, les blocs, les threads
	 static dim3 threadsPerBlock;
	 static int numBlocks;
	 static dim3 blockDim;
	 static dim3 blockIdx;
	 static dim3 threadIdx;
	 static dim3 dimGrid;
 
	 //Grandeurs d'entrées sur lesquelles on travaille pour l'instant 
	 static int **A;
	 static int m,n;
 
	//Initialisation
	threadsPerBlock=init_dim3(N,N,0); // On considére des blocs de threads N*N 2D 
	numBlocks=1; // 256 opérations à effectuer sur notre matrice N*N, un bloc est suffisant
	blockIdx=init_dim3(0,0,0); // On va travailler sur le premier bloc 0,0 en 2D 
	blockDim=init_dim3(N,N,0);
	dimGrid=init_dim3(M,M,0);
 
 
 
    while(1)
	{ 
 
		A=pos_in2.read();
		if(A==NULL)
		{
			goto end;
		}
 
        LAUNCHKERNEL2(numBlocks,threadsPerBlock,blockIdx,blockDim,A);
 
		//Debug
		for(i=0;i<N;i++)
		{
		 for(j=0;j<N;j++)
		 {
			 DEBUG_DATA("Kernel2 write pos_diag[%d][%d]=%d\n",i,j,A[i][j]);
		 }
		}
end:
			wait();
	}
 
   }
};

Je travaille sous VisualC++ 2008 et j'ai clairement une erreur lors de l'écriture du signal F4. Les résultats que je trouve pour F2 et F3 ne sont pas corrects, je pense que vu que c'est des signaux je ne copie pas dans les fichiers texte leur état initial ?

Lorsque je mets en commentaire le 3ème Kernel, les 2 permiers s'exécutent parfaitement. Mais je retrouve dans mes fichiers texte le même résultat pour F2 et F3 (résultat de l'opération sur la diagonale effectuée dans le kernel2).

J'espère que quelqu'un pourra m'aider, je ne suis pas très sur de comment allouer de la mémoire pour un signal travaillant sur des matrices d'entier, ça vient peut-être de là ?

En effet dans le thread de chacun de mes 2ème et 3ème kernel je me contente de venir lire l'entrée comme ceci :
Code : Sélectionner tout - Visualiser dans une fenêtre à part
A=pos_in2.read();
ou
Code : Sélectionner tout - Visualiser dans une fenêtre à part
A=pos_in3.read();
Enfin je vous donne le code de mon "common.h" où j'ai mis les fonctions relatives aux différents kernels& matrices

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
#define DEBUG_DATA(...) printf( __VA_ARGS__ );
 
 
#define N 16 // Matrice N*N représentant un bloc de threads
#define M 10 // Matrice M*M représentant une grille de bloc de threads
 
 
struct dim3 {
	unsigned int x;
	unsigned int y;
	unsigned int z;
};
 
 
//Allocation dynamique d'une matrice d'entiers n*n
int ** init_mat(int num,int n)
	{
		int **A;
		int i,j;
		A=(int**)calloc(n,sizeof(int*));		
		for(i=0;i<n;i++)
		{
			A[i] =(int*)calloc(n,sizeof(int));
		}
		for(i=0;i<n;i++)
		{
			for(j=0;j<n;j++)
			{
				A[i][j]=num;
			}
		}
		return A;
}
 
//Initialisation d'un élément dim3
dim3 init_dim3(unsigned int x, unsigned y, unsigned z)
{
	dim3 alpha;
	alpha.x=x;
	alpha.y=y;
	alpha.z=z;
	return alpha;
}
 
 
//Ecrire une matrice n*n dans un fichier texte
void ecrire_mat(char* nom, int **A, int n) 
{
    FILE* fichier = NULL;
	int i,j;
 
    fichier = fopen(nom, "w");
 
    if (fichier != NULL)
    {
 
		for(i=0;i<n;i++)
		{
			for(j=0;j<n;j++)
			{
	  		fprintf(fichier, "%d ",A[i][j]);
			}
			fprintf(fichier,"\n");
		}
		fclose(fichier);
	}
}
Merci d'avance pour votre aide si vous voulez des informations complémentaires je me tiens à votre disposition!

Bien cordialement,

Pascal