CUDA简介——CUDA内存模式

1. 引言

前序博客:

  • CUDA简介——基本概念
  • CUDA简介——编程模式
  • CUDA简介——For循环并行化
  • CUDA简介——Grid和Block内Thread索引

在这里插入图片描述

CUDA内存模式,采用分层设计,是CUDA程序与正常C程序的最大不同之处:

  • Thread-Memory Correspondence
  • Block-Memory Correspondence
  • Grid-Memory Correspondence

总体的CUDA内存模式为:
在这里插入图片描述

  • 1)Registers & Local Memory:声明在Kernel内的常规变量。小空间变量存储于Register中,大空间变量存储于Local Memory中。因Local Memory操作速度慢,应尽量避免使用Local Memory。
  • 2)Shared Memory:允许同一Block内的Threads间相互通信。
  • 3)Constant Memory:用于存储Kernel执行期间不变的数据。会缓存到On-Chip memory中,可大大降低Kernel执行期间Global Memory的通信吞吐量。
  • 4)Global Memory:用于存储与Host交互的数据。

2. Thread-Memory Correspondence

在这里插入图片描述

Thread-Memory Correspondence,即Threads等价为Local Memory (and Registers):

  • 其范围为:对相应Thread是private的,对其它Threads来说是不可访问的。
  • 其生命周期为:Thread。当Thread执行完成,与该Thread相关的任何local memory (and Registers) 都将自动释放。
  • 不过,local memory和registers存在完全不同的性能特性。

3. Block-Memory Correspondence

在这里插入图片描述
Block-Memory Correspondence,即Blocks等价为Shared Memory:

  • 其范围为:同一Block内的每个Thread均可访问。
  • 其生命周期为:Block。当Block执行完成,其shared memory内容都将自动释放。

4. Grid-Memory Correspondence

在这里插入图片描述
Grid-Memory Correspondence,即Grids等价为Global Memory:

  • 其范围为:所有 Grids内的每个Thread均可访问。
  • 其生命周期为:Host代码内的整个main()程序。或可通过在Host代码中手工调用cudaFree(...)来释放。

5. Device内存模式

在这里插入图片描述

  • Host:由CPU及机器内存等组成。
  • Device:由GPU及其DRAM等组成。【Device图上的绿色方格,表示的是一个CUDA core。】

在这里插入图片描述

Device的DRAM中,有:

  • Global Memory物理空间
  • Local Memory物理空间。需注意,此处的Local,并不是指其物理位置;此处的Local是指该内存空间的scope(范围)和 lifetime(生命周期)。

而Device的GPU中,有:

  • Registers物理空间
  • Shared Memory物理空间

Device图上的绿色方格,表示的是一个CUDA core。Device上的CUDA cores组合在一起,成为streaming Multiprocessor(简称为SM)。
Device图上的黄色方格,表示SM。黄色方格组合在一起为CUDA cores集合。

位于SM上的内存,称为:

  • “On-Chip” Device memory。因此,Registers和Shared Memory均对应为 “On-Chip” Device memory。因Registers和Shared Memory均物理存在与GPU的streaming Multiprocessor中。

非SM上的内存,称为::

  • “Off-Chip” Device memory。因其并不存在与GPU之上。对应,Global Memory和Local Memory均为“Off-Chip” Device memory。也即Device上的DRAM为 “Off-Chip” Device memory。

以NVIDIA GPU Geforce Titan 物理布局为例:
在这里插入图片描述

  • 上图蓝色框所示,为Device的DRAM,均为 “Off-Chip” Device memory。
  • 上图绿色框,即为实际的GPU,对应为“On-Chip” Device memory。

理解Blocks如何映射到SM,是设计kernel的基本要求,从而获得优化的GPU计算性能。

5.1 内存速度

不同的内存空间,其带宽和延迟各不相同:
在这里插入图片描述

  • on-chip memory操作速度,要快于off-chip memory。

5.2 Global Memory访问

在这里插入图片描述
在这里插入图片描述

Global Memory访问方式有:

  • cudaMalloc()
  • cudaMemset()
  • cudaMemCopy()
  • cudaFree()

无法避免不使用Global Memory,因必须使用Global Memory空间,来将数据由host传递到device。不过,应尽量减少Global Memory通讯量,因为其速度很慢。

Global Memory的优势在于:

  • 其通常很大。如计算机内存为8GB或16GB,而Titan和Tesla k40,均有6GB的global DRAM。

5.3 Registers and Local Memory

在这里插入图片描述
Kernel中声明的变量,存储于Register中:

  • 对应为On-Chip Device memory。
  • 为最快的内存形式。

