CUDA入门之统一内存

原文来自CUDA 编程入门之统一内存

🎬个人简介:一个全栈工程师的升级之路!
📋个人专栏:高性能(HPC)开发基础教程
🎀CSDN主页 发狂的小花
🌄人生秘诀:学习的本质就是极致重复!

目录

What Unified Memory Delivers

Simpler Programming and Memory Model

Performance Through Data Locality

Unified Memory or Unified Virtual Addressing?

Example: Eliminate Deep Copies

Example: CPU/GPU Shared Linked Lists

Unified Memory with C++

A Bright Future for Unified Memory


借助 CUDA 6,NVIDIA 引入了 CUDA 平台历史上最引人注目的编程模型改进之一,即统一内存。在当今典型的 PC 或集群节点中,CPU 和 GPU 的内存在物理上是不同的,并由 PCI-Express 总线分开。在 CUDA 6 之前,程序员就是这样看待事物的。CPU 和 GPU 之间共享的数据必须在两个内存中分配,并由程序在它们之间显式复制。这给 CUDA 程序增加了很多复杂性。

统一内存创建了一个在 CPU 和 GPU 之间共享的托管内存池,弥合了 CPU-GPU 鸿沟。CPU 和 GPU 都可以使用单个指针访问托管内存。关键是系统会自动在主机和设备之间迁移统一内存中分配的数据。

在这篇文章中,作者将向您展示统一内存如何显著简化 GPU 加速应用程序中的内存管理。下图显示了一个非常简单的示例。两种代码都从磁盘加载文件,对其中的字节进行排序,然后在 CPU 上使用排序后的数据,然后释放内存。右侧的代码使用 CUDA 和统一内存在 GPU 上运行。唯一的区别是 GPU 版本启动内核(并在启动后同步),并使用新的 API 为统一内存中的加载文件分配空间cudaMallocManaged()

如果您以前编写过 CUDA C/C++,那么您无疑会被右侧代码的简洁性所震撼。请注意,我们分配了一次内存,并且我们有一个指向可以从主机和设备访问的数据的指针。我们可以直接从文件中读取分配,然后我们可以直接将指针传递给在设备上运行的 CUDA 内核。然后,等待内核完成后,我们可以再次从 CPU 访问数据。CUDA 运行时隐藏了所有复杂性,自动将数据迁移到访问它的地方。

What Unified Memory Delivers

程序员从统一内存中受益的主要方式有两种。

Simpler Programming and Memory Model

统一内存通过使设备内存管理成为优化而不是要求,降低了在 CUDA 平台上进行并行编程的门槛。借助统一内存,现在程序员可以直接开发并行 CUDA 内核,而不会陷入分配和复制设备内存的细节中。这将使学习为 CUDA 平台编程和将现有代码移植到 GPU 变得更简单。但这不仅适用于初学者。本文后面的示例展示了统一内存如何使复杂的数据结构更容易与设备代码一起使用,以及它与 C++ 结合时的强大功能。

Performance Through Data Locality

通过在 CPU 和 GPU 之间按需迁移数据,Unified Memory 可以在 GPU 上提供本地数据的性能,同时提供全局共享数据的易用性。此功能的复杂性被隐藏在 CUDA 驱动程序和运行时,确保应用程序代码更易于编写。迁移的重点是实现每个处理器的全带宽;250 GB/s 的 GDDR5 内存对于满足 Kepler GPU 的计算吞吐量至关重要。

重要的一点是,使用流并cudaMemcpyAsync有效地与数据传输重叠执行的精心调整的 CUDA 程序可能比仅使用统一内存的 CUDA 程序表现得更好。可以理解的是:CUDA 运行时从来没有像程序员那样拥有关于何时何地需要数据的信息!CUDA 程序员仍然可以访问显式设备内存分配和异步内存副本,以优化数据管理和 CPU-GPU 并发性。统一内存首先是一种生产力功能,它为并行计算提供了更平滑的入口,而不会剥夺 CUDA 为高级用户提供的任何功能。

Unified Memory or Unified Virtual Addressing?

