当前位置： 首页 > news >正文

【CUDA】Sgemm单精度矩阵乘法（下）

news 来源：原创 2025/6/28 2:16:15

目录

    • 前言
    • 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/

相关文章：

【CUDA】Sgemm单精度矩阵乘法（下）

目录 前言1. 优化技巧5&#xff1a;使用register模拟二级缓存&#xff08;内积转外积&#xff09;2. 优化技巧6&#xff1a;使用register模拟二级缓存 float43. 优化技巧7&#xff1a;global memory转置再存放shared memory4. 优化技巧8&#xff1a;使用double buffer加速矩阵…...

编程日记 2025/6/28 2:16:15

cursor 学习

参考&#xff1a;AI编程神器&#xff01;Cursor无限续杯&#xff01;白嫖白嫖&#xff01;&#xff01;&#xff01;...

编程日记 2025/6/23 1:37:13

学术论文的科研流程概述 视频会议记录

CCF-Talk SPP131期 浙江大学研究员彭思达的报告。 举例视频生成要多快好省。 提升代码能力&#xff1a;先明白基础的函数&#xff0c;可以复现一个网络。最好是实现一个操作系统。...

编程日记 2025/6/26 12:38:16

【Linux笔记】——Linux线程理解与分页存储的奥秘

&#x1f525;个人主页&#x1f525;&#xff1a;孤寂大仙V &#x1f308;收录专栏&#x1f308;&#xff1a;Linux &#x1f339;往期回顾&#x1f339;&#xff1a;【Linux笔记】——进程信号的捕捉——从中断聊聊OS是怎么“活起来”的 &#x1f516;流水不争&#xff0c;争的…...

编程日记 2025/6/25 11:40:40

ACM算法

在ACM模式下使用JavaScript/TypeScript获取输入值 在ACM编程竞赛或在线判题系统(如LeetCode、牛客网等)中&#xff0c;JavaScript/TypeScript需要特定的方式来获取输入值。以下是几种常见的获取输入的方法&#xff1a; 1. 使用Node.js的readline模块 这是最常见的处理ACM模式…...

编程日记 2025/6/23 13:34:16

家用或办公 Windows 电脑玩人工智能开源项目配备核显的必要性（含 NPU 及显卡类型补充）

一、GPU 与显卡的概念澄清 首先需要明确一个容易误解的概念&#xff1a;GPU 不等同于显卡。 显卡和GPU是两个不同的概念。 【概念区分】 在讨论图形计算领域时&#xff0c;需首先澄清一个常见误区&#xff1a;GPU&#xff08;图形处理单元&#xff09;与显卡&#xff08;视…...

编程日记 2025/6/27 2:11:12

FastByteArrayOutputStream和ByteArrayInputStream有什么区别

FastByteArrayOutputStream 和 ByteArrayInputStream 是两种完全不同的 Java I/O 类&#xff0c;它们的主要区别体现在 设计目的 和 使用场景 上。以下是详细对比&#xff1a; 1. 核心区别总结 特性FastByteArrayOutputStream (Spring框架)ByteArrayInputStream (JDK原生)所属…...

编程日记 2025/6/27 11:37:12

远程连接电脑的方法？异地远程桌面连接和三方软件实现

远程连接电脑&#xff0c;是指通过网络技术&#xff0c;在一台设备上操控另一台设备的电脑桌面&#xff0c;实现跨地域的操作和管理。在日常工作、技术支持、远程办公等场景中&#xff0c;远程连接电脑都发挥着重要作用。实现远程连接电脑主要有系统自带工具和第三方软件两种方…...

编程日记 2025/6/22 6:59:10

编程题 03-树2 List Leaves【PAT】

文章目录 题目输入格式输出格式输入样例输出样例 题解解题思路完整代码 编程练习题目集目录 题目 Given a tree, you are supposed to list all the leaves in the order of top down, and left to right. 输入格式 Each input file contains one test case. For each case, …...

编程日记 2025/5/31 9:15:34

数据预处理之数据平滑处理详解

信号数据收到噪声干扰&#xff0c;影响检测的准确性。数据平滑处理的关键步骤&#xff0c;旨在降低噪声同时保留信号特征。 1.1 移动平均&#xff08;Moving Average&#xff09; 原理&#xff1a;通过计算窗口内数据的平均值来平滑噪声&#xff0c;适用于快速去除高频噪声。…...

编程日记 2025/6/23 18:08:42

deepseek梳理java高级开发工程师算法面试题

