4.5 A TILED MATRIX MULTIPLICATION KERNEL

news2024/9/22 7:34:56

我们现在准备展示一个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.coloradmin.cn/o/1367030.html

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

相关文章

过滤器和拦截器

上篇文章我们学习了 Session 认证和 Token 认证,这篇我们来学习一下过滤器和拦截器,过滤器和拦截器在日常项目中经常会用到。 一、过滤器 1.1、理论概念 过滤器 Filter 是 JavaWeb 三大组件(Servlet、Filter、Listener)之一&am…

为什么我国的计算机教育那么差?

建议看看计算机科学速成课,一门很全面的计算机原理入门课程,短短10分钟可以把大学老师十几节课讲不清楚的东西讲清楚!整个系列一共41个视频,B站上有中文字幕版。 每个视频都是一个特定的主题,例如软件工程、人工智能、…

CnosDB容灾方案概述

本文主要介绍了跟容灾相关的关键技术以及技术整合后形成的几种具体方案,每种方案都在RTO、RPO、部署成本和维护成本等方面有自己的特点和区别,可以根据具体场景选择最合适的方案。 基本概念 RTO(Recovery Time Objective)&#x…

计算机操作系统进程同步(信号量pv专题)

文章目录 一 基本概念1.1 多道程序中的制约关系1.2 临界资源(Critical Resouce)1.3 三区:进入区、临界区、退出区 二 同步机制应遵循的原则三 信号量机制类型3.1 整型信号量3.2 记录型信号量3.3 AND型信号量3.4 信号量集 四 信号量的应用4.1 信号量实现进程互斥4.2 …

Qt读取文件对比:每次获取自定义的长度和使用系统的API,耗时对比

0. 前言 在编程过程中,经常遇到文件读写操作,太频繁了。每次也都写的不一样。 突发奇想,想测试下几种不同的读取文件的效率。 测试以下三种方式读取文件效率: 自定义读取文件耗时使用QFile类API读取文件耗时使用QTextStream类AP…

黑马程序员Java项目实战《瑞吉外卖》,轻松掌握springboot + mybatis plus开发核心技术的真java实战项目——第四部分

黑马程序员Java项目实战《瑞吉外卖》,轻松掌握springboot mybatis plus开发核心技术的真java实战项目——第四部分 1. 套餐管理1.1 新增套餐1.1.1 添加菜品数据回显 1.2 保存添加套餐1.3 套餐信息分页查询1.4 删除套餐1.5 需要自己单独实现的功能1.5.1 套餐管理的启…

qt.network.ssl: QSslSocket::connectToHostEncrypted: TLS initialization failed

方法一: 如果是https,改为http。 方法二: Qt 解决qt.network.ssl: QSslSocket::connectToHostEncrypted: TLS initialization failed问题-CSDN博客 其他:

Elasticsearch零基础实战

分享后可优化点(待完成) java es8 查询如何打印查询入参 ?(直接执行的json) es自定义分词器 如何实现? kibana 监控jvm分子分母是什么 ? es如何 改索引结构? 修改数据原理 分享…

JDBC-数据库连接池(druid)

一、背景 在介绍JDBC基本概念中,似乎Java程序每次与数据库交互都要通过驱动创建一个新的连接对象(Connection),再由连接对象创建一个可执行SQL的Statement对象(或PreparedStatement对象),操作完…

海康威视摄像头+服务器+录像机配置校园围墙安全侦测区域入侵侦测+越界侦测

一、适用场景 1、校园内,防止课外时间翻越围墙到校外、从校外翻越围墙到校内; 2、通过服务器摄像头的侦测功能及时抓图保存,为不安全因素提供数字化依据; 3、网络录像机保存监控视频,服务器保存抓拍到的入侵与越界&am…

学习笔记16——操作系统

学习笔记系列开头惯例发布一些寻亲消息,感谢关注! 链接:https://www.mca.gov.cn/lljz/indexdetail.html?idd0afa7f6f36946319a206d61937f9b63&type0&t10.11199120579373845 八股——操作系统一些基础知识整理 一个java程序对应一个…

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

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

*4.3 CUDA MEMORY TYPES

CUDA设备包含几种类型的内存,可以帮助程序员提高计算到全局内存的访问率,从而实现高执行速度。图4.6显示了这些CUDA设备内存。全局内存和恒定内存出现在图片的底部。主机可以通过调用API函数来写入(W)和读取(R&#xf…

PHP反序列化漏洞利用及修复,示例代码讲解

您提到的PHP反序列化漏洞是一个重要的网络安全问题。在我的网络安全工程师的角色下,我可以提供关于此问题的深入分析。 PHP反序列化漏洞通常发生在当不可信的数据被反序列化时。序列化是将数据结构或对象状态转换为可存储或可传输的格式的过程,而反序列…

快速幂算法总结

知识概览 快速幂可以在O(logk)的时间复杂度之内求出来的结果。 例题展示 快速幂 题目链接 活动 - AcWing 系统讲解常用算法与数据结构,给出相应代码模板,并会布置、讲解相应的基础算法题目。https://www.acwing.com/problem/content/877/ 代码 #inc…

机器学习 前馈神经网络

人工神经网络(Artificial Neural Network,ANN)是指一系列受生物学和神经科学启发的数学模型.这些模型主要是通过对人脑的神经元网络进行抽象,构建人工神经元,并按照一定拓扑结构来建立人工神经元之间的连接…

一文讲透使用Python绘制双纵轴线图

双纵轴线图主要用来展示两个因变量和一个自变量的关系,并且两个因变量的数值单位不同。具体来说,双纵轴线图是指在一幅图上有一个横轴和两个纵轴,适用于三个变量。两个纵轴分别表示一个变量,横轴变量同时适用于两个纵轴上的变量&a…

Selenium-java元素等待三种方式

第二种方式需要写在创建driver时的代码下面 第三种则是对每个定位元素进行配置

Linux基础知识总结

目录 一、Linux权限设置 更改文件属性 chgrp - 更改文件属组 chown - 更改文件所有者,也可以同时更改文件所属组。 chmod - 更改文件属性 二、Linux文件与目录管理 处理目录的常用命令 ls(list files)- 列出目录及文件名 cd&#xff…

【Linux】Linux系统编程——Linux命令解析器

【Linux】Linux系统编程——Linux命令解析器 什么是Linux 命令解析器? Linux 命令解析器,通常被称为 shell,是 Linux 操作系统中的一个关键组件。它充当用户和系统内核之间的接口,允许用户通过输入命令来控制和管理操作系统和应…