0
  • 聊天消息
  • 系统消息
  • 评论与回复
登录后你可以
  • 下载海量资料
  • 学习在线课程
  • 观看技术视频
  • 写文章/发帖/加入社区
创作中心

完善资料让更多小伙伴认识你,还能领取20积分哦,立即完善>

3天内不再提示

CUDA 6中的统一内存模型

Linux阅码场 来源:Linuxer 作者:Linuxer 2020-07-02 14:08 次阅读

CUDA介绍

CUDA(Compute Unified Device Architecture,统一计算设备架构)是由NVIDIA公司于 2006 年所推出的一种并行计算技术,是该公司对于GPGPU( General-purpose computing on graphics processing units, 图形处理单元上的通用计算 )技术的正式命名。通过此技术,用户可在GPU上进行通用计算,而开发人员可以使用C语言来为CUDA架构编写程序 。相比CPU,拥有CUDA技术的GPU成本不高,但计算性能很突出。本文中提到的是2014年发布的CUDA6, CUDA6最重要的新特性就是支持统一内存模型(Unified Memory)。

注:文中经常出现“主机和设备”,本文的“主机”特指CPU、“设备”特指GPU。

CUDA 6中的统一内存模型

NVIDIA在CUDA 6中引入了统一内存模型 ( Unified Memory ),这是CUDA历史上最重要的编程模型改进之一。在当今典型的PC或群集节点中,CPU和GPU的内存在物理上是独立的,并通过PCI-Express总线相连。在CUDA6之前, 这是程序员最需要注意的地方。CPU和GPU之间共享的数据必须在两个内存中都分配,并由程序直接地在两个内存之间来回复制。这给CUDA编程带来了很大难度。

统一内存模型创建了一个托管内存池(a pool of managed memory),该托管内存池由CPU和GPU共享,跨越了CPU与GPU之间的鸿沟。CPU和GPU都可以使用单指针访问托管内存。关键是系统会自动地在主机和设备之间迁移在统一内存中分配的数据,从而使那些看起来像CPU内存中的代码在CPU上运行,而另一些看起来像GPU内存中的代码在GPU上运行。

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

CUDA 6中的统一内存模型

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

统一内存模型提供了什么

统一内存模型为程序员提供了两大捷径

简化编程、简化内存模型

统一内存模型通过使设备内存管理(device memory management)成为一项可选的优化,而不是一项硬性的要求,从而降低了CUDA平台上并行编程的门槛。借助统一内存模型,程序员现在可以直接开发并行的CUDA内核,而不必担心分配和复制设备内存的细节。这将降低在CUDA平台上编程的学习成本,也使得将现有代码移植到GPU的工作变得容易。但这些好处不仅有利于初学者。我在本文后面的示例中将展示统一内存模型如何使复杂的数据结构更易于与设备代码一起使用,以及它与C++结合时的强大威力。

通过数据局部性原理提高性能

通过在CPU和GPU之间按需迁移数据,统一内存模型可以满足GPU上本地数据的性能需求,同时还提供了易于使用的全局共享数据。这个功能的复杂细节被 CUDA驱动程序和运行时隐藏了,以确保应用程序代码更易于编写。迁移的关键是从每个处理器获得全部带宽。250 GB / s的GDDR5内存对于保证开普勒( Kepler )GPU的计算吞吐量至关重要。

值得注意的是, 一个经过精心调优的CUDA程序,即使用流(streams)和 cudaMemcpyAsync来有效地将执行命令与数据传输重叠的程序,会比仅使用统一内存模型的CUDA程序更好 。可以理解的是:CUDA运行时从来没有像程序员那样提供何处需要数据或何时需要数据的信息!CUDA程序员仍然可以显式地访问设备内存分配和异步内存拷贝,以优化数据管理和CPU-GPU并发机制 。首先,统一内存模型提高了生产力,它为并行计算提供了更顺畅的入口,同时它又不影响高级用户的任何CUDA功能。

统一内存模型 vs 统一虚拟寻址?

自CUDA4起,CUDA就支持统一虚拟寻址(UVA),并且尽管统一内存模型依赖于UVA,但它们并不是一回事。UVA为 系统中的所有内存提供了单个虚拟内存地址空间,无论指针位于系统中的何处,无论在设备内存(在相同或不同的GPU上)、主机内存、或片上共享存储器。UVA也允许 cudaMemcpy在不指定输入和输出参数确切位置的情况下使用。UVA启用“零复制(Zero-Copy)” 内存,“零复制”内存是固定的主机内存,可由设备上的代码通过PCI-Express总线直接访问,而无需使用 memcpy。零复制为统一内存模型提供了一些便利,但是却没有提高性能,因为它总是通过带宽低而且延迟高的PCI-Express进行访问。

