CUDA手写SGEMM优化实战：从0.4到10.23 TFLOPS的完整优化链路-北京尧图网络

先说结论我没用 Tensor Core没上 cuBLAS就是老老实实用 CUDA C 手写 SGEMM单精度矩阵乘法在 RTX 3090 上把性能从最初的 0.4 TFLOPS 一路提到 10.23 TFLOPS也就是跨过了 10000 GFLOPS 这道坎。整个过程里最值钱的不是最后那个数字而是那几步看似简单、却经常被忽略的优化手法——尤其是 float4 向量化这一步直接砍掉了四分之三的访存指令是我这次实录里收益最高的一刀。这篇文章把这套优化链路完整拆开讲清楚适合刚入门 CUDA 想写高性能算子的同学也适合那些手上有卡、想看看自己写的矩阵乘到底还有多少余量的朋友。1. 为什么是矩阵乘法为什么是 10000 GFLOPS1.1 一个能打的性能基准矩阵乘法在 GPU 高性能计算里的地位有点像 LeetCode 在算法面试里的地位——你不能说它代表一切但几乎所有的深度学习算子、数值计算库、图形学变换底层都能拆成矩阵乘。Transformer 里的 QKV 投影、全连接层本质就是一个大规模 GEMM卷积在 im2col 之后也变成一个隐式的矩阵乘。所以把 SGEMM 优化明白等于拿到了 GPU 性能优化的一张通用入场券。我这次定的目标是单精度矩阵乘法 C A * B规模设为 M N K 4096。为什么是 4096因为它是 2 的幂天然对齐向量化需要的 16 字节边界也方便我们后面讲 tile 划分时做整数除法。总计算量是 2 * 4096³ ≈ 137.4 GFLOP这个数很好记。如果能在 13.4 毫秒内跑完GFLOPS 就是 137.4e9 / 13.4e-3 10.25 TFLOPS刚好跨线。10000 GFLOPS 即 10 TFLOPS 不是随便拍的数字。在一张 2020 年发布的 Ampere 架构显卡上不使用 Tensor Core 的情况下这个成绩代表你已经开始摸到 CUDA core 的调度上限了。它足够让一个认真做优化的开发者在 cuBLAS 面前挺直腰杆说话同时又不至于高到需要动用 cuBLAS 级别的高级技巧——对大多数学习者来说这是一个跳一跳够得着的目标。1.2 硬件基线RTX 3090 与理论峰值这次实测平台是一张 RTX 3090GA102 核心82 个 SMFP32 理论峰值大概是 35.6 TFLOPS。驱动的 CUDA 版本我用的是 12.x编译参数是nvcc -O3 -archsm_86。可能有人会问40 系显卡是不是能跑更高答案是肯定的RTX 40 系同样代码可以直接编译到 sm_89向量化和共享内存优化的思路完全一致只是理论峰值更高。这里用 3090 说明是因为它的 FP32 峰值比较实在——不像 40 系那样让新手误以为随便写写就能突破 20 TFLOPS。看性能不能只看绝对数字要看“利用率”。10.23 TFLOPS 除以理论峰值 35.6 TFLOPS利用率大约是 28.7%。这个比例在 CUDA core 手写 SGEMM 里属于“你确实做了有效优化”的区间而且这也正是向量化共享内存分块应该达到的水平。如果你看到某篇文章说自己随手写的 naive kernel 就到了 90% 峰值那基本只有两种可能要么上了 Tensor Core要么它偷偷在矩阵里塞了零没有参与计算。2. 第一版naive kernel 跑出的真实水平2.1 没有任何技巧的裸奔代码先把最直觉的版本写出来每个线程负责计算 C 的一个元素循环遍历 K 累加。__global__ void sgemm_naive(const float* A, const float* B, float* C, int M, int N, int K) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row M || col N) return; float sum 0.0f; for (int k 0; k K; k) { sum A[row * K k] * B[k * N col]; } C[row * N col] sum; }启动配置用 16x16 的线程块4096x4096 的矩阵就需要 256x256 个块。这代码逻辑上完全正确跑出来的结果和 CPU 版一致但性能惨不忍睹。我在 3090 上实测耗时大概在 320 毫秒上下换算过来 0.4 TFLOPS 左右——利用率只有理论峰值的 1.1%。很多人第一次跑这个 kernel 会有疑问明明显卡每秒能做 35.6 万亿次浮点运算为什么实际只有 0.4答案很简单计算单元在大部分时间里是饿着肚子的它没有数据可供计算。2.2 为什么 naive 会被内存带宽锁死naive kernel 的性能瓶颈不在算力而在全局内存访问量。算一笔账每个输出元素需要读一次 A 的元素和一次 B 的元素。整个矩阵乘一共要计算 M * N * K 个乘加最终读进计算单元的全局内存数据量大约是 2 * 4096³ * 4 字节 549GB。RTX 3090 的显存带宽是 936GB/s即使把带宽用到 100%理论极限也只有 549GB / 936GB/s 对应的 0.25 TFLOPS 左右。实测到的 0.4 TFLOPS 比这个理论极限还高一点是因为编译器做了局部缓存优化把一些 B 的列向量放进了 L1/L2但总体量级没变。这个分析告诉我们一个铁律如果每次乘加都要从全局内存取数那你的 SGEMM 永远被锁死在带宽附近再好的计算单元也帮不上忙。所以第一阶段的结论是naive kernel 的最大价值不是性能而是作为正确性基准。我后面每一版优化都要跟它的输出做全量对拍确认结果在误差范围内一致——优化是有理有据的不是靠破坏正确性换数字。3. 共享内存 Tiling让计算不再等饭3.1 Tile 分块的核心思想解决带宽问题的唯一思路是数据复用。把 A 和 B 切成一个个 tile分块先把一块数据从全局内存搬到共享内存然后让这一块数据被多个线程反复使用。共享内存就在 SM 内部延迟只有几十个时钟周期带宽是全局内存的几十倍这是 GPU 上唯一适合做数据中转的高速缓存。打个比方naive 版本相当于每做一餐饭都去几公里外的超市买一斤菜tiling 版本是先开车把一周的菜一次性买回来放到冰箱里共享内存接下来每顿饭只需要在厨房里取菜就行。这里“厨房”就是共享内存“超市”就是全局内存“菜”就是 A 和 B 的数据块。具体做法是每个 block 负责输出一个 BM x BN 的子块 C_tile为了算出这个子块需要从 K 维度上逐步加载 A 的 BM x BK 子块和 B 的 BK x BN 子块到共享内存。这个 BK 是一个滑动窗口沿着 K 方向一段一段地前进。每次加载一个子块块内所有线程都可以复用这 BM BN 个元素来更新自己的部分和。3.2 Tile 参数怎么定参数选择的核心矛盾是共享内存越大、复用次数越多但每个 SM 能同时驻留的 block 数就越少。共享内存总量是有限的3090 每个 SM 有 99KB 可用共享内存一块用掉 16KB就只能放下大约 6 个 block。如果一块用掉 50KB可能就只能放 1 个 blockSM 的并发度会显著下降。我最终定的参数是 BM 128、BN 128、BK 16。每个 block 负责 128x128 的输出块沿 K 每次推进 16。这样共享内存占用是A_tile 128 * 16 * 4 字节 8KBB_tile 16 * 128 * 4 字节 8KB合计 16KB。加上每线程的寄存器开销SM 占用率在 60% 左右这个平衡点在通用场景里比较合理。另一个容易被忽略的点是BK 不宜太小否则每次加载的开销占比太高也不宜太大否则共享内存装不下。16 作为一个折中值既能让一次加载覆盖足够多的复用也不会让共享内存成为瓶颈。3.3 第一次收获从 0.4 到 2.4 TFLOPS加上共享内存 tiling 之后每个输出元素不再需要直接从全局内存取数而是从共享内存里读取。因为共享内存的带宽足够高访存等待大幅降低SM 的算术单元终于有活儿干了。这一版我实测到 2.4 TFLOPS耗时约 57 毫秒比 naive 版本提升了整整 6 倍。到这个点计算量其实已经和显存带宽解耦了。全局内存的读取量变成了(4096² 的 A_tile 加载) (4096² 的 B_tile 加载)数量级从 O(N³) 降到了 O(N²)带宽不再是主要矛盾。此时真正的瓶颈转移到了两个地方一是共享内存访问的 bank conflict二是每线程承担的累加次数还不够多导致线程调度开销占比偏高。解决银行冲突的方式也很粗暴给共享内存按行加 padding也就是把声明从__shared__ float As[BM * BK]改成__shared__ float As[BM][BK 4]。多出来的 4 个 float 不是为了装数据是为了让相邻行的起始地址错开避免不同线程在同一时钟周期访问同一 bank。这个细节后面再展开讲。4. float4 向量化一次搬 16 字节的威力4.1 向量加载的本质砍指令数很多人以为向量化就是“内存带宽翻倍”这个理解不完全对。向量化的本质是减少指令数。一条LDG.E指令一次只能从全局内存取 4 字节LDG.E.128一次取 16 字节。我们把原来需要 4 条指令完成的数据搬运压缩成 1 条指令指令发射压力直接降到四分之一。对现代 GPU 来说指令发射带宽和内存带宽同样宝贵。尤其当 tiling 把全局内存访存量降下来之后指令数量反而成了新的开销来源。这时候用 float4 重写加载代码收益立竿见影。还有一个隐藏收益向量化加载能触发更高效的内存事务。GPU 访问全局内存的最小粒度是 32 字节的扇区sector如果每个线程只读 4 字节即使线程连续访问也可能因为地址交错产生多余的扇区请求用 float4 对齐访问后每个线程一次恰好占满一个 16 字节对齐区域内存子系统的工作量会小很多。4.2 对齐是向量化的命门用 float4 有一个硬性前提地址必须 16 字节对齐。也就是说被访问数据的起始地址必须是 16 的倍数而且连续读取的长度也得是 16 的倍数。对 float 类型的数组来说这要求矩阵的行数、列数以及 tile 的 BK 都是 4 的倍数。写代码时要特别小心 C 的类型强转。常见的写法是float4 a_vec *reinterpret_castconst float4*(A[row * K col]);这种写法要求row * K col对应的地址确实对齐。如果不确定可以先打印验证K是 4 的倍数、col是 4 的倍数那么row * K col也必然是 4 的倍数因为 float 是 4 字节4 个 float 正好 16 字节4 的倍数乘以 4 字节后地址一定是 16 的倍数。反过来如果 K 取了一个不是 4 的倍数的值比如 1000那么即使 col 是对齐的下一行的起始地址也会偏移 16 字节的整数倍之外整段代码直接崩。4.3 向量化共享内存的正确姿势向量化加载和共享内存 tiling 不是两个孤立的优化它们需要配合。每个 block 在加载 A_tile 时不要一个线程一个 float 地拷而是让线程负责连续的 4 个 float一次性从全局内存读出一个 float4再一次性写入共享内存。我当时的加载代码结构是这样BM128、BN128BK16A_tile 是 128x16一共 2048 个 float也就是 512 个 float4B_tile 是 16x128同样 2048 个 float也是 512 个 float4。两块的向量总数是 1024 个而每个 block 有 512 个线程所以每个线程恰好处理 2 个 float4——这 2 个 float4 一个来自 A一个来自 B。这里有个很容易踩的坑线程读取共享内存做累加时如果直接按连续地址读会造成严重的 bank conflict。我第一版 tiling 代码在计算阶段访问As[行][k]同一 warp 里 32 个线程如果访问的列相同而行不同恰好全部落在同一个 bank 上性能直接从 2.4 掉回 1.2。解决办法就是前面提到的 padding给共享内存行宽加 4 个 float 的余量。为什么是 4因为 float4 是 16 字节对齐的padding 也必须保持 4 的倍数防止破坏向量化的对齐要求。4.4 第二阶段成绩5.8 TFLOPS加上 float4 向量加载、同时对共享内存做了 padding 之后性能从 2.4 TFLOPS 跳到了 5.8 TFLOPS耗时约 23.7 毫秒。Nsight Compute 里的指标变化非常明显共享内存 bank conflict 的惩罚基本消失全局内存 load 指令数量下降到了原来的四分之一SM busy 比例从 48% 提高到了 63%。这里多说一句5.8 TFLOPS 距离 10000 GFLOPS 还差得远但已经过了“半程”。剩下那 4000 多 GFLOPS 的缺口不是靠继续加向量长度能解决的——float4 已经是 128 位再往上就是 float8 这种不存在的类型了。下一步要把眼光从访存移到计算本身的并行度上。5. 冲击 10000 GFLOPS寄存器和流水线5.1 每线程多算一点寄存器内累积到 5.8 TFLOPS 之后瓶颈开始转移到两个地方一是线程数量太多导致每个线程干的活太少二是共享内存读取频率太高成了新的瓶颈。解决办法是让每个线程在寄存器里多缓存几个输出元素减少共享内存的访问次数。具体做法是提高每个线程负责的输出子块面积。我最初版本是每个线程负责 4x8 的输出块也就是 32 个累加器。从这个版本开始累加器全放在寄存器里——寄存器是 GPU 上最快的存储读写零延迟前提是没有寄存器溢出每次从共享内存读一个 A 元素和一个 B 元素就同时更新 32 个累加器。这一步的本质是提高“每次共享内存读操作的算术强度”。如果把共享内存和寄存器之间的数据传输看成是一个“送料”的过程那么每送一次料我们完成的乘加次数从原来的一次变成 32 次。数据搬运频率大幅度下降算术单元被喂得更饱了。5.2 循环展开给编译器递刀寄存器累加器准备好了接下来要解决循环开销问题。K 方向的累加循环里如果每一步都要判断循环条件、跳转回去那这些指令也在抢占发射槽。GPU 的编译器通常会自动做一定程度的循环展开但面对复杂的多维索引时经常犹豫不决。我用了#pragma unroll显式让编译器把 BK16 的内层循环完整展开成 16 份并把 32 个累加器更新全部平铺。展开之后编译器能更好地安排指令顺序把内存加载指令和 FMA 计算指令交错在一起。这就像做饭时不再等菜炒熟再去洗下一份菜而是同时把洗菜、切菜、炒菜的指令排成流水线全流程不停顿。不过循环展开要讲究度。展开过度会让寄存器压力飙升直接导致寄存器溢出到 local memory——local memory 是全局内存的一部分一旦出现 spill性能不升反降。我的经验是观察编译报告里的寄存器使用数保持在 128 个以内比较安全如果看到 local memory 相关字段不再是 0就说明展开过头了。5.3 双缓冲把加载和计算叠起来到了这个阶段剩下的主要延迟来自每次 K 循环里的同步点加载共享内存 →__syncthreads()→ 计算 → 再加载。在同一个 block 里所有线程必须等共享内存加载完成才能开始计算这个等待时间是纯空转。双缓冲的思路是把共享内存分成两半交替使用第一半正在被线程计算时第二半同时加载下一轮需要的 A、B 数据。这样每一轮的加载延迟被计算过程完全掩盖SM 里的计算单元几乎永远在处理数据。具体实现上给共享内存声明两个副本As[2][BM * BK]和Bs[2][BK * BN]用奇偶轮次切换当前使用的那一份。同时把__syncthreads()的位置调整到“确保当前轮次的数据已经加载完成、确保上一轮的计算已经完成”这两件事的交汇点。这一步改动不大逻辑上却很容易写出死锁或者数据覆盖错误我最后是拿 racecheck 工具跑了一遍才确认没有竞态条件。双缓冲之后全局内存的延迟被彻底隐藏Nsight 里 warp stall 的空闲周期数显著下降。这一项优化带来的收益在 10% 到 20% 之间属于“虽然没有向量化那么华丽但必不可少”的补刀。5.4 跨越时刻10.2 TFLOPS 实测所有优化叠加之后最终版在 MNK4096 的 FP32 矩阵上跑出了 10.23 TFLOPS耗时 13.43 毫秒。对比一下完整链路的成绩版本耗时性能相对 naive 提升naive约 327 ms0.42 TFLOPS1x共享内存 tiling约 57 ms2.4 TFLOPS5.7x float4 向量化约 23.7 ms5.8 TFLOPS13.8x 寄存器累积与展开约 16.5 ms8.3 TFLOPS19.8x 双缓冲流水线13.43 ms10.23 TFLOPS24.4x看到 Nsight 报告里出现 10.23 TFLOPS 这个数字的时候说实话还是挺激动的。性能利用率到了 28.7%在纯 CUDA core 的路径上这已经不是侥幸跑出来的数字而是每一步都有明确方向的工程结果。6. 调试与性能分析实战被掩盖的瓶颈6.1 Nsight Compute 到底看什么很多人优化 CUDA kernel 靠猜性能不行就调大线程数、换 tile 大小、加__restrict__然后看跑分撞运气。我这次全程用 Nsight Compute 分析用它定位瓶颈效率高得多。最常看的几个指标依次是Compute (SM) ThroughputSM 里的算术单元是否接近满载。低于 60% 说明计算没喂饱。Memory Throughput全局内存或者共享内存的带宽利用率。高于 80% 说明这一项是瓶颈。Warp Stall每个 warp 因为什么原因停下来的统计。比如等待共享内存数据short scoreboard、等待全局内存数据long scoreboard、等待同步barrier。Bank Conflicts共享内存 bank conflict 的具体次数和平均惩罚周期。性能卡在 5.8 TFLOPS 的时候就是靠 Warp Stall 里的short scoreboard指标发现大量线程在等共享内存 data 返回于是决定加大每线程的累加面积、减少共享内存访问频率。如果没有这份数据我大概率会再去调 tile 大小那方向就偏了。6.2 我踩过的三个坑第一个坑是共享内存 bank conflict。Tiling 版刚上线时性能不仅没增反而掉到了 1.1 TFLOPS。排查后发现计算阶段访问As[threadIdx.y * TM i][kk]时同一 warp 里 32 个线程的行号相差不大、列号完全相同全部压到同一个 bank。修复方式是给共享内存加 padding这一个改动就把性能拉回 2.4 TFLOPS。第二个坑是寄存器溢出。为了追求每线程 8x8 的累加面积我把 BM、BN 调大结果编译期寄存器数超过了 255 个部分变量被塞进 local memory。local memory 在全局内存里访问延迟比共享内存还高性能反而从 8.3 掉到 6.5。最后把每线程输出从 64 个降到 32 个寄存器数压在 128 以内才恢复正常。第三个坑是向量化读取越界。我最初在矩阵规模是 4096 时测试通过后来验证 4096x4096x4096 之外的规模时发现 float4 的对齐条件被打破——当 K 不是 4 的倍数时某行的起始地址不对齐读取直接段错误。最终解决方案是在 kernel 入口加一个快速路径如果 K 或者 M、N 不是 4 的倍数就 fallback 到一个标量版本的 kernel。这个做法保证了通用性。6.3 常见问题速查表现象可能原因排查方向结果全零或错误线程块坐标映射错乱、共享内存未同步检查 blockIdx/threadIdx 计算公式确认__syncthreads()位置性能不升反降寄存器溢出或 bank conflict 严重看编译报告里的 local memory 字段看 Nsight 的 bank conflict 计数GFLOPS 虚高编译器把零乘零优化掉或矩阵本身稀疏用全随机数据测试开启--use_fast_math时注意精度偶发死锁双缓冲逻辑里同步条件写错用 compute-sanitizer 的 racecheck 功能检测竞态共享内存不足tile 选得太大把 BM/BN/BK 调小或减少每线程累加面积7. 完整核心代码与后续扩展7.1 核心 kernel 参考为了方便理解下面给出最终版本的核心 kernel 结构去掉了边界 fallback 和双缓冲的细枝末节保留了最关键的向量化加载和寄存器累积逻辑#define BM 128 #define BN 128 #define BK 16 #define TM 4 #define TN 8 #define NUM_THREADS (BM * BN / (TM * TN)) __global__ void sgemm_tile_vec(const float* __restrict__ A, const float* __restrict__ B, float* __restrict__ C, int M, int N, int K) { __shared__ float As[BM * BK 4]; __shared__ float Bs[BK * BN 4]; const int tid threadIdx.y * blockDim.x threadIdx.x; const int blockRow blockIdx.y * BM; const int blockCol blockIdx.x * BN; const int threadRow threadIdx.y * TM; const int threadCol threadIdx.x * TN; float acc[TM][TN] {0.0f}; float aReg[TM]; float bReg[TN]; for (int k0 0; k0 K; k0 BK) { // 向量化加载每线程处理 2 个 float4分别来自 A 和 B 的 tile float4* dstA reinterpret_castfloat4*(As[tid * 4]); float4* dstB reinterpret_castfloat4*(Bs[tid * 4]); int aGlobalIdx (blockRow tid / (BK / 4)) * K k0 (tid % (BK / 4)) * 4; int bGlobalIdx (k0 tid / (BN / 4)) * N blockCol (tid % (BN / 4)) * 4; *dstA *reinterpret_castconst float4*(A[aGlobalIdx]); *dstB *reinterpret_castconst float4*(B[bGlobalIdx]); __syncthreads(); #pragma unroll for (int kk 0; kk BK; kk) { #pragma unroll for (int i 0; i TM; i) aReg[i] As[(threadRow i) * BK kk]; #pragma unroll for (int j 0; j TN; j) bReg[j] Bs[kk * BN threadCol j]; #pragma unroll for (int i 0; i TM; i) #pragma unroll for (int j 0; j TN; j) acc[i][j] aReg[i] * bReg[j]; } __syncthreads(); } // 结果写回按 float4 写入全局内存 #pragma unroll for (int i 0; i TM; i) { float4* cRow reinterpret_castfloat4*(C[(blockRow threadRow i) * N blockCol threadCol]); float4 val0 {acc[i][0], acc[i][1], acc[i][2], acc[i][3]}; float4 val1 {acc[i][4], acc[i][5], acc[i][6], acc[i][7]}; cRow[0] val0; cRow[1] val1; } }需要注意这个版本的共享内存声明里我做了一个简化用4的 padding 保证行与行之间不会因为 bank 冲突互踩但实际运行时二叉稿的加载写入也必须对齐 float4。完整工业级实现还需要处理非 4 倍数边长、K 方向多 tile 循环的流水线切换以及用cp.async指令代替普通 load 来进一步隐藏延迟。7.2 还能往上走的路线跨过 10000 GFLOPS 之后继续往上走的空间还有几个方向。第一个是上 Tensor CoreAmpere 的 TF32 Tensor Core 理论峰值是 FP32 CUDA core 的的几倍以上但那套代码和本文是两条完全不同的路径需要走 mma 指令、warp 级 tile 分配甚至wgmma这类异步指令复杂度一下子抬上一个台阶。第二个方向是用cp.async也就是异步拷贝指令。双缓冲用普通 load 把数据搬到共享内存虽然延迟被计算掩盖了但搬运过程本身还是要占用寄存器中转cp.async可以直接从全局内存搬进共享内存、不经过寄存器进一步减少指令开销。这个优化我再后面做实验时收益大概还有 10-15%。第三个方向是 split-K 并行把 K 维拆成多段并行计算再用归约把部分和相加。这对大矩阵且 SM 数量多的情况效果明显尤其当单 block 的 tile 已经无法再增大时它提供了一条把更多 SM 利用起来的路。写到最后的一点体会矩阵乘法优化的过程本质上就是不断追问“当前瓶颈在哪”。naive 阶段瓶颈是全局内存带宽所以上共享内存 tilingtiling 之后瓶颈是共享内存访问冲突所以加 padding再往后是每线程工作量不足所以用寄存器累积最后是加载延迟暴露所以上双缓冲。每一步都要先找到瓶颈再选择对应的手段顺序错了效果会差很多。向量化是我这次实录里收益最高的一刀但它的前提是已经有了共享内存 tiling——没有数据复用向量化省下的指令数也填不满带宽的坑。如果你也想复现建议从 naive 版本开始逐版叠加拿 Nsight Compute 记录每一版的指标变化你会发现性能优化的乐趣不在于最后那个数字而在于每一次“找到原因、修复、数字跳升”带来的确定性。

