前言

自动驾驶之心推出的 《CUDA与TensorRT部署实战课程》，链接。记录下个人学习笔记，仅供自己参考

Note：关于 stream 流杜老师之前也讲过，感兴趣的可以看看 3.4.cuda运行时API-流的学习，异步任务的管理

本次课程我们来学习课程第二章—CUDA 编程入门，一起来学习流和事件

课程大纲可以看下面的思维导图

0. 简述

本小节目标：理解什么是 stream，cuda 编程中的显式隐式同步，以及如何利用多流进行隐藏访存和核函数执行延迟的调度

这节课程我们来讲第二章节第 4 小节，stream 和 event 即流和事件，event 事件我们在统计核函数时间时给大家稍微提到过，这个小节主要跟大家去讲 stream，stream 其实是 CUDA 编程中比较有意思的东西

我们在写一个程序时希望它能做到高并发性，我们可以利用多流一起去操作，从核函数外部以核函数为级别进行一个调度，我们不仅仅可以考虑核函数本身还可以考虑核函数和 memory copy 之间的一个并发，比如在核函数执行的过程中去进行 memory copy，当 memory copy 执行完之后你的核函数差不多也计算完成了，这样就可以把你 memory copy 产生的延迟给隐藏掉，或者把你核函数计算的时间给隐藏掉

这个其实是一个有意思的事情，如果利用好你的 stream 进行一个多流程序的编写的话其实你的速度提升还是挺有发展空间的，这个小节的目标就是一起来学习下如何利用 stream 来编写多流程序隐藏访存和核函数执行延迟的调度。

1. 执行一下我们的第九个CUDA程序

源代码获取地址：https://github.com/kalfazed/tensorrt_starter

这节课程的案例代码是 2.9-stream-and-event，如下所示：

这个案例代码多了 stream.cu 和 stream.hpp 两个文件，以及 CUDA 加速 gelu 激活函数的实现代码 gelu.cu 和 gelu.hpp，这个跟这个小节没有特别强的关联性，大家感兴趣的可以尝试下如何利用多流操作实现 gelu 和矩阵乘法的实现，这节课更多的的是利用 nsight 分析 stream.cu 核函数和 memcpy 之间的一个调度关系，我们使用不同的策略来观察它的调度关系是什么样子的

本节案例执行的效果如下

这里显示了单流以及多流的实验结果，可以根据 block size，grid size 进行内部修改，我们可以发现当使用一个流去执行这一套操作时花费的时间大概是 8.76ms，而我们如果使用 5 个流去执行这一套操作花费的时间只需要 4.01ms

这里的执行结果是韩军老师在 RTX3080 上测试得到的，大家真正在做部署的时候可能用到的一般都是 Jetson 这种边缘设备，那它的计算速度可能没有 RTX3080 这么快，有的时候性能还会卡在访存、带宽这些地方，所以为了平衡 memcpy 以及 kernel 的执行，这里在 kernel 内部设定了等待，如下图所示：

也就是说我们为了平衡不想让核函数跑得太快，所以在核函数里面其实没有做任何东西，只是做了一个 sleep，也就是说等待的一个核函数，占用了资源占用了空间但内部什么都不干

其次这里面有做了一个多流进行异步的计算流程，这个我们到时候在代码中再仔细看

这里给大家还加了一个关于 GELU 激活函数 CUDA 实现的扩展，为后面课程中搭建 Plugin 做铺垫。大家以后利用 TensorRT 进行模型部署的时候可能会有发现一些算子是 TensorRT 官方不支持的，那我们怎么办呢？我们就只能自己去写 CUDA 核函数来对这个算子进行一个加速，从而让 TensorRT 可以去识别它，我们可以利用它来做多流的一个测试

2. Stream是什么

我们讲了这么多那 stream 到底是什么呢？官方给的解释是 A sequence of operation that execute in issue-order in GPU，在 GPU 中 stream 就是一串指令以 issue-order 启动的顺序来执行，也就是说同一个流的执行顺序和各个 kernel 以及 memcpy operation 的启动顺序是一致的，但是，只要资源没有被占用，不同流之间的执行是可以 overlap 的，这个就是多流