Java高级工程师算法面试题与答案 一、数据结构与算法基础 1. 红黑树与AVL树比较 题目&#xff1a;详细说明红黑树和AVL树的区别及各自的适用场景&#xff0c;并用Java实现红黑树的插入操作。 答案&#xff1a; 区别对比&#xff1a; ┌─────────────────…...

编程日记 2025/6/25 11:42:46

【SSL证书系列】SSL证书工作原理解读

SSL&#xff08;Secure Sockets Layer&#xff09;及其继任者TLS&#xff08;Transport Layer Security&#xff09;是用于保护网络通信安全的加密协议。SSL证书是实现HTTPS协议的核心&#xff0c;其工作原理涉及加密技术、身份验证和信任机制。以下是其工作原理的详细分步解析…...

编程日记 2025/6/25 11:38:46

模板源码建站、定制建站和SaaS 建站有什么区别？企业建站应该怎么选？

最近遇到不少客户问&#xff0c;为什么现在做一个网站为什么从几百到几万的都有呀&#xff1f;市面上五花八门有模板源码建站、SaaS建站和定制建站我该怎么选&#xff1f;有什么区别&#xff1f;今天小编就跟大家一起来唠一唠&#xff0c;接下来我们就一起来看看吧&#xff01;…...

编程日记 2025/6/26 4:44:08

OpenCV进阶操作：人脸检测、微笑检测

文章目录 前言一、OpenCV如何实现人脸检测1、haar特征2、级联分类器3、级联分类器的使用 二、人脸检测、微笑检测 案例实现1、预处理2、加载分类器3、标注人脸4、运行结果&#xff1a;4、微笑检测 总结 前言 要实现人脸识别首先要判断当前图像中是否出现了人脸&#xff0c;这就…...

编程日记 2025/6/26 7:16:12

论文查询的ai工具 —— SCAICH

&#xff08;1&#xff09;SCAICH的项目背景 SCAICH是由Scihub Web3 Community孵化的技术产品。SCAICH是一个非盈利性的平台&#xff0c;模式上采用免费邀请码模式&#xff0c;采用捐赠和广告维持成本。产品将会面向世界上所有国家的学者。 &#xff08;2&#xff09;SCAICH产品…...

编程日记 2025/6/24 19:53:23

Python+大模型 day01

Python基础 计算机系统组成 基础语法 如:student_num 4.标识符要做到见名知意,增强代码的可读性 关键字 系统或者Python定义的,有特殊功能的字符组合 在学习过程中,文件名没有遵循标识符命名规则,是为了按序号编写文件方便查找复习 但是,在开发中,所有的Python文件名称必须…...

编程日记 2025/6/22 9:11:21

elasticsearch硬件与资源配置优化

以下是Elasticsearch硬件与资源配置优化的综合方案，结合最新实践与核心优化逻辑： 一、硬件选型优化 ‌存储设备‌ 优先选用SSD作为存储介质，其随机读取性能比机械硬盘高5-10倍，尤其适合文档检索类高并发场景。单节点存储控制在2TB以内，避免超过5TB导致查询性能下降和系统…...

编程日记 2025/6/22 7:14:27

C++ 在 Windows 的开发经验与解决方案

一、开发环境搭建 在 Windows 上进行 C 开发&#xff0c;主流的集成开发环境&#xff08;IDE&#xff09;有 Visual Studio 和 CLion。Visual Studio 是微软官方推出的强大开发工具&#xff0c;对 Windows 平台有着原生的支持&#xff0c;集成了编译器、调试器、代码编辑器等一…...

编程日记 2025/6/20 1:55:52

1669上什么课

1.题目描述 暑假来了&#xff0c;晶晶报了四门课来充实自己的暑假生活&#xff1b;周一上游泳&#xff0c;周三上编程&#xff0c;周五上阅读&#xff0c;周六上数学&#xff1b;其余时间没课。请从键盘读入今天是星期几&#xff0c;输出晶晶今天应该上什么课。 请注意&#…...

编程日记 2025/6/23 12:30:47

通过MCP让LLM调用系统接口

场景 MCP的出现大大丰富了LLM的功能&#xff0c;对于存量系统&#xff0c;我们希望能让模型调用已有的接口&#xff0c;以最小的成本让AI能够获取系统内部数据。因此我们开发了一个名为http-api-call的MCP Server&#xff0c;来支持模型到内部API的调用 实现方案 使用用标准…...

编程日记 2025/6/23 2:03:23

Java NIO 深度解析：突破传统IO的性能瓶颈