UVA不会像统一内存模型一样自动将数据从一个物理位置迁移到另一个物理位置。由于统一内存模型能够在主机和设备内存之间的各级页面自动地迁移数据,因此它需要进行大量的工程设计,因为它需要在CUDA运行时(runtime)、设备驱动程序、甚至OS内核中添加新功能。以下示例旨在让您领会到这一点。示例:消除深层副本

统一内存模型的主要优势在于,在访问GPU内核中的结构化数据时,无需进行深度复制(deep copies),从而简化了异构计算内存模型。如下图所示,将包含指针的数据结构从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 name cudaMalloc(&d_elem, sizeof(dataElem)); cudaMalloc(&d_name, namelen);

// Copy up each piece separately, including new “name” pointer value cudaMemcpy(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之间分享链表。唯一的选择是在零拷贝内存(被pin住的主机内存)中分配链表,这意味着GPU的访问受限于PCI-express性能。通过在统一内存模型中分配链表数据,设备代码可以正常使用GPU上的指针,从而发挥设备内存的全部性能。程序可以维护单链表,并且无论在主机或设备中都可以添加和删除链表元素。

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

Unified Memory with C++

统一内存模型确实在C++数据结构中大放异彩。C++通过带有拷贝构造函数(copy constructors)的类来简化深度复制问题。拷贝构造函数是一个知道如何创建类所对应对象的函数,拷贝构造函数为对象的成员分配空间并从其他对象复制值过来。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-referenceclass String : public Managed { int length; char *data;

public:// Unified memory copy constructor allows pass-by-value String (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 copiesclass dataElem : public Managed {public:int prop1;int prop2; String name;};

通过这些更改,C++的类将在统一内存中分配空间,并自动处理深度复制。我们可以像分配任何C++的对象那样在统一内存中分配一个 dataElem。

dataElem *data = new dataElem;

请注意,您需要确保树中的每个类都继承自 Managed,否则您的内存映射中会有一个漏洞。实际上,任何你想在CPU和GPU之间分享的内容都应该继承 Managed。如果你倾向于对所有程序都简单地使用统一内存模型,你可以在全局重载 new和 delete, 但这只在这种情况下有作用——你的程序中没有仅被CPU访问的数据(即程序中的所有数据都被GPU访问),因为只有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 reference kernel_by_ref《《《1,1》》》(*data);

// pass data to kernel by value -- this will create a copy kernel_by_val《《《1,1》》》(*data);}

多亏了统一内存模型,深度复制、按值传递和按引用传递都可以正常工作。统一内存模型为在GPU上运行C++代码提供了巨大帮助。

这篇文章的例子可以在Github上找到。

统一内存模型的光明前景

CUDA 6中关于统一内存模型的最令人兴奋的事情之一就是它仅仅是个开始。我们针对统一内存模型有一个包括性能提升与特性的长远规划。我们的第一个发行版旨在使CUDA编程更容易,尤其是对于初学者而言。从CUDA 6开始, cudaMemcpy()不再是必需的。通过使用 cudaMallocManaged(),您可以拥有一个指向数据的指针,并且可以在CPU和GPU之间共享复杂的C / C++数据结构。这使编写CUDA程序变得容易得多,因为您可以直接编写内核,而不是编写大量数据管理代码并且要维护在主机和设备之间所有重复的数据。您仍然可以自由使用 cudaMemcpy()(特别是 cudaMemcpyAsync())来提高性能,但现在这不是一项要求,而是一项优化。

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

声明:本文内容及配图由入驻作者撰写或者入驻合作网站授权转载。文章观点仅代表作者本人,不代表电子发烧友网立场。文章及其配图仅供工程师学习之用,如有内容侵权或者其他违规问题,请联系本站处理。 举报投诉
  • 数据
    +关注

    关注

    8

    文章

    6504

    浏览量

    87444
  • 内存
    +关注

    关注

    8

    文章

    2737

    浏览量

    72613
  • 编程
    +关注

    关注

    88

    文章

    3431

    浏览量

    92218
收藏 人收藏

    评论

    相关推荐

    物理内存模型的演变

    内存管理概述中,主要是以Linux v2.6.11为例进行分析的,但是计算技术在不断发展,新的存储架构、新的指令集架构、新的SoC架构等都对物理内存模型的抽象提出了更高要求。为此,必须抽象一种完全独立于硬件架构的物理
    的头像 发表于 02-25 10:35 210次阅读

    jvm内存模型内存结构

    JVM(Java虚拟机)是Java程序的运行平台,它负责将Java程序转换成机器码并在计算机上执行。在JVM中,内存模型内存结构是两个重要的概念,本文将详细介绍它们。 一、JVM内存
    的头像 发表于 12-05 11:08 341次阅读

    求解大型COMSOL模型需要多少内存

    求解大型COMSOL模型需要多少内存? COMSOL是一种非常强大的跨学科有限元分析软件,可以用于解决各种复杂的问题,包括流体力学、电磁学、热传递、结构力学等。但是,在处理大型模型时,COMSOL
    的头像 发表于 10-29 11:35 929次阅读

    CUDA核心是什么?CUDA核心的工作原理

    CUDA核心(Compute Unified Device Architecture Core)是NVIDIA图形处理器(GPU)上的计算单元,用于执行并行计算任务。每个CUDA核心可以执行单个线程的指令,包括算术运算、逻辑操作和内存
    发表于 09-27 09:38 4248次阅读
    <b class='flag-5'>CUDA</b>核心是什么?<b class='flag-5'>CUDA</b>核心的工作原理

    Armv8-A内存定序模型详解

    1. 本指南介绍由 Armv8-A 架构定义的内存定序模型,并介绍所提供的不同内存障碍;本指南还确定些需要明确定序的常见情况,以及如何使用内存
    发表于 08-02 11:03

    Armv8-M内存模型内存保护用户指南

    本指南概述了Armv8-M内存模型内存保护单元(MPU)在Cortex-M处理器实现。本指南使用示例来帮助解释这些概念 它介绍了。 本章概述以下主题: •
    发表于 08-02 08:12

    周三研讨会预告 | 从 CUDA 到 CV-CUDA:如何为自己定制高效的 CV 任务算子

    CUDA (Compute Unified Device Architecture)编程模型 ,利用 GPU 强大的并行计算能力,为计算机视觉任务带来了前所未有的加速效果。 为了能让 CV
    的头像 发表于 06-13 20:55 261次阅读
    周三研讨会预告 | 从 <b class='flag-5'>CUDA</b> 到 CV-<b class='flag-5'>CUDA</b>:如何为自己定制高效的 CV 任务算子

    简化内存模型的调试

    在这篇博客中,我将展示另一个功能,它使Verdi Protocol Analyzer成为调试内存协议问题和验证正确系统行为的首选工具。此外,该工具还可用于验证命令序列器以及 DUT 和内存模型之间的交互。我们今天要看的功能是将交
    的头像 发表于 05-26 16:34 486次阅读
    简化<b class='flag-5'>内存</b><b class='flag-5'>模型</b>的调试

    如何减少内存模型调试时间

    Synopsys 内存模型 (VIP) 与 Verdi 一起可提高整体调试效率。以下是紧密耦合的调试解决方案如何帮助解决一些痛点的示例:
    的头像 发表于 05-26 10:30 789次阅读
    如何减少<b class='flag-5'>内存</b><b class='flag-5'>模型</b>调试时间

    覆盖模型 – 填补内存VIP的漏洞

    Synopsys 内存模型 (VIP) 具有内置的验证计划、功能和定时覆盖模型,可加速覆盖收敛。提供覆盖模型是为了帮助跨配置设置、模式寄存器设置、功能和时序参数的多种组合运行完整的验证
    的头像 发表于 05-25 16:19 493次阅读
    覆盖<b class='flag-5'>模型</b> – 填补<b class='flag-5'>内存</b>VIP的漏洞

    CUDA编程共享内存

    共享内存是使用__shared__内存空间说明符分配的。
    的头像 发表于 05-19 15:32 751次阅读
    <b class='flag-5'>CUDA</b>编程共享<b class='flag-5'>内存</b>

    介绍CUDA编程模型CUDA线程体系

    CUDA 编程模型主要有三个关键抽象:层级的线程组,共享内存和栅同步(barrier synchronization)。
    的头像 发表于 05-19 11:32 1025次阅读
    介绍<b class='flag-5'>CUDA</b>编程<b class='flag-5'>模型</b>及<b class='flag-5'>CUDA</b>线程体系

    GPU平台生态,英伟达CUDA和AMD ROCm对比分析

    CUDA 除了是并行计算架构外,还是 CPU 和 GPU 协调工作的通用语言。在CUDA 编程模型中,主要有 Host(主机)和 Device(设备)两个概念,Host 包含 CPU 和主机
    的头像 发表于 05-18 09:57 1607次阅读
    GPU平台生态,英伟达<b class='flag-5'>CUDA</b>和AMD ROCm对比分析

    Instruct-UIE:信息抽取统一模型

    Instruct-UIE 统一了信息抽取任务训练方法,可以融合不同类型任务以及不同的标注规范,统一进行训练。针对新的任务需求,仅需要少量的数据进行增量式学习,即可完成模型的升级。
    的头像 发表于 04-25 10:46 1007次阅读
    Instruct-UIE:信息抽取<b class='flag-5'>统一</b>大<b class='flag-5'>模型</b>

    RT1060注册内存被TFLite Micro模型的其他函数覆盖了怎么解决?

    您好,我遇到了些 TFLM 操作注册在其他操作中被覆盖的错误。这只发生在 INT8/UINT8 量化模型,这意味着相同的模型但在浮点数
    发表于 04-20 07:18