AI部署篇 | CUDA學(xué)習(xí)筆記2：矩陣乘法與GPU優(yōu)化(附CUDA代碼)

1學(xué)習(xí)筆記2——矩陣相乘與共享內(nèi)存

1、矩陣乘法 CPU 實(shí)現(xiàn)

CPU程序通過三層循環(huán)實(shí)現(xiàn)：

void?matrixMulCpu(float*?A,?float*?B,?float*?C,?int?width){
????float?sum?=?0.0f;
????for(int?i?=?0;?i?????????for(int?j?=?0;?j?????????????for(int?l?=?0;?l?????????????????sum?+=?A[i?*?width?+?l]?*?B[l?*?width?+?j];
????????????}
????????????C[i?*?width?+?j]?=?sum;
????????????sum?=?0.0f;
????????}
????}
}

• 通過上面CPU代碼的實(shí)驗(yàn)觀察可以看出，總共的計(jì)算次數(shù)為：

• 時(shí)間復(fù)雜度為：

2、GPU實(shí)現(xiàn)矩陣乘法

獲得 C 矩陣的計(jì)算方法都是相同的，只不過使用的是矩陣 A、B 不同的元素來進(jìn)行計(jì)算，即不同數(shù)據(jù)的大量相同計(jì)算操作，這種計(jì)算是特別適合使用GPU來計(jì)算，因?yàn)镚PU擁有大量簡單重復(fù)的計(jì)算單元，通過并行就能極大的提高計(jì)算效率。

在 GPU 中執(zhí)行矩陣乘法運(yùn)算操作：

1. 在 Global Memory 中分別為矩陣 A、B、C 分配存儲(chǔ)空間；
2. 由于矩陣 C 中每個(gè)元素的計(jì)算均相互獨(dú)立，NVIDIA GPU 采用的 SIMT (單指令多線程)的體系結(jié)構(gòu)來實(shí)現(xiàn)并行計(jì)算的, 因此在并行度映射中，讓每個(gè) thread 對應(yīng)矩陣 C 中1個(gè)元素的計(jì)算；
3. 執(zhí)行配置 (execution configuration)中 gridSize 和 blockSize 均有 x(列向)、y(行向)兩個(gè)維度，其中，

CUDA的kernel函數(shù)實(shí)現(xiàn)如下：

每個(gè) thread 需要執(zhí)行的 workflow 為：

• 從矩陣 A 中讀取一行向量 (長度為width) ==> A[Row * width + i]
• 從矩陣 B 中讀取一列向量 (長度為width（圖中為height）) ==> B[i * width + Col]
• 對這兩個(gè)向量做點(diǎn)積運(yùn)算 (單層 width 次循環(huán)的乘累加)==> A[Row * width + i] * B[i * width + Col]
• 最后將結(jié)果寫回矩陣 C。==> C[Row * width + Col] = Pervalue

//核函數(shù)的具體實(shí)現(xiàn)
__global__?void?matMul_GlobalKernel(int?*A,int?*B,int?*C,int?width){
???int?bx?=?blockIdx.x;
???int?by?=?blockIdx.y;
???int?tx?=?threadIdx.x;
???int?ty?=?threadIdx.y;

???int?Col?=?bx?*?blockDim.x?+?tx;
???int?Row?=?by?*?blockDim.y?+?ty;

???int?perValue?=?0;
???for(int?i?=?0;?i????????perValue?+=?A[Row?*?width?+?i]?*?B[i?*?width?+?Col];
???}
???C[Row?*?width?+?Col]?=?Pervalue;
}

下面來分析一下該 kernel 函數(shù)中 A、B、C 三個(gè)矩陣對 Global memory 的讀取和寫入情況：

讀取 Global Memory：

• 對于矩陣 C 中每一個(gè)元素計(jì)算，需要讀取矩陣 A 中的一行元素；

  對于矩陣 C 中同一行的 width 個(gè)元素，需要重復(fù)讀取矩陣 A 中同一行元素 width 次；

• 對于矩陣 C 中每一個(gè)元素計(jì)算，需要讀取矩陣 B 中的一列元素；

  對于矩陣 C 中同一列的 width 個(gè)元素，需要重復(fù)讀取矩陣 B 中同一列元素 width 次；

寫入 Global Memory：

• 矩陣 C 中的所有元素只需寫入一次

由此可見：