一、Java NIO 核心价值与演进历程 1.1 传统IO的局限性 Java传统的BIO（Blocking I/O）模型在应对高并发场景时存在显著缺陷： 线程资源浪费：每个连接需要独立线程处理上下文切换开销：线程数增加导致CPU调度成本指数级增长吞吐量瓶颈：受限于线程池大小和操作系统限制响应延…...

编程日记 2025/6/25 18:25:08

AI-02a5a5.神经网络-与学习相关的技巧-权重初始值

权重的初始值 在神经网络的学习中&#xff0c;权重的初始值特别重要。实际上&#xff0c;设定什么样的权重初始值&#xff0c;经常关系到神经网络的学习能否成功。 不要将权重初始值设为 0 权值衰减&#xff08;weight decay&#xff09;&#xff1a;抑制过拟合、提高泛化能…...

编程日记 2025/6/23 10:41:46

sqlalchemy库详细使用

SQLAlchemy 是 Python 中最强大、最受欢迎的 ORM&#xff08;对象关系映射&#xff09;库&#xff0c;它允许你使用 Python 对象来操作数据库&#xff0c;而不需要直接编写 SQL 语句。同时&#xff0c;它也提供了对底层 SQL 的完全控制能力&#xff0c;适用于从简单脚本到大型企…...

编程日记 2025/6/24 16:50:49

最短路和拓扑排序知识点

1、在一个有权无向图中&#xff0c;如果顶点b到顶点a的最短路径长度是10&#xff0c;顶点c与顶点b之间存在一条长度为3的边。&#xff08;c与a的最短路径长度不超过13&#xff1b;c与a的最短路径不小于7&#xff09; 2、我们用一个有向图来表示航空公司所有航班的航线。最适合…...

编程日记 2025/6/25 8:02:16

【Alist+RaiDrive挂载网盘到本地磁盘】

1.安装准备 安装RaiDrive RaiDrive - 像 USB 驱动器一样安装云存储 安装alist 安装方式请查看官网: AList文档 2.启动Alist(docker) docker官网 Install | Docker EngineDocker Desktop | Docker Docs 运行容器 docker run -d --restartalways -v /home/alist:/opt/alist/…...

编程日记 2025/6/21 21:20:48

达梦数据库 【-6111: 字符串转换出错】问题处理

达梦数据库 【-6111: 字符串转换出错】问题处理 问题背景问题分析问题总结 问题背景 今天在更新数据库某一个值属性的时候&#xff0c;执行更新语句报错提示 -6111: 字符串转换出错&#xff0c;但是自己检查了sql语句&#xff0c;只是一个简单的sql&#xff0c;并没有需要字符…...

编程日记 2025/6/22 19:53:06

Java的多线程笔记

创建一个线程的方法有多种&#xff0c;比如可以继承Thread类或者实现Runnable接口&#xff0c;结论是实现Runnable接口比前者更加优越。 二者代码对比 Java 不支持多继承&#xff0c;如果你继承了 Thread 类&#xff0c;就不能再继承其他类&#xff0c;实现 Runnable 接口后&am…...

编程日记 2025/6/27 16:18:37

学习51单片机01（安装开发环境）

新学期新相貌.......哈哈哈&#xff0c;我终于把贪吃蛇结束了&#xff0c;现在我们来学stc51单片机&#xff01; 要求&#xff1a;c语言的程度至少要到函数&#xff0c;指针尽量&#xff01;如果c语言不好的&#xff0c;可以回去看看我的c语言笔记。 1.开发环境的安装&#x…...

编程日记 2025/6/24 19:38:47

互联网协议的多路复用、Linux系统的I/O模式

目录 1. 互联网协议栈-多路复用 1.1. 应用层的多路复用 2.2. 传输层的多路复用 3.3. 网络层的多路复用 2. Linux系统的I/O模式 2.1. I/O 2.2. Socket 2.3. 从网卡到操作系统 2.4. Socket 编程模型 2.5. I/O多路复用 2.6. 阻塞/非阻塞、同步/异步 2.7. Question 1. …...

编程日记 2025/6/27 19:39:42

vue中，created和mounted两个钩子之间调用时差值受什么影响

在 Vue 中&#xff0c;created 和 mounted 是两个生命周期钩子&#xff0c;它们之间的调用时差主要受以下几个因素影响&#xff1a; &#x1f7e2; 1. 模板复杂度与渲染耗时&#xff08;最主要因素&#xff09; mounted 的触发时间是在组件的 DOM 被挂载之后&#xff08;也就是…...

编程日记 2025/6/24 1:37:10

软件设计师考试《综合知识》计算机编码考点分析——会更新软设所有知识点的考情分析，求个三连