太大不适于Register的数组,将存储在Local Memory中:

  • 对应为Off-Chip Device memory。
  • 由编译器控制。
  • “Local”是指范围,而不是位置。此处“Local”,是指相对每个Thread的Local Memory。
    • 每个Thread均有其自己的private local memory和registers,对其它Thread不可访问。
  • 应尽量避免,因Local Memory为最慢的内存形式之一。Registers为最快的内存形式。
    • 因此,设计目标为避免将更多信息存储于local变量中,以免超过register的存储空间。
    • register space为稀缺硬件资源。应更好地规划充分利用。

5.4 Shared Memory

在这里插入图片描述
借助Shared Memory:

  • 支持同一Block内的Threads之间相互通信:
    • 同步方式通信
  • Shared Memory为非常特殊的内存,对实现计算性能和正确性至关重要:
    • Shared Memory处理速度快,仅次于Registers。因其为On-Chip device memory。
    • 支持Block内的Threads相互通信,可将其看成是用户定义的L1 Cache,可用作“scratch-pad(高速暂存存储器)”内存。后续将介绍Shared Memory和L1 Cache关系密切。

在这里插入图片描述

使用__shared___关键字来表示分配的为shared memory:
在这里插入图片描述

5.5 Constant Memory

Constant Memory为Device Memory的特殊区域:

  • 用于存储Kernel执行过程中不变的数据。
  • 对Kernel来说是只读的。
  • Constant Memory为Off-Chip Device Memory。
  • Constant Memory 积极缓存到 On-Chip Memory中。

在这里插入图片描述
Constant Memory的思想在于:

  • GPU没有很大的cache。从而可使用constant memory来实现很简单的cache类型。
  • Constant Memory可以很大,因其实际位于Device DRAM中。所有Threads均可访问Constant Memory,但其为只读内存。
  • 对于需频繁访问,但Kernel执行过程中不变的数据,可使用Constant Memory。
  • Constant Memory是在off-chain DRAM硬件中实现的,但实际其内容会积极缓存到On-Chip Memory中。因此,使用Constant Memory,可大大降低Kernel执行期间Global Memory的通信吞吐量。

参考资料

[1] Intro to CUDA (part 5): Memory Model

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

如若内容造成侵权/违法违规/事实不符,请联系我们进行投诉反馈qq邮箱809451989@qq.com,一经查实,立即删除!

相关文章

《数字中台建设总体方案》

《数字中台建设总体方案》 制定数字中台战略规划:制定符合企业实际情况的数字中台战略规划,明确建设目标、重点任务和时间表。确定数字中台架构:根据企业业务需求和特点,确定数字中台的架构,包括技术架构、应用架构和数…

vFW搭建IRF

正文共:2328字 40图,预估阅读时间:5 分钟 IRF(Intelligent Resilient Framework,智能弹性架构)技术通过将多台设备连接在一起,虚拟化成一台设备,集成多台设备的硬件资源和软件处理能…

Selenium自动化测试技巧还不知道吗?

1、前言 与以前瀑布式开发模式不同,现在软件测试人员具有使用自动化工具执行测试用例套件的优势,而以前,测试人员习惯于通过测试脚本执行来完成测试。 但自动化测试的目的不是完全摆脱手动测试,而是最大程度地减少手动运行的测试…

玻色量子计算应用奖公布!MathorCup大赛圆满落幕

2023年8月15日,中国优选法统筹法与经济数学研究会主办的2023年第十三届MathorCup高校数学建模挑战赛圆满落幕。为了培养优秀学子的创新意识及运用数学方法和量子计算技术解决实际问题的能力,推动量子计算的实际落地应用,北京玻色量子科技有限…

计算机网络:可靠数据传输(rdt)、流水协议、窗口滑动协议

文章目录 前言一、Rdt1.Rdt1.02.Rdt2.03.Rdt2.14.Rdt2.25.Rdt3.0 二、流水线协议1.滑动窗口(slide window)协议发送窗口接收窗口正常情况下的2个窗口互动异常情况下GBN的2个窗口互动异常情况下SR的2窗口互动GBN协议和SR协议的异同 2.小结 总结 前言 Rdt…

Linux之重谈文件和c语言文件接口

重谈文件 文件 内容 属性, 所有对文件的操作都是: a.对内容操作 b.对属性操作 关于文件 一: 即使文件的内容为空,该文件也会在磁盘上也会占空间,因为文件不仅仅只有内容还有文件对应的属性,文件的内容会占用空间, 文件的属性也…

Leetcode每日一题学习训练——Python3版(从二叉搜索树到更大和树)

版本说明 当前版本号[20231204]。 版本修改说明20231204初版 目录 文章目录 版本说明目录从二叉搜索树到更大和树理解题目代码思路参考代码 原题可以点击此 1038. 从二叉搜索树到更大和树 前去练习。 从二叉搜索树到更大和树 给定一个二叉搜索树 root (BST),请…

麦田医学CEO参加2023年度苏州市青年创业标兵授奖仪式

