Co-authored-by: kuanzeren <kuanzeren@tencent.com>
This commit is contained in:
111qqz
2022-09-05 14:39:00 +08:00
committed by GitHub
parent cf00d15589
commit 7a948656f1

View File

@@ -288,7 +288,7 @@ __global__ void gemmKernel(const float * A,
### 提高计算强度
计算强度Compute Intensity指计算指令数量与访存指令数量的比值在现代GPU中往往有大量计算单元但只有有限的访存带宽程序很容易出现计算单元等待数据读取的问题因此提高计算强度是提升程序性能的一条切实有的指导思路。对于之前实现的GPU核函数我们可以粗略计算其计算强度在$K$次循环的内积计算中,对矩阵$A$与矩阵$B$的每次读取会计算一次浮点乘法与浮点加法因此计算强度为1——两次浮点运算除以两次数据读取。之前的版本是每个线程负责处理矩阵$C$的一个元素——计算矩阵$A$的一行与矩阵$B$的一列的内积,我们可以通过使每个线程计算$C$更多的元素——计算矩阵$A$的多行与矩阵$B$的多列的内积——从而提升计算强度。具体地,如果在$K$次循环的内积计算中一次读取矩阵$A$中的$m$个元素和矩阵$B$中的$n$个元素,那么访存指令为$m+n$条,而计算指令为$2mn$条,所以计算强度为$\frac{2mn}{m+n}$,因此可以很容易发现提高$m$和$n$会带来计算强度的提升。
计算强度Compute Intensity指计算指令数量与访存指令数量的比值在现代GPU中往往有大量计算单元但只有有限的访存带宽程序很容易出现计算单元等待数据读取的问题因此提高计算强度是提升程序性能的一条切实有的指导思路。对于之前实现的GPU核函数我们可以粗略计算其计算强度在$K$次循环的内积计算中,对矩阵$A$与矩阵$B$的每次读取会计算一次浮点乘法与浮点加法因此计算强度为1——两次浮点运算除以两次数据读取。之前的版本是每个线程负责处理矩阵$C$的一个元素——计算矩阵$A$的一行与矩阵$B$的一列的内积,我们可以通过使每个线程计算$C$更多的元素——计算矩阵$A$的多行与矩阵$B$的多列的内积——从而提升计算强度。具体地,如果在$K$次循环的内积计算中一次读取矩阵$A$中的$m$个元素和矩阵$B$中的$n$个元素,那么访存指令为$m+n$条,而计算指令为$2mn$条,所以计算强度为$\frac{2mn}{m+n}$,因此可以很容易发现提高$m$和$n$会带来计算强度的提升。
我们在上一个代码例子中对全局内存的访问与存储都是借助 `float` 指针完成的,具体到硬件指令集上实际是使用指令 `LDG.E` 与 `STG.E` 完成的。我们可以使用128位宽指令`LDG.E.128` 与 `STG.E.128` 一次读取多个 `float` 数。使用宽指令的好处是一方面简化了指令序列使用一个宽指令代替四个标准指令可以节省十几个指令的发射周期这可以为计算指令的发射争取到额外的时间另一方面128比特正好等于一个cache line的长度使用宽指令也有助于提高cache line的命中率。但我们并不提倡在一切代码中过度追求宽指令的使用开发者应当将更多的时间关注并行性设计和局部数据复用等更直接的优化手段。