2019-2023年真题深度解析与备考策略 分值占比分析 75分中编码相关分值分布与核心考点 年份编码相关题量分值占总分比例核心考点20232题2分2.67%补码表示范围、IEEE 754偏移量20223题3分4.00%原码/反码比较、浮点数规格化20211题1分1.33%补码表示-1的能力20202题2分2.67%移码…...

编程日记 2025/6/26 1:06:41

剖析提示词工程中的递归提示

递归提示：解码AI交互的本质，构建复杂推理链 递归提示的核心思想，正如示例所示，是将一个复杂任务分解为一系列更小、更易于管理、逻辑上前后关联的子任务。每个子任务由一个独立的提示来驱动，而前一个提示的输出（经过必要的解析和转换）则成为下一个提示的关键输入。这种…...

编程日记 2025/6/25 16:32:41

【SSL证书系列】https双向认证中客户端认证的原理

HTTPS双向认证&#xff08;也称为双向SSL/TLS认证&#xff09;是一种增强安全性的机制&#xff0c;其中客户端和服务器都需要验证彼此的数字证书&#xff0c;以确保双方身份的真实性。以下是其核心原理和步骤的详细解析&#xff1a; 一、双向认证的核心目标 双向身份验证&#…...

编程日记 2025/6/24 21:41:59

map格式可以接收返回 fastjson2格式的数据 而不需要显示的转换

Fastjson2 JSONObject 与 Map 的关系 Fastjson2 的 JSONObject 类定义如下&#xff1a; public class JSONObject extends JSON implements Map<String, Object>, Cloneable {// 实现了 Map 接口的所有方法&#xff08;put、get、keySet 等&#xff09; }解释&#xff…...

编程日记 2025/6/27 14:01:37

NHANES稀有指标推荐：PWI

文章题目&#xff1a;Association between plain water intake and the risk of osteoporosis among middle-aged and elderly people in the United States: a cross-sectional study DOI&#xff1a;10.3389/fnut.2025.1527771 中文标题&#xff1a;美国中老年人白开水摄入与…...

编程日记 2025/6/27 11:24:10

CN 第二章 应用层-单选题

非并行TCP连接 HTTP非持续连接 假定在同一Web服务器上的某HTML文件引用了3个非常小的对象&#xff08;例如图片&#xff09;。忽略传输时延&#xff0c;往返时延为RTT&#xff0c;不考虑连接释放时间&#xff0c;采用非并行TCP连接的HTTP非持续连接方式将该页面完整接收下来需…...

编程日记 2025/5/26 9:16:42

游戏引擎学习第279天:将实体存储移入世界区块

黑板讲解&#xff1a;为什么使用SOA&#xff08;结构体数组&#xff09;而不是AOS&#xff08;数组结构体&#xff09;来构建实体系统 我们在构建游戏实体系统时&#xff0c;探讨了使用结构体数组&#xff08;SOA, Struct of Arrays&#xff09;而不是结构体组成的数组&#x…...

编程日记 2025/6/26 23:56:18

zabbix7.2最新版本 nginx自定义监控(三) 设置触发器

安装zabbix-get服务 在zabbix-server端口安装zabbix-get服务 [rootlocalhost ~]# dnf install -y zabbix-get Last metadata expiration check: 1:55:49 ago on Wed 14 May 2025 09:24:49 AM CST. Dependencies resolved. Package Architectur…...

编程日记 2025/6/23 5:46:40

解密企业级大模型智能体Agentic AI 关键技术：MCP、A2A、Reasoning LLMs- OpenAI AGI 五阶段

解密企业级大模型智能体Agentic AI 关键技术：MCP、A2A、Reasoning LLMs- OpenAI AGI 五阶段 然后第三个阶段就是agent，注意这里面的agent和我们说应用程序开发的这个agent是一个不同的概念。AI just can take actions autonomously自动的去执行一些动作。但大家像今天我们看到…...

编程日记 2025/6/25 1:24:42

Flink实时统计任务CPU异常排查与解决方案

一、核心原因分析 ‌资源配置不合理‌ ‌CPU核数与并行度不匹配‌：TaskManager的taskmanager.numberOfTaskSlots设置过高，导致单个节点负载过载（如32核节点设置2个slot被多个任务占用，总需求超过物理CPU核数）。‌内存与CPU分配不均‌：内存不足引发频繁GC，间接导致CPU利…...

编程日记 2025/6/25 22:44:40

Vue3指令（二）--v-text、v-html数据渲染，计算属性

