CUDA优化:最大化内存吞吐量(官方文档翻译)

2023-10-27

毕业设计要翻译技术资料 3000 字,这里找了英伟达 CUDA TOOLKIT DOCUMENTATION 的 5.3 节“最大化内存吞吐量来”翻译一下,供参考,并希望此文对诸位的 CUDA 程序优化有所帮助。

 

 

5.3. 最大化内存吞吐量

最大化应用程序总内存吞吐量的第一步当是最大限度地减少低带宽的数据传输。

 

这意味着最小化主机(内存)和设备(显存)之间的数据传输,因为正如主机和设备间的数据传输中详述的那样——这样的数据传输的带宽远远低于全局内存和设备之间的数据传输。

 

这也意味着通过尽可能通过使用片上内存(on-chip memory):共享内存和缓存(即在计算能力大于等于 2.x 的设备上可用的L1 缓存和 L2 缓存,及所有设备上可用的纹理缓存(texture cache) 和常量缓存(constant cache))以最大限度地减少全局内存和设备之间的数据传输。

 

共享内存等价于“用户管理的缓存”:应用程序显式地分配和访问它。如 CUDA 运行时 所述,一个典型的编程模式是将来自设备内存的数据组织编排到共享内存中:换句话说,让一个块的每一个线程:

  1. 将数据从设备内存加载到共享内存,
  2. 与块的所有其他线程同步,以便每个线程可以安全地读取由不同线程填充的共享内存位置,
  3. 在共享内存中处理数据,
  4. 必要时再次同步,以确保共享内存已与结果一起更新,
  5. 将结果写回设备内存。

 

对于某些应用程序(例如那些全局内存访问模式依赖于数据的程序),传统的硬件管理缓存更适合利用数据的局域性。如计算能力-3-x、计算能力 7.x 和计算能力 8.x所述,对于计算能力 3.x、7.x 和 8.x 的设备, L1缓存和共享内存使用的是相同的空间,并且每次内核调用都可配置用于 L1 与共享内存的比值。

 

内核访问的吞吐量可能因不同内存的访问模式而异。因此,最大化内存吞吐量的下一步是根据 设备内存访问 中描述的最佳内存访问模式尽可能最佳地组织内存访问。这种优化对于全局内存访问尤为重要,因为与片上内存的带宽和算术指令吞吐量相比,全局内存带宽较低,因此未经优化的全局内存访问通常对性能有很高的(负面)影响。

 

5.3.1. 主机和设备之间的数据传输

应用程序应努力最大限度地减少主机和设备之间的数据传输。实现此目的的一种方法是将更多代码(计算过程)从主机移动到设备,即使这意味着运行没有展现出足够的并行性(以获得最高效能)的内核函数。你可以在设备内存中创建、在设备上运算、并销毁,而无需在主机产生映射或复制到主机内存的中间数据结构。

 

此外,由于每次传输的经常性开销(overhead),将许多次小的传输组合成单个大的数据传输中总是比单独地进行每次传输效果更好。

 

在具有前端总线(front-side bus)的系统上,使用页锁定主机内存(page-locked host memory)中描述的页面锁定主机内存可实现主机和设备之间的数据传输的更高性能。

 

此外,在使用映射的页面锁定内存(映射内存)时,无需分配任何设备内存,或明确在设备和主机内存之间拷贝数据。每次核函数访问映射内存时,都会隐式执行数据传输。若要获得最大性能,这些内存访问必须与访问全局内存一样聚合(将小访问聚合成大访问)(请参阅设备内存访问 )。假设它们这些映射内存仅读或写一次,则使用映射的页面锁定内存,相较于设备和主机内存之间的显式地拷贝,可能带来性能的提升。

 

在设备内存和主机内存实质上相同的集成系统中,主机和设备内存之间的任何拷贝都是多余的,应改为使用映射的页面锁定内存。应用可以通过检查集成设备属性(见设备枚举)是否等于 1 来查询设备是否为集成设备。

 

 

5.3.2. 设备内存访问

获取可地址指示的内存(即全局、局部、共享、常数或纹理内存)的指令可能需要多次重新发布,具体取决于线程束(warp)内线程的内存地址的分布。分布如何以这种方式影响指令吞吐量,取决于每种类型的内存,这将在以下部分进行描述。例如,对于全局存储器,一般来说,地址越分散,吞吐量就越低。

 

