7.3 CONSTANT MEMORY AND CACHING

news2024/12/23 1:28:57

掩模数组M在卷积中的使用方式有三个有趣的属性。首先,M阵列的大小通常很小。大多数卷积掩模在每个维度上都少于10个元素。即使在3D卷积的情况下,掩码通常也只包含少于1000个元素。其次,在内核执行过程中,M的内容不会改变。第三,所有线程都需要访问掩码元素。更好的是,所有线程都以相同的顺序访问M元素,从M[0]开始,并通过图7.6.中for循环的迭代一次移动一个元素。这两个属性使掩码数组成为恒定内存和缓存的绝佳候选(图7.7)。
在这里插入图片描述
正如我们在第5章(性能考虑因素)中讨论的那样,CUDA编程模型允许程序员在常量内存中声明一个变量。与全局内存变量一样,常量内存变量对所有线程块也是可见的。主要区别在于,在内核执行期间,线程不能更改常量内存变量。此外,恒定内存的大小相当小,目前为64KB。

为了使用常量内存,主机代码需要以与全局内存变量不同的方式分配和复制常量内存变量。要在常量内存中声明M数组,主机代码将其声明为全局变量,如下所示:

#define MAX_MASK_WIDTH 10
__constant__ float M[MAX_MASK_WIDTH];

这是一个全局变量声明,应该在源文件中的任何函数之外。关键字__constant__(每边两个下划线)告诉编译器,数组M应放入设备常量内存中。

假设主机代码已经在带有Mask_Width元素的主机内存中的掩码M_h数组中分配并初始化了掩码。M_h的内容可以在设备常量存储器中传输到M,如下所示:
在这里插入图片描述
请注意,这是一个特殊的内存复制函数,它通知CUDA运行时,在内核执行期间,复制到常量内存的数据不会更改。一般来说,cudaMemcpyToSymble()函数的使用如下:在这里插入图片描述
其中dest是指向常量内存中目标位置的指针,src是指向主机内存中源数据的指针,大小是要复制的字节数。

内核函数作为全局变量访问常量内存变量。因此,它们的指针不需要作为参数传递给内核。我们可以修改内核以使用常量内存,如图7.8.所示。请注意,内核看起来与图7.6.中几乎相同。唯一的区别是,M不再通过作为参数传入的指针访问。它现在作为主机代码声明的全局变量进行访问。请记住,全局变量的所有C语言范围规则都适用于这里。如果hostl代码和内核代码在不同的文件中,内核代码文件必须包含相关的外部声明信息,以确保M的声明对内核可见。

与全局内存变量一样,恒定内存变量也位于DRAM中。然而,由于CUDA运行时知道常量内存变量在内核执行期间不会被修改,它指示硬件在内核执行期间积极缓存常量内存变量。为了了解恒定内存使用的好处,我们首先需要更多地了解现代处理器内存和缓存层次结构。在这里插入图片描述
正如我们在第5章中讨论的,性能考虑,DRAM的长延迟和有限的带宽几乎是所有现代处理器的主要瓶颈。为了减轻内存瓶颈的影响,现代处理器通常使用片上缓存存储器或缓存,以减少需要从主存储器(DRAM)访问的变量数量,如图7.9所示。
在这里插入图片描述
CUDA共享内存或一般的暂存不同,缓存对程序是“透明”的。也就是说,为了使用CUDA共享内存,程序需要将变量声明为__shared_ _,并显式地将全局内存变量移动到共享内存变量中。另一方面,在使用缓存时,程序只需访问原始变量。处理器硬件将自动在缓存中保留一些最近或最常用的变量,并记住其原始DRAM地址。当稍后使用其中一个保留的变量时,硬件将从其地址中检测出该变量的副本在缓存中可用。然后,变量的值将从缓存中提供,无需访问DRAM。

内存的大小和内存的速度之间存在权衡。因此,现代处理器通常使用多个级别的缓存。这些缓存级别的编号约定反映了与处理器的距离。最低级别,L1或1级,是直接连接到处理器核心的缓存。它在延迟和带宽方面都以非常接近处理器的速度运行。然而,L1缓存体积较小,通常在16KB到64KB之间。L2缓存更大,范围在128KB到1MB之间,但可能需要数十个周期才能访问。它们通常在多个处理器内核或CUDA设备中的SM之间共享。在今天的一些高端处理器中,甚至有L3缓存,大小可以为几MB。