• 對 A 矩陣重復(fù)讀取 width 次，共計(jì) 次 32 bit Global Memory Load操作；
• 對 B 矩陣重復(fù)讀取 width 次，共計(jì) 次32 bit Global Memory Load操作；
• 對 C 矩陣共計(jì) 次 32 bit Global Memory Store操作。

在上述分析中是將每一個(gè) 32 bit 元素 (或者說每個(gè)thread)對 Global memory 的訪問 (access)獨(dú)立對待的；但實(shí)際情況是如此嗎？

假設(shè)矩陣規(guī)模為 width=32，執(zhí)行配置 blockSize=(32, 32, 1)，gridSize=(1, 1, 1)，使用上述的 Kernel 函數(shù)進(jìn)行計(jì)算，在 NVVP 中 Memory Bandwidth Analysis 結(jié)果如下：

按照前面的計(jì)算方式，Global Memory Load 次數(shù)為次，Store 次數(shù)為 1024 次；而 NVVP 顯示的讀取次數(shù)為 5120 次；寫入次數(shù)為 128 次, 這到底是為什么呢？

別有洞天之 Warp

GPU 編程中最重要的概念之一是 warp，每個(gè) warp 包含 32 個(gè) thread，而 GPU 的指令發(fā)射是以 warp 為最小單元的。當(dāng) warp 中的一條指令發(fā)射一次，稱為 1 次 “transaction”，重復(fù)發(fā)射一次, 稱為 1 次 “reply”。

對于 Global Memory 的訪問，warp 內(nèi)的指令需要幾次 transaction，或者說是否會(huì)發(fā)生 reply，取決于地址對齊及可合并訪問的情況。

Global Memory 的讀/寫訪問均是以 32 Byte 為單元的，稱為 1 個(gè) segment，即 1 transaction 可訪問 32 Byte 數(shù)據(jù) (假設(shè) L1 cache 為 non-caching)。假設(shè) 1 個(gè) warp 中的每個(gè) thread 需要訪問 1 個(gè) 32 bit (=4 Byte)數(shù)據(jù)，且訪問地址是 32 Byte 對齊的，則總共需要訪問 4 個(gè) segments，即產(chǎn)生 4 次 transaction (即 3 次 reply)指令發(fā)射。

接下來重新分析矩陣乘法中Global Memory訪問的情況：

• Global Memory Load：對于 1 個(gè) warp 中的 32 個(gè) thread，在每 1 次循環(huán)中，需要讀取矩陣 A 同一個(gè)元素 (1 次 transaction)，以及矩陣 B 連續(xù)的 32 個(gè)元素 (假設(shè)是理想的可合并訪問的，至少需要 4 次 transaction)，共發(fā)生 5 次 transaction (注意，并不是前文的 32+32 次)。K 次循環(huán)總共需要 k×5次 transactions。對于 ? ?個(gè) thread, 共有 個(gè) warp，總共的 Global Memory Load Transaction 數(shù)目為： (注意，并不是前文的 ).
• Global Memory Store：矩陣 C 的寫入過程中，每個(gè) warp 中的 32 thread 可連續(xù)寫入 32 個(gè) 32 bit 元素 (4次 transaction)，對于個(gè) thread，共有個(gè) warp，總共的 Global Memory Store Transaction 數(shù)目為：次 transaction。

對于前面驗(yàn)證矩陣, 其規(guī)模為width=32, 執(zhí)行配置blockSize=(32,32,1), gridSize=(1,1,1),

• Global Memory Load Transaction數(shù)目為: width×width÷32×width×5=32×32÷32×32×5=5120

分析結(jié)果與 NVVP 中 GPU 實(shí)際執(zhí)行結(jié)果是完全吻合的。

width*width*width\Blocksize 對計(jì)算性能的影響

結(jié)果分析：

• 隨著矩陣規(guī)模增大，計(jì)算性能不斷提升，到達(dá)峰值后又略有下降。在矩陣規(guī)模較小時(shí)，由于block數(shù)量不夠多，無法填滿所有SM單元，此時(shí)的性能瓶頸為Latency Bound（由于低Occupancy導(dǎo)致GPU計(jì)算資源的利用率低，延遲無法被很好的隱藏）；隨著矩陣規(guī)模增加，block數(shù)量增加，每個(gè)SM中的active warps數(shù)量隨之增大，此時(shí)Latency不再是性能瓶頸，轉(zhuǎn)而受限于Memory Bound（過多的高延遲、低帶寬的全局內(nèi)存訪問），在無法提升memory訪問效率的情況下，性能無法進(jìn)一步提升；

