行列の乗算CUDA

Question

私はいくつかのウェブサイトから読んでおり、NVIDA'sコードをガイドとして使用していますが、私はまだ間違った答えを得ています。

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 

しかし、それは間違っています。メインは、しかし、これは私のサンプル出力であるIがAとBの両方のために2×2行列を実行すると言うその後、結果の行列Cを表示するサイズをユーザに求めます、そしてAとBが表示されます。それは次のようになります。

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にゼロパディングするなど）、結果を確認できるはずです。