在大规模并行处理器中使用缓存的一个主要设计问题是缓存一致性,当一个或多个处理器内核修改缓存数据时,就会出现缓存一致性。由于L1缓存通常只直接连接到其中一个处理器内核,因此其他处理器内核不容易观察到其内容的变化。如果修改后的变量在不同处理器内核上运行的线程之间共享,这会导致问题。需要一个缓存一致性机制,以确保其他处理器内核的缓存内容得到更新。在大规模并行处理器中提供缓存一致性既困难又昂贵。然而,它们的存在通常简化了并行软件开发。因此,现代CPU通常支持处理器内核之间的缓存一致性。虽然现代GPU提供两级缓存,但它们通常没有缓存一致性,以最大限度地利用可用的硬件资源,以增加处理器的算术吞吐量。

恒定内存变量在大规模并行处理器中使用缓存中起着有趣的作用。由于它们在内核执行期间没有更改,因此在内核执行期间没有缓存一致性问题。因此,硬件可以积极缓存L1缓存中的常量变量值。此外,这些处理器中的缓存设计通常经过优化,以向大量线程广播值。因此,当warp中的所有线程访问相同的恒定内存变量时,就像M的情况一样,缓存可以提供大量的带宽来满足线程的数据需求。此外,由于M的大小通常很小,我们可以假设所有M元素都有效地从缓存中访问。因此,我们可以简单地假设**没有在M访问上花费DRAM带宽。**通过使用恒定内存和缓存,我们有效地将浮点算术与内存访问的比率增加了一倍,达到2。

事实证明,对输入N阵列元素的访问也可以从较新的GPU中的缓存中受益。我们将在第7.5节中回到这一点。

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

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

相关文章

移动端对大批量图片加载的优化方法(三)

移动端对大批量图片加载的优化方法(三)Flutter 本篇主要从Flutter开发中可以使用到的对大批量图片加载的优化方法进行整理。 1.合适的图片格式 详情请参考移动端对大批量图片加载的优化方法(一)。 2.缓存机制 在Flutter中&am…

2023年山东省职业院校技能大赛高职组信息安全管理与评估—内存取证解析

内存取证 1、从内存中获取到用户admin的密码并且破解密码,以Flag{admin,password} 形式提交(密码为 6 位); volatility -f 1.vmem imageinfo 操作系统我们一般取第一个就可以了

vulhub中的Apache SSI 远程命令执行漏洞

Apache SSI 远程命令执行漏洞 1.cd到ssi-rce cd /opt/vulhub/httpd/ssi-rce/ 2.执行docker-compose up -d docker-compose up -d 3.查看靶场是否开启成功 dooker ps 拉取成功了 4.访问url 这里已经执行成功了,注意这里需要加入/upload.php 5.写入一句话木马 &…

Qt实现简单的分割窗口

最近在学习一些关于Qt的新知识,今天来讲述下我学习到的窗口分割,如果有不正确的,大家可以指正哦~ 首先,先看一下实现之后的简单效果吧!省的说的天花乱坠,大家却不知道说的是哪个部分。 功能实现 整体demo…

阿里云RDMA通信库XRDMA论文详解

RDMA(remote direct memory access)即远端直接内存访问,是一种高性能网络通信技术,具有高带宽、低延迟、无CPU消耗等优点。RDMA相比TCP在性能方面有明显的优势,但在编程复杂度上RDMA verbs却比TCP socket复杂一个数量级。 开源社区和各大云厂…

如何在iOS手机上查看应用日志

引言 在开发iOS应用过程中,查看应用日志是非常重要的一项工作。通过查看日志,我们可以了解应用程序运行时的状态和错误信息,帮助我们进行调试和排查问题。本文将介绍两种方法来查看iOS手机上的应用日志,并提供相应的操作步骤。 …

Python学习之路——文件部分【书接上回】

