CUDA stream利用CUDA流重叠计

本文涉及的产品
数据传输服务 DTS,数据同步 small 3个月
推荐场景:
数据库上云
数据传输服务 DTS,数据迁移 small 3个月
推荐场景:
MySQL数据库上云
数据传输服务 DTS,同步至SelectDB 1个月
简介: CUDA stream利用CUDA流重叠计

利用CUDA流重叠计算和数据传输

CUDA中有一个重要的概念是 流(stream). 其实它代表着一系列的指令的执行队列. 这个执行队列就像他的名字一样, 有着固定的执行顺序(就像河流只能向一个方向固定的河道流淌一样).

而这条河的源头就是主机线程(host), 它开启了这个执行队列. 同样的, 也可能这座高山开启了不同的河流, 我们的主机线程(host)也可能启动了不同的执行队列. 或者多个主机线程(多座高山), 开启了多个stream(河流).

我更愿意理解流是更高一个层次的并行手段, 相对于thread, block 和 grid, 它的层级更高, 也更独立.

thread, block 和 grid其实都可以算作 kernel内的并行层次, 而流(stream)是kernel外的并行层次.

流(stream)分为默认流(或者叫做NULL流)非默认流.

默认流指的是你不显示声明,创建或指定的操作队列. 在任何CUDA程序中只要是你调用了kernel或调用相关的CUDA函数, 并且没指定他们运行在哪个流中, 那么他们会自动的被安排在默认流中.

或者你可以这么想, 当你开始CUDA程序的时候, 你调用的那些方法或函数就已经在默认流中了. 当你创建了新的流并把那些函数方法放在新的流中的时候, 他们才从默认流中解脱, 有了自己新的执行队列.

而你创建的那些新的流, stream0, stream1, stream2…就是非默认流.

CUDA流最重要的功能就是并发执行

这里提到的并发执行简单点说就是, 将要计算的数据或要处理的资源分割成多个部分, 让每个流分别处理.

这个流的概念也是其他基于CUDA的加速库使用最多的CUDA内容. 比如cuBLAS, cuFFT, TensorRT, cuDNN中的句柄概念就可以和CUDA流结合使用, 并发处理.

接下来我们分析一下这段代码:

for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
                    size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
    MyKernel<<<100, 512, 0, stream[i]>>>
          (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
    cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
                    size, cudaMemcpyDeviceToHost, stream[i]);

在这段代码中:

  • 第一个for循环将数据从CPU内存传输给GPU显存中. 而传输的过程是将input分割成两部分. 分别放在stream[i]中进行传输.
  • 在第二个for循环中, MyKernel<<<100, 512, 0, stream[i]>>>分别在stream[i]中执行.
  • 第三个for循环分别在两个stream中讲计算完成的结果传输给CPU内存.

这里需要注意的是:

  • 这里传输的函数使用的是cudaMemcpyAsync, 它是一个异步传输的函数. 简单点说, 可以同时进行多个传输任务, 他们之间并不是一起执行的.
  • 我们在使用上述代码的时候, 需要手动的进行同步.

我们需要让程序知道这些流(执行队列)是否已经完成, 而显式的同步流的方法有以下几种:

  • cudaDeviceSynchronize() 一直等待,直到所有主机线程的所有流中的所有先前命令都完成。
  • cudaStreamSynchronize() 将流作为参数并等待,直到给定流中的所有先前命令都已完成。 它可用于将主机与特定流同步,允许其他流继续在设备上执行。
  • cudaStreamWaitEvent() 将流和事件作为参数(有关事件的描述,请参阅事件),并在调用 cudaStreamWaitEvent() 后使添加到给定流的所有命令延迟执行,直到给定事件完成。
  • cudaStreamQuery() 为应用程序提供了一种方法来了解流中所有前面的命令是否已完成。

重叠计算

理解了上面的内容, 我们就可以简单的理解了重叠计算的内容.

如上图所示, 我们同时进行数据传输和数据计算. 那么我们就将数据传输的时间隐藏在数据计算的过程中, 或者将数计算的时间隐藏在数据传输的过程中.

如果主机线程在它们之间发出以下任一操作,则来自不同流的两个命令不能同时运行:

  • 页面锁定的主机内存分配,
  • 设备内存分配,
  • 设备内存设置,
  • 两个地址之间的内存拷贝到同一设备内存,
  • 对 NULL 流的任何 CUDA 命令,
  • 计算能力 3.x 和计算能力 7.x 中描述的 L1/共享内存配置之间的切换。

对于支持并发内核执行且计算能力为 3.0 或更低的设备,任何需要依赖项检查以查看流内核启动是否完成的操作:

  • 仅当从 CUDA 上下文中的任何流启动的所有先前内核的所有线程块都已开始执行时,才能开始执行;
  • 阻止所有以后从 CUDA 上下文中的任何流启动内核,直到检查内核启动完成。