全局内存

全局内存在设备内存中,可通过 32、64 或 128 个字节的规格进行内存访问。这些内存规格必须天然地对齐:只有与其大小对齐的 32、64 或 128 字节(即其第一个地址是其大小的倍数)的设备内存段才能通过内存事务进行读取或写入。

 

当线程束执行访问全局内存的指令时,它会根据每个线程访问的字大小和所有线程访问的内存地址的分布,将线程束内线程的内存访问汇合成这些内存事务中的一个或多个。一般来说,传输次数越多,未被使用但被线程访问的字越多,从而相应地降低了指令的吞吐量。例如,如果一个32字节的内存访问被每个线程用4字节访问完成,则吞吐量缩减为原来的八分之一。

 

需要多少访问以及最终影响多少吞吐量因设备的计算能力不同。计算能力 3.x计算能力 5.x计算能力 6.x计算能力 7.x 和计算能力 8.x 提供了有关处理各种计算能力的全局内存访问方式的更多详细信息。

 

因此,要最大限度地提高全局内存吞吐量,必须通过:

  1. 遵循基于计算能力 3.x 计算能力 5.x计算能力 6.x计算能力 7.x 和计算能力 8.x 的最佳访问模式
  2. 使用符合下面大小和对齐要求部分中详细说明的大小和对齐要求的数据类型,
  3. 在某些情况下,例如,在访问下面的二维矩阵部分中描述的二维矩阵时,应修补数据。

 

尺寸和对齐要求

 

全局内存指令支持读取或写入大小为 1、2、4、8 或 16 字节的字。如果数据类型大小为 1、2、4、8 或 16 字节且数据自然对齐(即其地址是该大小的倍数),则(通过变量或指针)对存储于全局内存中的数据的任何访问都可编译为单个全局内存指令。

 

如果此大小和对齐要求未实现,访问将被编译为多个指令,并使用交错访问模式,以防止这些指令完全结合。因此,我们建议存储在全局内存中的数据,都符合此要求。

 

对于 内置矢量类型 而言,程序自动实现其对齐要求。

 

对于结构体,通过可以使用__align__ (8) 或__align__ (16) 的对齐指示,编译器将使之满足大小和对齐要求,例如

 

struct __align__(8) {

    float x;

    float y;

};

或者

struct __align__(16) {

    float x;

    float y;

float z;

};

存储在全局内存中的变量的任何地址,或由驱动程序或运行时 API 的内存分配函数返回的地址始终与至少 256 字节对齐。

 

读取不自然对齐的 8 字节或 16 字节字会产生不正确的结果(相差几个字),因此必须特别小心地保证这些类型的任何值或值矩阵的起始地址的对齐。一个容易忽略这种情况的典型案例是使用一些自定义的全局内存分配方案,即将多个的分配(多次调用 cudaMalloc())或 cuMemAlloc() 替换为可分区为多个矩阵的单个大块内存的分配,在这种情况下,每个矩阵的起始地址的偏移与块的起始地址的偏移一致。

 

二维矩阵

常见的全局内存访问模式是,当每个索引线程 (tx,ty) 使用以下地址访问位于类型* 的地址在 BaseAddress 的宽度 width 的二维矩阵的一个元素时(其中“类型”符合 最大化利用 中描述的要求):

BaseAddress + width * ty + tx

要使这些访问完全结合,线程格的宽度和矩阵的宽度必须是线程数大小的倍数。

特别地,这意味着,宽度不是此大小的倍数的数组,如果实际分配的宽度补足到此大小的最接近的倍数,并且按行相应地填充后,将可被更有效地访问。参考手册中描述的 cudaMallocPitch() 和 cuMemAllocPitch() 函数和相关内存拷贝函数允许程序员编写非硬件依赖的代码来分配符合这些限制的矩阵。

 

局部内存

局部内存存取只在某些自动变量出现时存在,其中自动变量在 变量内存空间指示 提及。编译器可能放置在局部内存中的自动变量有:

  1. 无法确定它们与常量大小的矩阵,
  2. 消耗太多的寄存器空间大结构体或矩阵,
  3. 任何使得内核使用超出可用的寄存器数量的变量(这也称为寄存器溢出)。

 

