企业官网建设流程全解析

深圳专业网站排名公司,儿童早教网站模板,wordpress数据库越来越大,ppt要怎么做网站NVCC版本#xff1a;12关于GPU你需要了解的那些事从硬件视角看GPU之前在CUDA并行规约那篇文章中提到过#xff0c;在进行CUDA开发时#xff0c;我们是以Grid#xff0c;Block和Thread三级的层次结构来组织线程的#xff0c;那么这三者是如何对应到具体的硬件实现的呢…NVCC版本12关于GPU你需要了解的那些事从硬件视角看GPU之前在CUDA并行规约那篇文章中提到过在进行CUDA开发时我们是以GridBlock和Thread三级的层次结构来组织线程的那么这三者是如何对应到具体的硬件实现的呢从硬件视角下看一张显卡里有一个GPUGPU内部有多个流式多处理器Streaming Multiprocessor以下简称SM如下图所示image接下来把视角转向SM内部每个SM有多个处理器线程就是在这些处理器上具体执行的。除此之外每个SM还有一块共享内存区域之前文章里提到的共享内存Shared Memory以下简称SMEM就是在这个区域这个区域只能是SM内部的处理器访问。SM内部的每个处理器又有着只能是自己访问的寄存器REGS。从这里也可以总结出GPU上内存访问的速度排序寄存器REGS是最快的但是只能线程自身访问其次是SMEM但是只能是Block内的线程访问最慢的是GMEMGMEM就是我们通过cudaMalloc申请到的内存所有线程均可访问。那么上述的线程层级架构又是怎么和这个硬件架构相对应的呢 而且Warp在其中又是怎么体现的呢这就得从线程调度的角度来看一个内核被启动的过程了。在启动内核时我们会指定GridDim和BlockDim这就使得内核有了一定数量的线程块Block需要执行每个Block里有若干个Thread。在进行调度时GPU会以Block为单位把一个Block分给一个SM这时候一个SM可能会被分到多个Block。接下来就是SM的工作了一个Block里连续的32个线程为一个WarpSM会以Warp为单位进行调度即SM会选择32个连续的线程然后放到32个处理器上运行。上述过程如下图所示。image全局内存访问合并同一个Warp里的线程有很多有意思的特性对这些特性加以利用就能够达成不错的优化效果全局内存访问合并Global Memory Coalescing就是其中之一。这个特性是如果一个Warp里的线程访问的内存恰好是连续的32个4B的浮点数那么GPU就只会做一次长度为128B的访存操作 把128B的数据读取之后分发给32个线程。这里参考资料作者的精美的手绘图可以很形象地说明这一点image一些约定在开始正式实现前首先需要把一些容易混淆的设定给明确了。本文默认所有矩阵都是行优先存储的本文中的x都是指行下标y都是指列下标如下图所示image注本文所有的内核实现都只是在大小为4096的方阵下进行了正确性验证如果要适配任意形状需要考虑很多corner case这有点偏离主线了所以本文就暂时不做这方面的适配工作了。V42cuBLAS这里先放出cuBLAS实现的性能数据供后续比较和参考imageV1Naive Kernel对于矩阵乘法一个最朴素的想法就是让每个线程都计算C中的一个元素所以只需要一个Block使用Thread的x和y表示要计算的C的元素坐标然后2个for循环计算即可。这个想法没问题只是在实现的时候由于一个Block里面最多有1024个线程所以需要进行一次分块。具体而言可以把C分成若干个的块Tile每个Tile交给一个Block进行计算如下图所示image对于每个线程而言首先需要根据计算出当前线程需要处理的C的坐标然后用2个for循环计算结果并写回即可。至于布局每个Block里自然是个线程而Grid则是需要用向上取整的除法来分配尽量多的Block。注理清楚每个线程需要计算哪些C的元素是不被后面更复杂的分块绕晕的关键最终得到的源代码如下所示__global__ void MatmulKernelV1(const scalar_t *a, const scalar_t *b, scalar_t *out, uint32_t M, uint32_t N, uint32_t P) {uint32_t x blockIdx.x * blockDim.x threadIdx.x;uint32_t y blockIdx.y * blockDim.y threadIdx.y;if (x M y P) {scalar_t tmp 0;for (uint32_t k 0; k N; k) {// out[i][j] a[i][k] * b[k][j]tmp a[x * N k] * b[k * P y];}out[x * P y] tmp;}}void MatmulCoreV1(const scalar_t *a, const scalar_t *b, scalar_t *out, uint32_t M, uint32_t N, uint32_t P) {dim3 grid(std::ceil(M / 32.0), std::ceil(P / 32.0), 1);dim3 block(32, 32, 1);MatmulKernelV1grid, block(a, b, out, M, N, P);}实验数据最终的性能数据如下所示image性能只有cuBLAS的0.39%可以说是相当拉垮了。理论分析这里先插入一段理论分析来分析一下Naive Kernel可能的性能瓶颈在哪。首先计算一下理论最快的运行时间进行一个4096方阵的乘法所需要浮点运算次数为因为C有个元素每个元素需要进行4096次乘法和加法大约为137GFLO而内存读取最低需要约134MB写入需要约67MB。实验用的显卡浮点数计算性能为870GFLOPS显存带宽为29GB/s所以理论上计算最快需要157ms访存共需要6.9ms也就是说理论上来讲矩阵乘法应该是计算瓶颈的。但是我们的Naive Kernel似乎并不是这样的下面来详细分析一下。在计算次数上如果不考虑计算下标的开销那它的计算次数就是和理论最低值相等的在内存访问上实际上每个线程都会访问次全局内存GMEM如果这些访问没有经过任何优化那么这个内核一共就会有的访存需要耗时18.9s已经是计算的120倍了所以很显然目前的首要任务是优化内存访问。注这里内存访问事实上并没有计算的那么多因为有一些Warp层的自动优化这个后面马上会提到V2Global Memory Coalescing这里可以像防止Bank Conflict那样通过调整每个线程负责的区域来实现Coalescing。我们首先分析一下V1里面每个线程都在计算C的哪一个元素。通过代码可以知道线程计算的C的x坐标就是threadIdx.xy坐标就是threadIdx.y并且threadIdx是x先变化的所以第一个线程计算的是(0, 0)第二个是(0, 1)以此类推如下图所示用背景色来区分T0和T1加载的数据image可以发现一个Warp里计算的其实是C中的某一列那么同一时刻Warp里的线程访问的A一定不是连续的所以访问A的部分一个Warp需要次访存而访问B的部分由于一个Warp里的线程在同一时刻访问的都是同一个B所以这里只会有一次访存开销那么访问B总共就会有4096次访存。这里如果我们让一个Warp计算C中的一行会怎么样呢那么访问A就只需要4096次访存但是访问B的时候在同一时刻线程们访问的数据是连续的此时就可以触发Global Memory Coalescing把32次访存压缩为1次如下图所示image实验数据image这里仅仅是对换一下x和y性能就提升了接近8倍。但是和cuBLAS相比还是有不小的差距目前也还只有cuBLAS性能的2%。关于实现方式具体实现时只需要把x和y对换一下就行了。参考资料作者在这里的实现是取消了threadIdx.y这个维度然后把x维度的大小扩展为了1024之后在线程内部根据threadIdx.x以及BlockSize来计算当前线程对应到C的坐标如下所示constexpr uint32_t BLOCK_SIZE 32;__global__ void MatmulKernelV2(const scalar_t *a, const scalar_t *b, scalar_t *out, uint32_t M, uint32_t N, uint32_t P) {// 相当于首先把OUT分成若干个32*32的BlockV1和V2都是如此它们的区别在于Block内部的分配方式// 这里blockIdx.x * BLOCK_SIZE, blockIdx.y * BLOCK_SIZE 就是在定位Block的起始x和yconst uint32_t x blockIdx.x * BLOCK_SIZE threadIdx.x / BLOCK_SIZE;const uint32_t y blockIdx.y * BLOCK_SIZE threadIdx.x % BLOCK_SIZE;......个人认为这种实现肯定是不如直接对换x和y的因为这种实现引入了除法和求余数这种非常昂贵的操作但是实际测试下来两种实现的性能是几乎一致的。于是我看了下两种实现对应的PTX汇编如下图所示imageimage发现两者指令数量几乎都差不多可能是因为这里BlockSize为2的整数次幂的关系吧这种特性使得编译器可以对求余数指令做优化一般的求余数指令是需要用除法减法来实现的。如下图所示可以看到在把BlockSize换成34之后确实多出来了一条sub指令。image在把BlockSize改成31然后实际运行后确实出现了性能的损失。image