但是这里有两个点需要大家注意：

• PCIe 是共享的，所以 memcpy 只能够在同一个时间执行一个
• SM 计算资源是有限的，所以如果计算资源占满了，多流和单流是差不多的

上图是 NVIDIA 讲座中有关 stream 的一张图，首先 Serial 代表的是 Memcpy(H2D)、Kernel、Memcpy(D2H) 三个指令的串行；2-way concurrency 代表的是我们可以把 Kernel、D2H 分割成多个小块放在不同的 stream 中执行，这样我们整个执行时间减少的同时还没有打破它们之间的一个依赖关系；3-way concurrency 代表的是 H2D 这个操作我们也可以分块，host 端往 device 传输一点数据，我们核函数就做一部分运算，然后把数据在拷贝回去，同理另外的流也是这样，它的执行效率会比串行快很多；4-way concurrency 代表的是 GPU 和 CPU 可以异步，GPU 指令比如 Memcpy 或者 Kernel 发送完后 CPU 就闲置了，这是我们可以让 CPU 去做其它的一些事情；最后 4+ way concurrency 代表的是如果每个小块 Kernel 执行的操作太多，我们还可以进一步的进行细分来进行并行处理

到目前为止我们的程序中并没有出现 stream，当我们写 Memcpy 或者核函数的时候并没有指定流，核函数的参数我们就指定了 Grid 和 Block 两个，是不是意味着我们没有使用 stream 呢？其实不是的，当我们不指定核函数以及 Memcpy 的 stream 时，CUDA 会使用默认流（default stream），你所有的操作都在这个默认流里面按照它指令发行的顺序去依次执行

上图是 Nsight 工具的一个截图，从图中可以看到我们先做了一个 Memcpy(H2D)，然后做了一个 SleepKernel 最后做了一个 Memcpy(D2H)，这个完全就是串行的。我们还可以看到 Memcpy(H2D) 和 Memcpy(D2H) 它们两个是不同颜色的，这也就是意味着 H2D 和 D2H 是在同一个 stream 中的不同队列中，不同队列是不是意味着我们可以让这两个指令进行并行呢

同样的，我们在默认流中一次执行两个核函数在 Nsight 上的体现也是完全依次执行的，这个是默认流。但是我们如果想显式的指定流进行操作时我们该怎么做呢？

如上图所示，如果我们要显式指定流，我们核函数调用的时候需要添加一个 stream 的参数，此外在 Memcpy 时我们需要使用异步的数据拷贝 API 即 cudaMemcpyAsync 同时在参数中显式指定我们的流，还有一个值得注意的点是我们显式指定流时如果需要在 Host 端分配空间需要使用 cudaMallocHost 函数，多流操作时需要这个 API 函数来分配 Host 内存，它会在 Host 端给你显式的分配一个 pinned memory 页锁定内存

在上图中我们可以看到多流操作在 Nsight 上是怎么体现的，左边是计算资源被占满的时候，这个如果你前面的核函数已经把所有计算资源都占满了，下一个核函数其实并不能够继续利用你的计算资源，所以我们只能去等待，等前面的核函数执行得差不多了在开始执行；右边是计算资源没有被占满的时候，这时我们就可以利用多流进行并行执行了，前面核函数还没执行完我们就开始执行下一个核函数，依此类推。

这里很容易有个误区需要大家注意，多流执行并不是说我们所有得核函数都是在同一个时间一起执行的，因为核函数启动需要时间，它中间有一个延迟，这个是我们不能忽略的，所以实际上它们并不是完全在同一个时间启动的。另外每个流内部还是按照指令依次执行的，也就是 Memcpy(H2D)、SleepKernel、Memcpy(D2H)

我们刚才讲 cudaMallocHost 分配的是 pinned memory 页锁定内存，那大家可能对操作系统中的内存可能忘记了，我们这边稍微扩展下。内存我们从广义上可以分为以下两个：

• Pageable memory：可分页内存
• Pinned memory/Page-locked memory：页锁定内存