对 PTX 装配代码的检查(通过使用 -ptx 或 -keep 选项进行编译获得)将展示某一变量是否在第一个编译阶段被放置在局部内存中,因为它会被标记上 .local 助记符并被通过 ld.local 核st.local 助记符访问。即使没有存在于局部内存,如果发现它在所处的计算架构中消耗太多的寄存器空间,后续的编译阶段仍可能使之变为局部内存:使用 cuobjdump 对cubin对象的检查将判断是否是这种情况。此外,在使用  --ptxas-options=-v 选项进行编译时,编译器会报告每个内核 (lmem) 的局部内存总用量。请注意,某些数学函数具有可能访问局部内存的实现。

 

局部内存空间位于设备内存中,因此局部内存访问具有与全局内存访问相同的高延迟和低带宽的特性,并且受制于 设备内存访问 中描述的存储器合并的类似要求。但是,局部内存被组织为连续的 32 位字被连续的线程 ID 访问。因此,只要线程束中的所有线程访问一致的相对地址(例如,矩阵变量中的相同索引、结构体中的相同成员),访问就完全合并在一起。

 

在某些计算能力设备3.x 的设备上,局部内存访问始终以与全局内存访问相同的方式缓存在 L1 和 L2 中(参见计算能力 3.x)。

 

在计算能力 5.x 和 6.x 的设备上,局部内存访问始终以与全局内存访问相同的方式缓存在 L2 中(参见计算能力 5.x 和计算能力 6.x)。

 

共享内存

 

由于共享内存是片上存储器,因此与局部或全局内存相比,带宽要大得多,延迟也低得多。

为了实现高带宽,共享内存被划分为大小相等的内存模块,称为"库(bank)",可同时访问。因此,对于 n 地址在 n 个不同的内存库中提出的任何内存读写请求都可以同时进行响应,从而产生整体带宽,其带宽是单个模块带宽的 n 倍。

 

然而,如果存储器请求的两个地址位于同一内存库中,则存在库冲突,访问必须序列化。硬件根据需要将带有库冲突的内存请求拆分为尽可能多的独立无冲突请求,将吞吐量减少到等于独立内存请求数。如果单独的内存请求数为n,则初始内存请求被定义为 n 路库冲突。

 

因此,要获得最大的性能,了解内存地址如何映射到内存库非常重要,以便安排内存请求,从而最大限度地减少行内存库冲突。这些在计算能力 3.x计算能力 5.x计算能力 6.x计算能力 7.x 和计算能力 8.x 分别被详述。

 

常量内存

 

常量内存空间位于设备内存中,并缓存在常量缓存中。

 

一个请求被分割成与初始请求中不同的内存地址一样多的单独的请求,从而将吞吐量减少到等于单独请求数。

 

产生的请求将在常量缓存命中(cache hit)发生时的数据吞吐时响应,否则在设备内存吞吐时相应。

 

纹理和表面内存

 

纹理和表面内存空间位于设备内存中,并缓存在纹理缓存中,因此,一个纹理/表面读取仅仅在缓存未击中(cache miss)时消耗一次设备内存读取,否则只消耗一次纹理缓存读取。纹理缓存是针对二维区域优化的,因此读取二维空间中相邻的纹理或表面地址的相同线程束中的线程将实现最佳性能。此外,它专为恒定延迟的流获取而设计;缓存命中可降低 DRAM 带宽需求,但无法降低获取延迟。

 

通过纹理或表面获取读取设备内存相较于从全局或常量内存中读取设备内存这些好处:

  1. 如果内存读取不遵循全局或常量内存读取必须遵循才能获得良好的性能的访问模式,只要纹理/表面读取中有区域性,也可以实现更高的带宽:
  2. 计算部分由专用单元在内核之外执行:
  3. 打包的数据可以在单个操作中广播到不同的变量:
  4. 8 位和 16 位整数输入数据可在 [0.0、1.0] 或 [-1.0、 1.0] (参见纹理内存)范围内可选转换为 32 位浮点值。

 

 

本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)

