__shared__如何在下面的代碼中工作？

Question

我不是9號和10號線;該指數通過公式Col + (m*TILE_WIDTH + ty)*Width使用和計算。

有人可以幫助我理解此代碼，即使用__shared__？

__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) 
{ 
__shared__float Mds[TILE_WIDTH][TILE_WIDTH]; 
__shared__float Nds[TILE_WIDTH][TILE_WIDTH]; 
3. int bx = blockIdx.x; int by = blockIdx.y; 
4. int tx = threadIdx.x; int ty = threadIdx.y; 
// Identify the row and column of the Pd element to work on 
5. int Row = by * TILE_WIDTH + ty; 
6. int Col = bx * TILE_WIDTH + tx; 
7. float Pvalue = 0; ; 
// Loop over the Md and Nd tiles required to compute the Pd element 
8. for (int m = 0; m < Width/TILE_WIDTH; ++m) { 
// Coolaborative loading of Md and Nd tiles into shared memory 
9.Mds[ty][tx] = Md[Row*Width + (m*TILE_WIDTH + tx)]; 
10.Nds[ty][tx] = Nd[Col + (m*TILE_WIDTH + ty)*Width]; 
11.__syncthreads(); 
11. for (int k = 0; k < TILE_WIDTH; ++k) 
12.Pvalue += Mds[ty][k] * Nds[k][tx]; 
13. Synchthreads(); 
    } 
    Pd[Row*Width+Col] = Pvalue; 
    } 

Answer 1

__shared__內存是GPU的快速（但很小）的片上資源。

要乘以的矩陣從全局存儲器中開始（Md和Nd）。線10和11：

Mds[ty][tx] = Md[Row*Width + (m*TILE_WIDTH + tx)]; // line 10 
Nds[ty][tx] = Nd[Col + (m*TILE_WIDTH + ty)*Width]; // line 11 

每一個矩陣的「瓦片」（方形子部分）要被相乘（或者Md或Nd）負載到一個共享存儲器複製（Mds或Nds）。單行代碼可以加載整個「tile」的原因是因爲threadblock的所有線程都執行該行代碼。結果，數據的線程塊大小「塊」或「塊」從全局移動到共享內存。

一旦它在共享內存中，實際乘法將在第14行完成。由於第14行在共享內存中運行，而不是全局內存，並且由於塊中相鄰線程之間存在數據重用，因此總體乘法操作運行速度更快，因爲可以比全局內存更快地訪問共享內存。

提供了類似的代碼和描述in the programming guide。