CUDA 自 CUDA 4 起就支持统一虚拟寻址 (UVA),虽然统一内存依赖于 UVA,但它们并不是一回事。UVA 为系统中的所有内存提供一个单一的虚拟内存 地址空间 ,并允许从 GPU 代码访问指针,无论它们位于系统中的哪个位置,无论是设备内存(在相同或不同 GPU 上)、主机内存、或片上共享内存。它还允许cudaMemcpy在不指定输入和输出参数的确切位置的情况下使用。UVA 启用“零拷贝”内存,它是固定的主机内存,可通过 PCI-Express 直接由设备代码访问,无需memcpy. 零拷贝提供了统一内存的一些便利,但没有提供性能,因为它总是以 PCI-Express 的低带宽和高延迟访问。

UVA 不会像统一内存那样自动将数据从一个物理位置迁移到另一个物理位置。因为统一内存能够在主机和设备内存之间的单个页面级别自动迁移数据,所以它需要大量的工程来构建,因为它需要 CUDA 运行时、设备驱动程序甚至操作系统内核中的新功能。以下示例旨在让您了解这可以实现什么。

Example: Eliminate Deep Copies

统一内存的一个关键优势是通过在访问 GPU 内核中的结构化数据时消除对深拷贝的需求来简化异构计算内存模型。将包含指针的数据结构从 CPU 传递到 GPU 需要执行“深度复制”,如下图所示。

以下面的 struct 为例dataElem

struct dataElem {int prop1;int prop2;char *name;
}

要在设备上使用这个结构,我们必须复制结构本身及其数据成员,然后复制结构指向的所有数据,然后更新结构副本中的所有指针。这导致以下复杂代码,只是为了将数据元素传递给内核函数。

void launch(dataElem *elem) {dataElem *d_elem;char *d_name;int namelen = strlen(elem->name) + 1;// Allocate storage for struct and namecudaMalloc(&d_elem, sizeof(dataElem));cudaMalloc(&d_name, namelen);// Copy up each piece separately, including new “name” pointer valuecudaMemcpy(d_elem, elem, sizeof(dataElem), cudaMemcpyHostToDevice);cudaMemcpy(d_name, elem->name, namelen, cudaMemcpyHostToDevice);cudaMemcpy(&(d_elem->name), &d_name, sizeof(char*), cudaMemcpyHostToDevice);// Finally we can launch our kernel, but CPU & GPU use different copies of “elem”Kernel<<< ... >>>(d_elem);
}

可以想象,在 CPU 和 GPU 代码之间共享复杂数据结构所需的额外主机端代码对生产力有重大影响。在统一内存中分配我们的dataElem结构消除了所有多余的设置代码,只剩下内核启动,它在与主机代码相同的指针上运行。这是一个很大的进步!

void launch(dataElem *elem) {kernel<<< ... >>>(elem);
}

但这不仅仅是代码复杂性的重大改进。统一记忆还使以前无法想象的事情成为可能。让我们看另一个例子。

Example: CPU/GPU Shared Linked Lists

链表是一种非常常见的数据结构,但由于它们本质上是由指针组成的嵌套数据结构,因此在内存空间之间传递它们是非常复杂的。如果没有统一内存,在 CPU 和 GPU 之间共享一个链表是无法管理的。唯一的选择是在零拷贝内存(固定主机内存)中分配列表,这意味着 GPU 访问仅限于 PCI-express 性能。通过在统一内存中分配链表数据,设备代码可以在 GPU 上正常跟随指针,充分发挥设备内存的性能。该程序可以维护单个链表,并且可以从主机或设备中添加和删除列表元素。

将具有现有复杂数据结构的代码移植到 GPU 曾经是一项艰巨的任务,但统一内存使这变得容易得多。作者希望统一内存能够为 CUDA 程序员带来巨大的生产力提升。

Unified Memory with C++