• 不同的blockSize對性能的影響不大（這里僅限于討論8*8、16*16、32*32三種情況）。究其原因，是因?yàn)檫x擇的幾種block維度設(shè)計(jì)（每行分別有8/16/32個(gè)thread），對1個(gè)warp內(nèi)訪問Global Memory時(shí)（Load或Store）transaction的數(shù)量沒有變化。

3、Shared Memory 優(yōu)化矩陣乘法

雖然 warp 內(nèi)對 Global Memory 的訪問均已最大的實(shí)現(xiàn)了合并訪問，但在 A、B 矩陣的讀取操作中仍然有很多重復(fù)訪問，例如：

對于矩陣 A 的讀取操作，通過合并訪問（32 個(gè) thread 訪問 Global Memory 的同一地址，合并為一次訪問），實(shí)際重復(fù)讀取次數(shù)是(n/32)；

對于矩陣 B 的讀取操作，通過合并訪問（8 個(gè) thread 訪問 32 Byte 數(shù)據(jù)可合并為一次），實(shí)際重復(fù)讀取次數(shù)為(m/8)次。

在不改變這種數(shù)據(jù)讀取方式的前提下又如何優(yōu)化性能呢？

在GPU中除了 Global Memory 還有 Shared Memory，這部分 Memory 是在芯片內(nèi)部的，相較于 Global Memory 400~600 個(gè)時(shí)鐘周期的訪問延遲，Shared Memory 延時(shí)小 20-30 倍、帶寬高 10 倍，具有低延時(shí)、高帶寬的特性。因此性能優(yōu)化的問題可以轉(zhuǎn)變?yōu)槿绾卫?Shared Memory 代替 Global Memory 來實(shí)現(xiàn)數(shù)據(jù)的重復(fù)訪問。

如上圖所示，使用 Shared Memory 優(yōu)化 Global Memory 訪問的基本思想是充分利用數(shù)據(jù)的局部性。

讓一個(gè) block 內(nèi)的 thread 先從 Global Memory 中讀取子矩陣塊數(shù)據(jù)（大小為 BLOCK_SIZE × BLOCK_SIZE）并寫入 Shared Memory 中；在計(jì)算時(shí)，從 Shared Memory 中（重復(fù)）讀取數(shù)據(jù)做乘累加，從而避免每次都到 Global 中取數(shù)據(jù)帶來的高延遲影響。接下來讓子矩陣塊分別在矩陣 A 的行向以及矩陣 B 的列向上滑動(dòng)，直到計(jì)算完所有 width 個(gè)元素的乘累加。使用 Shared Memory 優(yōu)化后的 kernel 代碼如下所示：

//核函數(shù)的具體實(shí)現(xiàn)
__global__?void?matmul_ShareMemory(int?*M,int?*N,int?*P,int?width){
????__shared__?float?Mds[BLOCK_SIZE][BLOCK_SIZE];
????__shared__?float?Nds[BLOCK_SIZE][BLOCK_SIZE];

????int?bx?=?blockIdx.x;
????int?by?=?blockIdx.y;
????int?tx?=?threadIdx.x;
????int?ty?=?threadIdx.y;

????int?Col?=?bx?*?BLOCK_SIZE?+?tx;
????int?Row?=?by?*?BLOCK_SIZE?+?ty;

????int?Pervalue?=?0;
????//有多少個(gè)BLOCK_SIZE，每個(gè)循環(huán)計(jì)算一個(gè)塊的大小
????for(int?i?=?0;i?????????Mds[ty][tx]?=?M[Row?*?width?+?(i?*?BLOCK_SIZE?+?tx)];
????????Nds[ty][tx]?=?N[Col?+?(i?*?BLOCK_SIZE?+?ty)?*?width];
????????__syncthreads();

????????//BLOCK_SIZE相乘
????????for(int?k?=?0;k?????????????Pervalue?+=?Mds[ty][k]?*?Nds[k][tx];
????????__syncthreads();
????}
????P[Row?*?width?+?Col]?=?Pervalue;
}

• 1 個(gè) block 即看做 1 個(gè)子矩陣 C，且為方陣；
• 讀取子矩陣 A 和子矩陣 B 的 Shared Memory 的大小均等于子矩陣 C 的維度大?。?/section>
• 子矩陣 A 在矩陣 A 的行向上移動(dòng) width/BLOCK_SIZE 次，子矩陣 B 在矩陣 B 的列向上移動(dòng) width / BLOCK_SIZE 次；
• 每個(gè) thread 的計(jì)算過程，由原來的單層 k 次循環(huán)，變?yōu)榱藘蓪友h(huán)：外層循環(huán)次數(shù)為 width / BLOCK_SIZE（假設(shè)能夠整除），其任務(wù)是從 Global Memory 讀取數(shù)據(jù)到 Shared Memory；內(nèi)存循環(huán)次數(shù)為 BLOCK_SIZE，其任務(wù)是讀取 Shared Memory 中的數(shù)據(jù)做乘累加計(jì)算；
• 有 2 次 __syncthreads()操作：第一次表示同步 Shared Memory 中數(shù)據(jù)寫入之后，在進(jìn)入計(jì)算之前，保證block內(nèi)所有Shared Memory中數(shù)據(jù)已更新；第二次表示同步計(jì)算之后以及 Shared Memory 寫入之前，保證 block 內(nèi)所有 thread 的計(jì)算已完成，可以進(jìn)行 Shared Memory 的數(shù)據(jù)更新。