上面左边的图展示的是逻辑内存、虚拟内存以及物理内存之间的关系，我们简单梳理下它们之间的关系：(from chatGPT)

逻辑内存（Logical Memory）

逻辑内存也被称为应用程序内存，指的是应用程序视角下的内存。对于程序而言，它们操作的地址空间（即逻辑地址）是连续的，这提供了一个简单、连续的内存使用视角，无需关心底层物理内存的分配和布局。这种抽象允许程序员在不了解底层硬件细节的情况下进行编程。

虚拟内存（Virtual Memory）

虚拟内存是一种内存管理技术，它为每个程序提供了一个似乎独立的内存空间，即虚拟地址空间。这个空间通常远大于物理内存的大小，它通过在物理内存和磁盘之间动态交换数据（分页机制）来实现。虚拟内存的关键优点是它隔离了不同程序的地址空间，增加了系统的稳定性和安全性，并且通过延迟加载和内存共享等技术，提高了内存使用的效率。

物理内存（Physical Memory）

物理内存指的是计算机硬件（RAM）实际提供的存储空间。操作系统的任务之一是管理物理内存的分配给各个程序和系统进程。由于物理内存的大小是有限的，操作系统需要采用一系列策略（如分页、分段）来高效利用这些资源。

三者之间的关系

• 逻辑内存到虚拟内存：程序在运行时操作的是逻辑地址，这些地址通过操作系统的内存管理单元（MMU）映射到虚拟地址空间。这一过程对程序透明，即程序不需要知道其逻辑地址如何映射到虚拟地址。
• 虚拟内存到物理内存：虚拟内存地址通过页表（操作系统维护的结构）映射到物理内存地址。当程序访问的数据不在物理内存中时（缺页），操作系统会从磁盘中加载该数据到物理内存，这一过程称为分页交换（Paging）。
• 页锁定内存（Pinned/Page-locked Memory）：某些情况下，为了避免性能损失，特定的内存区域可以被“锁定”，使得这部分内存不会被交换到磁盘上。这对于需要快速响应或高性能的应用程序非常重要，如实时系统或高性能计算（HPC）。

通过这种层次化的管理，操作系统能够提供一个既高效又安全的内存使用环境，使得不同的程序能够同时运行而互不干扰。

上面右图是 NVIDIA 官网上关于 GPU 访问 Host 端 Memory 的过程，我们看到如果 GPU 想要访问 Pinned Memory 的话可以直接通过 DMA（Direct Memory Access）控制器进行数据传输，而无需 CPU 干预。由于这部分内存是页锁定的，操作系统保证它不会被换出到磁盘，这意味着 GPU 可以直接访问这块内存区域，大大加快了数据传输速度。

而相反如果 GPU 想要访问 Pageable Memory 的话，需要先将数据复制到一个内部的 Pinned Memory 缓冲区，然后才能通过 DMA 传输到 GPU，这增加了额外的复制开销和延迟。这也是为什么多流程序中我们分配的是 Pinned Memory 而非 Pageable Memory 的原因之一

此外 Pageable Memory 可能会被换出磁盘，而我们的 stream 是异步的，必须把数据计算完之后写回到之前的地方，不能让它消失，所以这个时候需要使用 Pinned Memory

所以大家在看很多 CUDA 程序时其实很少会看到使用 Malloc，更多的是使用 cudaMallocHost 来分配 Host 上的内存空间

关于内存杜老师之前也讲过，大家感兴趣的可以看看：3.3.cuda运行时API-内存的学习，pinnedmemory，内存效率问题

3. Streams实验(单流vs多流)

下面有几个单流和多流对比的实验，我们一起来看下：

我们设计了一个矩阵乘法的核函数，在上图中它被设置 256，此外 Grid Size 和 Block Size 都被设置为 4x4，我们对比了单个 stream 和 5 个 stream 执行 1 memcpy，5 kernel 所使用的时间，结果表明多流比单流要快 2.32 倍，这是因为这个时候 Grid 和 Block 都很小，GPU 上有空余的资源，所以多流有效。单流串行执行时 SM Warp 占用率不高，而多流并行执行时 SM Warp 占用率非常高

