Stencil计算GPU性能建模

最新推荐文章于 2024-06-09 02:34:19 发布

zongy17

最新推荐文章于 2024-06-09 02:34:19 发布

阅读量1.5k

点赞数 2

分类专栏：高性能计算文章标签： c语言 hpc gpgpu cuda 数学建模

本文链接：https://blog.csdn.net/weixin_43614211/article/details/122106716

版权

高性能计算专栏收录该内容

14 篇文章 12 订阅

订阅专栏

Stencil计算GPU性能建模

3D Blocking 建模
2D Tiling 建模

机器：NVidia Tesla V100单卡

此文承接专栏中Stencil-GPU一文。由于作业2中7点stencil使用了naïve版本的3D blocking实现，仅做了参数调优；27点stencil使用了2D tiling的优化，故本次作业对两种实现分别做建模。模型中所用到的GPU硬件信息（各级内存带宽、延迟等）按照Github仓库上的基准进行测试，并参照了Elias Konstantinidis等写的文章A quantitative performance evaluation of fast on-chip memories of GPUs，及其代码做对比参考。测试得到的硬件信息，与Zhe Jia等写的技术报告Dissecting the NVidia Turing T4 GPU via Microbenchmarking中提到的Tesla GV100的相关数值在差异不大，可以认为是靠谱的。

3D Blocking 建模

本部分主要采用了Huayou Su等写的文章An analytical GPU performance model for 3D stencil computations from the angle of data traffic。

原文中建模的kernel

作业kernel核心代码

区别在于原文的u_old使用老版GPU的RODC（read-only-data-cache），而作业的local_g使用shared_memory，但两者类似。由于stencil计算往往是memory-bounded的，模型只从数据传输的角度分析执行时间。

模型输入

分块大小 $B_x,B_y,B_z$ ，问题规模 $n$

模型参数

主要是机器参数

#SM=80：流处理器（Stream Multiprocessor）的数目
#MAX_threads_per_SM=2048：每个SM最多运行的线程数
coalesced_size=32 Bytes：共享内存载入数据时的聚合传输大小
L2_cacheLine_size=64 Bytes：L2 cache的缓存行大小
L2_cache_size=6144*1024 Bytes：L2 cache的大小，6144 KB
$BW_{GM}$ =771.4 GB/s：全局内存的可达带宽
$BW_{L_2}$ =2066.5 GB/s：L2 cache的可达带宽
$BW_{SMX}$ =11860 GB/s：SM片上共享内存的可达带宽
$\epsilon$ =0.01：描述L2 cache缺失率的常数

模型假设

执行时间是三层数据传输中的耗时较大者：

$\geq max(\frac{V_{GM}}{BW_{BM}}, \frac{V_{L_2}}{BW_{L_2}}, \frac{V_{SMX}}{BW_{SMX}})$

其中 $V_{GM}$ 表示从全局内存中读或写到全局内存中的数据总量， $BW_{GM}$ 表示全局内存到L2 cache的可达带宽。同理， $V_{L_2}$ 和 $BW_{L_2}$ 表征 L2 cache和流处理器（SM）片上的共享内存之间的数据传输总量和可达带宽； $V_{SMX}$ 和 $BW_{SMX}$ 表征片上共享内存和线程寄存器之间的数据传输总量和可达带宽。

以下逐个考察这三项。下面由于公式较多，本人一开始是写在word上的，再重新type一遍有点麻烦，所以直接截图了。
在这里插入图片描述

模型效果

这是一个纯分析型模型，不需要训练或回归等方法得知模型参数。直接用于7点stencil的 $n$ =256,384,512的各个并行度（即GPU的分块数）下的时间预测，效果如下。其中横坐标( $B_x$ , $B_y$ , $B_z$ )表示三维分块各个维度的大小。右边坐标轴表示每对（测试，模型预测）值的相对误差。