Shared Memory 性能分析

• 隨著矩陣規(guī)模增大，計(jì)算性能不斷提升，到達(dá)峰值后又略有下降——這與未優(yōu)化前使用 Global Memory 時(shí)的性能分析結(jié)果一致；

• 不同 blockSize 對性能影響很大。外層循環(huán)中從 Global Memory 讀取數(shù)據(jù)寫入到 Shared Memory 時(shí)，無論是讀取 A 矩陣或 B 矩陣，當(dāng) warp 的組織形式（行x列）為 8x4 或 16x2，對于 Global Memory 的 load 操作，均可實(shí)現(xiàn) 32B 的合并訪問（8 thread x 4B 及 4 次 transactions）。而對于 Shared Memory 的 Store 操作，則會(huì)出現(xiàn) Bank Conflict 導(dǎo)致 reply 的發(fā)生；同樣地，內(nèi)層循環(huán)中讀取 Shared Memory 時(shí)，當(dāng) warp 的組織形式為 8x4 或 16x2 時(shí)，則會(huì)出現(xiàn) Bank Conflict，導(dǎo)致 Shared memory 讀取時(shí)的 reply，從而影響性能。

4、Register 優(yōu)化矩陣乘法

前面的算法設(shè)計(jì)中，每個(gè)線程只計(jì)算了矩陣 C 中的一個(gè)元素，每個(gè)線程每個(gè)內(nèi)層循環(huán)需要從子矩陣 A 和子矩陣 B 中各讀取一個(gè) 4 Byte 的元素（共取 8 Byte 數(shù)據(jù)執(zhí)行2次浮點(diǎn)運(yùn)算），實(shí)際上可以讓每個(gè)線程讀取一組 Shared Memory 數(shù)據(jù)后（放入寄存器中），計(jì)算更多的元素，從而減少 Shared Memory 的訪問。

//?using?ILP?2?to?improve?the?performance
__global__?void?matrixMulSharedILPkernel(float*?A,?float*?B,?float*?C,?int?width){
????int?row?=?blockIdx.y?*?blockDim.y?*?2?+?threadIdx.y;
????int?col?=?blockIdx.x?*?blockDim.x?+?threadIdx.x;
????float?val[2]?=?{0.0f};

????__shared__?float?shTileA[BLOCK_SIZE][BLOCK_SIZE];
????__shared__?float?shTileB[BLOCK_SIZE][BLOCK_SIZE];

????int?iter?=?(width?+?BLOCK_SIZE?-?1)?/?BLOCK_SIZE;
????for(int?i?=?0;?i?????????//?read?data?from?global?memory?to?shared?memory
????????shTileA[threadIdx.y][threadIdx.x]=A[row?*?width+i*BLOCK_SIZE+threadIdx.x];
????????shTileA[threadIdx.y+16][threadIdx.x]=A[(row+16)*width+i*BLOCK_SIZE+threadIdx.x];
????????shTileB[threadIdx.y][threadIdx.x]=B[(i*BLOCK_SIZE+threadIdx.y)*width+col];
????????shTileB[threadIdx.y+16][threadIdx.x]=B[(i*BLOCK_SIZE+threadIdx.y+16)*width+col];
????????__syncthreads();

????????for(int?j?=?0;?j?????????????val[0]?+=?shTileA[threadIdx.y][j]?*?shTileB[j][threadIdx.x];
????????????val[1]?+=?shTileA[threadIdx.y?+?16][j]?*?shTileB[j][threadIdx.x];
????????}
????????__syncthreads();
????}
????C[row?*?width?+?col]?=?val[0];
????C[(row?+?16)?*?width?+?col]?=?val[1];
}

