CUDA内存模型

核函数性能并不只与线程束的执行有关。

CUDA内存模型概述

GPU和CPU内存模型的主要区别是,CUDA编程模型能将内存层次结构更好地呈现给用户,能让我们显示的控制它的行为。

对程序员来说,一般有两种类型的存储器:

  • 可编程的:需要显示的控制哪些数据存放在可编程内存中
  • 不可编程的:不能决定数据的存放位置,程序自动生成存放位置

CUDA内存模型提供了多种可编程的内存类型:

  • 寄存器
  • 共享内存
  • 本地内存
  • 常量内存
  • 纹理内存
  • 全局内存

图中所示为这些内存空间的层次结构,每个都有不同的作用域、生命周期和缓存行为。一个核函数中的线程都有自己私有的本地内存。一个线程块有自己的共享内存,对同一线程块内的所有线程可见,其内容持续线程块的整个生命周期。所有线程都可以访问全局内存。所有线程都能访问的只读空间有:常量内存和纹理内存。对一个应用程序来说,全局内存、常量内存和纹理内存中的内容具有相同的生命周期。

  1. 全局内存(Global Memory)

    • 特性:全局内存是GPU中所有线程共享的内存空间,可以被所有线程访问。
    • 优点:全局内存容量较大,适合存储大量数据。
    • 缺点:全局内存访问延迟较高,需要通过内存控制器访问,因此访问速度相对较慢。
    • 适用场景:适用于需要在不同线程之间共享数据的情况,但对访问速度要求不高的情况。
  2. 常量内存(Constant Memory)

    • 特性:常量内存是只读的内存空间,所有线程都可以读取其中的数据,但不能写入。
    • 优点:常量内存具有高速缓存,访问速度较快。
    • 缺点:常量内存容量有限(通常为64KB),且只能在程序启动时初始化。
    • 适用场景:适用于存储只读数据,且需要高速访问的情况,如常数表、查找表等。
  3. 纹理内存(Texture Memory)

    • 特性:纹理内存是一种特殊的只读内存,具有缓存和插值功能,适合于图像处理等需要纹理特性的应用。
    • 优点:纹理内存具有缓存和插值功能,可以提高某些类型数据的访问速度。
    • 缺点:容量相对较小,不能直接写入。
    • 适用场景:适用于需要纹理特性的数据访问,如图像处理、光线追踪等。

寄存器

寄存器是GPU上运行速度最快的内存空间。核函数中声明的变量通常储存在寄存器中。寄存器对每个线程来说是私有的。与核函数生命周期相同。

寄存器和线程不是一对一的,费米架构一个线程可以有63个寄存器,开普勒架构可以有255个寄存器。

如果,一个核函数使用了超过硬件限制数量的寄存器,则会用本地内存替代。这种寄存器溢出的情况会降低性能。

本地内存

核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。

  • 在编译时,使用未知索引引用的本地数组
  • 可能会占用大量寄存器空间的较大本地结构体或数组
  • 任何不满足核函数寄存器限定条件的变量

共享内存

在核函数中使用如下修饰符修饰的变量存放在共享内存中

__shared__

每个SM都有一定数量的由线程块分配的共享内存。共享内存在核函数的范围内声明,其生命周期伴随着整个线程块。

共享内存是线程之间相互通信的基本方式。访问共享内存必须同步使用如下调用,

void __syncthreads();

SM中的一级缓存和共享内存都使用64KB的片上内存,通过静态划分,在运行时也可以进行动态配置。

片上内存(On-chip Memory)是指集成在芯片(例如GPU、CPU等处理器)内部的内存单元。与传统的外部内存(如RAM)相比,片上内存具有更低的访问延迟和更高的带宽,因为它们与处理器核心更加接近,减少了数据传输的距离和延迟。

常量内存

常量内存保存在设备内存中,每个SM都有自己专门用于缓存常量内存的缓存。

常量变量用以下修饰符修饰:

__constant__

常量变量必须在全局空间内和所有核函数之外进行声明。常量内存是静态声明的,并对同一编译单元内的所有核函数可见。

核函数只能从常量内存中读取数据,因此,常量内存必须在主机端由以下函数初始化

cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);

将count个字节从src指向的内存复制到symbol指向的内存中,这个变量存放在设备的全局内存或常量内存中。

线程束中的所有线程从相同的地址读取数据时,常量内存表现最好。比如,所有线程都需要读取一个公式的系数,对不同的数据根据同样的系数做相同的计算,这时,系数存放在常量内存中最好。

纹理内存

  • 特性:纹理内存是一种特殊的只读内存,具有缓存和插值功能,适合于图像处理等需要纹理特性的应用。
  • 优点:纹理内存具有缓存和插值功能,可以提高某些类型数据的访问速度。
  • 缺点:容量相对较小,不能直接写入。
  • 适用场景:适用于需要纹理特性的数据访问,如图像处理、光线追踪等。

