CODEWITHSUNDEEP
matrixcudamatrix-multiplication
user990246
user990246

Reputation:

Matrix multiplication with CUDA using multiple columns

The code below computes C=A*B, how would I expand it to do multi tiled multiplication? One row of A is multiplied with one column of B on every one loop. My question is how would I modify it such that one row of A is multiplied with multiple columns of B so that we avoid reloading them again for the same row.

(code taken from: Dynamic matrix multiplication with CUDA)

// CUDA Kernel
__global__ void matrixMul( float* C, float* A, float* B, int wA, int wB,size_t block_size)
 {
int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;

int aBegin = wA * block_size * by;
int aEnd   = aBegin + wA - 1;
int aStep  = block_size;

int bBegin = block_size * bx;

int bStep  = block_size * wB;
float Csub=0;

for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
{
    extern __shared__ float As[];
    extern __shared__ float Bs[];
    extern __shared__ float smem[];

    smem[ty*block_size+tx] = A[a + wA * ty + tx];

    smem[block_size*block_size+ty*block_size+tx]  = B[b + wB * ty + tx];

    __syncthreads();

    for (int k = 0; k < block_size; ++k)
        Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ;

    __syncthreads();
}

int c = wB * block_size * by + block_size * bx;
C[c + wB * ty + tx] = Csub;
}

Thanks

Upvotes: 1

Views: 1538

Answers (1)

Brendan Wood
Brendan Wood

Reputation: 6440

There is a book by Kirk and Hwu that goes into great depth with developing an efficient matrix multiplication kernel:

http://www.amazon.com/Programming-Massively-Parallel-Processors-Hands/dp/0123814723

There have also previously been several questions here on Stack Overflow regarding tiled matrix multiplication on CUDA. See the following:

Matrix Multiplication CUDA

Non Square Matrix Multiplication in CUDA

Upvotes: 1

Related Questions