这个时候我们修改下 Grid Size 和 Block Size 如上图所示，我们扩大了 Block Size 可以让 Warp 的填充率提升，但计算资源还是有很多，所以多流的速度提升和之前的一样

我们继续看我们提高矩阵的 Size 如上图所示，可以看看到此时单流的 SM Warp 占用率有所提升但依旧没有被占满，所以多流和单流的加速比是 2.47 和之前的差不多

既然如此我们让矩阵 Size 继续扩大让它占满 SM Warp 如上图所示，这个时候我们发现当计算资源已经占满的时候，核函数不能够再 overlapping 了，只能够一个一个的去等待，此时多流和单流的加速比只有 1.33 勉勉强强进行一个加速

大部分情况下我们会通过各种调参来设计核函数让它尽可能的把计算资源占满，此时多流带来的加速在核函数与核函数之间体现得就不是非常明显了，这时多流的优势就体现在核函数和内存拷贝之间了。

4. 如何隐藏延迟(memory)

我们下面来看当核函数资源占满时如何利用多流隐藏 memory 延迟，如下图所示：

在上图中红色代表一个流的指令，绿色代表另一个流的指令，我们提供了两种方案，方案 1 是按照正常的方式去调度，按部就班先执行 H2D1、H2D2 然后执行 KERNEL1 最后执行 D2H1、D2H2。我们可以看到在方案 1 中第二个流中 D2H2 的指令被放在了第一个流中 D2H1 的指令之后，这其实就有点浪费了，因为 D2H2 这个数据拷贝指令和前面的指令没有任何的依赖，因此我们可以把 D2H2 放在最一开始，减少 D2D2 Queue 的等待时间，也就是上面右图方案 2 中所做的

5. 如何隐藏延迟(kernel)

我们下面来看当核函数资源没有被占满时如何利用多流隐藏 kernel 延迟，如下图所示：

上图中我们有两个流，每个流有两个 Kernel 核函数，方案 1 是按照正常的方式去调度，但是我们发现 KernelA1 启动后其实是有空余的，因此我们可以在 KernelA1 执行等待时 launchKernelA2，也就是方案 2 所做的事情，KernelA1 和 KernelA2 几乎是同一时间启动的，这样就可以减少方案 1 中正常调度所空闲的时间

我们再来看另外一种情况，两个流中 KernelA1 和 KernelB2 所需要占据的计算资源更多，方案 1 是按照正常的方式调度，空闲出来了 KernelA1 执行的时间，方案 2 则是在 KernelA1 执行等待的时间启动 KernelA2， 这时可以减少一部分空余时间，方案 3 则是在 KernelA2 执行完后立即启动 KernelB2，不用等到 KernelA1 执行完，这样可以把 KernelB2 的时间完全节省下来，这是一个比较好的调度方式

6. 如何隐藏延迟(kernel+memory)

最后我们来看下如何同时隐藏 memory 和 kernel 延迟，如下图所示：

上图不同的颜色代表不同流的指令，图中左边和右边展示了两种不同的调度执行情况，左边的调度是先执行完所有的 H2D 后执行所有的 Kernel 最后执行所有的 D2H 操作，而右边的则是按流中的指令执行，先执行 H2D1、Kernel1、D2H1，依此类推

我们看到如果是左边的调度顺序，D2H1 需要等到 KERNEL3 之后再启动，这个时间就存在空闲，因为它和前面的指令没有任何关系，所以我们最好的方式是右边这种，在 KERNEL2 启动之后立即启动 D2H1，这个就是一个比较完美的调度执行过程

7. 代码分析

这里我们简单过一下 2.9-stream-and-event 案例代码，我们先从 main.cpp 开始，代码如下所示：

