緩存分塊是一種內(nèi)存優(yōu)化技術(shù)，主要用于提高數(shù)據(jù)的局部性（Locality），以減少緩存未命中（Cache Miss）的次數(shù)。在現(xiàn)代計(jì)算機(jī)體系結(jié)構(gòu)中，處理器（CPU）的速度通常比內(nèi)存快得多。因此，如果CPU在處理數(shù)據(jù)時(shí)需要頻繁地等待數(shù)據(jù)從內(nèi)存中加載，就會(huì)大大降低程序的執(zhí)行效率。Cache Tiled技術(shù)通過將數(shù)據(jù)分割成較小的塊（Tiles），并確保這些小塊能夠完全裝入CPU的高速緩存（Cache），來減少這種等待時(shí)間。

CUDA編程中，用于優(yōu)化內(nèi)存訪問模式，以減少全局內(nèi)存(DRAM)訪問次數(shù)并提高內(nèi)存帶寬的利用率。它的核心思想是將數(shù)據(jù)分成小塊（稱為“tiles”或“blocks”），這樣每個(gè)塊可以完全加載到共享內(nèi)存中。共享內(nèi)存是一種CUDA核心內(nèi)的高速緩存內(nèi)存，其訪問速度比全局內(nèi)存快得多。

基本原理

見啥使用DRAM，也就是全局內(nèi)存。轉(zhuǎn)而多用L1 Cache。緩存分塊是有的時(shí)候數(shù)據(jù)太多了，每次只能加載一部分。

• 減少內(nèi)存延遲：通過將數(shù)據(jù)加載到共享內(nèi)存中，可以減少對全局內(nèi)存的訪問次數(shù)，從而減少延遲。
• 提高內(nèi)存帶寬利用率：將數(shù)據(jù)劃分為小塊后，可以更有效地利用內(nèi)存帶寬。
• 協(xié)同工作：多個(gè)線程可以協(xié)作加載一個(gè)Tile，然后從共享內(nèi)存中高效讀取數(shù)據(jù)。

實(shí)現(xiàn)步驟

1. 定義Tile的大小：確定目標(biāo)內(nèi)存以及GPU的共享內(nèi)存大小。計(jì)算index用于加載到共享內(nèi)存。
2. 加載數(shù)據(jù)到共享內(nèi)存：在CUDA核心中，多個(gè)線程協(xié)作將全局內(nèi)存中的數(shù)據(jù)加載到共享內(nèi)存。
3. 同步線程：確保所有數(shù)據(jù)都加載到共享內(nèi)存后，再進(jìn)行處理。
4. 處理數(shù)據(jù)：從共享內(nèi)存讀取數(shù)據(jù)，進(jìn)行計(jì)算。
5. 將結(jié)果寫回全局內(nèi)存：如果需要，將處理后的數(shù)據(jù)寫回到全局內(nèi)存。

Coding

TILE_WIDTH是一個(gè)預(yù)定義的常量，它定義了Tile的大小。

__syncthreads() 是一個(gè)同步原語，用于確保一個(gè)線程塊內(nèi)的所有線程都達(dá)到這一點(diǎn)后才能繼續(xù)執(zhí)行。這在使用共享內(nèi)存時(shí)尤其重要，因?yàn)樗_保在所有線程開始讀取共享內(nèi)存中的數(shù)據(jù)之前，所有的寫入操作都已完成。

#define TILE_WIDTH  16*16*4  // b c bit 定義每個(gè)Tile的寬度// CUDA核心函數(shù)，用于矩陣乘法
__global__ void MatrixMulKernel(float* Md, float* Nd, float* Pd, int Width) {__shared__ float Mds[TILE_WIDTH][TILE_WIDTH]; // 定義共享內(nèi)存，用于存儲(chǔ)Md的一個(gè)Tile__shared__ float Nds[TILE_WIDTH][TILE_WIDTH]; // 定義共享內(nèi)存，用于存儲(chǔ)Nd的一個(gè)Tileint bx = blockIdx.x;  // 獲取當(dāng)前塊的x坐標(biāo)int by = blockIdx.y;  // 獲取當(dāng)前塊的y坐標(biāo)int tx = threadIdx.x; // 獲取當(dāng)前線程在塊中的x坐標(biāo)int ty = threadIdx.y; // 獲取當(dāng)前線程在塊中的y坐標(biāo)// 計(jì)算Pd矩陣中的行號和列號int Row = by * TILE_WIDTH + ty;int Col = bx * TILE_WIDTH + tx;float Pvalue = 0; // 初始化計(jì)算值// 遍歷Md和Nd矩陣的Tile，計(jì)算Pd矩陣的元素for (int m = 0; m < Width/TILE_WIDTH; ++m) {// 協(xié)作加載Md和Nd的Tile到共享內(nèi)存Mds[ty][tx] = Md[Row*Width + (m*TILE_WIDTH + tx)];Nds[ty][tx] = Nd[Col + (m*TILE_WIDTH + ty)*Width];__syncthreads(); // 確保所有線程都加載完畢// 計(jì)算Tile內(nèi)的乘積并累加到Pvaluefor (int k = 0; k < TILE_WIDTH; ++k) {Pvalue += Mds[ty][k] * Nds[k][tx];}__syncthreads(); // 確保所有線程都計(jì)算完畢}// 將計(jì)算結(jié)果寫入Pd矩陣Pd[Row*Width + Col] = Pvalue;
}