CUDA优化:最大化内存吞吐量(官方文档翻译) 的相关文章

  • __syncthreads() 死锁

    如果只有部分线程执行 syncthreads 会导致死锁吗 我有一个这样的内核 global void Kernel int N int a if threadIdx x
  • 设置最大 CUDA 资源

    我想知道是否可以设置 CUDA 应用程序的最大 GPU 资源 例如 如果我有一个 4GB GPU 但希望给定的应用程序只能访问 2GB 如果它尝试分配更多 就会失败 理想情况下 这可以在进程级别或 CUDA 上下文级别上设置 不 目前没有允
  • Yocto for Nvidia Jetson 由于 GCC 7 而失败 - 无法计算目标文件的后缀

    我正在尝试将 Yocto 与 meta tegra 一起使用 https github com madisongh meta tegra https github com madisongh meta tegra 为 Nvidia Jets
  • TensorRT 多线程

    我正在尝试使用 python API 来使用 TensorRt 我试图在多个线程中使用它 其中 Cuda 上下文与所有线程一起使用 在单个线程中一切正常 我使用 docker 和 tensorrt 20 06 py3 图像 onnx 模型和
  • 有没有一种有效的方法来优化我的序列化代码?

    这个问题缺乏细节 因此 我决定创建另一个问题而不是编辑这个问题 新问题在这里 我可以并行化我的代码吗 还是不值得 https stackoverflow com questions 17937438 can i parallelize my
  • 无法在内存位置找到异常源:cudaError_enum

    我正在尝试确定 Microsoft C 异常的来源 test fft exe 中 0x770ab9bc 处的第一次机会异常 Microsoft C 异常 内存位置 0x016cf234 处的 cudaError enum 我的构建环境是 I
  • 内联 PTX 汇编代码强大吗?

    我看到一些代码示例 人们在 C 代码中使用内联 PTX 汇编代码 CUDA工具包中的文档提到PTX很强大 为什么会这样呢 如果我们在 C 代码中使用这样的代码 我们会得到什么好处 内联 PTX 使您可以访问未通过 CUDA 内在函数公开的指
  • CUDA 中指令重放的其他原因

    这是我从 nvprof CUDA 5 5 获得的输出 Invocations Metric Name Metric Description Min Max Avg Device Tesla K40c 0 Kernel MyKernel do
  • 如何运行和理解CUDA Visual Profiler?

    我已经设置了 CUDA 5 0 并且我的 CUDA 项目运行良好 但我不知道如何使用 Visual Profiler 分析我的 CUDA 项目 如何运行它 我还需要安装更多吗 又该如何做呢 我的电脑使用Window 7 64位 CUDA 5
  • cudaDeviceScheduleBlockingSync 和 cudaDeviceScheduleYield 之间有什么区别?

    正如这里所说 如何减少 CUDA 同步延迟 延迟 https stackoverflow com questions 11953722 how to reduce cuda synchronize latency delay 等待设备结果有
  • Bank 在字长方面存在冲突

    我读过一些关于共享内存的好文章 但我对银行冲突有初步疑问 据说 如果线程 1 和线程 2 从存储体 0 访问字 0 则不存在存储体冲突 但如果他们访问不同的单词 就会出现银行冲突 但我的问题是不同的单词如何可以驻留在一个银行中 由于bank
  • CUDA Thrust 库中counting_iterators 的用途和用法

    我很难理解counting iterator在 CUDA 的推力库中 它的目的是什么以及如何使用 它在其他编程语言 例如 C 中也可用吗 计数迭代器只是一个迭代器 它从每次迭代器递增时前进的序列中返回下一个值 最简单的例子是这样的 incl
  • 使用 cudamalloc()。为什么是双指针?

    我目前正在浏览有关的教程示例http code google com p stanford cs193g sp2010 http code google com p stanford cs193g sp2010 学习CUDA 演示的代码 g
  • 无法在 CUDA 中找到 1 到 100 数字的简单和?

    我正在研究使用 CUDA 的图像处理算法 在我的算法中 我想使用 CUDA 内核找到图像所有像素的总和 所以我在cuda中制作了内核方法 来测量16位灰度图像的所有像素的总和 但我得到了错误的答案 所以我在cuda中编写了一个简单的程序来查
  • CUDA 和 Eigen 的成员“已声明”错误

    我只是 CUDA 和 Nsight 的初学者 希望利用出色的 GPU 性能进行线性代数运算 例如 CUBLAS 我在以下人员的帮助下编写了很多自定义代码Eigen http eigen tuxfamily org index php tit
  • 使用推力来处理 CUDA 类中的向量?

    我对 C 类的推力的适用性有疑问 我正在尝试实现一个类对象 该对象接收顶点的 x y z 坐标作为 ver1 ver2 和 ver3 然后 分配给一个三角形并计算面积和法向量 然而 我不太明白如何创建一类推力向量 这是我从文件中读取的顶点坐
  • 如何在没有 nvcc 的情况下在编译时获取 CUDA 工具包版本?

    我在 cpp 文件中对 cuSPARSE 库进行了一些调用 这些调用在旧工具包中不可用 为了支持使用旧工具包的系统 我想使用编译器指令编译不同的代码部分 特别是 我想使用旧工具包的 CSR 格式矩阵和新工具包的 BSR 格式矩阵来求解稀疏三
  • 布尔实现的atomicCAS

    我想弄清楚是否存在错误答案 https stackoverflow com a 57444538 11248508 现已删除 关于Cuda like的实现atomicCAS for bool是 答案中的代码 重新格式化 static inl
  • 完全禁用 NVCC 优化

    我正在尝试测量 GPU 上的峰值单精度触发器 为此我正在修改 PTX 文件以在寄存器上执行连续的 MAD 指令 不幸的是 编译器正在删除所有代码 因为它实际上没有做任何有用的事情 因为我没有执行任何数据的加载 存储 是否有编译器标志或编译指
  • 如何使用 Visual Studio 2008 调试 CUDA 内核代码?

    嘿 我正在使用带有 CUDA 3 2 的 Visual Studio 2008 我正在尝试调试具有此签名的函数 MatrixMultiplication Kernel lt lt

