CUDA学習の二:shared_memory使用、マトリクス乗算

CUDAでshared_を使用memoryは演算を加速することができ,行列乗算では体現される.
行列C=A*B,通常の演算ではC[i,j]=A[i,:*B[:,j]を用いて結果を算出する.しかし、CPU上でこの演算を完了するには、A[m,n],B[n,k]を設定すると、Cマトリクスはm*kであり、全体として、m*n*k乗算、m*(b-1)*k加算、シリアル実行が必要であり、全体の複雑度はO(m*n*k)である.
マトリックスクラス:

1 class Matrix
2 {
3 public:
4     int cols;   // x
5     int rows;   // y
6     float *data;  //  ，    
7 }

CPU上のプログラム、1つの3層の循環

for(int i =0;i< C.rows;i++)
    {
        for(int j =0;j< C.cols;j++)
        {
            float *a = A.data;
            float *b = B.data;
            for(int k=0;k<A.cols;k++)
                C.data[i*C.cols+j]+=a[i*A.cols+k] * b[k*B.cols+j];
        }
    }
}

GPUで加速し、CUDAで実現したいと思います.kernelと書きます.

__global__ void matrixMulKernel(const Matrix A, const Matrix B, Matrix C)
{  
    // Each thread computes one element of C  
    // by accumulating results into Cvalue  
    float Cvalue = 0;  
    int row = blockIdx.y * blockDim.y + threadIdx.y;  
    int col = blockIdx.x * blockDim.x + threadIdx.x;  
    for (int e = 0; e < A.cols; ++e)  
    Cvalue += A.data[row * A.cols + e]* B.data[e * B.cols + col];  
    C.data[row * C.cols + col] = Cvalue;  
}  

この場合、計算プロセスは並列であるが、A、Bマトリクスにアクセスする場合、同時にアクセスできないため、主な時間はメモリ読み取りに費やされ、各スレッドはAの1行、Bの1列を読み取り、Cの対応値を計算する.したがってglobal memoryからn回A,m回Bを読む必要がある.時間複雑度はO(m+n)次メモリアクセスとk次乗算である.
 
実際にはshared memoryを使う方法もありますが、ここではA、Bマトリクスをblocksizeに従ってサブマトリクスsubA[blocksize][blocksize]、subB[blocksize][blocksize]に分けます.サブマトリクスを__に設定shared__. thread block内のすべてのthreadsはshared memoryを共用している(読み書き可能).するとAはglobal memoryからn/block_しか読まなかったsize回、Bはm/blockしか読めませんでした_size次;時間的複雑さは、O(m/block_size+n/block_size)二次メモリアクセス、およびk次乗算である.時間の複雑さをさらに減らす.コードは次のとおりです.

__global__ void matrixMulKernel(const float *A, const float *B, float *C,int Aw ,int Bw)
{
    const int bs = CUDA_LG::block_size;
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int bx = blockIdx.x;
    int by = blockIdx.y;

    int aBlockFisrt = by * bs * Aw ; 
    int aBlockStep  = bs ; 
    int aBlockLast  = by * bs * Aw  + Aw - 1 ; 
    int bBlockFisrt = bx * bs ;
    int bBlockStep  = bs * Bw ;
    
    float subC=0;

    for(int a = aBlockFisrt,int b = bBlockFisrt; a <= aBlockLast ;a+=aBlockStep,b+=bBlockStep )
    {
        //    shared memory    
        __shared__ float  subA[bs][bs];
        __shared__ float  subB[bs][bs];

        subA[ty][tx] = A[a + ty * Aw + tx];
        subB[ty][tx] = B[b + ty * Bw + tx];

        __syncthreads(); 

        for(int i = 0;i<bs;i++)
        {
            subC += subA[ty][i] * subB[i][tx];
        }    

        __syncthreads();
    }
    C[ by*bs*Bw + bx*bs + ty * Bw +tx] = subC;

}

 
参照sample_6.5\0_SimplematrixMulプログラム.コメント詳細
Rachel zhangのブログCUDA学習シリーズの2つを参照してください.http://blog.csdn.net/abcjennifer/article/details/42528569