在這個(gè)示例中，MatrixMulKernel 是用于矩陣乘法的CUDA核心。它使用了兩個(gè)共享內(nèi)存數(shù)組Mds和Nds來存儲(chǔ)兩個(gè)輸入矩陣的Tile。每個(gè)線程塊處理輸出矩陣Pd的一個(gè)Tile。線程塊中的每個(gè)線程共同工作，加載輸入矩陣的相應(yīng)部分到共享內(nèi)存，然后使用這些數(shù)據(jù)來計(jì)算輸出矩陣的一個(gè)元素。
__syncthreads() 出現(xiàn)在兩個(gè)關(guān)鍵位置:

1. 加載數(shù)據(jù)到共享內(nèi)存之后：這里的 __syncthreads() 確保了所有線程都完成了對共享內(nèi)存的寫入操作。即使這個(gè)寫入操作是在 for 循環(huán)中完成的，我們也需要確保每個(gè)線程都完成了當(dāng)前迭代的加載操作，才能安全地開始使用這些共享內(nèi)存中的數(shù)據(jù)進(jìn)行計(jì)算。
2. 計(jì)算Tile內(nèi)的乘積并累加到Pvalue之后：第二個(gè) __syncthreads() 確保了所有線程都完成了當(dāng)前Tile的計(jì)算。在開始處理下一個(gè)Tile之前，這是必要的，因?yàn)橄乱粋€(gè)Tile的計(jì)算可能依賴于共享內(nèi)存中的新數(shù)據(jù)。

在這兩種情況下，__syncthreads() 的作用是確保所有線程在繼續(xù)執(zhí)行之前都達(dá)到同一點(diǎn)。

對比原始矩陣乘法的代碼：

__global__ void MatrixMulSimple(float* A, float* B, float* C, 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 += A[Row * Width + k] * B[k * Width + Col];}C[Row * Width + Col] = Pvalue;}
}

變量存儲(chǔ)類別 關(guān)鍵字總結(jié)

用于指定變量的存儲(chǔ)類別，這些關(guān)鍵字決定了變量的存儲(chǔ)位置以及如何在不同線程和線程塊之間共享：

關(guān)鍵字	描述	作用域	生命周期
device	用于在GPU的全局內(nèi)存中聲明變量。	所有線程	應(yīng)用程序執(zhí)行期間
global	用于定義在主機(jī)上調(diào)用但在設(shè)備上執(zhí)行的函數(shù)（即CUDA核心函數(shù)）。	-	-
host	用于定義在主機(jī)上調(diào)用并執(zhí)行的函數(shù)。	-	-
shared	用于聲明位于共享內(nèi)存中的變量。	同一個(gè)線程塊內(nèi)的線程	線程塊的執(zhí)行期間
constant	用于聲明位于常量內(nèi)存中的變量。	所有線程	應(yīng)用程序執(zhí)行期間
managed	用于聲明在主機(jī)和設(shè)備之間共享的統(tǒng)一內(nèi)存變量。	所有線程和主機(jī)	應(yīng)用程序執(zhí)行期間

• __device__：這些變量存儲(chǔ)在設(shè)備的全局內(nèi)存中，可以被所有線程訪問，但訪問延遲較高。
• __global__：定義的是CUDA核心函數(shù)，這種函數(shù)可以從主機(jī)（CPU）調(diào)用并在設(shè)備（GPU）上異步執(zhí)行。
• __host__：定義的是常規(guī)的C++函數(shù)，僅在主機(jī)上執(zhí)行。
• __shared__：聲明的變量位于共享內(nèi)存中，這是一種較快的內(nèi)存類型，但僅在同一個(gè)線程塊內(nèi)的線程之間共享。
• __constant__：用于聲明常量內(nèi)存中的變量，這種內(nèi)存對于所有線程來說是只讀的，訪問速度快，但空間有限。
• __managed__：Unified Memory（統(tǒng)一內(nèi)存）中的變量，可以被GPU和CPU共同訪問，CUDA運(yùn)行時(shí)負(fù)責(zé)管理內(nèi)存的遷移。