麦田医学CEO参加2023年度苏州市青年创业标兵授奖仪式 2023年12月4日,麦田(苏州)医学科技有限公司(下简称麦田医学)首席执行官(CEO)李任远同志受邀参加在苏州广电总台举行的2023年度苏州大学生创…

如何在亚马逊平台上找到最畅销的产品

这是您掌握在全球最大的在线市场上销售艺术的首选资源。在本博客中,我们将深入探讨在英国亚马逊上查找最畅销产品的秘密。在浩瀚的选择海洋中导航可能会令人不知所措,但不要害怕——有了正确的策略,您就可以发现利润丰厚的机会并提高销售额。…

C++ day53 最长公共子序列 不相交的线 最大子序和

题目1:1143 最长公共子序列 题目链接:最长公共子序列 对题目的理解 返回两个字符串的最长公共子序列的长度,如果不存在公共子序列,返回0,注意返回的是最长公共子序列,与前一天的最后一道题不同的是子序…

uniapp 在线预览PDF

1、官网地址&#xff1a; https://mozilla.github.io/pdf.js/getting_started/ 2、解压文件到static中 3、定义查看组件FilePreview <template><view><web-view :src"allUrl"></web-view></view> </template><script> e…

(04730)电路分析基础之电阻、电容及电感元件

04730电子技术基础 语雀&#xff08;完全笔记&#xff09; 电阻元件、电感元件和电容元件的概念、伏安关系&#xff0c;以及功率分析是我们以后分析电 路的基础知识。 电阻元件 电阻及其与温度的关系 电阻 电阻元件是对电流呈现阻碍作用的耗能元件&#xff0c;例如灯泡、…

mongoose学习记录

mongoose安装和连接数据库 npm i mongoose导入mongoose const mongoose require(mongoose) mongoose.set("strictQuery",true)连接数据库 mongoose.connect(mongodb:127.0.0.1:27017/test)设置回调 mongoose.connection.on(open,()>{console.log("连接成…

弱口令防护和网站防盗链有什么用

弱口令防护主要针对用户账户的安全。弱口令是指容易被猜测或破解的密码&#xff0c;如常见的密码、简单的数字序列或字典中的单词等。弱口令防护的目的是防止恶意用户或攻击者通过猜测或暴力破解密码的方式获取合法用户的账户权限。通过实施强密码策略、密码复杂度要求和账户锁…

centos 7.9 二进制部署 kubernetes 1.27.7

文章目录 1. 预备条件2. 基础配置2.1 配置root远程登录2.2 配置主机名2.3 安装 ansible2.4 配置互信2.5 配置hosts文件2.6 关闭防firewalld火墙2.7 关闭 selinux2.8 关闭交换分区swap2.9 修改内核参数2.10 安装iptables2.11 开启ipvs2.12 配置limits参数2.13 配置 yum2.14 配置…

【JavaEE进阶】 Spring核⼼与设计思想

文章目录 &#x1f332;Spring 是什么&#xff1f;&#x1f384;什么是IoC呢&#xff1f;&#x1f388;传统程序开发&#x1f388;传统程序开发的缺陷&#x1f388;如何解决传统程序的缺陷&#xff1f;&#x1f388;控制反转式程序开发&#x1f388;对⽐总结规律 &#x1f340;…

十二、FreeRTOS之FreeRTOS任务相关API函数

本节需要掌握以下内容&#xff1a; 1&#xff0c;FreeRTOS任务相关API函数介绍&#xff08;熟悉&#xff09; 2&#xff0c;任务状态查询API函数实验&#xff08;掌握&#xff09; 3&#xff0c;任务时间统计API函数实验&#xff08;掌握&#xff09; 4&#xff0c;课堂总结…

【鸿蒙应用开发】开发环境搭建及IDE安装使用

1.下载安装包 安装包下载地址&#xff1a; 点击跳转下载页面 可以根据自己的操作系统选择对应版本下载。 本文以Windows安装为例&#xff0c;Mac安装方式相同 2. 安装 下载好后&#xff0c;打开安装包&#xff0c;进入安装界面&#xff1a; 点击Next&#xff0c;进入安…

vue创建项目,使用可视化界面安装插件

安装项目&#xff1a; vue create vue-app 选择默认配置就行&#xff0c;也可以按需选择自定义配置 vue ui通过可视化管理项目 通过可视化安装全家桶插件

Spring Cloud 配置 Druid(二)

不废话&#xff0c;直接上代码&#xff0c; Nacos搭建的微服务&#xff0c;可以看Spring Cloud 配置 Nacos&#xff08;一&#xff09;-CSDN博客 一&#xff0c;pom文件 spring-cloud-starter-alibaba-nacos-discovery 和 spring-cloud-starter-openfeign 都是基于spring-cl…