注意， kernel launch 時(shí)的 blocksize 需要變化為：blockSize.y = BLOCK_SIZE / 2。而 gridSize 不變。

上面的 kernel 函數(shù)中，注意觀察內(nèi)層循環(huán)：讓 1 個(gè) thread 分別從子矩陣 A 中讀取 2 個(gè)數(shù)據(jù)，從子矩陣 B 中讀取 1 個(gè)數(shù)據(jù)（注意2次取數(shù)據(jù)是同一地址?。缓笸瑫r(shí)計(jì)算2個(gè)元素 val[0] 和 val[1]。此時(shí)，通過讀取 4B*3 個(gè)數(shù)據(jù)，實(shí)現(xiàn)了2次乘加共4次計(jì)算。減少了 shared memory 中子矩陣B一半的數(shù)據(jù)訪問。

下面詳細(xì)分析一下上述代碼對 Shared Memory 的訪問情況：

Shared Memory Store：每 1 次外層循環(huán)中對矩陣 A 和矩陣 B 有 2 次 load，總的 thread 個(gè)數(shù)減少了一半。但是實(shí)際上總的 load transactions 次數(shù)沒有變化（假設(shè) no bank conflict）。

Shared Memory Load：在每 1 次內(nèi)層循環(huán)中，1 個(gè) warp 內(nèi)的 32 個(gè) thread 需要從 shTileA 讀取同 2 個(gè)元素，需要 2 次 Shared Memory Load Transactions，再從 shTileB 中讀取連續(xù)的 32 個(gè)元素（假設(shè)沒有 Bank Conflict，需要1次 Shared Memory Load Transactions）（注意val[0]和val[1]的計(jì)算中，shTileB 的地址是一樣的），即總共需要 3 次 Shared Memory Load Transactions。 次循環(huán)總共需要 次 Shared Memory Load Transactions。對于 個(gè) threads，共有個(gè) warp，總共的 Shared Memory Load Transactions 數(shù)目為：。對比優(yōu)化前的 Shared Memory Load Transactions 數(shù)目 。

當(dāng)然，我們還可以繼續(xù)對 kernel 函數(shù)進(jìn)行優(yōu)化，讓每個(gè) thread 計(jì)算的元素個(gè)數(shù)從 2 個(gè)提高到 4/8/16/32 個(gè)，對比測試結(jié)果如下（m=n=k=1024, BLOCK_SIZE=32*32）：

此時(shí)測出的峰值性能相比于優(yōu)化前提升了約 82.5%。但是，為什么繼續(xù)增大每個(gè) thread 計(jì)算元素個(gè)數(shù)后，性能反而下降呢？

一方面是繼續(xù)增大該指標(biāo)后，每個(gè) block 中 thread 的個(gè)數(shù)是在減少的，但是每個(gè) block 中需要的 Shared Memory 數(shù)量沒有減少。這將導(dǎo)致由于 Block 總數(shù)受限，降低 SM 中的 active threads 數(shù)量，即降低了 Occupancy。另外，每個(gè) thread 計(jì)算更多元素會(huì)使用更多的 Registers。而每個(gè) thread 中 Registers 的個(gè)數(shù)又會(huì)反過來影響 SM 中 active threads 的個(gè)數(shù)，進(jìn)而影響 Occupancy。

具體 Shared Memory 或者 Registers 哪個(gè)會(huì)影響 Occupancy，取決于哪個(gè)指標(biāo)先到達(dá)閾值。在沒有換來同等指令集并行的情況下，Occupancy 的減少會(huì)導(dǎo)致計(jì)算性能受限。

注：Shared_2代碼中，每個(gè)thread計(jì)算 8 個(gè)元素。

上圖為優(yōu)化前后 3 個(gè)版本CUDA程序的性能差異，從圖中可以得出：

• 在句子規(guī)模為 的情況下，第三個(gè)版本的方法達(dá)到的峰值性能超過 7T；

• 隨著矩陣規(guī)模的增加，計(jì)算性能也逐漸增加；

• 通過利用 Shared Memory 和寄存器能有效的降低 IO 帶寬對性能的影響，從而更加高效的利用 GPU 的硬件計(jì)算資源。

完整代碼如下：

#include?
#define?BLOCK_SIZE?5