目录 &#xff08;一&#xff09;数据渲染 1.插值表达式渲染数据 1.1实战案例 1.1.1代码&#xff1a; 1.1.2实现截图&#xff1a; 2.使用v-text和v-html来渲染数据 2.1实战案例&#xff1a; 2.1.1代码&#xff1a; 2.1.2实现截图&#xff1a; &#xff08;二&#xff…...

编程日记 2025/6/26 1:22:49

【深入Spring系列】源码级深入剖析SpringBoot如何实现自动装载

1. SpringBoot自动装载 Spring Boot 实现“自动装载”&#xff08;Auto Configuration&#xff09;是其最核心、最强大的功能之一&#xff0c;使得开发者可以快速搭建项目而无需进行复杂的 XML 配置。这一机制的底层实现主要依赖于 Spring Framework 的条件注解机制 和 Spring…...

编程日记 2025/6/26 18:14:42

【AI News | 20250514】每日AI进展

AI Repos 1、ocr-workbench OCR Workbench 是一款使用 AI&#xff08;Gemini 或 Tesseract&#xff09;进行文档光学字符识别&#xff08;OCR&#xff09;并生成 Markdown 或 HTML 转录的开源 Web 应用。它专为处理需要大量编辑的 OCR 文本而设计&#xff0c;特别是老旧文档。…...

编程日记 2025/6/21 17:00:19

嵌入式设计模式基础--C语言的继承封装与多态

继承&#xff0c;封装和多态是OOP的三大核心特性&#xff0c;它们共同构了面向对象的基础.但嵌入式开发中大量的使用到的却是C语言这种面向过程的语言&#xff0c;那么我们就需要了解如何在C中使用设计模式的思想做功能开发。要了解设计模式&#xff0c;我们就需要先搞清楚 继承…...

编程日记 2025/6/26 21:07:06

【python爬虫】python+selenium实现Google Play Store应用信息爬虫+apk下载

实验要求&#xff1a;利用pythonselenium实现Google Play Store应用信息爬虫apk下载。 其中&#xff1a; 1、热门应用列表包含200个app&#xff0c;需要点击右侧按钮滑动产生下一页数据&#xff0c;所以需要Selenium来控制页面操作。 2、每个应用的爬虫信息包括&#xff1a;ap…...

编程日记 2025/6/24 13:49:49

RPC协议及库介绍

一.RPC介绍 RPC(Remote Procedure Call)&#xff0c;远程过程调用协议&#xff0c;客户端在不知道调用细节的情况下&#xff0c;调用存在于远程计算机上的某个对象&#xff0c;就像调用本地应用程序中的对象一样&#xff0c;即允许像调用本地服务一样调用远程服务。 RPC框架的…...

编程日记 2025/6/24 4:56:41

【教程】Docker更换存储位置

转载请注明出处&#xff1a;小锋学长生活大爆炸[xfxuezhagn.cn] 如果本文帮助到了你&#xff0c;欢迎[点赞、收藏、关注]哦~ 目录 背景说明 更换教程 1. 停止 Docker 服务 2. 创建新的存储目录 3. 编辑 Docker 配置文件 4. 迁移已有数据到新位置 5. 启动 Docker 服务 6…...

编程日记 2025/6/25 13:55:49

vue3实现JSON格式化和JSONPath提取功能

功能简介 1、JSON数据的格式化 2、通过JSONPath语法对格式化后的数据匹配提取 基础环境参考 vue3flasksqlite前后端项目实战 包安装 npm install jsonpath src/views/JsonFormat.vue <template><div class"json-formatter-container"><el-card cla…...

编程日记 2025/6/25 1:28:20

【springcloud学习(dalston.sr1)】服务消费者通过restTemplate来访问服务提供者（含源代码）（五）

该系列项目整体介绍及源代码请参照前面写的一篇文章​​​​​​【springcloud学习(dalston.sr1)】项目整体介绍&#xff08;含源代码&#xff09;&#xff08;一&#xff09; 一般情况下&#xff0c;我们远程调用服务&#xff0c;可以用restTemplate来进行http请求的访问。接…...

编程日记 2025/6/24 10:08:25

在 Angular 中， `if...else if...else`

在 Angular 中&#xff0c;模板语法本身并不直接支持 if...else if...else 这样的多条件分支结构。不过&#xff0c;你可以通过使用 *ngIf 指令结合其else模板功能来实现类似的效果。下面是如何模拟if...else if...else逻辑的方法&#xff1a; 示例&#xff1a;实现if...else …...

编程日记 2025/6/25 22:11:38