7.3 CONSTANT MEMORY AND CACHING

掩模数组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.mzph.cn/news/625310.shtml

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

相关文章

启动Vue项目,报错:‘vue-cli-service‘ 不是内部或外部命令,也不是可运行的程序

前言: 最近在打开一个Vue项目的时候,打开之后输入命令行:npm run serve之后发现,报错:vue-cli-service 不是内部或外部命令,也不是可运行的程序,以下是解决方案: 报错图片截图&…

HNU-算法设计与分析-实验3

算法设计与分析实验3 计科210X 甘晴void 202108010XXX 目录 文章目录 算法设计与分析<br>实验31 用Dijkstra贪心算法求解单源最短路径问题问题重述证明模板&#xff1a;Dijkstra算法代码验证算法分析 1【扩展】 使用堆优化的Dijkstra原因代码算法分析验证 2 回溯法求解…

运筹说 第98期|无约束极值问题

上一期我们一起学习了关于非线性规划问题的一维搜索方法的相关内容&#xff0c;本期小编将带大家学习非线性规划的无约束极值问题。 下面&#xff0c;让我们从实际问题出发&#xff0c;学习无约束极值问题吧&#xff01; 一、问题描述及求解原理 1 无约束极值问题的定义 无约…

ArkUI-X跨平台已至,何需其它!

运行环境 DevEco Studio&#xff1a;4.0Release OpenHarmony SDK API10 开发板&#xff1a;润和DAYU200 自从写了一篇ArkUI-X跨平台的文章之后&#xff0c;好多人都说对这个项目十分关注。 那么今天我们就来完整的梳理一下这个项目。 1、ArkUI-X 我们之前可能更多接触的…

登录验证

目录 会话技术 Cookie Session JWT JWT生成 JWT校验 会话技术 会话 打开浏览器&#xff0c;访问web服务器的资源&#xff0c;会话建立&#xff0c;直到有一方断开连接&#xff0c;会话结束。在一次会话中可以包含多次请求与响应 会话跟踪 一种维护浏览器的方法 服务器需要…

性能测试jmeter

选的这些怎么添加 在一个列表里面 方法调用${__time(YMD)} 两个下划线&#xff0c;后跟函数名&#xff0c;小括号内是输入参数&#xff0c;整个用大括号包裹。 注意POST一定要在消息体数据里面写,不能再参数里面 否则报错:loginOut,没cookie等

VueCli-自定义创建项目

参考 1.安装脚手架 (已安装可以跳过) npm i vue/cli -g2.创建项目 vue create 项目名 // 如&#xff1a; vue create dn-demo键盘上下键 - 选择自定义选型 Vue CLI v5.0.8 ? Please pick a preset:Default ([Vue 3] babel, eslint)Default ([Vue 2] babel, eslint) > M…

小迪安全第二天

文章目录 一、Web应用&#xff0c;架构搭建二、web应用环境架构类三、web应用安全漏洞分类总结 一、Web应用&#xff0c;架构搭建 #网站搭建前置知识 域名&#xff0c;子域名&#xff0c;dns,http/https,证书等 二、web应用环境架构类 理解不同web应用组成角色功能架构 开发…

显示CPU架构的有关信息 lscpu

文章目录 显示CPU架构的有关信息 lscpu默认实例更多信息 显示CPU架构的有关信息 lscpu Linux的CPU设备查看器。lscpu命令用来显示cpu的相关信息。 lscpu从sysfs和/proc/cpuinfo收集cpu体系结构信息&#xff0c;命令的输出比较易读 。 命令输出的信息包含cpu数量&#xff0c;线…

tensorflow报错: DNN library is no found

错误描述 如上图在执行程序的时候&#xff0c;会出现 DNN library is no found 的报错 解决办法 这个错误基本上说明你安装的 cudnn有问题&#xff0c;或者没有安装这个工具。 首先检测一下你是否安装了 cudnn 进入CUDA_HOME下&#xff0c;也就是进入你的cuda的驱动的安装目…

个人数据备份方案分享(源自一次悲惨经历)

