Discussion :

[pepito3364]

Je cherche a trouver les limites sur le nombre de block et de thread que je peux utiliser pour un calcul (ici multiplication matricielle).

Actuellement, je teste dans l'ordre plusieurs choses avant d'appeler le kernel:
- la taille mémoire de mes tableaux VS la ram max de la carte
- la bonne allocation mémoire des objets de cuda
- le nombre de block x et y VS la taille max possible
- le nombre de thread par block VS le nombre max

Si toute ces conditions sont ok je lance le calcul.
A l'intérieur du kernel je test aussi les valeurs des indices par rapport à la taille des matrices.

Sauf que pour des matrices 2000 par 2000 je récupère un écran noir.

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
 
__global__ void MatrixMulKernel2(double * m1, double * m2, double * Pd, int m, int n, int p)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
 
	if (i < m)
	{
		int j = blockIdx.y * blockDim.y + threadIdx.y;
		if (j < p)
		{
			double Pvaleur = 0;
 
			for (int k = 0; k < n; ++k)
				Pvaleur  += m1[k  * m + i] * m2[j * n + k];
 
			Pd[j * m + i] = Pvaleur;
		}
	}
}
 
 
int MatrixMulOnDevice(cudaDeviceProp &cDP, double * matrice1, double * matrice2, double * matriceResult, int m, int n, int p)
{
	//calcul de la taille des matrices
	double *cuda_matrix1;
	double *cuda_matrix2;
	double *cuda_matrixResult;
 
	int size1 = sizeof(double) * m * n;
	int size2 = sizeof(double) * n * p;
	int size3 = sizeof(double) * m * p;
 
	cudaError_t cudaStatus = cudaMalloc((void**)&cuda_matrix1, size1);
	if (cudaStatus != cudaSuccess)
		return cudaStatus;
	cudaStatus = cudaMalloc((void**)&cuda_matrix2, size2);
	if (cudaStatus != cudaSuccess)
		return cudaStatus;
	cudaStatus = cudaMalloc((void**)&cuda_matrixResult, size3);
	if (cudaStatus != cudaSuccess)
		return cudaStatus;
 
	cudaStatus = cudaMemcpy(cuda_matrix1, matrice1, size1, cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess)
		return cudaStatus;
	cudaStatus = cudaMemcpy(cuda_matrix2, matrice2, size2, cudaMemcpyHostToDevice);
	if (cudaStatus != cudaSuccess)
		return cudaStatus;
 
	int BLOCK_WIDTH = 16;
 
	dim3 dimGrid((m / BLOCK_WIDTH) + ( (m % BLOCK_WIDTH) == 0 ? 0 : 1), (p / BLOCK_WIDTH) + ( (p % BLOCK_WIDTH) == 0 ? 0 : 1));
 
	if (dimGrid.x < cDP.maxGridSize[0] && dimGrid.y < cDP.maxGridSize[1])
	{
		dim3 dimBlock(BLOCK_WIDTH, BLOCK_WIDTH);
		if (BLOCK_WIDTH * BLOCK_WIDTH < cDP.maxThreadsPerBlock)
		{
			MatrixMulKernel2<<<dimGrid, dimBlock>>>(cuda_matrix1, cuda_matrix2, cuda_matrixResult, m, n, p);
			cudaMemcpy(matriceResult, cuda_matrixResult, size3, cudaMemcpyDeviceToHost);
		}
		else
			return -2;
	}
	else
		return -1;
 
	cudaFree(cuda_matrix1);
	cudaFree(cuda_matrix2);
 
	cudaFree(cuda_matrixResult);
 
	return 0;
}

Dans mon exemple, j'arrive avec dimGrid = (125, 125, 1).

J'ai plusieurs questions:
- Est-ce que je test mal des choses ou j'en oublie certaines?
- Comment récupèrer le nombre max de block?
- Si mes matrices sont trop grosses pour la carte, faut-il appeler le kernel dans une boucle pour n'envoyer que des "petits" morceaux?

Merci de votre aide, je débute avec cuda et ça à l'air vraiment intéressant.

[pepito3364]

J'ai finalement trouvé ce qui provoqué l'écran noir.
C'est le timeout d'un wtachdog qui fait que au bout d'environ 5s de calcul, la carte est coupé.

Il n'est pas conseillé d'utiliser cuda lorsque la carte graphique est utilisé pour l'affichage principal.