统一内存确实在 C++ 数据结构中大放异彩。C++ 通过使用带有复制构造函数的类来简化深复制问题。复制构造函数是一个知道如何创建类的对象、为其成员分配空间以及从另一个对象复制其值的函数。C++ 还允许重载 new 和 delete 内存管理操作符。这意味着我们可以创建一个基类,我们将其称为 Managed,它在重载的 new 运算符中使用 cudaMallocManaged(),如下面的代码所示。

class Managed {
public:void *operator new(size_t len) {void *ptr;cudaMallocManaged(&ptr, len);cudaDeviceSynchronize();return ptr;}void operator delete(void *ptr) {cudaDeviceSynchronize();cudaFree(ptr);}
};

然后我们可以让我们的String类从该类继承Managed,并实现一个复制构造函数,为复制的字符串分配统一内存。

// Deriving from “Managed” allows pass-by-reference
class String : public Managed {int length;char *data;public:// Unified memory copy constructor allows pass-by-valueString (const String &s) {length = s.length;cudaMallocManaged(&data, length);memcpy(data, s.data, length);}// ...
};

同样,我们让我们的dataElem类继承Managed

// Note “managed” on this class, too.
// C++ now handles our deep copies
class dataElem : public Managed {
public:int prop1;int prop2;String name;
};

通过这些更改,C++ 类在统一内存中分配它们的存储,并且自动处理深拷贝。我们可以像任何 C++ 对象一样在统一内存中分配一个 dataElem。

dataElem *data = new dataElem;

请注意,您需要确保树中的每个类都继承自Managed,否则您的内存映射中有一个漏洞。实际上,您可能需要在 CPU 和 GPU 之间共享的所有内容都应该继承Managed. 如果您更喜欢简单地对所有内容使用统一内存,则可以在全局范围内重载newdelete但这仅在您没有纯 CPU 数据时才有意义,否则数据将不必要地迁移。

现在,当我们将对象传递给内核函数时,我们有一个选择;在 C++ 中很正常,我们可以按值传递或按引用传递,如下面的示例代码所示。

// Pass-by-reference version
__global__ void kernel_by_ref(dataElem &data) { ... }// Pass-by-value version
__global__ void kernel_by_val(dataElem data) { ... }int main(void) {dataElem *data = new dataElem;...// pass data to kernel by referencekernel_by_ref<<<1,1>>>(*data);// pass data to kernel by value -- this will create a copykernel_by_val<<<1,1>>>(*data);
}

多亏了统一内存,深拷贝、按值传递和按引用传递都可以正常工作。这为在 GPU 上运行 C++ 代码提供了巨大的价值。

这篇文章中的示例可在 Github 上找到。

A Bright Future for Unified Memory

CUDA 6 中统一内存最令人兴奋的事情之一是它只是一个开始。他们围绕统一内存​​计划了一个很长的改进和功能路线图。他们的第一个版本旨在使 CUDA 编程更容易,特别是对于初学者。从 CUDA 6 开始,cudaMemcpy()不再需要。通过使用cudaMallocManaged(),您有一个指向数据的指针,并且您可以在 CPU 和 GPU 之间共享复杂的 C/C++ 数据结构。这使得编写 CUDA 程序变得更加容易,因为您可以直接编写内核,而不是编写大量数据管理代码并维护所有数据的重复主机和设备副本。您仍然可以自由使用cudaMemcpy()(尤其是 cudaMemcpyAsync())性能,但它现在不是一种要求,而是一种优化。

CUDA 的未来版本可能会通过添加数据预取和迁移提示来提高使用统一内存的应用程序的性能。他们还将增加对更多操作系统的支持。他们的下一代 GPU 架构将带来多项硬件改进,以进一步提高性能和灵活性。

🌈我的分享也就到此结束啦🌈
如果我的分享也能对你有帮助,那就太好了!
若有不足,还请大家多多指正,我们一起学习交流!
📢未来的富豪们:点赞👍→收藏⭐→关注🔍,如果能评论下就太惊喜了!
感谢大家的观看和支持!最后,☺祝愿大家每天有钱赚!!!欢迎关注、关注!

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

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

相关文章

Python元组(Tuple)深度解析!

