上篇的方法是在全局存儲區中，這樣對取數據時速度回很慢，影響性能，而設備中線程對塊中的共享存儲區中數據讀取時速度是很快的，並且在全局存儲區中進行讀取時，有很多數組元素的重復讀取。因此，先將需要計算的數組數據讀取到共享存儲區中，再利用共享存儲區中的數據進行計算，就會提高性能。
但由於每個塊的共享存儲區的存儲空間一般很小，以本人8400MG為例，只有16KB，因此在一個塊內需要的數據量大時，有必要對數據進行分塊，分塊進行計算。
1.分塊策略
比如上篇中的：
網格維度：(width/TILE_WIDTH,width/TILE_WIDTH) (64,64)
塊維度：(TILE_WIDTH,TILE_WIDTH) (16,16)
計算如圖中Pd中的一個塊需要Md和Nd虛線包圍的數據。Pd一個塊中的所需的數據大小：TILE_WIDTH*width*2*4/1024 (KB) 2因為Md和Nd兩個數組，4因為float。計算得128MB。顯然大大超過了16KB，這時就需要采用分塊的方式計算。(注意：分塊計算的方式在大數據處理中是很常用的方法)
分塊大小：可以先嘗試嘗試，就以TILE_WIDTH*TILE_WIDTH的矩形塊為一小塊，進行分塊分階段計算。計算需要的存儲器大小為16*16*2*4/1024=2KB，可以，根據上圖也很清晰，思路明了(這是關鍵)。
2.源程序
__shared__關鍵字說明變量存儲在共享存儲區內，共享存儲區內的數據對塊內的線程是共享的。在線程內申明的變量作用范圍知識線程。
每個塊中含有TILE_WIDTH*TILE_WIDTH個線程，而每個小塊中Md和Nd中的數據元素為TILE_WIDTH*TILE_WIDTH，因此塊中的一個線程將Md和Nd中的一個元素加載進共享存儲器。如下代碼第21、22行所示(tx,ty)處的線程分別將Md和Nd中的一個元素加載進共享存儲器中。
采用此方法，kernel函數就要改變，如下：
1 __global__ static void MatrixMulKernel(const float* Md,const float* Nd,float* Pd,int Width) 2 { 3 //共享存儲器保存從從全局存儲器中加載的數據 4 __shared__ float Mds[TILE_WIDTH][TILE_WIDTH]; 5 __shared__ float Nds[TILE_WIDTH][TILE_WIDTH]; 6 7 //計算Pd和Md中元素的行索引 8 int bx=blockIdx.x; 9 int by=blockIdx.y; 10 int tx=threadIdx.x; 11 int ty=threadIdx.y; 12 //Pd的行和列 13 int Row = by*TILE_WIDTH+ty; 14 int Col = bx*TILE_WIDTH+tx; 15 16 float Pvalue = 0.0; 17 //第k小塊 18 for (int k=0;k<Width/TILE_WIDTH;k++) 19 { 20 //通過協作把Md和Nd的塊加載到共享存儲器中 21 Mds[ty][tx]=Md[Row*Width+k*TILE_WIDTH+tx]; 22 Nds[ty][tx]=Nd[(k*TILE_WIDTH+ty)*Width+Col]; 23 __syncthreads(); //等待塊中其他線程同步 24 25 for(int m=0;m<TILE_WIDTH;m++) 26 Pvalue +=Mds[ty][m]*Nds[m][tx]; 27 __syncthreads(); //等待其他線程計算完，因為Pvalue要用到下一個塊的計算 28 } 29 //每個線程負責計算P中的一個元素 30 Pd[Row*Width+Col]=Pvalue; 31 }
里面比較重要的是線程同步，用__syncthreads()函數，完成一個塊內線程的同步的功能，因為要讓線程將需要的數據都加在進共享存儲區內才能進行計算。
3.測試結果
CPU計算比較耗時，就不測試了，可以看到采用此方法比上篇中的性能要提高了不少。
4.不同分配策略比較
上面只是用到了2KB的共享存儲器，可以考慮將分塊數據更大一點，即將上面Md和Nd中四個小塊合並為一個小塊
Md一個小塊大小：(TILE_WIDTH,TILE_WIDTH*4)
Nd一個小塊的大小：(TILE_WIDTH*4,TILE_WIDTH)
上面代碼就要進行些變化，主要是共享存儲器分配大小以及循環變量的循環條件。
測試結果如下：
可見比前面要慢些。這里還可以利用其它分塊策略進行測試，從而找到最適合的分塊策略，分塊是同時也要兼顧物理存儲限制及存儲器訪問模型。