矩阵乘法CUDA

Question

我一直在阅读几个网站，甚至使用NVIDA's代码作为指导，但我仍然得到了错误的答案。主会要求大小的用户，并显示A和B则显示生成的矩阵C.但是说我运行一个2×2矩阵A和B，这是我的样本输出：

Matrix A 
0.000000 8.000000 
2.000000 2.000000 


Matrix B 
3.000000 1.000000 
5.000000 7.000000 


Matrix C (Results) 
0.000000 9.000000 
7.000000 4.000000 

但是，这是不正确。它应该是：

40.000 56.000 
16.000 16.000 

我改变了它从小数到整数，这样它会更容易检查，我发现它不正确。我不明白为什么它不正确，尤其是即使我从代码示例中正确使用它。

#ifndef _MATRIXMUL_KERNEL_H_ 
#define _MATRIXMUL_KERNEL_H_ 

#include <stdio.h> 

// Thread block size 
#define BLOCK_SIZE 16 
#define TILE_SIZE 16 



// CUDA Kernel 
__global__ void matrixMul(float* C, float* A, float* B, int wA, int wB) 
{ 
    // Block index 
    int bx = blockIdx.x; 
    int by = blockIdx.y; 

// Thread index 
int tx = threadIdx.x; 
int ty = threadIdx.y; 

// Index of the first sub-matrix of A processed 
// by the block 
int aBegin = wA * BLOCK_SIZE * by; 

// Index of the last sub-matrix of A processed 
// by the block 
int aEnd = aBegin + wA - 1; 

// Step size used to iterate through the 
// sub-matrices of A 
int aStep = BLOCK_SIZE; 

// Index of the first sub-matrix of B processed 
// by the block 
int bBegin = BLOCK_SIZE * bx; 

// Step size used to iterate through the 
// sub-matrices of B 
int bStep = BLOCK_SIZE * wB; 
float Csub=0; 
// Loop over all the sub-matrices of A and B 
// required to compute the block sub-matrix 
for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
{ 
    // Declaration of the shared memory array As 
    // used to store the sub-matrix of A 
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 

    // Declaration of the shared memory array Bs 
    // used to store the sub-matrix of B 
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; 

    // Load the matrices from global memory 
    // to shared memory; each thread loads 
    // one element of each matrix 
    As[ty][tx] = A[a + wA * ty + tx]; 
    Bs[ty][tx] = B[b + wB * ty + tx]; 

    // Synchronize to make sure the matrices 
    // are loaded 
    __syncthreads(); 

    // Multiply the two matrices together; 
    // each thread computes one element 
    // of the block sub-matrix 
    for (int k = 0; k < BLOCK_SIZE; ++k) 
     Csub += As[ty][k] * Bs[k][tx]; 

    // 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 the block sub-matrix to device memory; 
// each thread writes one element 
int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx; 
C[c + wB * ty + tx] = Csub; 
} 

#endif // #ifndef _MATRIXMUL_KERNEL_H_ 

主机代码：

//perform the calculation 
    //setup execution parameters 
    dim3 threads(BLOCK_SIZE, BLOCK_SIZE); 
    dim3 grid(c.colSize/threads.x, c.rowSize/threads.y); 

    // execute the kernel 
    matrixMul<<< grid, threads >>>(deviceMatrixC, deviceMatrixA, deviceMatrixB, a.colSize, b.colSize); 

感谢您的帮助， 丹

Answer 1

您使用隐式的代码要求矩阵的大小是块大小的圆形倍数（16×16在这种情况下）。内积计算一次处理图块宽度，而不检查存储器访问是否超出边界。出于这个原因，2x2矩阵将无法工作。

如果您尝试运行带有16x16输入的内核（例如将您的2x2外壳填充为16x16），则应该能够确认结果。