一、书接上回 上个博客我说过,为什么最开始的时候一定要将文件内的中文的逗号替换为英文的逗号,接下来,请看(其实想一想,感觉没必要,不过也是好的,总要练练手的嘛) def func03(st…

Spring——基于注解的AOP配置

基于注解的AOP配置 1.创建工程 1.1.pom.xml <?xml version"1.0" encoding"UTF-8"?> <project xmlns"http://maven.apache.org/POM/4.0.0"xmlns:xsi"http://www.w3.org/2001/XMLSchema-instance"xsi:schemaLocation"…

基于DNA的密码学和隐写术综述

摘要 本文全面调研了不同的脱氧核糖核酸(DNA)-基于密码学和隐写术技术。基于DNA的密码学是一个新兴领域,利用DNA分子的大规模并行性和巨大的存储容量来编码和解码信息。近年来,由于其相对传统密码学方法的潜在优势,如高存储容量、低错误率和对环境因素的抗性,该领域引起…

怎么把图片表格转成word表格?教你简单转换

怎么把图片表格转成word表格&#xff1f;在我们的日常工作中&#xff0c;经常会遇到需要将图片表格转换成Word表格的情况。这样不仅可以方便我们编辑、整理数据&#xff0c;还能提高工作效率。那么&#xff0c;如何将图片表格快速、准确地转换成Word表格呢&#xff1f;下面就为…

openGauss学习笔记-190 openGauss 数据库运维-常见故障定位案例-服务启动失败

文章目录 openGauss学习笔记-190 openGauss 数据库运维-常见故障定位案例-服务启动失败190.1 服务启动失败190.1.1 问题现象190.1.2 原因分析190.1.3 处理办法 openGauss学习笔记-190 openGauss 数据库运维-常见故障定位案例-服务启动失败 190.1 服务启动失败 190.1.1 问题现…

国芯科技荣膺高工智能汽车“年度车规MCU高成长供应商”,加速产品精准化系列化布局

2023年12月13—15日&#xff0c;2023&#xff08;第七届&#xff09;高工智能汽车年会在上海召开&#xff0c;大会以“寻找拐点”为主题&#xff0c;通过超80场主题演讲及多场圆桌对话&#xff0c;为智能汽车赛道参与者「备战2024」提供全方位的决策支持。 作为汽车电子芯片领…

DNS解析原理和k8s DNS 实践

1. 问题背景 1.1 域名解析异常 近期开发的一个功能&#xff0c;需要在k8s集群容器环境中调用公司内部api&#xff0c;api提供了内网域名&#xff0c;解析内网域名异常导致请求超时&#xff0c;因此梳理了下DNS的知识点。 可以先看到下面&#x1f447;这段配置&#xff0c;修…

基于net6的asp.net core webapi项目打包为docker镜像,并推送至私有镜像仓库harbor中

基于net6的asp.net core webapi项目打包为docker镜像&#xff0c;并推送至私有镜像仓库harbor中 0、环境说明1、打包步骤1.1 创建Asp.net core WebApi项目1.2 在Asp.net core WebApi项目根目录下创建Dockerfile文件1.3 在子系统Ubuntu20.04.4中通过docker build生成docker镜像1…

资源分享:GIS插件-海怪工具箱(附下载链接)

工具详情&#xff1a; ArcGIS海怪工具箱&#xff0c;可直接在ArcGIS中安装使用。 功能&#xff1a; 1、按属性编码&#xff1a;根据字段编码&#xff0c;相同的字段值同一个编码&#xff0c;然后依次升序。 可以手动指定编码样式&#xff0c;比如输入1&#xff0c;那么后续编…

three.js实现旋转棱锥效果

three.js实现旋转棱锥效果 图例 步骤 创建一个四面棱锥添加贴图render函数中旋转&#xff0c;修改高度 代码 <template><div class"app"><div ref"canvesRef" class"canvas-wrap"></div></div> </template…

Python——欢迎来到吱昂张游乐园

欢迎来到吱昂张游乐园&#xff01;&#xff01;&#xff01; 凡是身高小于120或者您的vip等级大于三级的皆可免费游玩。 那我们接下来就来设计一下以上的规则叭 print("欢迎来到吱昂张游乐园") if int(input("输入您的身高&#xff1a;"))>120:print…

Python print 高阶玩法

Python print 高阶玩法 当涉及到在Python中使用print函数时&#xff0c;有许多方式可以玩转文本样式、字体和颜色。在此将深入探讨这些主题&#xff0c;并介绍一些print函数的高级用法。 1. 基本的文本样式与颜色设置 使用ANSI转义码 ANSI转义码是一种用于在终端&#xff0…

62.网游逆向分析与插件开发-游戏增加自动化助手接口-游戏公告类的C++还原

内容来源于&#xff1a;易道云信息技术研究院VIP课 上一个内容&#xff1a;游戏红字公告功能的逆向分析-CSDN博客 码云地址&#xff08;master分支&#xff09;&#xff1a;https://gitee.com/dye_your_fingers/sro_-ex.git 码云版本号&#xff1a;0888e34878d9e7dd0acd08ef…

HarmonOS 日期选择组件(DatePicker)

本文 我们一起来看基础组件中的 DatePicker 这就是 日程开发中的日期组件 他可以创建一个日期的范围 并创建一个日期的滑动选择器 这里 我们先写一个组件的骨架 Entry Component struct Index {build() {Row() {Column() {}.width(100%)}.height(100%)} }然后 我们先在Column组…