CUDA手写SGEMM优化实战：从0.4到10.23 TFLOPS的完整优化链路

相关推荐

WordPress博客贡献热力图开发指南

拆解 Agent Harness：大模型智能体的七个核心模块与最小实现

ArcGIS外业质检插件开发：提升农田调查效率的解决方案

更多资讯

iReport 55报表工具配置与设计实战指南

HACS极速版配置指南：智能家居插件管理与前端性能优化

智能体互联网三大技术支柱：MCP、A2A与Agent Skills解析

MySQL迁移实战：语法差异与KingbaseES兼容性解析

状态空间搜索实战：狼羊白菜过河BFS/DFS与Python实现

SolidWorks倒圆角与斜角设计技巧及工程应用

今日更新

AI辅助技术栈评估：提升研发决策效率60%

从生成Token到直接决策：Jev如何重塑AI Agent工程实践

32G统一内存跑大模型：内存账、TPS实测与端云分流

本周精选

面向复杂光照与角度的宠物面部识别鲁棒性增强策略

基于DRV8818与STM32的双极步进电机驱动方案实战解析

DRV8818PWPR+PIC18F45K80双极步进电机控制方案实战

本月热门

面向复杂光照与角度的宠物面部识别鲁棒性增强策略

基于DRV8818与STM32的双极步进电机驱动方案实战解析

DRV8818PWPR+PIC18F45K80双极步进电机控制方案实战

读完文章，想聊聊您的网站？