目录 1. 什么是元组&#xff1f; 2. 创建元组 3.访问元组 4.元组的运算 5.修改元组不可行 6.元组的应用场景 前面的博客里&#xff0c;我们详细介绍了列表&#xff08;List&#xff09;这一种数据类型&#xff0c;现在我们来讲讲与列表相似的一种数据类型&#xff0c;元组…

使用Docker安装运行RabbitMQ---阿里云服务器

目录 0、阿里云没开端口的得要去安全组规则去添加&#xff1a; 1、下载RabbitMQ镜像&#xff1a; 2、查看镜像是否下载成功&#xff0c;得到docker镜像id&#xff1a; 3、运行RabbitMQ: 4、查看RabbbitMQ容器是否启动成功&#xff1a; 5、启动RabbitMQ中的插件管理 6、访…

基于YOLOv8/YOLOv7/YOLOv6/YOLOv5的人群密度检测系统(Python+PySide6界面+训练代码)

摘要&#xff1a;开发人群密度检测系统对于公共安全等领域具有关键作用。本篇博客详细介绍了如何运用深度学习构建一个人群密度检测系统&#xff0c;并提供了完整的实现代码。该系统基于强大的YOLOv8算法&#xff0c;并对比了YOLOv7、YOLOv6、YOLOv5&#xff0c;展示了不同模型…

Python 读取写入excel文件

使用Python读取和写入excel的xlsx、xls文件 目录 读取xlsx文件 安装三方库 引入三方库 读取数据 打开文件 表名 最大行数 最大列数 读取一张表 读取整个文件 返回xls整体内容 安装三方包 读取内容 写入xls文件 引入三方库 创建文件并写入数据 报错及解决 报错…

KBL610-ASEMI开关电源专用KBL610

编辑&#xff1a;ll KBL610-ASEMI开关电源专用KBL610 型号&#xff1a;KBL610 品牌&#xff1a;ASEMI 封装&#xff1a;KBL-4 最大重复峰值反向电压&#xff1a;1000V 最大正向平均整流电流(Vdss)&#xff1a;6A 功率(Pd)&#xff1a;中小功率 芯片个数&#xff1a;4 …

项目解决方案:视频监控接入和录像系统设计方案(下)

目 录 1.概述 2. 建设目标及需求 2.1建设总目标 2.2 需求描述 ​2.3 需求分析 3.设计依据与设计原则 3.1设计依据 3.2 设计原则 4.建设方案设计 4.1系统方案设计 4.2组网说明 5.产品介绍 5.1视频监控综合资源管理平台介绍 5.2视频录像服务器和存储 5.2.…

MySQL--优化(索引--聚簇和非聚簇索引)

MySQL–优化&#xff08;索引–聚簇和非聚簇索引&#xff09; 定位慢查询SQL执行计划索引 存储引擎索引底层数据结构聚簇和非聚簇索引索引创建原则索引失效场景 SQL优化经验 一、聚簇索引 聚簇索引&#xff1a;将数据存储与索引放到了一块&#xff0c;索引结构的叶子节点保存…

ChatGPT无法登录,提示我们检测到可疑的登录行为,将阻止进一步的尝试。请与管理员联系

1. 问题描述 之前本来已经连续稳定使用ChatGPT好几个月了&#xff0c;但是今天尝试登录ChatGPT的时候&#xff0c;却提示&#xff1a;我们检测到可疑的登录行为&#xff0c;将阻止进一步的尝试。请与管理员联系。 此外&#xff0c;我还在网上看到了一些相关的消息&#xff0c;…

3D资产管理

3D 资产管理是指组织、跟踪、优化和分发 3D 模型和资产以用于游戏、电影、AR/VR 体验等各种应用的过程。 3D资产管理也称为3D内容管理。 随着游戏、电影、建筑、工程等行业中 3D 内容的增长&#xff0c;实施有效的资产管理工作流程对于提高生产力、减少错误、简化工作流程以及使…

【Android 内存优化】KOOM 快手开源框架线上内存监控方案-源码剖析