//基礎(chǔ)版本
__global__?void?matMul_GlobalKernel(int?*A,int?*B,int?*C,int?width){
????int?bx?=?blockIdx.x;
????int?by?=?blockIdx.y;
????int?tx?=?threadIdx.x;
????int?ty?=?threadIdx.y;

????int?Col?=?bx?*?blockDim.x?+?tx;
????int?Row?=?by?*?blockDim.y?+?ty;

????int?perValue?=?0;
????for(int?i?=?0;?i?????????perValue?+=?A[Row?*?width?+?i]?*?B[i?*?width?+?Col];
????}
????C[Row?*?width?+?Col]?=?Pervalue;
}

//優(yōu)化1
__global__?void?matmul_ShareMemory(int?*M,int?*N,int?*P,int?width){
????__shared__?float?Mds[BLOCK_SIZE][BLOCK_SIZE];
????__shared__?float?Nds[BLOCK_SIZE][BLOCK_SIZE];

????int?bx?=?blockIdx.x;
????int?by?=?blockIdx.y;
????int?tx?=?threadIdx.x;
????int?ty?=?threadIdx.y;

????int?Col?=?bx?*?BLOCK_SIZE?+?tx;
????int?Row?=?by?*?BLOCK_SIZE?+?ty;

????int?Pervalue?=?0;
????//有多少個(gè)BLOCK_SIZE，每個(gè)循環(huán)計(jì)算一個(gè)塊的大小
????for(int?i?=?0;i?????????Mds[ty][tx]?=?M[Row?*?width?+?(i?*?BLOCK_SIZE?+?tx)];
????????Nds[ty][tx]?=?N[Col?+?(i?*?BLOCK_SIZE?+?ty)?*?width];
????????__syncthreads();

????????//BLOCK_SIZE相乘
????????for(int?k?=?0;k?????????????Pervalue?+=?Mds[ty][k]?*?Nds[k][tx];
????????__syncthreads();
????}
????P[Row?*?width?+?Col]?=?Pervalue;
}

//優(yōu)化2
__global__?void?matrixMul_SharedILPkernel(float*?A,?float*?B,?float*?C,?int?width){
????int?row?=?blockIdx.y?*?blockDim.y?*?2?+?threadIdx.y;
????int?col?=?blockIdx.x?*?blockDim.x?+?threadIdx.x;
????float?val[2]?=?{0.0f};

????__shared__?float?shTileA[BLOCK_SIZE][BLOCK_SIZE];
????__shared__?float?shTileB[BLOCK_SIZE][BLOCK_SIZE];

????int?iter?=?(width?+?BLOCK_SIZE?-?1)?/?BLOCK_SIZE;
????for(int?i?=?0;?i?????????//?read?data?from?global?memory?to?shared?memory
????????shTileA[threadIdx.y][threadIdx.x]=A[row?*?width+i*BLOCK_SIZE+threadIdx.x];
????????shTileA[threadIdx.y+16][threadIdx.x]=A[(row+16)*width+i*BLOCK_SIZE+threadIdx.x];
????????shTileB[threadIdx.y][threadIdx.x]=B[(i*BLOCK_SIZE+threadIdx.y)*width+col];
????????shTileB[threadIdx.y+16][threadIdx.x]=B[(i*BLOCK_SIZE+threadIdx.y+16)*width+col];
????????__syncthreads();

????????for(int?j?=?0;?j?????????????val[0]?+=?shTileA[threadIdx.y][j]?*?shTileB[j][threadIdx.x];
????????????val[1]?+=?shTileA[threadIdx.y?+?16][j]?*?shTileB[j][threadIdx.x];
????????}
????????__syncthreads();
????}
????C[row?*?width?+?col]?=?val[0];
????C[(row?+?16)?*?width?+?col]?=?val[1];
}

完整代碼見【集智書童】知識星球

星球主要內(nèi)容包括：經(jīng)典論文分享、可落地方法復(fù)現(xiàn)、模型部署、Transformer部署與改進(jìn)、電子資源分享等內(nèi)容。

以下是星球的部分內(nèi)容，持續(xù)更新中......

掃描下方二維碼，我們一起學(xué)習(xí)吧！??！

參考

[1]. 矩陣乘法的 CUDA 實(shí)現(xiàn)、優(yōu)化及性能分析

2推薦閱讀

AI部署篇 | CUDA學(xué)習(xí)筆記1：向量相加與GPU優(yōu)化(附CUDA C代碼)

掃描下方二維碼關(guān)注【集智書童】公眾號，獲取更多實(shí)踐項(xiàng)目源碼和論文解讀，非常期待你我的相遇，讓我們以夢為馬，砥礪前行！