文章目录 1 起源2 备份架构2.1 生活照片2.2 生活录音2.3 微信文件2.4 工作文件2.5 笔记、影视音乐、书籍 3 使用工具介绍3.1 小米云服务3.2 中国移动云盘3.3 小米移动硬盘&#xff08;1T&#xff09;3.4 FreeFileSync 4 总结 1 起源 本文的灵感源于我个人的一次不幸遭遇&#…

领域驱动设计——DDD领域驱动设计进阶

摘要 进阶篇主要讲解领域事件、DDD 分层架构、几种常见的微服务架构模型以及中台设计思想等内容。如何通过领域事件实现微服务解耦&#xff1f;、怎样进行微服务分层设计&#xff1f;、如何实现层与层之间的服务协作&#xff1f;、通过几种微服务架构模型的对比分析&#xff0…

记一个有关 Vuetify 组件遇到的一些问题

Vuetify 官网地址 所有Vuetify 组件 — Vuetify 1、Combobox使用对象数组 Combobox 组合框 — Vuetify items数据使用对象数组时&#xff0c;默认选中的是整个对象&#xff0c;要对数据进行处理 <v-comboboxv-model"defaultInfo.variableKey":rules"rules…

基于springboot体育场馆运营管理系统源码

基于springboot体育场馆运营管理系统源码330 -- MySQL dump 10.13 Distrib 5.7.31, for Linux (x86_64) -- -- Host: localhost Database: springboot3cprm -- ------------------------------------------------------ -- Server version 5.7.31/*!40101 SET OLD_CHARACT…

网络安全全栈培训笔记(53-WEB攻防-通用漏洞CRLF注入URL重定向资源处理拒绝服务)

第53天 WEB攻防-通用漏洞&CRLF注入&URL重定向&资源处理拒绝服务 知识点&#xff1a; 1、CRLF注入-原理&检测&利用 2、URL重定向-原理&检测&利用 3、Web拒绝服务-原理&检测&利用 #下节预告&#xff1a; 1、JSONP&CORS跨域 2、域名安全…

嵌入式软件工程师面试题——2025校招社招通用(十八)

说明&#xff1a; 面试群&#xff0c;群号&#xff1a; 228447240面试题来源于网络书籍&#xff0c;公司题目以及博主原创或修改&#xff08;题目大部分来源于各种公司&#xff09;&#xff1b;文中很多题目&#xff0c;或许大家直接编译器写完&#xff0c;1分钟就出结果了。但…

共识算法介绍

文章目录 共识算法Paxos 算法三种角色一致性提交算法prepare 阶段accept 阶段commit 阶段 CAP 定理BASE 理论Zookeeper 算法实现三类角色三个数据三种模式四种状态消息广播算法Leader选举算法 共识算法 Paxos 算法 Paxos 算法是莱斯利兰伯特(Leslie Lamport)1990 年提出的一种…

基于Java (spring-boot)的社团管理系统

一、项目介绍 系统管理员的功能概述&#xff1a; ①用户管理 a.注册用户账户 当一个新用户注册时&#xff0c;用户填写基本信息并上传。用户基本信息包括账号、 姓名、密码、手机、地址等信息。 b.用户信息管理 管理员可以查看系统所有用户的基本信息&#xff0c;并修改和…

面试官常问问题:Java中的128陷阱详解

看这样两段代码&#xff0c;思考结果返回的是什么 Integer num1 100; Integer num2 100; System.out.println(num1 num2);Integer num3 128; Integer num4 128; System.out.println(num3 num4); 揭晓答案&#xff1a;第一段代码的结果是true&#xff0c;第二段代码的结…

数据结构学习 jz59 滑动窗口的最大值

关键词&#xff1a;排序 大顶堆 双端队列 题目&#xff1a; 望远镜中最高的海拔 方法一&#xff1a;维护一个辅助队列。 方法二&#xff1a;大顶堆。 我还在主站 239 写了找最小值的方法。 方法一&#xff1a;最优解 这个方法和jz30维护一个非严格递减的辅助栈是基本一样的…