随机推荐

  • Windows混音器API使用

    1 首先用mixerGetNumDevs 函数获取系统中的混音器设备的数量 一般 机器上都至少有一个混音器设备 声卡 如果机器上没有连接其它的音频设备 那么也就只有声卡这一个混音器设备 我的机器上接有一个名为USB EMP Audio De
  • go-redis 框架基本使用

    文章目录 redis使用场景 下载框架和连接redis 1 安装go redis 2 连接redis 字符串操作 有序集合操作 流水线 事务 1 普通事务 2 Watch redis使用场景 缓存系统 减轻主数据库 MySQL 的压力 计数
  • 自适应控制设计(二)

    自适应控制设计 二 自适应控制基本思想 一文主要介绍了自适应控制设计的基本思路 但是针对控制率的设计没有具体说明 这里针对反馈控制率的设计步骤进行具体介绍 控制器设计基本思想 对于任何一个动态系统 我们都可以根据Lyapunov稳定性设计其
  • C++MFC编程之按钮控件Button、Radio Button和Check Box

    钮控件包括命令按钮 Button 单选按钮 Radio Button 和复选框 Check Box 等 命令按钮就是我们前面多次提到的狭义的按钮控件 用来响应用户的鼠标单击操作 进行相应的处理 它可以显示文本也可以嵌入位图 单选按钮使用时
  • 解决el-select下拉框多选在赋完值之后,不能对tag和已选中的值取消掉

    这种是原先的写法
  • 如何在Linux环境创建GRE Tunnel

    Question I want to connect to remote networks by using a GRE tunnel How can I create a GRE tunnel between two end points
  • ini配置文件读写操作入门

    ini配置文件读写操作入门 ini文件 Initialization file 这种类型的文件中通常存放的是一个程序的初始化信息 ini文件由若干个节 Section 组成 每个Section由若干键 Key 组成 每个Key可以赋相应的值
  • 大点云的可以用opencv和pcl结合粗配准

    比如 可以把点云转换为 tif 通过一部分截取8位图像 进行sift 将得到的内点序号序列 转换为三维坐标序列 然后将三维坐标转换为点云sift cloud1和sift cloud2 这时候 由于是一一对应的 即在pcl中 query in
  • 目标跟踪综述 (持续更新)

    这几天对目标跟踪挺感兴趣的 但是在CSDN和知乎上面找的相关介绍资料都看的一知半解 所以自己找了一篇 2022 04 26 发表的综述文章作下笔记学习下 目录 一 基于相关滤波的目标跟踪算法 1 相关滤波视频目标跟踪算法的框架 2 相关滤波
  • c#特性(Attribute)与反射(Reflection)学习

    概念 特性 Attribute 用于在运行时传递程序中各种元素 比如类 方法 结构 枚举 组件等 的行为信息的声明性标签 放置在他所修饰的元素前面用 包裹 用于添加元数据 如编译器指令和注释 描述 方法 类等其他信息 可以使用预定义的特性或
  • 通过SpringBoot生成微信小程序二维码,跳转指定页面

    以下通过两种方法实现生成微信小程序二维码保存 通过华为存储obs服务 通过 IO流 字符流的使用 读取字符流 字符流写入 微信小程序获取二维码参数 onLoad function options console log options 方式
  • qt 如何在另一个线程更新控件状态

    在 Qt 中 如果要在另一个线程中更新控件的状态 有以下几种方法可以考虑 使用信号和槽 在另一个线程中发射信号 连接到控件的槽函数 在槽函数中更新控件的状态 使用事件队列 在另一个线程中调用 QCoreApplication postEve
  • 手机投屏不是全屏怎么办_手机投屏到竖放的电视

    手机投屏到竖放的电视 今天有网友给我发了手机投屏电视 但是手机横屏之后 电视机还是竖屏 其实这个很正常 因为当手机与电视机处于镜像投屏模式的时候 因为显示比例的关系 电视机是不会满屏的 只能以竖屏模式在中间显示一部分画面 就像上面的图一样
  • Maven Install 报错:To see the full stack trace of the errors, re-run Maven with the -e switch

    博主 在eclipse打包项目 将war包部署到 linux 上 install 时的报错 蓝色标志 出现这个error信息 说明仓库里有些对应的jar包没下载完全 试过非常多办法就是不行 最后 处理方法是去到本地仓库里 把对应的jar包先
  • 【干货】日志管理与分析(一)——日志收集及来源

    对广大IT工作者 尤其是运维和安全人员来说 日志 是一个再熟悉不过的名词 日志从哪来 机房中的各种软件 系统 防火墙 和硬件 交换机 路由器等 都在不断地生成日志 IT安全业界的无数实践告诉我们 健全的日志记录和分析系统 是系统正常运营 优
  • [C++]抽象工厂模式

    抽象工厂模式 Abstract Factory Pattern 是围绕一个超级工厂创建其他工厂 该超级工厂又称为其他工厂的工厂 这种类型的设计模式属于创建型模式 它提供了一种创建对象的最佳方式 在抽象工厂模式中 接口是负责创建一个相关对象的
  • KETTLE WEB管理控制台设计

    KETTLE WEB管理控制台设计 系统效果图 1 资源配置管理 1 1 用例图 1 2 用例叙述 1 2 1 添加资源库用例 用例名称 添加资源库 前置条件 操作员在启动这个用例之前 必须先执行过 登录 用例 主要事件流 1 当用户选择添
  • 清除input的缓存

    有过表单设计经验的朋友肯定知道 当我们在浏览器中输入表单信息的时候 往往input文本输入框会记录下之前提交表单的信息 以后每次只要双击 input文本输入框就会出现之前输入的文本 这样有时会觉得比较方便 但有时也会暴露用户的隐藏数据 那么
  • oracle简单的备份

    企业之中这种备份肯定用不到 只是闲着没事简单备份一下自己的数据库 方便以后调数据 查阅 一 备份 1 打开cmd 切到oracle的bin目录 2 备份 exp scott 123456 file E oracle备份 scott back
  • CUDA优化:最大化内存吞吐量(官方文档翻译)

    毕业设计要翻译技术资料 3000 字 这里找了英伟达 CUDA TOOLKIT DOCUMENTATION 的 5 3 节 最大化内存吞吐量来 翻译一下 供参考 并希望此文对诸位的 CUDA 程序优化有所帮助 5 3 最大化内存吞吐量 最大