Nvidia的”护城河” | Waang’s Blog

type

status

date

slug

summary

本文含少量AI生成内容，请注意鉴别

cuda.ipynb

26.6KB

矩阵乘法

考虑,利用计算密度为. CPU端串行程序来看,计算访存比为

如下内容参考了这个实现：

cuda_code/matrix at main · xgqdut2016/cuda_code

easy cuda code. Contribute to xgqdut2016/cuda_code development by creating an account on GitHub.

https://github.com/xgqdut2016/cuda_code/tree/main/matrix

v1&2.二维线程块实现矩阵乘法使用shared memory

针对当前线程块所有线程可见。仍然存在重复读取的情况，如果重复的不是global memory，而是shared memory。

v3.一个线程处理多个元素

计算访存比为

v4.优化矩阵load至shared memory过程

线程索引的重新映射

v5.float4类型访存

减少内存事务：单次内存操作可加载/存储 4 个元素，提升带宽利用率。

注意问题：循环展开次数增加，寄存器压力过大（寄存器爆炸），可能引发性能下降，所以需要线程块参数调优，以平衡寄存器使用。

v6.bank conflict

Bank Conflict 定义

共享内存分为 32 个 Bank，每个 Bank 宽度为 4 字节。同一 Warp（32 线程）访问不同 Bank 时性能最佳，否则产生 Bank Conflict。

原始 SA 访问的冲突问题

C 元素索引：[8×threadIdx.x+index.q][8×threadIdx.y+index.v]。

SA 索引公式：

SB 索引公式：

固定 index.q 和 index.v 时，SA 的访问模式导致同一 Warp 的线程访问相同 Bank，引发 Bank Conflict。

优化方法

转置 SA 的加载方式：将 SA 的共享内存布局从 [128][BK] 调整为 [8][128]。调整后 SA 的索引模式分散到不同 Bank，消除 Bank Conflict。

代码示意（转置 SA 后的访问逻辑）：

v7.降低shared memory的读取

💡

SA 和 SB 的转置布局使得对 index_k 的访问是非连续的，但对 index_q 和 index_v 是连续的，可以借助寄存器和float4进行优化

核心思想：内积转换成外积

外积方式将分块乘法分解为多个外积的累加：

可以通过计算实现

v8.流水并行

重叠全局内存加载与计算操作，减少内存访问延迟，提升 GPU 计算效率。

原始模式

加载 data[0] → 计算 data[0] → 加载 data[BK-1] → 计算 data[BK-1]

顺序执行加载和计算，存在空闲等待时间.

流水并行模式

加载 data[0] → [加载 data[1] + 计算 data[0]] → ... → [加载 data[BK-1] + 计算 data[BK-2]] → 计算 data[BK-1]

v9.cuBLAS库

cuBLAS 是 NVIDIA 提供的 GPU 加速线性代数库，支持矩阵乘法等基础运算。其核心公式为C=α⋅op(A)⋅op(B)+β⋅C，

αβ：标量系数。

op(X)：对矩阵X的转置操作（可选不转置、转置或共轭转置）。

需注意列优先存储与 C 语言行优先的差异！ v10.Tensor Core

warp读取矩阵：tensor core计算过程中，left_frag和right_frag在内存中不连续，这时再结合二位索引重排。

PTX

a low-level Parallel Thread eXecution virtual machine and instruction set architecture (ISA),直面意思是低级并行线程执行虚拟机和指令集架构。PTX是上承GPU编程语言CUDA C++，下启GPU硬件SASS指令，可以借助NVRTC实现运行时优化，某些层面上来说可以称之为GPU设备无关代码，因此PTX可以理解为”CUDA IR“。