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

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/606797.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

Redis原理篇(Dict的收缩扩容机制和渐进式rehash)

Dict(即字典) Redis是一种键值型数据库,其中键与值的映射关系就是Dict实现的。 Dict通过三部分组成:哈希表(DictHashTable),哈希节点(DictEntry),字典(Dict&#xff09…

书生·浦语大模型全链路开源体系 学习笔记 第二课

基础作业: 使用 InternLM-Chat-7B 模型生成 300 字的小故事(需截图)。熟悉 hugging face 下载功能,使用 huggingface_hub python 包,下载 InternLM-20B 的 config.json 文件到本地(需截图下载过程&#xf…

tiktok云手机有用吗?用哪个好?

很多做独立站的跨境卖家都会搭配一些社媒平台给自己引流带货,比如说目前很火的TikTok,这也是目前比较有效的一种引流方式。本文将介绍tiktok运营方法以及如何用tiktok云手机规避运营风险。 TikTok是个不错的风口,不过我们在国内想要运营好Tik…

数环通12月产品更新:新增数据表相关功能、优化编辑器,15+应用进行更新

为了满足用户不断增长的需求,我们持续努力提升产品的功能和性能,以更好地支持用户的工作。 数环通12月的最新产品更新已经正式发布,带来了一系列强大的功能,以提升您的工作效率和系统的可靠性。 更新快速预览 新增&优化功能&a…

软考-软件设计师 知识点整理(一篇就过了 建议收藏)

文章目录 一 计算机组成CPU寻址方式校验码奇偶校验码(只能检一位错,并且不能纠错)循环冗余校验码CRC(只能检错,不能纠错)海明码 计算机体系结构分类Flynn分类法(理论存在:多指令单数…

Java LeetCode篇-二叉搜索树经典解法(实现:二叉搜索树的最近公共祖先、根据前序遍历建树等)

🔥博客主页: 【小扳_-CSDN博客】 ❤感谢大家点赞👍收藏⭐评论✍ 文章目录 1.0 判断合法 1.1 使用遍历方式实现验证二叉搜索树 1.2 使用递归方式实现验证二叉搜索树 2.0 求范围和 2.1 使用非递归实现二叉搜索树的范围和 2.2 使用递归方式实现…

腾讯云com域名注册1元一年,非常可以!

腾讯云com域名注册优惠价格1元首年,条件是企业新用户,个人新用户注册com域名是33元首年,第二年续费价格85元一年。活动 txybk.com/go/domain-sales 活动打开如下图: 腾讯云com域名注册优惠价格 腾讯云com域名注册原价是85元一年&a…

已解决 ValueError: Setting an array element with a sequence. 问题

博主猫头虎的技术世界 🌟 欢迎来到猫头虎的博客 — 探索技术的无限可能! 专栏链接: 🔗 精选专栏: 《面试题大全》 — 面试准备的宝典!《IDEA开发秘籍》 — 提升你的IDEA技能!《100天精通Golang》…

C++类与对象基础(6)

(注:本篇文章介绍部分内容时,需要用到上盘文章中日期类的代码,文章链接如下:C类与对象基础(5)——日期类的实现-CSDN博客​​​​​​) 目录 1. 运算符重载的相关补充: 1.1流运算符重载出现的问题&#x…

李沐-《动手学深度学习》--03-注意力机制

一、注意力机制 1 . 注意力提示 1)框架 **随意:**跟随自己的想法的,自主的想法,例如query **不随意:**没有任何偏向的选择,例如 Keys 如何得到 k v q 2)Nadaraya-Watson核回归 就是一个so…

《2024 AIGC 应用层十大趋势白皮书》:近屿智能OJAC带您一起探索AI未来

Look!👀我们的大模型商业化落地产品📖更多AI资讯请👉🏾关注Free三天集训营助教在线为您火热答疑👩🏼‍🏫 近日国际知名咨询机构IDC发布《2024 AIGC 应用层十大趋势白皮书》的发布&am…

Spring 动态数据源事务处理

在一般的 Spring 应用中,如果底层数据库访问采用的是 MyBatis,那么在大多数情况下,只使用一个单独的数据源,Spring 的事务管理在大多数情况下都是有效的。然而,在一些复杂的业务场景下,如需要在某一时刻访问不同的数据库,由于 Spring 对于事务管理实现的方式,可能不能达…

二叉树OJ练习(二)

1. 二叉树的最近公共祖先 题目描述: ​ 题解: 1.p或者q其中一个等于root,那么root就是最进公共祖先 2.p和q分布在root的左右两侧,那么root就是最进公共祖先 3.p和q在root的同一侧,就是要遍历这棵树,遇到p或者q返回 ​…

一款好的葡萄酒关键在哪里?

除了易于种植,赤霞珠还因其独特的口感、难以置信的味道和质量而闻名。这种葡萄主要用于中高端干红葡萄酒,通常表现出成熟的黑色水果味道,带有辛辣和泥土气息。 在橡木桶中陈酿后,赤霞珠表现极佳。随着葡萄酒的陈年,橡木…

【金猿人物展】数元灵科技CEO朱亚东:何以数智化

‍ 朱亚东 本文由数元灵科技CEO朱亚东撰写并投递参与“数据猿年度金猿策划活动——2023大数据产业年度趋势人物榜单及奖项”评选。 大数据产业创新服务媒体 ——聚焦数据 改变商业 在大数据经济的高速发展下,数据已经成为第5生产要素。打造以数据驱动为中心的标准化…

腾讯云免费服务器申请1个月攻略,亲测可行教程

腾讯云免费服务器申请入口 https://curl.qcloud.com/FJhqoVDP 免费服务器可选轻量应用服务器和云服务器CVM,轻量配置可选2核2G3M、2核8G7M和4核8G12M,CVM云服务器可选2核2G3M和2核4G3M配置,腾讯云服务器网txyfwq.com分享2024年最新腾讯云免费…

NUXT3学习笔记

1.邂逅SPA、SSR 1.1 单页面应用程序 单页应用程序 (SPA) 全称是:Single-page application,SPA应用是在客户端呈现的(术语称:CSR(Client Side Render)) SPA的优点 只需加载一次 SPA应用程序只需…

(二)Explain使用与详解

explain中的列 sql语句: EXPLAIN SELECT * from user WHERE userId=1340; 执行结果: 1. id列 id列的编号是 select 的序列号,有几个 select 就有几个id,并且id的顺序是按 select 出现的顺序增长的。 id列越大执行优先级越高,id相同则从上往下执行,id为NULL最后执行…

Chrome您的连接不是私密连接或专用连接

方法一: 在当前页面用键盘输入 thisisunsafe ,不是在地址栏输入,就直接敲键盘就行了因为Chrome不信任这些自签名ssl证书,为了安全起见,直接禁止访问了,thisisunsafe 这个命令,说明你已经了解并…

富文本编辑器

富文本:带样式,多格式的文本,在前端一般使用标签配合内联样式实现 富文本编辑器(Rich Text Editor,简称 RTE)是一种用户可以使用来创建格式化的文本内容的界面组件。它通常可以嵌入到网页或应用程序中&…