#include <stdio.h>
#include <cuda_runtime.h>#include "utils.hpp"
#include "timer.hpp"
#include "gelu.hpp"
#include "stream.hpp"int seed;void sleep_test(){Timer timer;int width     = 1<<10; // 4,096int size      = width * width;float low     = -1.0;float high    = 1.0;int blockSize = 16;int taskCnt   = 5;bool statMem  = true;char str[100];/* 初始化 */float* src_host;float* tar_host;cudaMallocHost(&src_host, size * sizeof(float));cudaMallocHost(&tar_host, size * sizeof(float));seed += 1;initMatrixSigned(src_host, size, low, high, seed);LOG("Input size is %d", size);/* GPU warmup */timer.start_gpu();SleepSingleStream(src_host, tar_host, width, blockSize, taskCnt);timer.stop_gpu();/* 1 stream，处理一次memcpy，以及n个kernel */blockSize = 16;timer.start_gpu();SleepSingleStream(src_host, tar_host, width, blockSize, taskCnt);timer.stop_gpu();std::sprintf(str, "sleep <<<(%2d,%2d), (%2d,%2d)>>>, %2d stream, %2d memcpy, %2d kernel", width / blockSize, width / blockSize, blockSize, blockSize, 1, 1, taskCnt);timer.duration_gpu(str);/* n stream，处理一次memcpy，以及n个kernel */timer.start_gpu();SleepMultiStream(src_host, tar_host, width, blockSize, taskCnt);timer.stop_gpu();std::sprintf(str, "sleep <<<(%2d,%2d), (%2d,%2d)>>>, %2d stream, %2d memcpy, %2d kernel", width / blockSize, width / blockSize, blockSize, blockSize, taskCnt, 1, taskCnt);timer.duration_gpu(str);
}void matmul_test() {/** 大家试着在这里对matmul计算做一个多流的计算看看整体延迟的改变* 可以观测到相比于kernel的计算, memcpy的延迟会很小*/
}void gelu_test() {/** 大家试着在这里对gelu计算做一个多流的计算看看整体延迟的改变* 可以观测到相比于memcpy的计算, kernel的延迟会很小*/
}int main(){cudaDeviceProp prop;cudaGetDeviceProperties(&prop, 0);// 需要先确认自己的GPU是否支持overlap计算if (!prop.deviceOverlap) {LOG("device does not support overlap");} else {LOG("device supports overlap");}sleep_test();// matmul_test();// gelu_test();// 这里供大家自由发挥。建议花一些在这里做调度的练习。根据ppt里面的方案实际编写几个测试函数。举几个例子在这里// e.g. 一个stream处理: H2D, 多个kernel，D2H。之后多个stream进行overlap// e.g. 一个stream处理: H2D, 大kernel，小kernel, D2H。之后多个stream进行overlap// e.g. 一个stream处理: H2D, 大kernel, H2D, 小kernel, D2H。之后多个stream进行overlap// e.g. 一个stream处理: H2D, 小kernel, H2D, 大kernel, D2H。之后多个stream进行overlap// e.g. 一个stream处理: H2D, 另外几个流分别只处理kernel, 和D2H。之后所有stream进行overlap// e.g. 一个stream处理: H2D(局部), kernel(局部), D2H(局部)。之后所有stream进行overlapreturn 0;
}

这段代码是一个多流 CUDA 程序的示例，主要展示了如何使用 CUDA 流（Streams）来并行化数据传输和 Kernel 执行，以提高 GPU 利用率和整体程序性能。以下是对程序的主要部分的逐步分析：(from chatGPT)

程序结构

程序包含三个测试函数（sleep_test、matmul_test、gelu_test），但只有 sleep_test 函数被实现和调用。这个函数演示了使用单个流和多个流执行内存复制（memcpy）和 Kernel 计算操作的性能差异。

CUDA Streams

CUDA 流允许异步执行 Kernel 计算和内存传输，可以在不同的流之间实现操作的并行。这对于提高资源利用率和程序性能非常关键。

sleep_test 函数

