【CUDA】Sgemm單精度矩陣乘法（下）

目錄

    • 前言
    • 1. 優化技巧5：使用register模擬二級緩存（內積轉外積）
    • 2. 優化技巧6：使用register模擬二級緩存 + float4
    • 3. 優化技巧7：global memory轉置再存放shared memory
    • 4. 優化技巧8：使用double buffer加速矩陣乘法
    • 結語
    • 下載鏈接
    • 參考

前言

學習 UP 主 比飛鳥貴重的多_HKL 的 【CUDA】Sgemm單精度矩陣乘法（已完結~） 視頻，記錄下個人學習筆記，僅供自己參考😄

refer 1：【CUDA】Sgemm單精度矩陣乘法（已完結~）

refer 2：https://github.com/tpoisonooo/how-to-optimize-gemm/cuda

refer 3：https://chatgpt.com/

1. 優化技巧5：使用register模擬二級緩存（內積轉外積）

我們接著上篇文章來講解 sgemm 的優化

前面在 v2 版本中我們通過分塊的方式將數據從 global memory 放置在 shared memory 中，大大減少了訪存所需要的時延，這里，我們進一步考慮從 shared memory 到 register 的過程，如下圖所示：

在這里插入圖片描述

Note：圖片來自于：深入淺出GPU優化系列：GEMM優化（一）

通過寄存器來模擬二級緩存，可以將內積形式轉換為外積形式，如下圖所示：

在這里插入圖片描述

上圖左邊展示的是經典內積的形式，通過三層循環，每次計算 C_tile[m][n] 時，從 A_tile[m][k]（藍色行）與 B[k][n]（藍色列）中加載值，計算后逐步累加

上圖右邊改寫為按 k 為最外層的循環，變為外積實現。先固定一個 k，將 A_tile[:,k]（藍色列向量）和 B_tile[k,:] 加載到寄存器中，對 C_tile 的一個子塊進行外積更新，相當于更新一個 rank-1 子矩陣，這樣可以就減少對全體 A_tile/B_tile 數據的重復加載，起到模擬二級緩存（寄存器暫存）的作用

下圖展示了利用 register 內積轉外積的整體流程：

在這里插入圖片描述

在將數據從 global memory 加載到 shared memory 之后，還需要進一步加載到 register 中，接著通過外積計算方式逐步累加得到最終的結果

下圖還對比了 v4 和 v5 版本的差異，主要體現在寄存器的使用以及內積轉外積的實現：

在這里插入圖片描述

代碼如下：

template<unsigned int BLOCK_SIZE, unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v5_register_outer_product(float* A, float* B, float* C, const int M, const int N, const int K) {int row = blockIdx.y * blockDim.y + threadIdx.y;int col = (blockIdx.x * blockDim.x + threadIdx.x) * NUM_PER_THREAD;extern __shared__ float shared_mem[];float* A_tile = shared_mem;float* B_tile = shared_mem + BLOCK_SIZE * BLOCK_SIZE;constexpr int REG_NUM = NUM_PER_THREAD / 2;float A_reg[REG_NUM] = {0.0f};float B_reg[REG_NUM] = {0.0f};float sum[REG_NUM * REG_NUM] = {0.0f};// re-arrange the layoutint tid = threadIdx.y * blockDim.x + threadIdx.x;int ctx = tid % (BLOCK_SIZE / REG_NUM);int cty = tid / (BLOCK_SIZE / REG_NUM);for(int k_base = 0; k_base < K; k_base += BLOCK_SIZE){// load A_tile from global memory to shared memoryint a_col = k_base + threadIdx.x * NUM_PER_THREAD;FLOAT4(A_tile[threadIdx.y * BLOCK_SIZE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(A[row * K + a_col]);// load B_tile from global memory to shared memoryint b_row = k_base + threadIdx.y;FLOAT4(B_tile[threadIdx.y * BLOCK_SIZE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(B[b_row * N + col]);__syncthreads();// use register to compute the sum of A_tile * B_tile for(int k = 0; k < BLOCK_SIZE; ++k){A_reg[0] = A_tile[(cty * REG_NUM) * BLOCK_SIZE + k];A_reg[1] = A_tile[(cty * REG_NUM + 1) * BLOCK_SIZE + k];B_reg[0] = B_tile[k * BLOCK_SIZE + ctx * REG_NUM];B_reg[1] = B_tile[k * BLOCK_SIZE + ctx * REG_NUM + 1];for(int i = 0; i < REG_NUM; ++i){for(int j = 0; j < REG_NUM; ++j){sum[i * REG_NUM + j] += A_reg[i] * B_reg[j];}}}__syncthreads();        }// write the result to Cfloat* C_start = C + blockIdx.y * blockDim.y * N + blockIdx.x * blockDim.x * NUM_PER_THREAD;for(int i = 0; i < REG_NUM; ++i){for(int j = 0; j < REG_NUM; ++j){C_start[(cty * REG_NUM + i) * N + ctx * REG_NUM + j] = sum[i * REG_NUM + j];}}
}

下面是該代碼的詳細分析：(from ChatGPT)

1. 核函數簽名和參數

template<unsigned int BLOCK_SIZE, unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v5_register_outer_product(float* A, float* B, float* C, const int M, const int N, const int K)
  • 模板參數：
    • BLOCK_SIZE：分塊（TILE）的大小，也就是每次只加載 BLOCK_SIZE * BLOCK_SIZE 大小的元素到共享內存中，在示例中 BLOCK_SIZE = 16
    • NUM_PER_THREAD：每個線程處理的元素數量，在示例中 NUM_PER_THREAD = 4 也就是每個線程負責 C 矩陣中 2x2 大小的元素，

2. 線程索引計算

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = (blockIdx.x * blockDim.x + threadIdx.x) * NUM_PER_THREAD;
  • row：當前線程處理的 A 矩陣的行索引
  • col：當前線程處理的 B 矩陣的起始列索引，注意由于每個線程處理 4 個連續元素，因此需要乘以 NUM_PER_THREAD

3. 共享內存分配

extern __shared__ float shared_mem[];
float* A_tile = shared_mem;
float* B_tile = shared_mem + BLOCK_SIZE * BLOCK_SIZE;
  • 共享內存布局：
    • 動態共享內存 shared_mem 被劃分為兩部分：
      • A_tile：前 BLOCK_SIZE * BLOCK_SIZE = 256 個 float，用于緩存 A 矩陣的分塊
      • B_tile：后 256 個 float，用于緩存 B 矩陣的分塊
    • 總共享內存需求：512 個 float（2KB）

4. 寄存器變量聲明

constexpr int REG_NUM = NUM_PER_THREAD / 2;
float A_reg[REG_NUM] = {0.0f};
float B_reg[REG_NUM] = {0.0f};
float sum[REG_NUM * REG_NUM] = {0.0f};
  • REG_NUM：
    • NUM_PER_THREAD = 4，所以 REG_NUM = 2
    • 這個設計意味著每個線程將處理 2x2=4 個輸出元素
  • 寄存器變量：
    • A_reg[2]：緩存從 A_tile 加載的數據
    • B_reg[2]：緩存從 B_tile 加載的數據
    • sum[4]：累加 2x2 輸出塊的部分和

5. 數據布局重排

int tid = threadIdx.y * blockDim.x + threadIdx.x;
int ctx = tid % (BLOCK_SIZE / REG_NUM);
int cty = tid / (BLOCK_SIZE / REG_NUM);
  • tid 計算：
    • 線性化的線程 ID，當前線程在 block 中的全局索引，范圍 0~63（4x16 線程塊）
  • ctx 和 cty 計算：
    • BLOCK_SIZE / REG_NUM = 16 / 2 = 8
    • ctx = tid % 8：線程在 x 方向的邏輯索引（0~7）
    • cty = tid / 8：線程在 y 方向的邏輯索引（0~7）
    • 這種重排將 64 個線程組織為 8x8 的網格，每個線程負責 2x2 的輸出塊，如下圖所示

在這里插入圖片描述

6. 主計算循環

for(int k_base = 0; k_base < K; k_base += BLOCK_SIZE)
  • 循環結構：
    • 在 K 維度上分塊處理，步長為 BLOCK_SIZE = 16
    • 對于 K = 512，共需要 512 / 16 = 32 次迭代

6.1 A_tile 加載

