4.5 A TILED MATRIX MULTIPLICATION KERNEL

2024-01-09 02:12

本文主要是介绍4.5 A TILED MATRIX MULTIPLICATION KERNEL,希望对大家解决编程问题提供一定的参考价值,需要的开发者们随着小编来一起学习吧!

我们现在准备展示一个tiled矩阵乘法内核,该内核使用共享内存来减少对全局内存的流量。图中4.16显示的内核。实施图4.15.中所示的阶段。在图4.16中,第1行和第2行声明Mds和Nds为共享内存变量。回想一下,共享内存变量的范围是一个块。因此,将为每个块创建一对Mds和Nds,并且一个块的所有线程都可以访问相同的Mds和Nds。这很重要,因为块中的所有线程都必须能够访问加载的M和N元素由同行进入Mds和Nds,以便他们可以使用这些值来满足他们的输入需求。
在这里插入图片描述
第3行和第4行将threadIdx和blockIdx值保存到自动变量中,从而保存到寄存器中,以便快速访问。回想一下,自动标量变量被放置在寄存器中。它们的范围在每个线程中;即一个tx、ty、bx和by的私有版本由运行时系统为每个线程创建,并将驻留在线程可访问的寄存器中。它们使用threadIdx和blockIdx值初始化,并在线程生命周期内多次使用。一旦线程结束,这些变量的值将不复存在。

第5行和第6行确定了线程要生成的P元素的行和列索引。该代码假设每个线程负责计算一个P元素。如第6行所示,水平(x)位置或由线程生成的P元素的列索引可以计算为bxTILE_WIDTH+ tx,因为每个块都涵盖了水平维度中的TILE_WIDTH元素。块bx中的线程在它之前会有bx线程块,或(bxTILE_WIDTH)线程;它们涵盖了P的bxTILE_WIDTH元素。同一块中的另一个tx线程将覆盖另一个tx元素。因此,带有bx和tx的线程应该负责计算x索引为bxTILE_WIDTH+ tx的P元素。这个水平索引保存在线程的变量Col中,图4.17.中也说明了。
在这里插入图片描述

在图4.14中,由block1,0的thread0,1计算的P元素的x索引为02+ 1= 1。同样,y索引可以通过byTILE_WIDTH+ ty计算。此垂直索引保存在线程的变量行中。因此,每个线程计算Col列和Row行的P元素,如图4.17所示。.回顾图4.14中的例子,由block1,0的线程1.0计算的P元素的y索引,”为1*2+ 0=2。因此,由此线程计算的P元素是P2.1。

图4.16中的第8行。标志着循环的开始,循环贯穿计算P元素的所有阶段。循环的每个迭代都对应于图4.15.中所示计算的一个阶段。Ph变量表示点积已经完成的阶段数。每个阶段使用一个M的图块和一个N个元素的图块。因此,在每个阶段开始时,前几个阶段都处理了M和N元素的ph*TILE_WIDTH对。

在每个阶段,第9行将适当的M元素加载到共享内存中。由于我们已经知道要由线程处理的M行和N列,我们现在讨论M的列索引和N的行索引。如图4.17所示。每个区块都有TILE_WIDTH 线程将协作将TILE_ WIDTH M元素加载到共享内存中。因此,我们只需要分配每个线程来加载一个M元素,这可以使用blockldx和threadIdx方便地完成。要加载的M元素部分的起始列索引是ph*TILE_WIDTH。因此,一个简单的方法是让每个线程加载一个tx(threadldx.x值)位置远离该起点的元素。