• 初始化：分配页锁定内存（Pinned Memory）来存储源（src_host）和目标（tar_host）矩阵，并初始化源矩阵。页锁定内存用于加速主机到设备（H2D）和设备到主机（D2H）的数据传输。
• GPU Warmup：通过执行 SleepSingleStream 函数进行GPU预热，以确保后续的性能测试不受初始化延迟的影响。
• 单流执行：使用单个 CUDA 流执行一次内存复制和多个 Kernel 计算。这演示了没有流并行时的基线性能。
• 多流执行：使用多个 CUDA 流（数量等于任务数taskCnt）并行执行一次内存复制和多个 Kernel 计算。每个流处理一部分任务，旨在展示通过并行化可以提高性能。

CUDA Device Properties

在 main 函数中，首先检查 CUDA 设备是否支持操作重叠（Overlap）。操作重叠是指设备能否同时进行 Kernel 计算和数据传输（H2D或D2H）。如果设备支持重叠，那么使用多流技术可以显著提高性能。

性能测量

程序使用自定义的 Timer 类来测量 GPU 操作的执行时间，通过对比单流和多流的执行时间来展示多流并行的性能优势。

总结

这个示例程序主要展示了如何利用 CUDA 流来并行化数据传输和 Kernel 执行，以及如何测量和对比单流与多流执行模式的性能差异。通过并行化操作，可以显著提高GPU的工作效率和整个程序的执行速度，尤其是在处理大量数据和复杂计算时。在实际应用中，开发者可以根据具体的计算任务和硬件特性来选择最适合的并行策略，以达到最佳的性能优化。

我们接着来看核心的 stream.cu 代码如下所示：