需要依赖检查的操作包括与正在检查的启动相同的流中的任何其他命令以及对该流的任何 cudaStreamQuery() 调用。 因此,应用程序应遵循以下准则来提高并发内核执行的潜力:

  • 所有独立操作都应该在依赖操作之前发出,
  • 任何类型的同步都应该尽可能地延迟。
相关实践学习
自建数据库迁移到云数据库
本场景将引导您将网站的自建数据库平滑迁移至云数据库RDS。通过使用RDS,您可以获得稳定、可靠和安全的企业级数据库服务,可以更加专注于发展核心业务,无需过多担心数据库的管理和维护。
Sqoop 企业级大数据迁移方案实战
Sqoop是一个用于在Hadoop和关系数据库服务器之间传输数据的工具。它用于从关系数据库(如MySQL,Oracle)导入数据到Hadoop HDFS,并从Hadoop文件系统导出到关系数据库。 本课程主要讲解了Sqoop的设计思想及原理、部署安装及配置、详细具体的使用方法技巧与实操案例、企业级任务管理等。结合日常工作实践,培养解决实际问题的能力。本课程由黑马程序员提供。
目录
相关文章
|
并行计算 TensorFlow 调度
推荐场景GPU优化的探索与实践:CUDA Graph与多流并行的比较与分析
RTP 系统(即 Rank Service),是一个面向搜索和推荐的 ranking 需求,支持多种模型的在线 inference 服务,是阿里智能引擎团队沉淀多年的技术产品。今年,团队在推荐场景的GPU性能优化上又做了新尝试——在RTP上集成了Multi Stream,改变了TensorFlow的单流机制,让多流的执行并行,作为增加GPU并行度的另一种选择。本文详细介绍与比较了CUDA Graph与多流并行这两个方案,以及团队的实践成果与心得。
|
并行计算 异构计算
CUDA streamCUDA流的基本概念
CUDA streamCUDA流的基本概念
2414 0
CUDA streamCUDA流的基本概念
|
存储 编解码 对象存储
Python提取指定时间、经度与纬度的NC数据
【2月更文挑战第15天】本文介绍基于Python语言的netCDF4库,读取.nc格式的数据文件,并提取指定维(时间、经度与纬度)下的变量数据的方法~
501 3
Python提取指定时间、经度与纬度的NC数据
|
机器学习/深度学习 存储 并行计算
一篇就够:高性能推理引擎理论与实践 (TensorRT)
本文分享了关于 NVIDIA 推出的高性能的深度学习推理引擎 TensorRT 的背后理论知识和实践操作指南。
13916 9
一篇就够:高性能推理引擎理论与实践 (TensorRT)
|
网络协议 网络安全 开发工具
|
缓存 测试技术 API
解锁开源模型高性能服务:SGLang Runtime 应用场景与实践
SGLang 是一个用于大型语言模型和视觉语言模型的推理框架。
|
11月前
|
机器学习/深度学习 人工智能 并行计算
【AI系统】GPU 架构回顾(从2010年-2017年)
自1999年英伟达发明GPU以来,其技术不断革新。本文概述了从2010年至2024年间,英伟达GPU的九代架构演变,包括费米、开普勒、麦克斯韦、帕斯卡、伏特、图灵、安培、赫柏和布莱克韦尔。这些架构不仅在游戏性能上取得显著提升,还在AI、HPC、自动驾驶等领域发挥了重要作用。CUDA平台的持续发展,以及Tensor Core、NVLink等技术的迭代,巩固了英伟达在计算领域的领导地位。
421 1
|
缓存 人工智能 并行计算
diffusers SD推理加速方案的调研实践总结(1)
diffusers SD推理加速方案的调研实践总结
433 13
|
Kubernetes 关系型数据库 网络架构
ray集群部署vllm的折磨
概括如下: 在构建一个兼容多种LLM推理框架的平台时,开发者选择了Ray分布式框架,以解决资源管理和适配问题。然而,在尝试集成vllm时遇到挑战,因为vllm内部自管理Ray集群,与原有设计冲突。经过一系列尝试,包括调整资源分配、修改vllm源码和利用Ray部署的`placement_group_bundles`特性,最终实现了兼容,但依赖于非官方支持的解决方案。在面对vllm新版本和Ray部署的`reconfigure`方法问题时,又需权衡和调整实现方式。尽管面临困难,开发者认为使用Ray作为统一底层仍具有潜力。
|
机器学习/深度学习
一文看懂卷积运算(convolution)与互相关运算(cross-correlation)的区别
一文看懂卷积运算(convolution)与互相关运算(cross-correlation)的区别