这种情况由第9行表示,其中每个线程加载M[Rowwidth + phTILE_WIDTH + txJ,其中线性化索引与行索引行和列索引ph*TILE_WIDTH + tx形成。由于Row的值是ty的线性函数,每个TILE_WIDTH2线程都会将一个唯一的M元素加载到共享内存中。这些线程将一起加载图4.17.中M的暗方子集。读者应该使用图4.14中的示例。和图4.15验证单个线程的地址计算是否正确。

第11行中的屏障_syncthreads()确保所有线程在任何线程向前移动之前都已完成将M和N的tile加载到Mds和Nds中。然后,第12行的循环在这些tile元素的基础上执行点积的一个阶段。Thready.tx的循环进度如图4.17所示,M和N元素沿箭头的访问方向,箭头标有k,第12行的循环变量。这些元素将从Mds和Nds访问**,Mds和Nds是包含这些M和N元素的共享内存阵列**。第14行中的屏障__syncthreads()确保所有线程在进入下一个迭代并从下一个tile加载元素之前,所有线程都已完成使用共享内存中的M和N元素。通过这种方式,没有一个线程会过早地加载元素并破坏其他线程的输入值。

从8行到14行的嵌套环路说明了一种称为 strip-mining 的技术,该技术需要一个长期运行的环路并将其分阶段。每个阶段都由一个内部循环组成,该循环执行原始循环的多次连续迭代。原始循环成为一个外部循环,其作用是迭代调用内部循环,以便原始循环的所有迭代都按照原始顺序执行。通过在内部循环之前和之后添加屏障同步,我们强制同一块中的所有线程将其工作完全集中在其输入数据的一部分上。Strip mining可以通过在数据并行程序中tile来创建所需的阶段。

点积的所有阶段完成后,执行将退出第8行的循环。所有线程都通过使用线性化索引计算的 Row和Col写入其P元素。

tile算法提供了巨大的好处。对于矩阵乘法,全局内存访问减少了TILE_WIDTH的倍数。如果使用16 x 16的tile,我们可以将全局内存访问量减少16倍。这将计算与全局内存的访问率从1提高到16。这种改进允许CUDA设备的内存带宽支持接近其峰值性能的计算速率;例如,具有150 GB/s全局内存带宽的设备可以接近((150/4)*16)=600 GFLOPS!

虽然tile矩阵乘法内核的性能改进令人印象深刻,但它包括一些简化的假设。首先,假设矩阵的宽度是线程块宽度的倍数。这种假设阻止了内核正确处理任意大小的矩阵。第二个假设是矩阵是平方矩阵,这在现实生活中并不总是正确的。在下一节中,我们将介绍一个带有边界检查的内核,以消除这些假设。

这篇关于4.5 A TILED MATRIX MULTIPLICATION KERNEL的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



http://www.chinasem.cn/article/585613

相关文章

Linux_kernel驱动开发11

一、改回nfs方式挂载根文件系统         在产品将要上线之前,需要制作不同类型格式的根文件系统         在产品研发阶段,我们还是需要使用nfs的方式挂载根文件系统         优点:可以直接在上位机中修改文件系统内容,延长EMMC的寿命         【1】重启上位机nfs服务         sudo service nfs-kernel-server resta

笔记整理—内核!启动!—kernel部分(2)从汇编阶段到start_kernel

kernel起始与ENTRY(stext),和uboot一样,都是从汇编阶段开始的,因为对于kernel而言,还没进行栈的维护,所以无法使用c语言。_HEAD定义了后面代码属于段名为.head .text的段。         内核起始部分代码被解压代码调用,前面关于uboot的文章中有提到过(eg:zImage)。uboot启动是无条件的,只要代码的位置对,上电就工作,kern

[论文笔记]LLM.int8(): 8-bit Matrix Multiplication for Transformers at Scale

引言 今天带来第一篇量化论文LLM.int8(): 8-bit Matrix Multiplication for Transformers at Scale笔记。 为了简单,下文中以翻译的口吻记录,比如替换"作者"为"我们"。 大语言模型已被广泛采用,但推理时需要大量的GPU内存。我们开发了一种Int8矩阵乘法的过程,用于Transformer中的前馈和注意力投影层,这可以将推理所需

欧拉系统 kernel 升级、降级

系统版本  cat  /etc/os-release  NAME="openEuler"VERSION="22.03 (LTS-SP1)"ID="openEuler"VERSION_ID="22.03"PRETTY_NAME="openEuler 22.03 (LTS-SP1)"ANSI_COLOR="0;31" 系统初始 kernel 版本 5.10.0-136.12.0.

[Linux Kernel Block Layer第一篇] block layer架构设计

目录 1. single queue架构 2. multi-queue架构(blk-mq)  3. 问题 随着SSD快速存储设备的发展,内核社区越发发现,存储的性能瓶颈从硬件存储设备转移到了内核block layer,主要因为当时的内核block layer是single hw queue的架构,导致cpu锁竞争问题严重,本文先提纲挈领的介绍内核block layer的架构演进,然

Kernel 中MakeFile 使用if条件编译

有时需要通过if  else来选择编译哪个驱动,单纯的obj-$(CONFIG_)就不是很方便,下面提供两种参考案例: 案例一: 来源:drivers/char/tpm/Makefileifdef CONFIG_ACPItpm-y += tpm_eventlog.o tpm_acpi.oelseifdef CONFIG_TCG_IBMVTPMtpm-y += tpm_eventlog.o

笔记整理—内核!启动!—kernel部分(1)驱动与内核的关系

首先,恭喜完成了uboot部分的内容整理,其次补充一点,uboot第一部分和第二部分的工作不是一定的,在不同的版本中,可能这个初始化早一点,那个的又放在了第二部分,版本不同,造成的工作顺序不同,但终归是要完成基本内容初始化并传参给kernel的。         那么至于驱动与内核的关系,用一张图来说明最适合不过:         驱动位于OS层的中下层与硬件相接。驱动是内

73. Set Matrix Zeros

题目: 解答: 提供了两种解题思路: 第一种,使用两个数组,分别标记每一行、每一列是否有0的存在,然后再去更新二维数组。 第二种,使用两个变量brow,bcol分别标记第0行,第0列是否存在0,然后使用每一行、每一列的第一个单元存储是否该行、该列存在0. 代码: class Solution {public:// 方法一void setZeroes(vector<vector<i

finalshell 4.5.x在m1mac闪退

使用过程中会出现突然闪退,尤其在定位生产打开一堆窗口的情况下,绝绝子 闪退崩溃日志: Thread 116 Crashed:: Java: pool-4-thread-280 libsystem_kernel.dylib 0x18e926600 __pthread_kill + 81 libsystem_pthread.dylib

Error: label vector and instance matrix must be double的解决方法

在使用uci下载的数据时,建模时出现这个错误的解决方法 首先现在UCI上面下载数据 然后右键另存为就行了。这样我们就从UCI里面下载到了训练数据 在matlab 点 导入数据,数据类型要记得选第二个, 如果选择最后一个table就会出现这个问题 最后附上代码 %%之前先import wine.date IMPORTED DATA 设为Numeric Matrix (数值矩