#include <cuda_runtime.h>
#include <utils.hpp>// #define MAX_ITER 10000     // memcpy == kernel / 10    (kernel执行的太快看不出来overlapping)
#define MAX_ITER 100000    // memcpy == kernel / 100   (开始能够看出来kernel的overlapping)
// #define MAX_ITER 10000000   // memcpy == kernel / 10000  (可以非常清楚的看到kernel的Overlapping)
#define SIZE 32// 为了能够体现延迟，这里特意使用clock64()来进行模拟sleep
// 否则如果kernel计算太快，而无法观测到kernel在multi stream中的并发
// 大家根据自己的情况需修改sleep的时间
__global__ void SleepKernel(int64_t num_cycles)
{int64_t cycles = 0;int64_t start = clock64();while(cycles < num_cycles) {cycles = clock64() - start;}
}/* 1 stream，处理一次memcpy，以及n个kernel */
void SleepSingleStream(float* src_host, float* tar_host, int width, int blockSize, int count) 
{int size = width * width * sizeof(float);float *src_device;float *tar_device;CUDA_CHECK(cudaMalloc((void**)&src_device, size));CUDA_CHECK(cudaMalloc((void**)&tar_device, size));;for (int i = 0; i < count ; i++) {for (int j = 0; j < 1; j ++) CUDA_CHECK(cudaMemcpy(src_device, src_host, size, cudaMemcpyHostToDevice));dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);SleepKernel <<<dimGrid, dimBlock >>> (MAX_ITER);CUDA_CHECK(cudaMemcpy(src_host, src_device, size, cudaMemcpyDeviceToHost));}CUDA_CHECK(cudaDeviceSynchronize());cudaFree(tar_device);cudaFree(src_device);
}/* n stream，处理一次memcpy，以及n个kernel */
void SleepMultiStream(float* src_host, float* tar_host,int width, int blockSize, int count) 
{int size = width * width * sizeof(float);float *src_device;float *tar_device;CUDA_CHECK(cudaMalloc((void**)&src_device, size));CUDA_CHECK(cudaMalloc((void**)&tar_device, size));/* 先把所需要的stream创建出来 */cudaStream_t stream[count];for (int i = 0; i < count ; i++) {CUDA_CHECK(cudaStreamCreate(&stream[i]));}for (int i = 0; i < count ; i++) {for (int j = 0; j < 1; j ++) CUDA_CHECK(cudaMemcpyAsync(src_device, src_host, size, cudaMemcpyHostToDevice, stream[i]));dim3 dimBlock(blockSize, blockSize);dim3 dimGrid(width / blockSize, width / blockSize);/* 这里面我们把参数写全了 <<<dimGrid, dimBlock, sMemSize, stream>>> */SleepKernel <<<dimGrid, dimBlock, 0, stream[i]>>> (MAX_ITER);CUDA_CHECK(cudaMemcpyAsync(src_host, src_device, size, cudaMemcpyDeviceToHost, stream[i]));}CUDA_CHECK(cudaDeviceSynchronize());cudaFree(tar_device);cudaFree(src_device);for (int i = 0; i < count ; i++) {// 使用完了以后不要忘记释放cudaStreamDestroy(stream[i]);}}/* n stream，处理一次memcpy，以及n个kernel */

这段CUDA代码演示了如何使用单个流和多个流（CUDA Streams）来管理内存复制和 Kernel 执行的过程，通过调整 Kernel 执行时间，观察不同配置下的性能差异和操作重叠（Overlapping）效果。这是一种常见的技术，用于提高GPU资源利用率和加速CUDA程序的执行。代码主要由两个关键部分组成：SleepSingleStream和SleepMultiStream函数。(from chatGPT)

SleepSingleStream 函数

这个函数使用单个默认流（Implicit Default Stream）来顺序执行内存复制和 Kernel 执行操作。具体步骤如下：

• 首先，为源（src_device）和目标（tar_device）设备内存分配空间。
• 使用循环，对每一个任务：
  • 通过 cudaMemcpy 函数将数据从主机内存（src_host）复制到设备内存（src_device）。
  • 配置 Kernel 的执行参数（dimBlock和dimGrid）并执行 SleepKernel，模拟延迟。
  • 将结果从设备内存复制回主机内存。
• 执行完所有任务后，调用 cudaDeviceSynchronize 以确保所有 CUDA 操作都已完成。
• 释放设备内存。

这个过程中，所有操作都是在单个默认流中顺序执行的，不会有操作重叠。

SleepMultiStream 函数

与 SleepSingleStream 不同，SleepMultiStream 函数演示了如何使用多个 CUDA 流来并行执行内存复制和 Kernel 执行操作。步骤如下：

• 同样先为源和目标设备内存分配空间。
• 创建与任务数（count）相同数量的CUDA流。
• 对每个任务：
  • 使用 cudaMemcpyAsync 将数据异步复制到设备内存，指定使用对应的流（stream[i]）。
  • 配置 Kernel 执行参数并在指定流中异步执行 SleepKernel。
  • 使用 cudaMemcpyAsync 异步将结果复制回主机内存，同样指定使用对应的流。
• 使用 cudaDeviceSynchronize 等待所有流中的操作完成。
• 释放设备内存和销毁创建的CUDA流。

在 SleepMultiStream 函数中，通过将每个任务分配给不同的流并异步执行，可以实现数据复制和 Kernel 计算的操作重叠，从而提高了GPU的利用率和程序的总体执行效率。这种方法尤其适用于当 Kernel 计算时间短于数据传输时间或当有大量独立任务需要并行处理的情况。

MAX_ITER 的作用

通过定义 MAX_ITER，代码可以调整 SleepKernel 执行的时长。通过增加MAX_ITER的值，可以模拟更长时间的 Kernel 执行，从而更容易观察到多流并发执行时 Kernel 操作的重叠效果。这对于理解和优化CUDA程序中的并行性和性能至关重要。

总结

这个示例代码展示了单流和多流策略在CUDA编程中的应用，以及如何通过控制 Kernel 执行时间来观察和利用操作重叠，从而提高程序性能。通过合理利用CUDA流，可以显著提升复杂CUDA应用的执行效率和资源利用率。

总结

本次课程我们重点讲解了 CUDA 编程中的 stream 流，并讲解了如何利用多流来隐藏 Memory 和 Kernel 执行的延迟，我们可以利用 Nsight 可视化工具来对比分析单流和多流执行的结果

OK，以上就是第 4 小节有关 CUDA 中的 stream 与 event 的全部内容了，下节我们来学习双线性插值，敬请期待😄

参考

• NVIDIA: Stream and concurrency webinar
• 3.3.cuda运行时API-内存的学习，pinnedmemory，内存效率问题
• 3.4.cuda运行时API-流的学习，异步任务的管理