文章目录 前言OOMMonitorInitTask.INSTANCE.initOOMMonitor.INSTANCE.startLoopsuper.startLoopcall() LoopState.Terminate dumpAndAnalysisdumpstartAnalysisService回到startLoop方法总结 前言 这篇文章主要剖析KOOM的Java层源码设计逻辑。 使用篇请看上一篇: 【Android …

基于el-tree实现懒加载穿梭条

一、关键代码 <template><div><!-- 左侧待选列表 --><div class"left-box"><p>待选列表</p><el-input placeholder"输入关键词过滤" v-model"leftFilterText" clearable/><el-treeref"tree…

Windows系统安装MongoDB并结合内网穿透实现公网访问本地数据库

文章目录 前言1. 安装数据库2. 内网穿透2.1 安装cpolar内网穿透2.2 创建隧道映射2.3 测试随机公网地址远程连接 3. 配置固定TCP端口地址3.1 保留一个固定的公网TCP端口地址3.2 配置固定公网TCP端口地址3.3 测试固定地址公网远程访问 前言 MongoDB是一个基于分布式文件存储的数…

颜色检测python项目

注意&#xff1a;本文引用自专业人工智能社区Venus AI 更多AI知识请参考原站 &#xff08;[www.aideeplearning.cn]&#xff09; 什么是颜色检测&#xff1f; 颜色检测是检测任何颜色名称的过程。很简单不是吗&#xff1f;嗯&#xff0c;对于人类来说&#xff0c;这是一项极…

荔枝派zero驱动开发06:GPIO操作(platform框架)

参考&#xff1a; 正点原子Linux第五十四章 platform设备驱动实验 一张图掌握 Linux platform 平台设备驱动框架 上一篇&#xff1a;荔枝派zero驱动开发05&#xff1a;GPIO操作&#xff08;使用GPIO子系统&#xff09; 下一篇&#xff1a;更新中… 概述 platform是一种分层思…

static的用法和作用

从三个方面来讲述static的用法和作用&#xff0c;分别是静态局部变量和静态函数以及静态成员 1. 在一个函数中定义静态局部变量的时候&#xff0c;该静态局部变量的生命周期将贯穿整个程序的运行期&#xff0c;而不会像局部变量一样函数运行结束的时候就被销毁。 #include &…

华为北向网管NCE开发教程(1)闭坑选接口协议

华为北向网管NCE开发教程&#xff08;1&#xff09;闭坑选接口协议 华为北向网管NCE开发教程&#xff08;2&#xff09;REST接口开发 华为北向网管NCE开发教程&#xff08;3&#xff09;CORBA协议开发 华为北向网管NCE开发教程&#xff08;4&#xff09;&#xff08;源代码接口…

网络地址转换协议NAT

网络地址转换协议NAT NAT的定义 NAT&#xff08;Network Address Translation&#xff0c;网络地址转换&#xff09;是1994年提出的。当在专用网内部的一些主机本来已经分配到了本地IP地址&#xff08;即仅在本专用网内使用的专用地址&#xff09;&#xff0c;但现在又想和因…

如何打造定制化企业内训系统?企培源码开发实战教学

在当今竞争激烈的商业环境中&#xff0c;企业内训系统的建立和定制化已成为提高企业竞争力和员工素质的关键。本文将探讨如何打造定制化企业内训系统&#xff0c;并介绍企培源码开发的实战教学。 第一步&#xff1a;需求分析与规划 建立定制化的企业内训系统之前&#xff0c…

echarts绘制雷达图

<template><div><div>【云端报警风险】</div><div ref"target" class"w-full h-full" stylewidth&#xff1a;200px;height:300px></div></div> </template><script setup> import { ref, onMounte…

【漏洞复现】Salia PLCC cPH2 远程命令执行漏洞(CVE-2023-46359)

0x01 漏洞概述 Salia PLCC cPH2 v1.87.0 及更早版本中存在一个操作系统命令注入漏洞&#xff0c;该漏洞可能允许未经身份验证的远程攻击者通过传递给连接检查功能的特制参数在系统上执行任意命令。 0x02 测绘语句 fofa&#xff1a;"Salia PLCC" 0x03 漏洞复现 ​…