纹理特性的数据访问指的是使用CUDA中的纹理内存(Texture Memory)来访问数据的一种特殊方式。纹理内存在访问数据时具有缓存和插值等特性,适用于某些类型的数据访问需求,例如图像处理、光线追踪等。纹理内存的特性包括以下几个方面:

  1. 缓存:纹理内存具有缓存功能,可以将最近访问的数据缓存在内存中,以提高后续访问相同数据的速度。这种缓存机制有助于减少对全局内存的访问次数,从而提高访问速度。

  2. 插值:纹理内存支持插值功能,可以对访问的数据进行插值计算,以获取平滑的插值结果。这在图像处理等需要对像素进行插值的应用中特别有用,可以减少图像处理过程中的锯齿状边缘或伪影现象。

  3. 边界处理:纹理内存支持边界处理功能,可以在访问超出边界的数据时进行特定的处理,例如将边界外的像素值设置为固定值或通过重复边界像素值来填充。

  4. 硬件优化:纹理内存的访问方式经过硬件优化,可以提高数据访问的效率,并充分利用GPU的并行计算能力。

全局内存

全局内存是GPU中最大、延迟最高且最常使用的内存。

可以使用如下修饰符静态声明一个全局内存变量

__device__

在主机端使用cudaMalloc函数分配全局内存,使用cudaFree函数释放全局内存。全局内存可以存在于应用程序的整个生命周期中,可以被所有核函数的所有线程访问。从多个线程访问全局变量时,由于线程之间不能跨线程块同步,不同线程块内的多个线程并发地修改全局内存变量可能会出现问题。

GPU缓存

与CPU缓存一样,GPU缓存是不可编程的内存。在GPU上有四种缓存

  • 一级缓存
  • 二级缓存
  • 只读常量缓存
  • 只读纹理缓存

每个SM都有一个一级缓存,所有的SM共享一个二级缓存。一级缓存和二级缓存都被用来存储本地内存和全局内存中的数据。

实践

静态声明一个全局变量,并且在一个核函数中修改值,在主机端获取修改后的值。

__device__ float num;__global__ void test_num() {printf("the num is %f\n", num);num += 0.3;
}int main(){float a = 3.14;cudaMemcpyToSymbol(num, &a, sizeof(float));test_num << <1, 1 >> > ();cudaMemcpyFromSymbol(&a, num, sizeof(float));cudaDeviceReset();cout << a;return 0;
}

这里面cudaMemcpyToSymbol函数在VS里可能会标红

显示num需要一个地址,其实并不是。在CUDA 编程指南手册里,

实际上动手做一下也会发现传地址结果会不同。

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

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

相关文章

JSP和tomcat

JSP&#xff08;JavaServer Pages&#xff09;是一种用于开发动态Web内容的技术&#xff0c;它允许开发者将Java代码嵌入到HTML页面中。JSP页面在服务器端被解析并转换成Servlet&#xff0c;然后由Servlet容器&#xff08;比如Tomcat&#xff09;执行。JSP允许开发者在页面中使…

YOLO系列改进,自研模块助力涨点

目录 一、原理 二、代码 三、添加到YOLOv5中 一、原理 论文地址:

截取视频第一帧当做封面

看了好多处理视频的框架 比如ffmpeg&#xff0c;很多都需要依赖安装第三方插件&#xff0c;比较麻烦&#xff0c;找到一个内嵌进去不需要额外安装的&#xff1a;jcodec 一 首先代码中添加依赖 <!--视频生成预览图用--><dependency><groupId>org.jcodec</…

企业职能部门定岗定编项目如何做?

某国家级高新技术开发区成立于上世纪90年代&#xff0c;位于南方某市&#xff0c;地处三省交界处&#xff0c;是直接由该省委省政府创办的全省首批高新技术开发区。该开发区面积达到300平方公里&#xff0c;辖区人口30万人。历经数十年发展&#xff0c;在省委省政府的高度重视下…

QT:核心控件-QWidget

文章目录 控件enableobjectNamegeometrysetWindowTitleopacitycursorFonttooltipstyleSheet 控件 什么是控件&#xff1f; 如上所示&#xff0c;就是控件&#xff0c;而本篇要做的就是对于这些控件挑选一些比较有用的常用的进行讲解分析 在QT的右侧&#xff0c;会有对应的空间…

unity制作app(1)--登录 注册 界面

把学到的知识投入到生产中反而是一件简单的事情&#xff01; 1.调整canvas的形状&#xff0c;这里和camera没有任何关系! overlay&#xff01; 2.既然自适应&#xff0c;空间按钮的位置比例就很重要了&#xff01; game窗口中新增720*1280的分辨率&#xff01; 3.再回到can…

【论文阅读】ViTAE:Vision transformer advanced by exploring intrinsic inductive bias

ViTAE:Vision transformer advanced by exploring intrinsic inductive bias 论文地址摘要&#xff1a;简介&#xff1a;3 方法论3.1 重温视觉变压器3.2 ViTAE3.3 缩减单元3.4 Normal cell3.5 模型细节 4 训练4.1 Implementation details4.2 Comparison with the state-of-the-…

高速收发器(GTX)文章导航

关于FPGA的开发&#xff0c;大致可以分为算法和接口两类&#xff0c;其中接口又可以分为高速和低速两类&#xff0c;像UART、I2C、SPI、SDRAM、HDMI、LVDS等等&#xff0c;都被归为低速接口类别&#xff0c;最高线速率不过几百Mbps&#xff0c;使用FPGA的普通IO即可实现数据收发…