int a_col = k_base + threadIdx.x * NUM_PER_THREAD;
FLOAT4(A_tile[threadIdx.y * BLOCK_SIZE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(A[row * K + a_col]);
  • 加載邏輯：
    • a_col：當前處理的 A 矩陣列索引
    • 使用 FLOAT4 寬指令一次加載 A 矩陣中 4 個連續的 float
    • 寫入共享內存 A_tile 中，按行主序排序

Note：關于索引的計算博主在 v4 版本中已經講過了，這邊就不再贅述了

6.2 B_tile 加載

int b_row = k_base + threadIdx.y;
FLOAT4(B_tile[threadIdx.y * BLOCK_SIZE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(B[b_row * N + col]);
  • 加載邏輯：
    • b_row：當前處理的 B 矩陣行索引
    • 同樣使用 FLOAT4 寬指令加載 B 矩陣的 4 個連續元素
    • 寫入共享內存 B_tile 中，按行主序排序

6.3 同步

__syncthreads();
  • 確保所有線程完成共享內存的加載后才開始計算

6.4 寄存器緩存計算

for(int k = 0; k < BLOCK_SIZE; ++k){A_reg[0] = A_tile[(cty * REG_NUM) * BLOCK_SIZE + k];A_reg[1] = A_tile[(cty * REG_NUM + 1) * BLOCK_SIZE + k];B_reg[0] = B_tile[k * BLOCK_SIZE + ctx * REG_NUM];B_reg[1] = B_tile[k * BLOCK_SIZE + ctx * REG_NUM + 1];
  • 寄存器加載：
    • 每次迭代處理 K 維度的一個元素（k=0 到 15）
    • 從 A_tile 加載兩行到 A_reg（由 cty 決定）
    • 從 B_tile 加載兩列到 B_reg（由 ctx 決定）
    • 這種訪問模式確保了合并的內存訪問
  • 外積計算：
    for(int i = 0; i < REG_NUM; ++i){for(int j = 0; j < REG_NUM; ++j){sum[i * REG_NUM + j] += A_reg[i] * B_reg[j];}}
}
  • 計算模式：
    • 這是典型的外積計算：A 的列向量（2 元素）與 B 的行向量（2 元素）相乘，得到 2x2 的矩陣
    • 結果累加到 sum 數組中
    • 共進行 BLOCK_SIZE = 16 次外積累加

博主繪制了一個草圖來簡要說明整個流程，如下圖所示（以 (ctx,cty)=(0,0) 線程為例）：

在這里插入圖片描述

每個線程處理 A_tile 兩行數據與 B_tile 兩列數據相乘，每次加載 A_tile 兩個數據和 B_tile 兩個數據到寄存器，循環 BLOCK_SIZE 次

6.5 第二次同步

__syncthreads();
  • 確保所有線程完成當前塊的計算后再加載下一塊

7. 結果寫回

float* C_start = C + blockIdx.y * blockDim.y * N + blockIdx.x * blockDim.x * NUM_PER_THREAD;
for(int i = 0; i < REG_NUM; ++i){for(int j = 0; j < REG_NUM; ++j){C_start[(cty * REG_NUM + i) * N + ctx * REG_NUM + j] = sum[i * REG_NUM + j];}
}
  • C_start 計算：
    • 計算當前線程塊對應的的 C 矩陣起始位置
    • blockIdx.y * blockDim.y * N：當前線程塊在 y 方向的偏移
    • blockIdx.x * blockDim.x * NUM_PER_THREAD：當前線程塊在 x 方向的偏移
  • 寫回邏輯：
    • 每個線程將其累加的 2x2 結果塊寫回全局內存
    • 使用 cty 和 ctx 確定寫入位置，保持與計算時相同的布局

這里博主在自己實現時發現始終不能很好的定位到 C 的全局索引，繞著繞著就把自己給繞暈了

那其實 UP 主的方法非常的有效，我們先定位到當前 block 對應的矩陣 C 的位置，然后再來處理具體 block 中每個 thread 的部分，這樣我們就只要關注每個 block 的索引計算就行了，會簡單不少

它也可以寫成下面的這種形式：

for(int i = 0; i < REG_NUM; ++i){for(int j = 0; j < REG_NUM; ++j){int c_row = blockIdx.y * blockDim.y + cty * REG_NUM + i;int c_col = blockIdx.x * blockDim.x * NUM_PER_THREAD + ctx * REG_NUM + j;C[c_row * N + c_col] = sum[i * REG_NUM + j];}
}

這個核函數使用寄存器來模擬二級緩存，其中：

  • 寄存器緩存層次：
    • 第一級：共享內存緩存全局內存數據
    • 第二級：寄存器緩存共享內存數據
    • 這種層次結構減少了共享內存的訪問壓力
  • 數據流分析：
    • 全局內存 ? 共享內存（通過 FLOAT4 寬加載）
    • 共享內存 ? 寄存器（標量加載）
    • 寄存器 ? 計算單元（高效計算）
  • 性能優化點：
    • 每個線程的 A_reg 和 B_reg 在循環中重復使用，減少共享內存訪問
    • 外積計算模式最大化數據復用率
    • 寄存器訪問完全無沖突，延遲極低

此外，這個核函數還將內積計算轉換為外積計算，博主沒有太理解內積轉外積過程，這邊簡要說明下：(from ChatGPT)

首先我們需要理解內積和外積這兩個基本概念，再來分析它們在矩陣乘法中的應用

1. 內積（點積）

數學定義：內積是兩個向量的乘積，結果是一個標量（單個數值）

對于兩個 n n n 維向量 a = [ a 1 , a 2 , … , a n ] \mathbf{a} = [a_1,a_2,\ldots,a_n] a=[a1?,a2?,…,an?] 和 b = [ b 1 , b 2 , … , b n ] \mathbf{b} = [b_1,b_2,\ldots,b_n] b=[b1?,b2?,…,bn?]

內積計算方式如下：

a ? b = a 1 b 1 + a 2 b 2 + … + a n b n = ∑ i = 1 n ( a i b i ) \mathbf{a} \cdot \mathbf{b} = a_1b_1+a_2b_2+\ldots+a_nb_n=\sum_{i=1}^{n}(a_ib_i) a?b=a1?b1?+a2?b2?+…+an?bn?=i=1∑n?(ai?bi?)

傳統矩陣乘法就是基于內積的：

for(int m = 0; m < M; m++){for(int n = 0; n < N; n++){float sum = 0;for(int k = 0; k < K; k++){sum += A[m * K + k] * B[k * N + n];}C[m * N + n] = sum;}
}

特點：

  • 每個輸出元素需要遍歷 K 維度
  • 內存訪問模式不連續（A 按行訪問，B 按列訪問）
  • 計算訪存比低

2. 外積

數學定義：外積是兩個向量的乘積，結果是一個矩陣

對于向量 u = [ u 1 , u 2 , … , u m ] \mathbf{u} = [u_1,u_2,\ldots,u_m] u=[u1?,u2?,…,um?] 和 v = [ v 1 , v 2 , … , v n ] \mathbf{v} = [v_1,v_2,\ldots,v_n] v=[v1?,v2?,…,vn?]

外積計算方式如下：

u × v T = [ u 1 u 2 ? u m ] [ v 1 v 2 ? v n ] = [ u 1 v 1 u 1 v 2 ? u 1 v n u 2 v 1 u 2 v 2 ? u 2 v n ? ? ? ? u m v 1 u m v 2 ? u m v n ] \mathbf{u} \times \mathbf{v}^T = \begin{bmatrix} u_1 \\ u_2 \\ \vdots \\ u_m \end{bmatrix} \begin{bmatrix} v_1 & v_2 & \cdots & v_n \end{bmatrix}= \begin{bmatrix} u_1 v_1 & u_1 v_2 & \cdots & u_1 v_n \\ u_2 v_1 & u_2 v_2 & \cdots & u_2 v_n \\ \vdots & \vdots & \ddots & \vdots \\ u_m v_1 & u_m v_2 & \cdots & u_m v_n \end{bmatrix} u×vT= ?u1?u2??um?? ?[v1??v2????vn??]= ?u1?v1?u2?v1??um?v1??u1?v2?u2?v2??um?v2???????u1?vn?u2?vn??um?vn?? ?

外積方法將矩陣乘法計算重構為多個小矩陣的乘積累加，在核函數中的具體表現為：

1. 數據分塊：

  • 將 A_tile 按列分塊（16x2 的小塊）
  • 將 B_tile 按行分塊（2x16 的小塊）

2. 計算單元：

// 外積計算核心代碼
A_reg[0] = A_tile[...]; // 加載A的一列2個元素
A_reg[1] = A_tile[...];B_reg[0] = B_tile[...]; // 加載B的一行2個元素
B_reg[1] = B_tile[...];// 計算2x2外積并累加
for(int i = 0; i < 2; i++)for(int j = 0; j < 2; j++)sum[i * 2 + j] += A_reg[i] * B_reg[j];

3. 數學表示：

每次迭代計算的是：

[ A 1 A 2 ] [ B 1 B 2 ] = [ A 1 B 1 A 1 B 2 A 2 B 1 A 2 A 2 ] \begin{bmatrix} A_1 \\ A_2 \\ \end{bmatrix} \begin{bmatrix} B_1 \ B_2 \end{bmatrix}= \begin{bmatrix} A_1 B_1 & A_1 B_2\\ A_2 B_1 & A_2 A_2\\ \end{bmatrix} [A1?A2??][B1??B2??]=[A1?B1?A2?B1??A1?B2?A2?A2??]

然后將這些 2x2 的小矩陣累加到最終結果中

與傳統內積方法相比，外積更適合 GPU，這主要是因為：

  • 更高的數據復用：
    • 在計算 2x2 外積時，A 的 2 個元素與 B 的 2 個元素產生 4 次乘加
    • 相比內積方法（1 個 A 元素 x 1 個 B 元素 = 1 次乘加），計算密度提高 4 倍
  • 更連續的內存訪問：
    • 外積方法中，A 按列訪問，B 按行訪問，都是連續內存訪問
    • 內積方法中，B 必須按列訪問，導致非連續訪問
  • 更適合 SIMT 架構：
    • GPU 擅長并行執行相同指令
    • 外積方法讓每個線程處理一個小矩陣，并行度高
    • 內積方法線程間計算模式差異大，不利于并行
  • 更高的計算強度：
    • 每次加載的數據參與更多計算操作
    • 更好地隱藏內存訪問延遲

性能和帶寬測試結果如下：

優化手段矩陣維度GridBlock耗時(us)Memory Throughout(%)DRAM Throughout(%)
v0_global_memory512x512(32,32)(16,16)471.7896.941.56
v1_shared_memory256x256(16,16)(16,16)82.1178.921.84
v2_shared_memory_sliding_window512x512(32,32)(16,16)362.5094.457.05
v3_increase_work_of_per_thread512x512(16,16)(16,16)204.2684.013.64
v4_using_float4512x512(32,32)(4,16)209.6091.993.44
v5_register_outer_product512x512(32,32)(4,16)206.1879.103.50

2. 優化技巧6：使用register模擬二級緩存 + float4

這個小節我們嘗試將 v4 和 v5 版本融合起來，既使用寄存器外積形式，又使用 float4 向量化加載，流程如下圖所示：

在這里插入圖片描述

由于我們需要使用 float4 來向量化加載，因此寄存器的數量我們各增加到 4 個，此時每個線程負責處理 4x4 大小的數據，如上圖所示。另外由于 B_tile 中的四個元素在內存中是連續的，因此可以使用 float4 加載，而對應的 A_tile 中的四個元素并不是連續存儲的，所以對于 A_tile 我們還是來一個個手動加載到寄存器中

代碼如下：

template<unsigned int NUM_PER_TILE, unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v6_register_outer_product_float4(float* A, float* B, float* C, const int M, const int N, const int K){int row = (blockIdx.y * blockDim.y + threadIdx.y) * NUM_PER_THREAD;int col = (blockIdx.x * blockDim.x + threadIdx.x) * NUM_PER_THREAD;extern __shared__ float shared_mem[];float* A_tile = shared_mem;float* B_tile = shared_mem + NUM_PER_TILE * NUM_PER_TILE;float A_reg[NUM_PER_THREAD] = {0.0f};float B_reg[NUM_PER_THREAD] = {0.0f};float sum[NUM_PER_THREAD * NUM_PER_THREAD] = {0.0f};for(int k_base = 0; k_base < K; k_base += NUM_PER_TILE){for(int i = 0; i < NUM_PER_THREAD; ++i){// load A_tile from global memory to shared memoryint a_col = k_base + threadIdx.x * NUM_PER_THREAD;FLOAT4(A_tile[(threadIdx.y * NUM_PER_THREAD + i) * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(A[(row + i) * K + a_col]);// load B_tile from global memory to shared memoryint b_row = k_base + threadIdx.y * NUM_PER_THREAD;FLOAT4(B_tile[(threadIdx.y * NUM_PER_THREAD + i) * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(B[(b_row + i) * N + col]);}__syncthreads();// use register to compute the sum of A_tile * B_tilefor(int k = 0; k < NUM_PER_TILE; ++k){A_reg[0] = A_tile[(threadIdx.y * NUM_PER_THREAD + 0) * NUM_PER_TILE + k];A_reg[1] = A_tile[(threadIdx.y * NUM_PER_THREAD + 1) * NUM_PER_TILE + k];A_reg[2] = A_tile[(threadIdx.y * NUM_PER_THREAD + 2) * NUM_PER_TILE + k];A_reg[3] = A_tile[(threadIdx.y * NUM_PER_THREAD + 3) * NUM_PER_TILE + k];FLOAT4(B_reg[0]) = FLOAT4(B_tile[k * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]);for(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}}}__syncthreads();        }// write the result to Cfloat* C_start = C + blockIdx.y * NUM_PER_TILE * N + blockIdx.x * NUM_PER_TILE;for(int i = 0; i < NUM_PER_THREAD; ++i){FLOAT4(C_start[(threadIdx.y * NUM_PER_THREAD + i) * N + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(sum[i * NUM_PER_THREAD]);}
}

下面是該代碼的詳細分析：(from ChatGPT)

1. 核函數簽名與參數

template<unsigned int NUM_PER_TILE, unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v6_register_outer_product_float4(float* A, float* B, float* C, const int M, const int N, const int K)
  • 模板參數：
    • NUM_PER_TILE：每個線程塊每次處理的分塊矩陣的大小，也就是將 NUM_PER_TILE * NUM_PER_TILE 大小的矩陣加載到共享內存中，相當于之前的 BLOCK_SIZE 變量，在當前示例中其值為 64
    • NUM_PER_THREAD：每個線程處理的元素數量，在當前示例中為 4
      • 注意，這里可能會引發歧義，實際上在當前核函數中每個線程處理的元素數量是 4x4=16，NUM_PER_THREAD 在這里表示的是行和列方向上各處理 4 個元素

2. 線程索引計算

int row = (blockIdx.y * blockDim.y + threadIdx.y) * NUM_PER_THREAD;
int col = (blockIdx.x * blockDim.x + threadIdx.x) * NUM_PER_THREAD;
  • row：當前線程處理的 A 矩陣的行索引，乘以 NUM_PER_THREAD 得到實際起始行
  • col：當前線程處理的 B 矩陣的列索引，乘以 NUM_PER_THREAD 得到實際起始列

3. 共享內存分配

extern __shared__ float shared_mem[];
float* A_tile = shared_mem;
float* B_tile = shared_mem + NUM_PER_TILE * NUM_PER_TILE;
  • 動態共享內存 shared_mem 被劃分為兩部分：
    • A_tile：前 NUM_PER_TILE * NUM_PER_TILE = 4096 個 float，用于緩存 A 矩陣的分塊
    • B_tile：后 4096 個 float，用于緩存 B 矩陣的分塊
  • 總共享內存需求：8192 個 float（32KB）

4. 寄存器變量聲明

float A_reg[NUM_PER_THREAD] = {0.0f};  // 4個寄存器緩存A數據
float B_reg[NUM_PER_THREAD] = {0.0f};  // 4個寄存器緩存B數據
float sum[NUM_PER_THREAD * NUM_PER_THREAD] = {0.0f}; // 16個寄存器存儲結果
  • 每個線程：
    • 緩存 4 個 A 元素和 4 個 B 元素
    • 累加 4x4=16 個部分和

5. 主計算循環

for(int k_base = 0; k_base < K; k_base += NUM_PER_TILE)
  • 循環結構：
    • 在 K 維度上分塊處理，步長為 NUM_PER_TILE = 64
    • 對于 K = 512，共需要 512 / 64 = 8 次迭代

5.1 數據加載階段

for(int i = 0; i < NUM_PER_THREAD; ++i){// 加載A_tileint a_col = k_base + threadIdx.x * NUM_PER_THREAD;FLOAT4(A_tile[(threadIdx.y * NUM_PER_THREAD + i) * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(A[(row + i) * K + a_col]);// 加載B_tileint b_row = k_base + threadIdx.y * NUM_PER_THREAD;FLOAT4(B_tile[(threadIdx.y * NUM_PER_THREAD + i) * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(B[(b_row + i) * N + col]);
}
  • A_tile 加載
    • 每個線程加載 4 個 FLOAT4（16 個 float）
    • 訪問模式：連續訪問 A 的 4 行
    • 使用 FLOAT4 實現向量化加載
  • B_tile 加載
    • 同樣加載 4 個 FLOAT4
    • 訪問模式：連續訪問 B 的 4 列

5.2 同步

__syncthreads();
  • 確保所有線程完成共享內存的加載后才開始計算

5.3 外積計算

for(int k = 0; k < NUM_PER_TILE; ++k){// 加載A的4個元素（一列）A_reg[0] = A_tile[(threadIdx.y * NUM_PER_THREAD + 0) * NUM_PER_TILE + k];A_reg[1] = A_tile[(threadIdx.y * NUM_PER_THREAD + 1) * NUM_PER_TILE + k];A_reg[2] = A_tile[(threadIdx.y * NUM_PER_THREAD + 2) * NUM_PER_TILE + k];A_reg[3] = A_tile[(threadIdx.y * NUM_PER_THREAD + 3) * NUM_PER_TILE + k];// 用FLOAT4加載B的4個元素（一行）FLOAT4(B_reg[0]) = FLOAT4(B_tile[k * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]);// 計算4x4外積for(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}}
}
  • 關鍵優化點：
    • 寄存器緩存：A_reg 和 B_reg 緩存共享內存數據
    • 外積計算：4 元素列向量 x 4 元素行向量 ? 4x4 矩陣
    • 向量化加載：B_reg 使用 FLOAT4 一次性加載 4 個元素
  • 計算過程：
    • 對每個 k（0~63）：
      • 從 A_tile 加載一列 4 個元素到 A_reg
      • 從 B_tile 加載一行 4 個元素到 B_reg（用 FLOAT4）
      • 計算 4x4 外積并累加到 sum

整個過程如下圖所示（以 thread(0, 0) 線程為例）：

在這里插入圖片描述

每個線程處理 A_tile 四行數據與 B_tile 四列數據相乘，每次加載 A_tile 四個數據和 B_tile 四個數據到寄存器，其中 B_tile 的四個數據內存連續，因此可以使用 float4 向量化加載，循環 NUM_PER_TILE 次

5.4 第二次同步

__syncthreads();
  • 確保所有線程完成當前塊的計算后再加載下一塊

6. 結果寫回

float* C_start = C + blockIdx.y * NUM_PER_TILE * N + blockIdx.x * NUM_PER_TILE;
for(int i = 0; i < NUM_PER_THREAD; ++i){FLOAT4(C_start[(threadIdx.y * NUM_PER_THREAD + i) * N + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(sum[i * NUM_PER_THREAD]);
}
  • 寫回優化：
    • 使用 FLOAT4 一次性寫回 4 個結果
    • 每個線程寫回 4x4=16 個結果元素
    • 寫回位置計算考慮了線程塊和線程的索引

和 v5 版本相比的差異如下：

  • 分塊大小
    • 之前：BLOCK_SIZE = 16
    • 現在：NUM_PER_TILE = 64
  • 線程組織
    • 之前：4x46=64 線程/塊
    • 現在：16x16=256 線程/塊
  • 外積規模
    • 之前：2x2 外積
    • 現在：4x4 外積

v6 版本的實現通過更大的分塊、更大的線程和更大的外積規模，進一步提高了計算效率和內存訪問信息，此外通過 float4 向量化加載 B 矩陣和寫回 C 矩陣進一步提高了吞吐量

性能和帶寬測試結果如下：

優化手段矩陣維度GridBlock耗時(us)Memory Throughout(%)DRAM Throughout(%)
v0_global_memory512x512(32,32)(16,16)471.7896.941.56
v1_shared_memory256x256(16,16)(16,16)82.1178.921.84
v2_shared_memory_sliding_window512x512(32,32)(16,16)362.5094.457.05
v3_increase_work_of_per_thread512x512(16,16)(16,16)204.2684.013.64
v4_using_float4512x512(32,32)(4,16)209.6091.993.44
v5_register_outer_product512x512(32,32)(4,16)206.1879.103.50
v6_register_outer_product_float4512x512(8,8)(16,16)84.9960.387.28

3. 優化技巧7：global memory轉置再存放shared memory

在 v6 版本中 A_tile 從共享內存到寄存器的加載我們并沒有使用 float4，因為它們之間的內存并不是連續的。那這里我們可以考慮在將 A 矩陣存入 shared memory 之前做一次轉置，這樣就可以也使用 float4 來處理 A_tile，如下圖所示：

在這里插入圖片描述

相比于之前的版本（v6）這里我們考慮在將矩陣 A 的元素從 global memory 加載到 shared memory 時做一個轉置，這樣我們在做外積計算時就可以直接取 A_tile 中的連續 4 個元素

加載流程如下：

在這里插入圖片描述

這個實現方式就是借用 4 個臨時的寄存器來完成轉置操作，首先通過 float4 向量化讀取 A 中的 4 個元素并存儲在臨時的 4 個寄存器中，接著將 4 個寄存器的值按照轉置的方式填充到 A_tile 共享內存中，然后依次循環其他加載的元素，最終它相當于把之前 A_tile 整個給轉置過來了

計算流程如下：

在這里插入圖片描述

這個實現在外積計算時會相對簡單些，因為 A_tile 是轉置存儲的，因此我們現在完全可以像加載 B_tile 一樣通過 float4 來加載 A_tile 了，所以在上圖中我們可以清晰的看到 A_tile 和 B_tile 的加載相同

代碼如下：

template<unsigned int NUM_PER_TILE, unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v7_A_smen_transpose(float* A, float* B, float* C, const int M, const int N, const int K){int row = (blockIdx.y * blockDim.y + threadIdx.y) * NUM_PER_THREAD;int col = (blockIdx.x * blockDim.x + threadIdx.x) * NUM_PER_THREAD;extern __shared__ float shared_mem[];float* A_tile = shared_mem;float* B_tile = shared_mem + NUM_PER_TILE * NUM_PER_TILE;float A_reg[NUM_PER_THREAD] = {0.0f};float B_reg[NUM_PER_THREAD] = {0.0f};float A_load_reg[NUM_PER_THREAD] = {0.0f};float sum[NUM_PER_THREAD * NUM_PER_THREAD] = {0.0f};for(int k_base = 0; k_base < K; k_base += NUM_PER_TILE){for(int i = 0; i < NUM_PER_THREAD; ++i){// col-major load A_tile from global memory to shared memoryint a_col = k_base + threadIdx.x * NUM_PER_THREAD;FLOAT4(A_load_reg[0]) = FLOAT4(A[(row + i) * K + a_col]);A_tile[(threadIdx.x * NUM_PER_THREAD + 0) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[0];A_tile[(threadIdx.x * NUM_PER_THREAD + 1) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[1];A_tile[(threadIdx.x * NUM_PER_THREAD + 2) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[2];A_tile[(threadIdx.x * NUM_PER_THREAD + 3) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[3];// load B_tile from global memory to shared memoryint b_row = k_base + threadIdx.y * NUM_PER_THREAD;FLOAT4(B_tile[(threadIdx.y * NUM_PER_THREAD + i) * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(B[(b_row + i) * N + col]);}__syncthreads();// use register to compute the sum of A_tile * B_tilefor(int k = 0; k < NUM_PER_TILE; ++k){FLOAT4(A_reg[0]) = FLOAT4(A_tile[k * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD]);FLOAT4(B_reg[0]) = FLOAT4(B_tile[k * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]);for(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}}}__syncthreads();        }// write the result to Cfloat* C_start = C + blockIdx.y * NUM_PER_TILE * N + blockIdx.x * NUM_PER_TILE;for(int i = 0; i < NUM_PER_THREAD; ++i){FLOAT4(C_start[(threadIdx.y * NUM_PER_THREAD + i) * N + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(sum[i * NUM_PER_THREAD]);}
}

v7 版本最核心的改進是對 A_tile 進行轉置存儲，使得從共享內存加載到寄存器時也能使用 FLOAT4 向量化加載，這種優化帶來了以下關鍵變化：

  • A_tile 內存布局重構：從行主序改為列主序
  • 加載模式改變：使用 FLOAT4 同時加載 A 和 B 的數據
  • 新增中間寄存器：A_load_reg 用于臨時存儲轉置數據

關鍵代碼分析如下：(from ChatGPT)

1. 新增寄存器變量

float A_load_reg[NUM_PER_THREAD] = {0.0f};  // 新增的臨時寄存器
  • 作用：臨時存儲從全局內存加載的 A 矩陣數據，用于轉置寫入共享內存
  • 必要性：實現從行優先到列優先的布局轉換

2. A_tile 加載邏輯重構（核心變化）

// v6版本（原始行主序加載）:
FLOAT4(A_tile[(threadIdx.y * NUM_PER_THREAD + i) * NUM_PER_TILE + threadIdx.x * NUM_PER_THREAD]) = 
FLOAT4(A[(row + i) * K + a_col]);// v7版本（轉置列主序加載）:
FLOAT4(A_load_reg[0]) = FLOAT4(A[(row + i) * K + a_col]);
A_tile[(threadIdx.x * NUM_PER_THREAD + 0) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[0];
A_tile[(threadIdx.x * NUM_PER_THREAD + 1) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[1];
A_tile[(threadIdx.x * NUM_PER_THREAD + 2) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[2];
A_tile[(threadIdx.x * NUM_PER_THREAD + 3) * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD + i] = A_load_reg[3];
  • 轉置操作解析：
    • 1. 用 FLOAT4 從全局內存連續加載 4 個元素到 A_load_reg
    • 2. 將這些元素分散存儲到共享內存的不同位置，實現轉置
    • 3. 存儲模式：(threadIdx.x * NUM_PER_THREAD + n) 決定列，(threadIdx.y * NUM_PER_THREAD + i) 決定行
  • 內存布局對比：
    • v6：A_tile 按行存儲，同一行的元素在內存中連續
    • v7：A_tile 按列存儲，同一列的元素在內存中連續

整個加載過程如下圖所示（以 thread(0, 0) 為例）：

在這里插入圖片描述

數據從 global memory 加載到寄存器中和之前保持一致，都是行主序加載。但是將數據從寄存器加載到 shared memory 中時是轉置加載的，也就是列主序加載，這樣我們在后續計算時也可以使用 float4 來加載 A_tile

3. 計算核心的優化（關鍵改進）

// v6版本（標量加載A）:
A_reg[0] = A_tile[(threadIdx.y * NUM_PER_THREAD + 0) * NUM_PER_TILE + k];
// ...（加載4個標量）// v7版本（FLOAT4加載A）:
FLOAT4(A_reg[0]) = FLOAT4(A_tile[k * NUM_PER_TILE + threadIdx.y * NUM_PER_THREAD]);
  • 優化效果：原來需要 4 次單獨加載，現在 1 次 FLOAT4 完成

我們舉個簡單的例子來說明下這一過程

假設矩陣 A 和矩陣 B 都是 4x4 大小，則二者相乘計算如下圖所示：

在這里插入圖片描述

其中 sum[0][0] 的計算結果如圖中所示，A 的一行和 B 的一列相乘

假設矩陣 A 和矩陣 B 都加載到了 A_tile 和 B_tile 中，且 A_tile 是列主序加載，則此時 sum[0][0] 的計算如下圖所示：

在這里插入圖片描述

由于 A_tile 是列主序存儲，因此可以和 B_tile 一樣通過 float4 向量化加載，加載到寄存器后由于是外積相乘，因此每次循環恰好計算相應元素的乘積，最終的結果也和前面保持一致

性能和帶寬測試結果如下：

優化手段矩陣維度GridBlock耗時(us)Memory Throughout(%)DRAM Throughout(%)
v0_global_memory512x512(32,32)(16,16)471.7896.941.56
v1_shared_memory256x256(16,16)(16,16)82.1178.921.84
v2_shared_memory_sliding_window512x512(32,32)(16,16)362.5094.457.05
v3_increase_work_of_per_thread512x512(16,16)(16,16)204.2684.013.64
v4_using_float4512x512(32,32)(4,16)209.6091.993.44
v5_register_outer_product512x512(32,32)(4,16)206.1879.103.50
v6_register_outer_product_float4512x512(8,8)(16,16)84.9960.387.28
v7_A_smem_transpose512x512(8,8)(16,16)118.2165.855.39

4. 優化技巧8：使用double buffer加速矩陣乘法

之前的版本中我們的 shared memory 只有一組，如下圖所示：

在這里插入圖片描述

這里考慮使用兩組，也就是 double buffer 的優化策略，其中一組先預填充，另一組異步加載，然后對預填充的緩沖區計算，這樣可以確保加載和計算重疊，有助于延遲隱藏，具體的實現流程我們還是來看代碼慢慢講解吧

代碼如下：

template<unsigned int BLOCK_SIZE_M,unsigned int BLOCK_SIZE_N,unsigned int BLOCK_SIZE_K,unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v8_double_buffer(float* A, float* B, float* C, const int M, const int N, const int K){float* A_start = A + blockIdx.y * BLOCK_SIZE_M * K;float* B_start = B + blockIdx.x * BLOCK_SIZE_N;// double bufferextern __shared__ float shared_mem[];int A_tile_per_buffer_size = BLOCK_SIZE_K * BLOCK_SIZE_M;int B_tile_per_buffer_size = BLOCK_SIZE_K * BLOCK_SIZE_N;float* A_tile = shared_mem;float* B_tile = shared_mem + 2 * A_tile_per_buffer_size;float A_reg[NUM_PER_THREAD] = {0.0f};float B_reg[NUM_PER_THREAD] = {0.0f};float A_load_reg[4] = {0.0f};float sum[NUM_PER_THREAD * NUM_PER_THREAD] = {0.0f};// re-arrange the layoutint tid = threadIdx.y * blockDim.x + threadIdx.x;   // 0~256int A_tile_tx = tid % (BLOCK_SIZE_K / 4);           // 0~1int A_tile_ty = tid / (BLOCK_SIZE_K / 4);           // 0~127int B_tile_tx = tid % (BLOCK_SIZE_N / 4);           // 0~31int B_tile_ty = tid / (BLOCK_SIZE_N / 4);           // 0~7// prefetch first tileFLOAT4(A_load_reg[0]) = FLOAT4(A_start[A_tile_ty * K + A_tile_tx * 4]);A_tile[(A_tile_tx * 4 + 0) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[0]; A_tile[(A_tile_tx * 4 + 1) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[1]; A_tile[(A_tile_tx * 4 + 2) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[2]; A_tile[(A_tile_tx * 4 + 3) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[3]; FLOAT4(B_tile[B_tile_ty * BLOCK_SIZE_N + B_tile_tx * 4]) = FLOAT4(B_start[B_tile_ty * N + B_tile_tx * 4]);__syncthreads();int buffer_idx = 1;for(int k_base = BLOCK_SIZE_K; k_base < K; k_base += BLOCK_SIZE_K){// prefetch next tileFLOAT4(A_load_reg[0]) = FLOAT4(A_start[A_tile_ty * K + A_tile_tx * 4 + k_base]);A_tile[buffer_idx * A_tile_per_buffer_size + (A_tile_tx * 4 + 0) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[0]; A_tile[buffer_idx * A_tile_per_buffer_size + (A_tile_tx * 4 + 1) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[1]; A_tile[buffer_idx * A_tile_per_buffer_size + (A_tile_tx * 4 + 2) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[2]; A_tile[buffer_idx * A_tile_per_buffer_size + (A_tile_tx * 4 + 3) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[3]; FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + B_tile_ty * BLOCK_SIZE_N + B_tile_tx * 4]) = FLOAT4(B_start[(B_tile_ty + k_base) * N + B_tile_tx * 4]);// toggle buffer indexbuffer_idx = buffer_idx ^ 1;// compute current tilefor(int k = 0; k < BLOCK_SIZE_K; ++k){// load A_tile and B_tile from shared memory to registerFLOAT4(A_reg[0]) = FLOAT4(A_tile[buffer_idx * A_tile_per_buffer_size + k * BLOCK_SIZE_M + threadIdx.y * NUM_PER_THREAD]);FLOAT4(A_reg[4]) = FLOAT4(A_tile[buffer_idx * A_tile_per_buffer_size + k * BLOCK_SIZE_M + threadIdx.y * NUM_PER_THREAD + 4]);FLOAT4(B_reg[0]) = FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + k * BLOCK_SIZE_N + threadIdx.x * NUM_PER_THREAD]);FLOAT4(B_reg[4]) = FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + k * BLOCK_SIZE_N + threadIdx.x * NUM_PER_THREAD + 4]);// use register to compute the sum of A_tile * B_tilefor(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}}}__syncthreads();}buffer_idx = buffer_idx ^ 1;for(int k = 0; k < BLOCK_SIZE_K; ++k){// compute the last tileFLOAT4(A_reg[0]) = FLOAT4(A_tile[buffer_idx * A_tile_per_buffer_size + k * BLOCK_SIZE_M + threadIdx.y * NUM_PER_THREAD]);FLOAT4(A_reg[4]) = FLOAT4(A_tile[buffer_idx * A_tile_per_buffer_size + k * BLOCK_SIZE_M + threadIdx.y * NUM_PER_THREAD + 4]);FLOAT4(B_reg[0]) = FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + k * BLOCK_SIZE_N + threadIdx.x * NUM_PER_THREAD]);FLOAT4(B_reg[4]) = FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + k * BLOCK_SIZE_N + threadIdx.x * NUM_PER_THREAD + 4]);for(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}}}    // write the result to Cfloat* C_start = C + blockIdx.y * BLOCK_SIZE_M * N + blockIdx.x * BLOCK_SIZE_N;for(int i = 0; i < NUM_PER_THREAD; ++i){FLOAT4(C_start[(threadIdx.y * NUM_PER_THREAD + i) * N + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(sum[i * NUM_PER_THREAD]);FLOAT4(C_start[(threadIdx.y * NUM_PER_THREAD + i) * N + threadIdx.x * NUM_PER_THREAD + 4]) = FLOAT4(sum[i * NUM_PER_THREAD + 4]);}
}

下面是該代碼的詳細分析：(from ChatGPT)

1. 核函數簽名與參數

template<unsigned int BLOCK_SIZE_M,unsigned int BLOCK_SIZE_N,unsigned int BLOCK_SIZE_K,unsigned int NUM_PER_THREAD>
__global__ void cuda_sgemm_v8_double_buffer(float* A, float* B, float* C, const int M, const int N, const int K)
  • 模板參數
    • BLOCK_SIZE_M = 128：每個線程塊處理的 M 維度大小
    • BLOCK_SIZE_N = 128：每個線程塊處理的 N 維度大小
    • BLOCK_SIZE_K = 8：每個線程塊處理的 K 維度大小
    • NUM_PER_THREAD：每個線程 x 方向和 y 方向分別處理的元素數量
  • 啟動參數：
    • block(16, 16)：每個線程塊包含 256 個線程
    • grid(4, 4)：整個網格包含 16 個線程塊
    • shared_mem_size：雙緩沖區所需共享內存

2. 輸入矩陣起始指針偏移

float* A_start = A + blockIdx.y * BLOCK_SIZE_M * K;
float* B_start = B + blockIdx.x * BLOCK_SIZE_N;
  • A_start：當前 block 塊處理的 A 矩陣的起始指針
  • B_start：當前 block 塊處理的 B 矩陣的起始指針

關于索引的計算，每個 block 要處理的是 A 中子矩陣 BLOCK_SIZE_M * K 與 B 中子矩陣 K * BLOCK_SIZE_N 的乘積。因此 A 矩陣的起始指針是 blockIdx.y * (BLOCK_SIZE_M * K)，其中 blockIdx.y 表示 y 方向上當前 block 的索引；B 矩陣的起始指針是 blockIdx.x * BLOCK_SIZE_N，其中 blockIdx.x 表示 x 方向上當前 block 的索引

在這里插入圖片描述

如上圖所示，block(2, 1) 處理的 A 矩陣和 B 矩陣的起始指針是 ? 的位置

3. 共享內存分配（雙緩沖）

extern __shared__ float shared_mem[];
int A_tile_per_buffer_size = BLOCK_SIZE_K * BLOCK_SIZE_M;  // 8*128=1024
int B_tile_per_buffer_size = BLOCK_SIZE_K * BLOCK_SIZE_N;  // 8*128=1024
float* A_tile = shared_mem;  // 緩沖區0的A_tile
float* B_tile = shared_mem + 2 * A_tile_per_buffer_size;  // 緩沖區0的B_tile
  • 分配共享內存用于矩陣 A 的兩個分塊緩沖區（A_tile）和矩陣 B 的兩個分塊緩沖區（B_tile）
  • 總共四個 tile 分塊，兩個用于預加載下一個塊（prefetch），兩個用于計算當前塊

共享內存布局為：

[A_tile_buffer_0, A_tile_buffer_1, B_tile_buffer_0, B_tile_buffer_1]

4. 寄存器初始化

float A_reg[NUM_PER_THREAD] = {0.0f};
float B_reg[NUM_PER_THREAD] = {0.0f};
float A_load_reg[4] = {0.0f};
float sum[NUM_PER_THREAD * NUM_PER_THREAD] = {0.0f};
  • 寄存器用于局部存儲：加載小塊、計算用值、累積結果

5. 線程布局重排（線程索引與 tile 索引轉換）

int tid = threadIdx.y * blockDim.x + threadIdx.x;  // 0~255
int A_tile_tx = tid % (BLOCK_SIZE_K / 4);         // BLOCK_SIZE_K=8 → 0~1
int A_tile_ty = tid / (BLOCK_SIZE_K / 4);         // 0~127
int B_tile_tx = tid % (BLOCK_SIZE_N / 4);         // BLOCK_SIZE_N=128 → 0~31
int B_tile_ty = tid / (BLOCK_SIZE_N / 4);         // 0~7
  • 將 256 個線程重新劃分為不同的加載工作組（如下圖所示）
  • A_tile 加載：128 個線程組（每個組 2 個線程）
  • B_tile 加載：8 個線程組（每個組 32 個線程）

在這里插入圖片描述

6. 預取一個數據塊

// prefetch first tile
FLOAT4(A_load_reg[0]) = FLOAT4(A_start[A_tile_ty * K + A_tile_tx * 4]);
A_tile[(A_tile_tx * 4 + 0) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[0]; 
A_tile[(A_tile_tx * 4 + 1) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[1]; 
A_tile[(A_tile_tx * 4 + 2) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[2]; 
A_tile[(A_tile_tx * 4 + 3) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[3]; FLOAT4(B_tile[B_tile_ty * BLOCK_SIZE_N + B_tile_tx * 4]) = FLOAT4(B_start[B_tile_ty * N + B_tile_tx * 4]);__syncthreads();
  • 使用 FLOAT4 宏從 global memory 讀取 A
  • A_tile 采用轉置存儲（列主序），同理 B_tile 也進行加載，實現計算階段的內存訪問連續性
  • 等待所有線程完成預加載

這里的重點是各個索引的計算，下面我們簡要分析下：(from ChatGPT)

6.1 A_tile 預取的索引計算

全局內存加載索引

FLOAT4(A_load_reg[0]) = FLOAT4(A_start[A_tile_ty * K + A_tile_tx * 4]);
  • A_start：當前線程塊對應的 A 矩陣起始指針
  • 索引分解：
    • A_tile_ty：范圍 0~127
    • A_tile_tx：范圍 0~1
  • 實際訪問：
    • 每個線程處理 4 個連續元素（FLOAT4）
    • 訪問位置：A_start + (A_tile_ty * K) + (A_tile_tx * 4)
    • 相當于：
      • 在行方向：A_tile_ty（0~127）
      • 在列方向：A_tile_tx * 4（0 或 4）
  • 線程分工：
    • 256 個線程分成 128 組（A_tile_ty = 0~127），每組 2 個線程（A_tile_tx = 0~1）
    • 每組線程負責加載 8 個連續元素（2 個 FLOAT4）

我們以第一個線程塊 block(0, 0) 為例來講解 A_tile 預取第一個數據塊部分的索引計算，如下圖所示：

在這里插入圖片描述

圖中灰色區域就是 block(0, 0) 線程塊需要加載的數據，總共 128x8 大小的元素數量，一個 block 包含 256 個線程，每個線程負責加載 4 個元素。線程索引的轉換在步驟 5 中來完成的，(16, 16)?(128, 2)

共享內存存儲索引（轉置關鍵）

A_tile[(A_tile_tx * 4 + n) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[n];
  • 存儲布局：
    • 將全局內存的行主序轉為共享內存的列主序
    • 公式：(col * BLOCK_SIZE_M) + row
  • 索引分解：
    • A_tile_tx * 4 + n：列索引（0~7）
    • A_tile_ty：行索引（0~127）
  • 轉置效果：
    • 全局內存中的行 A_tile_ty 變為共享內存中的列 A_tile_ty
    • 全局內存中的列 A_tile_tx * 4 + n 變為共享內存中的行 A_tile_tx * 4 + n

整個加載過程如下圖所示：

在這里插入圖片描述

6.2 B_tile 預取的索引計算

全局內存加載索引

... = FLOAT4(B_start[B_tile_ty * N + B_tile_tx * 4]);
  • B_start：當前線程塊對應的 B 矩陣的起始指針
  • 索引分解：
    • B_tile_ty：范圍 0~7
    • B_tile_tx：范圍 0~31
  • 實際訪問：
    • 每個線程處理 4 個連續元素（FLOAT4）
    • 訪問位置：B_start + (B_tile_ty * N) + (B_tile_tx * 4)
    • 相當于：
      • 在行方向：B_tile_ty（0~7）
      • 在列方向：B_tile_tx * 4（0~124）
  • 線程分工：
    • 256 個線程分成 8 組（B_tile_ty = 0~7），每組 32 個線程（B_tile_tx = 0~31）
    • 每組線程負責加載 128 個連續元素（32 個 FLOAT4）

我們以第一個線程塊 block(0, 0) 為例來講解 B_tile 預取第一個數據塊部分的索引計算，如下圖所示：

在這里插入圖片描述

圖中灰色區域就是 block(0, 0) 線程塊需要加載的數據，總共 8x128 大小的元素數量，一個 block 包含 256 個線程，每個線程負責加載 4 個元素。線程索引的轉換在步驟 5 中來完成的，(16, 16)?(32, 8)

共享內存存儲索引

B_tile[B_tile_ty * BLOCK_SIZE_N + B_tile_tx * 4] = ...
  • 存儲布局：
    • 保持行主序（不轉置）
    • 公式：(row * BLOCK_SIZE_N + col)
  • 索引分解：
    • B_tile_ty：行索引（0~7）
    • B_tile_tx：列索引（0~124）

整個加載過程如下圖所示：

在這里插入圖片描述

由于我們是行主序加載存儲，因此索引計算方式相比 A_tile 來說更加簡單

7. 主計算循環（雙緩沖核心）

int buffer_idx = 1;  // 初始緩沖區索引
for(int k_base = BLOCK_SIZE_K; k_base < K; k_base += BLOCK_SIZE_K){// 1. 預取下一塊到非活動緩沖區FLOAT4(A_load_reg[0]) = FLOAT4(A_start[A_tile_ty * K + A_tile_tx * 4 + k_base]);A_tile[buffer_idx * A_tile_per_buffer_size + (A_tile_tx * 4 + 0) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[0];// ...（寫入4個元素）FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + B_tile_ty * BLOCK_SIZE_N + B_tile_tx * 4]) = FLOAT4(B_start[(B_tile_ty + k_base) * N + B_tile_tx * 4]);// 2. 切換緩沖區索引buffer_idx = buffer_idx ^ 1;  // 0?1切換// 3. 計算當前塊（使用另一個緩沖區）for(int k = 0; k < BLOCK_SIZE_K; ++k){FLOAT4(A_reg[0]) = FLOAT4(A_tile[buffer_idx * A_tile_per_buffer_size + k * BLOCK_SIZE_M + threadIdx.y * NUM_PER_THREAD]);FLOAT4(A_reg[4]) = FLOAT4(/*... +4*/);  // 加載8個元素FLOAT4(B_reg[0]) = FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + k * BLOCK_SIZE_N + threadIdx.x * NUM_PER_THREAD]);FLOAT4(B_reg[4]) = FLOAT4(/*... +4*/);// 4. 8x8外積計算for(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}}}__syncthreads();
}
  • 流水線設計：
    • 當計算在使用 buffer0 時，異步加載下一塊到 buffer1
    • 下次迭代切換緩沖區，計算 buffer1 同時加載到 buffer0
  • 優勢：
    • 計算與內存傳輸重疊
    • 隱藏內存訪問延遲
    • 提高計算單元利用率

7.1 雙緩沖索引管理

int buffer_idx = 1;  // 初始緩沖區索引
for(int k_base = BLOCK_SIZE_K; k_base < K; k_base += BLOCK_SIZE_K){// ...預取和計算代碼...buffer_idx = buffer_idx ^ 1;  // 緩沖區切換
}
  • 初始值：buffer_idx = 1（因為第 0 個緩沖區已在預取階段填充）
  • 切換邏輯：buffer_idx ^ 1 在 0 和 1 之間切換
  • 雙緩沖工作流程：
    • 計算使用 buffer_idx 指向的緩沖區
    • 同時預取數據到 buffer_idx ^ 1 指向的緩沖區
    • 每次迭代后切換緩沖區

7.2 A_tile 預取索引（下一塊）

FLOAT4(A_load_reg[0]) = FLOAT4(A_start[A_tile_ty * K + A_tile_tx * 4 + k_base]);
A_tile[buffer_idx * A_tile_per_buffer_size + (A_tile_tx * 4 + n) * BLOCK_SIZE_M + A_tile_ty] = A_load_reg[n];

全局內存加載索引：

  • A_tile_ty * K：行偏移（0~127 行，每行跳 K 元素）
  • A_tile_tx * 4：列偏移（0 或 4）
  • k_base：當前 K 維度的基偏移（8，16，…，504）

實際訪問模式：

  • 每個線程加載全局內存中相隔 K 元素的 4 個連續 float
  • 整體訪問模式是跨步的但合并的（coalesced）

整個加載過程如下圖所示：

在這里插入圖片描述

那其實加載與緩沖區 0 的索引計算一樣，只是有一個 k_base 的偏移量，代表著處理下一個緩沖區

共享內存存儲索引：

  • buffer_idx * A_tile_per_buffer_size：選擇緩沖區（0 或 1024）
  • (A_tile_tx * 4 + n) * BLOCK_SIZE_M：列主序的列計算（0~7 * 128）
  • A_tile_ty：行索引（0~127）

在這里插入圖片描述

從寄存器到共享內存的加載如上圖所示，值得注意的是這里我們加載的是 buffer1 緩沖區，因此有一個 buffer_idx * A_tile_per_buffer_size 的偏移量存在

7.3 B_tile 預取索引（下一塊）

FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + B_tile_ty * BLOCK_SIZE_N + B_tile_tx * 4]) = 
FLOAT4(B_start[N * (B_tile_ty + k_base) + B_tile_tx * 4]);

全局內存加載索引：

  • B_tile_ty + k_base：行索引（0~7 + 8，16，…，504）
  • B_tile_tx * 4：列索引（0~124）
  • N * row + col：行主序訪問

訪問特點：

  • 每個線程加載 B 矩陣中連續的 4 個元素
  • 訪問模式是完全連續的

整個加載過程如下圖所示：

在這里插入圖片描述

和 A_tile 一樣，這里也有一個 k_base 的偏移量

共享內存存儲索引

  • buffer_idx * B_tile_per_buffer_size：選擇緩沖區（0 或 1024）
  • B_tile_ty * BLOCK_SIZE_N：行偏移（0~7 * 128）
  • B_tile_tx * 4：列偏移（0~124）

布局特點：

  • 保持行主序存儲
  • 與全局內存布局一致

在這里插入圖片描述

B_tile 緩沖區 1 從寄存器到共享內存的加載過程如上圖所示，同樣有一個 buffer_idx * B_tile_per_buffer_size 的偏移量存在

7.4 計算階段索引（當前塊）

buffer_idx = buffer_idx ^ 1;
for(int k = 0; k < BLOCK_SIZE_K; ++k){FLOAT4(A_reg[0]) = FLOAT4(A_tile[buffer_idx * A_tile_per_buffer_size + k * BLOCK_SIZE_M + threadIdx.y * NUM_PER_THREAD]);FLOAT4(A_reg[4]) = FLOAT4(/*...+4*/);FLOAT4(B_reg[0]) = FLOAT4(B_tile[buffer_idx * B_tile_per_buffer_size + k * BLOCK_SIZE_N + threadIdx.x * NUM_PER_THREAD]);FLOAT4(B_reg[4]) = FLOAT4(/*...+4*/);// ...外積計算...
}

A_tile 加載索引

  • buffer_idx * A_tile_per_buffer_size：選擇緩沖區
  • k * BLOCK_SIZE_M：K 維度偏移（0~7 * 128）
  • threadIdx.y * NUM_PER_THREAD：線程在 M 維度的偏移（0~15 * 8）

關鍵點：

  • 由于 A_tile 是轉置存儲的，這里實際上是按列連續訪問
  • 每次加載 8 個連續元素（2 個 FLOAT4）

B_tile 加載索引

  • buffer_idx * B_tile_per_buffer_size：選擇緩沖區
  • k * BLOCK_SIZE_N：K 維度偏移（0~7 * 128）
  • threadIdx.x * NUM_PER_THREAD：線程在 N 維度的偏移（0~15 * 8）

關鍵點：

  • 保持行主序訪問
  • 每次加載 8 個連續元素（2 個 FLOAT4）

從共享內存到寄存器的加載過程如下圖所示（以 thread(0,0) 為例）：

在這里插入圖片描述

7.5 外積計算索引

for(int i = 0; i < NUM_PER_THREAD; ++i){for(int j = 0; j < NUM_PER_THREAD; ++j){sum[i * NUM_PER_THREAD + j] += A_reg[i] * B_reg[j];}
}
  • A_reg 索引：i（0~7）
  • B_reg 索引：j（0~7）
  • sum 索引：i * 8 + j（0~63）

計算模式：

  • 8 元素 A 列向量 x 8 元素 B 行向量 ? 8x8 外積
  • 結果累加到 64 個局部和寄存器中

8. 處理最后一個數據塊

buffer_idx = buffer_idx ^ 1;  // 切換回最后一個緩沖區
for(int k = 0; k < BLOCK_SIZE_K; ++k){// 加載并計算最后一個塊FLOAT4(A_reg[0]) = FLOAT4(/*...*/);// ...（完整8x8外積計算）
}
  • 與前面一樣，只是不再 prefetch

9. 結果寫回

float* C_start = C + blockIdx.y * BLOCK_SIZE_M * N + blockIdx.x * BLOCK_SIZE_N;
for(int i = 0; i < NUM_PER_THREAD; ++i){FLOAT4(C_start[(threadIdx.y * NUM_PER_THREAD + i) * N + threadIdx.x * NUM_PER_THREAD]) = FLOAT4(sum[i * NUM_PER_THREAD]);FLOAT4(/*...+4*/) = FLOAT4(/*...+4*/);  // 寫回8個元素
}

9.1 輸出矩陣 C 的起始位置計算

float* C_start = C + blockIdx.y * BLOCK_SIZE_M * N + blockIdx.x * BLOCK_SIZE_N;
  • blockIdx.y 維度：
    • blockIdx.y * BLOCK_SIZE_M * N：計算當前線程塊在 M 維度的偏移
    • 每個線程塊處理 BLOCK_SIZE_M = 128 行
    • 乘以 N 得到正確的行偏移量（因為 C 是行主序）
  • blockIdx.x 維度：
    • blockIdx.x * BLOCK_SIZE_N：計算當前線程塊在 N 維度的偏移
    • 每個線程塊處理 BLOCK_SIZE_N = 128 行
  • 組合效果：
    • 定位到當前線程塊負責計算的 C 矩陣子塊的起始位置，和 A_start、B_start 類似

9.2 線程到輸出位置的映射

行索引計算

  • threadIdx.y * NUM_PER_THREAD + i：
    • threadIdx.y：線程在塊內的 y 坐標（0~15）
    • NUM_PER_THREAD = 8：每個線程負責 8 行
    • i：當前迭代（0~7）
    • 組合效果：0~127（覆蓋 BLOCK_SIZE_M = 128）

列索引計算

  • threadIdx.x * NUM_PER_THREAD：
    • threadIdx.x：線程在塊內的 y 坐標（0~15）
    • NUM_PER_THREAD = 8：每個線程負責 8 列
  • threadIdx.x * NUM_PER_THREAD + 4：
    • 額外的列偏移，用于處理每個線程的 8 列中的后 4 列

v8 版本通過雙緩沖共享內存、寄存器 blocking、數據預取（prefetching）與流水線方式來提高計算效率，其中：

  • 1. 雙緩沖共享內存：使用兩個緩沖區來重疊數據傳輸和計算
    • 一個緩沖區用于當前計算
    • 另一個緩沖區用于異步加載下一批數據
  • 2. 寄存器 blocking：每個線程使用寄存器緩存多個元素
  • 3. 數據預取：提前從全局內存加載下一 tile 的數據
  • 4. 流水線執行
    • 在計算當前分塊時， 異步加載下一個分塊
    • 通過 buffer_idx 在 0 和 1 之間切換，實現緩沖區輪換
  • 5. 性能優勢
    • 隱藏了全局內存訪問延遲
    • 計算和內存傳輸可以并行進行
    • 減少了線程等待時間

性能和帶寬測試結果如下：

優化手段矩陣維度GridBlock耗時(us)Memory Throughout(%)DRAM Throughout(%)
v0_global_memory512x512(32,32)(16,16)471.7896.941.56
v1_shared_memory256x256(16,16)(16,16)82.1178.921.84
v2_shared_memory_sliding_window512x512(32,32)(16,16)362.5094.457.05
v3_increase_work_of_per_thread512x512(16,16)(16,16)204.2684.013.64
v4_using_float4512x512(32,32)(4,16)209.6091.993.44
v5_register_outer_product512x512(32,32)(4,16)206.1879.103.50
v6_register_outer_product_float4512x512(8,8)(16,16)84.9960.387.28
v7_A_smem_transpose512x512(8,8)(16,16)118.2165.855.39
v8_double_buffer512x512(4,4)(16,16)135.7131.564.44

OK，以上就是 sgemm 各種優化的代碼實現了

結語

這篇文章中 sgemm 的一些優化技巧相比上篇文章來說復雜一些，博主經常被其中的索引計算搞破防，曾一度想放棄，不過靜下心來畫畫圖慢慢思考總是能理解的。在計算時可以先定位到當前 block 要處理的起始元素位置，然后思考 block 中每個 thread 負責處理幾個元素，都是怎么處理的，行和列索引分別是多少，這樣會相對簡單一些

OK，以上就是整篇文章的全部內容了

總的來說，跟隨 up 主一步步來實現還是能理解的，大家感興趣的可以多看看 up 主的視頻，還是非常不錯的🤗

下載鏈接

  • Sgemm 矩陣乘法代碼下載鏈接【提取碼：1234】

參考

  • 【CUDA】Sgemm單精度矩陣乘法（已完結~）
  • https://github.com/tpoisonooo/how-to-optimize-gemm/cuda
  • 深入淺出GPU優化系列：GEMM優化（一）
  • [施工中] CUDA GEMM 理論性能分析與 kernel 優化
  • cuda 入門的正確姿勢：how-to-optimize-gemm
  • CUDA 矩陣乘法終極優化指南
  • CUDA實現矩陣乘法的8種優化策略編程介紹
  • https://chatgpt.com/

本文來自互聯網用戶投稿，該文觀點僅代表作者本人，不代表本站立場。本站僅提供信息存儲空間服務，不擁有所有權，不承擔相關法律責任。
如若轉載，請注明出處：http://www.pswp.cn/bicheng/80992.shtml
繁體地址，請注明出處：http://hk.pswp.cn/bicheng/80992.shtml
英文地址，請注明出處：http://en.pswp.cn/bicheng/80992.shtml

如若內容造成侵權/違法違規/事實不符，請聯系多彩編程網進行投訴反饋email:809451989@qq.com，一經查實，立即刪除！

相關文章

【1000以內具有12個以上因子的整數并輸出它的因子】2021-12-27

【1000以內具有12個以上因子的整數并輸出它的因子】2021-12-27

緣由c語言輸入1000以內具有12個以上因子的整數 并輸出它的因子-編程語言-CSDN問答 int 求因子個數(int 數, int* 因子 { 0 }) {//緣由https://bbs.csdn.net/topics/399168406int 和 0, 求 1, 商 0, 含 0;//因子不含1和數本身while (求 < (商 數 / 求))if (!(數 % 求)…
閱讀更多...
C#中的dynamic與var：看似相似卻迥然不同

C#中的dynamic與var：看似相似卻迥然不同

在C#編程的世界里&#xff0c;var和dynamic這兩個關鍵字常常讓初學者感到困惑。它們看起來都在定義變量時省略了顯式類型聲明&#xff0c;但實際上它們的工作方式和應用場景有著天壤之別。今天&#xff0c;讓我們一起揭開這兩個關鍵字的神秘面紗。 var&#xff1a;編譯時的類型…
閱讀更多...
流速儀數據處理及流量斷面線繪制

流速儀數據處理及流量斷面線繪制

1 需求描述 在實際航道測量項目中&#xff0c;有測量斷面線流量流速的需求&#xff0c;得使用流速儀在現場進行測量&#xff0c;相關操作在之前已經寫了記錄。本次手冊記錄后期數據處理與流量線繪制&#xff0c;以該區域為例。 流速儀設備操作說明 2 規范要求 3 流量斷面表格…
閱讀更多...
購物車構件示例

購物車構件示例

通用購物車構件設計 注：代碼僅用于演示原理，不可用于生產環境。 一、設計目標 設計一個高度可復用的購物車構件，具備以下特點： 與具體業務系統解耦支持多種應用場景（商城、積分系統等）提供標準化接口易于集成和擴展二、核心架構設計 1. 分層架構 ┌─────────…
閱讀更多...
數據結構·字典樹

數據結構·字典樹

字典樹trie 顧名思義&#xff0c;在一個字符串的集合里查詢某個字符串是否存在樹形結構。 樹存儲方式上用的是結構體數組&#xff0c;類似滿二叉樹的形式。 模板 定義結構體和trie 結構體必須的內容&#xff1a;當前結點的字符&#xff0c;孩子數組可選&#xff1a;end用于查…
閱讀更多...
ES面試題系列「一」

ES面試題系列「一」

1、Elasticsearch 是什么&#xff1f;它與傳統數據庫有什么區別&#xff1f; 答案&#xff1a;Elasticsearch 是一個基于 Lucene 的分布式、開源的搜索和分析引擎&#xff0c;主要用于處理大量的文本數據&#xff0c;提供快速的搜索和分析功能。與傳統數據庫相比&#xff0c;E…
閱讀更多...
2025年6月一區SCI-不實野燕麥優化算法Animated Oat Optimization-附Matlab免費代碼

2025年6月一區SCI-不實野燕麥優化算法Animated Oat Optimization-附Matlab免費代碼

引言 近年來&#xff0c;在合理框架內求解優化問題的元啟發式算法的發展引起了全球科學界的極大關注。本期介紹一種新的元啟發式算法——不實野燕麥優化算法Animated Oat Optimization algorithm&#xff0c;AOO。該算法模擬了不實野燕麥的3種獨特行為&#xff0c;于2025年6月…
閱讀更多...
Agent Builder API - Agent Smith 擴展的后端服務（開源代碼）

Agent Builder API - Agent Smith 擴展的后端服務（開源代碼）

?一、軟件介紹 文末提供程序和源碼下載 Agent Builder API - Agent Smith 擴展的后端服務&#xff08;開源代碼&#xff09;手動設置&#xff1a;在本地計算機中克隆此存儲庫并啟動 python FAST API 服務器。&#xff08;可選&#xff09;安裝并設置 Mongo DB。Dev Container…
閱讀更多...
C及C++的SOAP協議庫

C及C++的SOAP協議庫

一.gSOAP gSOAP 是一個功能強大的開源工具包&#xff0c;專為 C 和 C 設計&#xff0c;用于快速開發基于 SOAP 協議的 Web 服務和客戶端。 1.協議支持 SOAP 版本&#xff1a;完整支持 SOAP 1.1/1.2 規范&#xff0c;包括消息格式、編碼規則和錯誤處理。 傳輸協議&#xff1a…
閱讀更多...
html5+css3實現傅里葉變換的動態展示效果(僅供參考)

html5+css3實現傅里葉變換的動態展示效果(僅供參考)

<!DOCTYPE html> <html lang"zh-CN"> <head><meta charset"UTF-8"><meta name"viewport" content"widthdevice-width, initial-scale1.0"><title>傅里葉變換的動態展示效果</title><sty…
閱讀更多...
ECharts中Map(地圖)樣式配置、漸變色生成

ECharts中Map(地圖)樣式配置、漸變色生成

前言 在日常開發中&#xff0c;ECharts 幾乎成了我們繪制數據圖表的標配工具&#xff0c;功能強大到幾乎無所不能。不過每次用的時候都要翻官方文檔查配置項&#xff0c;確實有點小繁瑣 &#x1f605; 為了提升效率&#xff0c;也方便以后快速復用&#xff0c;這里就整理記錄…
閱讀更多...
內存分配器ptmalloc2、tcmalloc、jemalloc，結構設計、內存分配過程詳解

內存分配器ptmalloc2、tcmalloc、jemalloc，結構設計、內存分配過程詳解

1. 引言 博主之前做過一個高并發內存池的項目實踐&#xff0c;在實踐中對于內存分配器的內存分配過程理解更加深刻了。在此期間&#xff0c;翻查了不少資料以及博客&#xff0c;發現源碼分享的博客不多&#xff0c;能生動完整的講述ptmalloc2、tcmalloc、jemalloc它們的結構設…
閱讀更多...
【擁抱AI】Deer-Flow字節跳動開源的多智能體深度研究框架

【擁抱AI】Deer-Flow字節跳動開源的多智能體深度研究框架

最近發現一款可以對標甚至可能超越GPT-Researcher的AI深度研究應用&#xff0c;Deer-Flow&#xff08;Deep Exploration and Efficient Research Flow&#xff09;作為字節跳動近期開源的重量級項目&#xff0c;正以其模塊化、靈活性和人機協同能力引發廣泛關注。該項目基于 La…
閱讀更多...
openfeign與dubbo調用下載excel實踐

openfeign與dubbo調用下載excel實踐

一、前言 openfeign和dubbo均是rpc框架 RPC&#xff08;Remote Procedure Call&#xff0c;遠程過程調用&#xff09;框架 是一種允許程序像調用本地方法一樣調用遠程服務器上函數的技術。它隱藏了底層網絡通信的復雜性&#xff0c;讓開發者可以專注于業務邏輯&#xff0c;實現…
閱讀更多...
解密企業級大模型智能體Agentic AI 關鍵技術：MCP、A2A、Reasoning LLMs-強化學習算法

解密企業級大模型智能體Agentic AI 關鍵技術：MCP、A2A、Reasoning LLMs-強化學習算法

解密企業級大模型智能體Agentic AI 關鍵技術&#xff1a;MCP、A2A、Reasoning LLMs-強化學習算法 現在我們的核心問題是有一些同學會知道要才能強化學習。為什么才能強化學習&#xff1f;是實現AGI。例如從這個其實你從第一階段開始以后&#xff0c;就是chatbot&#xff0c;這…
閱讀更多...
音頻分類的學習

音頻分類的學習

1.深度學習PyTorch入門-語音分類 https://blog.csdn.net/sinat_41787040/article/details/129795496 https://github.com/musikalkemist/pytorchforaudio https://github1s.com/musikalkemist/pytorchforaudio/blob/main/04%20Creating%20a%20custom%20dataset/urbansoundda…
閱讀更多...
美SEC主席：探索比特幣上市證券交易所

美SEC主席：探索比特幣上市證券交易所

作者/演講者&#xff1a;美SEC主席Paul S. Atkins 編譯&#xff1a;Liam 5月12日&#xff0c;由美國SEC加密貨幣特別工作組發起的主題為《資產上鏈&#xff1a;TradFi與DeFi的交匯點》系列圓桌會議如期舉行。 會議期間&#xff0c;現任美SEC主席Paul S. Atkins發表了主旨演講。…
閱讀更多...
Qt file文件操作詳解

Qt file文件操作詳解

1.引言 很多應用程序都具備操作文件的能力&#xff0c;包括對文件進行寫入和讀取&#xff0c;創建和刪除文件等等&#xff0c;甚至某些應用程序的就是為了操作文件&#xff0c;像WPS Office。基于此Qt框架中專門提供了對文件操作的類&#xff1a;QFile。 2.QFile文件操作 QF…
閱讀更多...
【測試開發知識儲備】之Jacoco（Java Code Coverage）

【測試開發知識儲備】之Jacoco（Java Code Coverage）

文章目錄 Jacoco是什么Jacoco的主要功能&#xff08;一&#xff09;多樣化覆蓋率指標分析&#xff08;二&#xff09; 豐富的報告生成&#xff08;三&#xff09;實時數據收集 Jacoco的工作原理&#xff08;一&#xff09;字節碼增強&#xff08;二&#xff09;測試執行與數據收…
閱讀更多...
Docker 介紹與使用

Docker 介紹與使用

Docker 文章目錄 Docker介紹與虛擬機的比較啟動速度占用資源 優勢更容易遷移更容易維護更容易擴展 使用場景持續集成提供可伸縮的云服務搭建微服務架構 鏡像與容器鏡像構成&#xff08;分層結構&#xff09;鏡像與容器的區別 安裝 Docker常用命令介紹鏡像相關容器相關 實戰&…
閱讀更多...
最新文章

Copyright @ 2022~2023