CUDA: Multiplication masortingce-masortingce en mosaïque avec mémoire partagée et taille de masortingce non multiple de la taille de bloc

J’essaie de me familiariser avec la programmation CUDA et de passer un bon moment. Je regarde actuellement ce pdf qui traite de la multiplication masortingcielle, faite avec et sans mémoire partagée. Le code complet pour les deux versions peut être trouvé ici . Ce code est presque identique à celui contenu dans les échantillons de multiplication de masortingce CUDA. Bien que la version à mémoire non partagée puisse fonctionner à n’importe quelle taille de masortingce, quelle que soit sa taille, la version à mémoire partagée doit fonctionner avec des masortingces multiples de la taille de bloc (définie à 4, la valeur par défaut étant 16) .

Un des problèmes suggéré à la fin du pdf est de le changer pour que la version à mémoire partagée puisse également fonctionner avec des non-multiples de la taille du bloc. Je pensais que ce serait une vérification d’index simple, comme dans la version non partagée:

int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if(row > A.height || col > B.width) return; 

Mais ça ne marche pas. Voici le code complet, moins la méthode principale (désolé), qui a été quelque peu modifiée par moi:

 void MatMul(const Masortingx A, const Masortingx B, Masortingx C) { // Load A and B to device memory Masortingx d_A; d_A.width = d_A.ssortingde = A.width; d_A.height = A.height; size_t size = A.width * A.height * sizeof(float); cudaError_t err = cudaMalloc(&d_A.elements, size); printf("CUDA malloc A: %s\n",cudaGetErrorSsortingng(err)); err = cudaMemcpy(d_A.elements, A.elements, size, cudaMemcpyHostToDevice); printf("Copy A to device: %s\n",cudaGetErrorSsortingng(err)); Masortingx d_B; d_B.width = d_B.ssortingde = B.width; d_B.height = B.height; size = B.width * B.height * sizeof(float); err = cudaMalloc(&d_B.elements, size); printf("CUDA malloc B: %s\n",cudaGetErrorSsortingng(err)); err = cudaMemcpy(d_B.elements, B.elements, size, cudaMemcpyHostToDevice); printf("Copy B to device: %s\n",cudaGetErrorSsortingng(err)); Masortingx d_C; d_C.width = d_C.ssortingde = C.width; d_C.height = C.height; size = C.width * C.height * sizeof(float); err = cudaMalloc(&d_C.elements, size); printf("CUDA malloc C: %s\n",cudaGetErrorSsortingng(err)); dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); dim3 dimGrid((B.width + dimBlock.x - 1) / dimBlock.x, (A.height + dimBlock.y-1) / dimBlock.y); MatMulKernel<<>>(d_A, d_B, d_C); err = cudaThreadSynchronize(); printf("Run kernel: %s\n", cudaGetErrorSsortingng(err)); // Read C from device memory err = cudaMemcpy(C.elements, d_C.elements, size, cudaMemcpyDeviceToHost); printf("Copy C off of device: %s\n",cudaGetErrorSsortingng(err)); // Free device memory cudaFree(d_A.elements); cudaFree(d_B.elements); cudaFree(d_C.elements); } // Get a masortingx element __device__ float GetElement(const Masortingx A, int row, int col) { return A.elements[row * A.ssortingde + col]; } // Set a masortingx element __device__ void SetElement(Masortingx A, int row, int col, float value) { A.elements[row * A.ssortingde + col] = value; } // Get the BLOCK_SIZExBLOCK_SIZE sub-masortingx Asub of A that is // located col sub-masortingces to the right and row sub-masortingces down // from the upper-left corner of A __device__ Masortingx GetSubMasortingx(Masortingx A, int row, int col) { Masortingx Asub; Asub.width = BLOCK_SIZE; Asub.height = BLOCK_SIZE; Asub.ssortingde = A.ssortingde; Asub.elements = &A.elements[A.ssortingde * BLOCK_SIZE * row + BLOCK_SIZE * col]; return Asub; } // Masortingx multiplication kernel called by MatMul() __global__ void MatMulKernel(Masortingx A, Masortingx B, Masortingx C) { // Block row and column int blockRow = blockIdx.y; int blockCol = blockIdx.x; int rowTest = blockIdx.y * blockDim.y + threadIdx.y; int colTest = blockIdx.x * blockDim.x + threadIdx.x; if (rowTest>A.height || colTest>B.width) return; // Each thread block computes one sub-masortingx Csub of C Masortingx Csub = GetSubMasortingx(C, blockRow, blockCol); // Each thread computes one element of Csub // by accumulating results into Cvalue float Cvalue = 0.0; // Thread row and column within Csub int row = threadIdx.y; int col = threadIdx.x; // Loop over all the sub-masortingces of A and B that are // required to compute Csub // Multiply each pair of sub-masortingces together // and accumulate the results for (int m = 0; m < (BLOCK_SIZE + A.width - 1)/BLOCK_SIZE; ++m) { // Get sub-matrix Asub of A Matrix Asub = GetSubMatrix(A, blockRow, m); // Get sub-matrix Bsub of B Matrix Bsub = GetSubMatrix(B, m, blockCol); // Shared memory used to store Asub and Bsub respectively __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; // Load Asub and Bsub from device memory to shared memory // Each thread loads one element of each sub-matrix As[row][col] = GetElement(Asub, row, col); Bs[row][col] = GetElement(Bsub, row, col); // Synchronize to make sure the sub-matrices are loaded // before starting the computation __syncthreads(); // Multiply Asub and Bsub together for (int e = 0; e < BLOCK_SIZE; ++e) { Cvalue += As[row][e] * Bs[e][col]; } // Synchronize to make sure that the preceding // computation is done before loading two new // sub-matrices of A and B in the next iteration __syncthreads(); } // Write Csub to device memory // Each thread writes one element SetElement(Csub, row, col, Cvalue); } 

choses notables que j’ai modifiées: j’ai ajouté une vérification dans MatMulKernel qui vérifie si notre thread actuel essaie de travailler sur un emplacement du C qui n’existe pas. Cela ne semble pas fonctionner. Bien que cela change le résultat, les modifications ne semblent pas avoir de motif autre que celui des entrées ultérieures (valeur x ou y plus élevée) semble être plus affecté (et j’obtiens beaucoup plus de résultats non entiers). J’ai également changé la méthode de calcul dimGrid donnée et la condition de boucle pour m dans MatMulKernel (avant que ce soit juste la largeur ou la hauteur divisée par la taille du bloc, ce qui semblait faux).

Même le guide de solutions que j’ai trouvé pour ce guide semble suggérer qu’il devrait simplement s’agir d’une simple vérification d’index. Je pense donc qu’il me manque quelque chose de fondamental.

Lorsque les dimensions de la masortingce ne sont pas des multiples des dimensions de la mosaïque, il peut arriver que certaines mosaïques ne couvrent que partiellement les masortingces. Les éléments de mosaïque situés en dehors des mosaïques ne se chevauchant pas complètement doivent être correctement mis à zéro. Ainsi, étendre votre code à des masortingces de taille arbitraire est facile, mais ne constitue pas une simple vérification d’index. Ci-dessous, je copie et colle ma version du kernel de multiplication masortingce-masortingce en mosaïque avec des masortingces de taille arbitraire.

 __global__ void MatMul(float* A, float* B, float* C, int ARows, int ACols, int BRows, int BCols, int CRows, int CCols) { float CValue = 0; int Row = blockIdx.y*TILE_DIM + threadIdx.y; int Col = blockIdx.x*TILE_DIM + threadIdx.x; __shared__ float As[TILE_DIM][TILE_DIM]; __shared__ float Bs[TILE_DIM][TILE_DIM]; for (int k = 0; k < (TILE_DIM + ACols - 1)/TILE_DIM; k++) { if (k*TILE_DIM + threadIdx.x < ACols && Row < ARows) As[threadIdx.y][threadIdx.x] = A[Row*ACols + k*TILE_DIM + threadIdx.x]; else As[threadIdx.y][threadIdx.x] = 0.0; if (k*TILE_DIM + threadIdx.y < BRows && Col < BCols) Bs[threadIdx.y][threadIdx.x] = B[(k*TILE_DIM + threadIdx.y)*BCols + Col]; else Bs[threadIdx.y][threadIdx.x] = 0.0; __syncthreads(); for (int n = 0; n < TILE_DIM; ++n) CValue += As[threadIdx.y][n] * Bs[n][threadIdx.x]; __syncthreads(); } if (Row < CRows && Col < CCols) C[((blockIdx.y * blockDim.y + threadIdx.y)*CCols) + (blockIdx.x * blockDim.x)+ threadIdx.x] = CValue; }