LeetCode刷题笔记第145题:二叉树的后序遍历

LeetCode刷题笔记第145题&#xff1a;二叉树的后序遍历 题目&#xff1a; 给定一棵二叉树的根节点 root &#xff0c;返回其节点值的后序遍历 。 想法&#xff1a; 后序遍历的是通过对树经过“左右根”的顺序进行遍历。使用递归的方式&#xff0c;先遍历左子树&#xff0c;…

Elasticsearch:理解近似最近邻 (ANN) 算法

作者&#xff1a;来自 Elastic Elastic Platform Team 如果你是在互联网出现之前长大的&#xff0c;你会记得找到新喜好并不总是那么容易。我们是在无意中听到收音机里的新乐队时发现他们的&#xff0c;是因为忘了换频道偶然看到一个新电视节目的&#xff0c;也是几乎完全依据游…

ClickHouse安装(成功安装)

1.下载安装包 下面通过阿里镜像&#xff08;https://mirrors.aliyun.com/clickhouse/rpm/lts/&#xff09;进行下载&#xff0c;下载哪里&#xff0c;自行指定。 # deb包下载使用如下4行 wget https://mirrors.aliyun.com/clickhouse/deb/pool/stable/clickhouse-client_22.8…

【DIY小记】用爬虫+clean-mark+zhihu-on-vscode搬运技术博客到知乎

今天灵光一闪&#xff0c;决定调研一下自己的技术博客&#xff0c;可以怎样方便迁移到其它社媒平台。想要达到的效果是&#xff0c;把自己在掘金的专栏&#xff1a;从1到∞精通Python&#xff0c;迁移到知乎上面去。 简单花了两三小时时间&#xff0c;找到一个比较快捷的方法&…

【ARMv8/v9 常见汇编指令学习 6.1 - armv8 右移指令 LSR 详细介绍】

请阅读【嵌入式开发学习必备专栏 】 文章目录 armv8 右移指令 LSR语法示例注意事项 armv8 右移指令 LSR 在ARMv8架构中&#xff0c;lsr&#xff08;Logical Shift Right&#xff09;指令是一种逻辑右移指令&#xff0c;用于将一个寄存器中的数值向右移动指定的位数。逻辑右移操…

git 的迁移

现象是gitlab经常会挂掉&#xff0c;linux会显示磁盘空间不足&#xff0c;实际上&#xff0c;我们linux某个目录的空间是4T。这个空间应该是足够的。猜测是gitlab的安装目录不对导致的空间不足。 1、查找原因 用rpm 安装gitlab会有自己的目录&#xff0c;很多安装文件会在opt…

MySql 导出导入(备份还原)

1&#xff0c;导出备份 要导出MySQL数据库中的数据&#xff0c;使用mysqldump命令。假设要导出名为mydatabase的数据库到名为backup.sql的文件中&#xff1a; mysqldump -u 用户名 -p 数据库名 > backup.sql 参数说明&#xff1a; -u mysql用户名称 -p 执行后会要求输入…

利用Python实现Smithwaterman算法

def smith_waterman(seq1, seq2, match_score3, mismatch_penalty-1, gap_penalty-2): # 初始化矩阵 matrix [[0] * (len(seq2) 1) for _ in range(len(seq1) 1)] max_score 0 max_i, max_j 0, 0 # 填充矩阵 for i in range(1, len(seq1) 1):for j in range(1, len(seq2…

【C++】学习笔记——string_4

文章目录 六、string类7. string类的模拟实现 未完待续 六、string类 7. string类的模拟实现 我们在上文简单实现了string类的构造函数。不知道大家有没有发现一个问题&#xff0c;我们在进行实现无参的构造函数时&#xff0c;初始化列表将 _str 初始化为 nullptr 了&#xf…

CRI-O的原理及应用详解(四)

本系列文章简介&#xff1a; 随着云计算和容器技术的迅猛发展&#xff0c;容器化应用已经成为现代软件开发和部署的重要趋势。在这一领域中&#xff0c;Kubernetes作为容器编排的领军者&#xff0c;为开发者提供了强大的容器管理和调度能力。然而&#xff0c;Kubernetes本身并不…

Unity LensFlare 入门

概述 在项目的制作过程中&#xff0c;太阳光的使用一定是不可缺少的部分&#xff0c;但是如果想实现真实太阳光眼睛看到的镜头炫光效果&#xff0c;那这部分的内容一定不要错过喔&#xff0c;接下来让我们来学习这部分的内容吧&#xff01; Hale(光环效果) Color&#xff1a;…

TiDB中的PD--元数据管理和调度中心

目录 PD 架构etcd 的使用TiDB 的 PD(Placement Driver)组件是整个分布式数据库系统的关键部分, 是整个集群的元数据管理和调度中心,负责存储集群的元数据和进行 Region 调度。 主要包括: 元数据存储:PD 存储了整个 TiDB 集群的元数据,包括 TiKV 集群的拓扑结构、Namespa…