可见这个分析模型在对较大规模的算例时，预测效果较好。对 $n$ =256,384,512的平均误差分别为40%，4%和4%。

对于27点stencil，预测效果如下图所示。对 $n$ =256,384,512的平均误差分别为29%，20%和22%。

2D Tiling 建模

本部分借鉴了一点Sunpyo Hong等写的文章An Analytical Model for a GPU Architecture with Memory-level and Thread-level Parallelism Awareness的思路。但原文实现实在过于精巧，工作量巨大，本作业没有忠实实现。

模型输入

分块大小 $B_x,B_y$ ，问题规模 $n$

模型参数

主要是机器参数：

#SM=80：流处理器（Stream Multiprocessor）的数目
#SP_per_SM=128：一个SM上的硬件核心数
#threads_per_warp=32：一个硬件调度warp内的线程数
coalesced_size=32 Bytes：共享内存载入数据时的聚合传输大小
$BW_{GM}$ =771.4 GB/s：全局内存的可达带宽
$BW_{L_2}$ =2066.5 GB/s：L2 cache的可达带宽
$BW_{SMX}$ =11860 GB/s：SM片上共享内存的可达带宽

模型假设

对于27点stencil的2D tiling的算法，一个线程块内的线程逐次读入全局内存的一层数据到共享内存上，然后转移到寄存器上，利用寄存器中的一层9个数做它相应有贡献的3层的计算。如下图所示。

程序的主体部分就是一个沿k方向的逐层循环。因此有
#threads = $B_x \times B_y$
#blocks = $\mathrm{ceil}(\frac{N_x}{B_x}) \times \mathrm{ceil} (\frac{N_y}{B_y})$

其中ceil函数向上取整是为估计不整除的情况。

仍然假定stencil计算是严重memory bounded的，只考虑数据传输的开销，忽略循环中computeStencil_planex_27()的计算开销。由于在shm2regs()和readBlockAndHalo()函数头都有__syncthreads()，因此在每个线程块内，都可以显式地将shm2regs()和readBlockAndHalo()两者分离，即时间上不重叠，可以直接加和。

$t=t_{\mathrm{shm2regs}} + t_{\mathrm{readBlk}} + t_{\mathrm{calc}} \approx t_{\mathrm{shm2regs}} + t_{\mathrm{readBlk}}$

分别考虑这两部分，shm2regs是数据从SM片上block-shared的共享内存到thread-private的寄存器的传输，readBlockAndHalo是数据从全局内存到SM片上共享内存的传输。注意后者一定是从全局内存，因为每次读入一块新的（k+2）平面，这块平面在之前未被访问过，一定在L2 cache中不命中，因此后续计算使用的是主存的带宽 $BW_{GM}$ 。

按照Sunpyo Hong等的分析，根据一个SM内的MWP（Memory Warp Parallelism）和CWP（Computational Warp Parallelism）的关系，有可能发生计算和通信的重叠，这也是GPU通过大量线程来隐藏访存延迟的特性。根据上述分析，shm2regs和readBlockAndHalo不重叠，因此分别考虑各自内部的访存重叠可能。

原文在这里的分析采用了测量聚合、非聚合访存的延迟、内存带宽等信息，并分析由反汇编NVIDIA PTX代码得到的指令流，来确定MWP，CWP，和计算的时钟周期数。这样的工作量对于一个作业而言显然太多了。因此这里采用待定系数的方法，假设SM在shm2regs和readBlockAndHalo内分别可以同时执行（重叠） $α$ 和 $β$ 个warp，如下图所示：（此处忽略了重叠时应错开的一小段由warp切换带来的延迟，认为它相比于整个访存过程的执行时间是足够小的）
在这里插入图片描述
再考虑到V100每个SM内有128个SP（Streaming Processor，相当于硬件核心），4个warp调度器，因此相当于四条流水线，可以同时执行warp的数目为#simWarp_per_SM=4，所以每条流水线上分到的warp数目为