CUDA共有メモリがタイル型行列乗算のグローバルメモリよりも遅いのはなぜですか？

Question

私は共有メモリの有無にかかわらず行列乗算コードをタイルしています。以下は、グローバルメモリを使用して行列の乗算である：

__global__ 
void MatrixMulKernel(float* M, float* N, float* P, int Width) 
{ 
int Row = blockIdx.y*blockDim.y + threadIdx.y; 
int Col = blockIdx.x*blockDim.x + threadIdx.x; 
if ((Row < Width) && (Col < Width)) { 
    float Pvalue = 0; 
    for (int k = 0; k < Width; ++k) 
    { 
     Pvalue += M[Row*Width + k] * N[k*Width + Col]; 
    } 

    P[Row*Width + Col] = Pvalue; 
} 
} 

以下は、共有メモリを使用して行列の乗算である：

__global__ 
void MatrixMulKernel(float* d_M, float* d_N, float* d_P, int Width) 
{ 
__shared__ float Mds[blockWidth][blockWidth]; 
__shared__ float Nds[blockWidth][blockWidth]; 
int tx = threadIdx.x; int ty = threadIdx.y; 
int bx = blockIdx.x; int by = blockIdx.y; 

int row = by * blockWidth + ty; 
int col = bx * blockWidth + tx; 
float pvalue = 0; 

for (int m = 0; m < Width/blockWidth; ++m) 
{ 
    Mds[ty][tx] = d_M[row * Width + m*blockWidth + tx]; 
    Nds[ty][tx] = d_N[(m*blockWidth + ty)*Width + col]; 
    __syncthreads(); 
    for (int k = 0; k < blockWidth; ++k) 
    { 
     pvalue += Mds[ty][k]*Nds[k][tx]; 
    } 
    __syncthreads(); 
} 
d_P[row*Width + col] = pvalue; 
} 

同じくらい私は速くする必要があり、共有メモリを使用して知っているが、この2つのコードを比較して、私はコードを見られるように共有メモリなしでは、1600 * 1600のマトリックスでは約2秒高速です。この速度の違いについての説明はありますか、何かが私のコードで間違っていますか？

私の先生は「プログラミング大量並列プロセッサ」ブックをメインのテキストリソースとして使用しています。

EDIT：カーネル用

構成：

int NumBlocks =ceil(Width/blockWidth); // blockWidth = 16 
dim3 dimGrid(NumBlocks, NumBlocks,1); // Width = 1600 
dim3 dimBlock(blockWidth, blockWidth,1); 
clock_t startGpuCalculation = clock(); 
MatrixMulKernel <<<dimGrid, dimBlock >>>(d_M, d_N, d_P, Width); 
cudaThreadSynchronize(); 
clock_t endGpuCalculation = clock(); 

Answer 1

私はデバッグモード（VS 2017 & CUDA 9）でプロジェクトを実行していました。私はリリースモードでコードを実行し、共有メモリはグローバルメモリよりもはるかに高速です。私の悪い。