前言

CUDA程序的并行层次主要有两个，一个是核函数内部的并行，一个是核函数外部的。合理使用CUDA流能实现核函数外部的并行。

CUDA流

核函数外部并行：
（1）核函数计算与数据传输之间的并行
（2）主机计算与数据传输之间的并行
（3）不同数据传输之间的并行
（4）核函数计算与主机计算之间的并行
（5）不同核函数之间的并行

CUDA流：一个CUDA流指的是由主机发出的在一个设备中执行的CUDA操作序列。CUDA流各个操作的次序是由主机控制的，按照主机发布的次序执行。两个不同的CUDA流中的操作不一定按照顺序执行，和线程一样，要么并发要么交错地执行。
任何CUDA操作都存在于某个CUDA流中，要么是默认流和空流，要么是指定地非空流。前面的CUDA程序没有指定的都是默认和空流中执行的。
非默认的CUDA流实在主机端产生与销毁的。

// 创建
cudaError_r cudaStreamCream_t(cudaStream_t*);  // 在yolov5_trt中使用过
// 销毁
cudaError_r cudaStreamDestroy(cudaStream_t);

由创建的代码可知，输入参数是cudaStream_t类型指针，返回值类型是一个错误代码。

不同CUDA流之间的并发

主机向某个CUDA流中发布一系列命令后必须马上获得程序的控制权，不用等待该CUDA流中的命令在设备中执行完毕。这样，就可以通过主机生成多个独立的CUDA流。
同时，CUDA运行时API还提供了两个函数：

// 这两个函数和等待和守护线程有点像
// 会强制阻塞主机，知道cuda流stream中的所有操作都执行完毕。
cudaError_r cudaStreamSynchronize(cudaStream_t stream);
// 不阻塞主机，只是检查CUDA流stream中的所有操作是否都执行完毕。
cudaError_r cudaStreamQuery(cudaStream_t stream);

在默认流中重叠主机与设备

合理利用主机与设备之间的执行次序。如前面的简单的数组求和的核函数：

   // 将某些数据从主机复制到设备上cudaMemcpy(d_x,h_x,M,cudaMemcpyHostToDevice);cudaMemcpy(d_y,h_y,M,cudaMemcpyHostToDevice);// 调用核函数在设备中进行计算,数组求和const int block_size = 128;  // 不同型号的GPU有线程限制，开普勒到图灵最大为1024const int gride_size = N/block_size;add<<<gride_size,block_size>>>(d_x,d_y,d_z);// 将某些数据从设备复制到主机上,这个数据传输函数隐式的起到了同步主机与设备的作用,所以后面用不用cudaDeviceSynchronize都可以cudaMemcpy(h_z,d_z,M,cudaMemcpyDeviceToHost);

从主机看，数据传输是同步的。在主机执行核函数之前的CUDA操作语句将在默认的CUDA流中按代码的顺序执行。在进行数据传输时，主机是闲置的，不能进行其他操作。
不同的是，在执行核函数时，核函数的启动是异步的。主机发出执行的命令，不会等核函数执行完毕，而会立刻的到程序的控制权，往下执行。执行到从设备到主机传输数据这条语句，该语句不会立即执行。因为这是默认流中的CUDA操作，必须等待前一个CUDA操作执行完毕才会开始执行。
从上面的分析可知，核函数启动异步，对默认流中CUDA的操作会阻塞，但是对主机中的程序不会进行阻塞。如果调用核函数下一句是主机中的某个计算任务，那么主机就会在设备执行核函数的同时去进行一些计算。这样主机和设备就可以同时进行计算。
当然，重叠主机与设备的计算要考虑两者计算的时间。一般来说设备函数的计算速度是主机函数的10倍左右。当设备函数执行完的时间与接下来的主机函数执行完的时间差不多时，两者擦汗不读时间结束，加速效果最好。当主机函数与设备函数的计算时间相差很多的情况下，设备函数占主要计算时间或者主机函数占主要时间，加速效果就差了。

用非默认CUDA流重叠多个核函数的执行

一个默认流可以实现主机计算与设备计算的并行，但是多个核函数之间的并行必须使用多个CUDA流。

重叠多个核函数的例子

使用非默认流时，核函数的执行配置中必须包含一个流对象。例如，核函数有如下三种调用方式：

add<<<gride_size,block_size>>>();
add<<<gride_size,block_size,shared>>>();
add<<<N_gride,N_block_size,N_shared,stream_id>>>();

其中，第一个调用方式在默认流中执行。第二个是使用了动态共享内存的，同样在默认流中执行。只有第三个调用方式，说明核函数在编号为stream_id的CUDA流中执行，而且可以使用动态共享内存，不使用动态共享内存的情况下，动态共享字节参数不能省略，必须为。
如修改原来的核函数中数组相加的程序：

// 多个非默认流for (int n = 0; n < num; ++n){int offset = n * N1;add<<<grid_size, block_size, 0, streams[n]>>>(d_x + offset, d_y + offset, d_z + offset);}
-----------------------------------------------------------------------------
void __global__ add(const real *d_x, const real *d_y, real *d_z)
{const int n = blockDim.x * blockIdx.x + threadIdx.x;if (n < N1){// 使用了for循环进行数组元素相加，这是为了后面的计时方便，不使用也行。for (int i = 0; i < 1000000; ++i){d_z[n] = d_x[n] + d_y[n];}}
}

和最初的代码，有两个区别：
（1）在调用核函数的时候，使用了for循环，表示使用非默认流中重叠多个核函数的执行。n，表示CUDA流的索引。offset 表示多个核函数中，每个核函数使用N1个线程。
（2）在每个核函数计算时，使用了for循环进行数组元素相加，这是为了后面的计时方便，不使用也行。
在测试中发现，使用多个流相对于使用一个流有了加速。但是，存在加速比的限制：
（1）当所有CUDA流中对应核函数的线程总数和超过某个值时，再增加流的数目就不会带来更高的加速比了。
（2）单个GPU中能够并发执行的核函数个数的上限，不同的GPU架构有不同的上限。

用非默认CUDA流重叠核函数的执行与数据传递

不可分页主机内存与异步的数据传输函数

要实现核函数执行与数据传输的并发（重叠），必须让这两个操作处于不同的非默认流，而且数据传输必须使用cudaMemcpy()函数的异步版本，即cudaMemcpyAsync()函数（前面tensorrt部署中有使用）。异步传输由GPU中的DMA直接实现，不需要主机参与。如果用同步的数据传输，主机无法再一个流中进行数据传输时，去另一个流调用核函数。这样核函数执行与数据传输的并发也就无法实现。

异步的数据传输函数，只比同步的多了一个参数，表示所在流的变量。在使用异步数据传输函数时，需要将主机内存定义为不可分页内存或者固定内存。不可分页主机内存分分配可以由以下两个函数中任何一个实现：

尝试使用不同的流执行不同的操作来提升性能。

一般一个cuda程序要有主机向设备进行数据传输（H2D），核函数的调用（KER），设备向主机进行数据传输（D2H）。这三个操作在一个CUDA流中执行的顺序：
s t r e a m 0 : H 2 D − > K E R − > D 2 H \qquad \qquad \qquad \qquad stream0 : H2D -> KER -> D2H stream0:H2D−>KER−>D2H
如果简单的将3个操作放入3个不同的流中：
s t r e a m 0 : H 2 D \qquad \qquad \qquad \qquad stream0 : H2D stream0:H2D
s t r e a m 1 : − > K E R \qquad \qquad \qquad \qquad stream1 : \quad \qquad -> KER stream1:−>KER
s t r e a m 2 : − > D 2 H \qquad \qquad \qquad \qquad stream2: \qquad \qquad \qquad \qquad -> D2H stream2:−>D2H
这样操作并不能带来性能提升。必须创造出在逻辑上可以并发执行的CUDA操作，可以将以上3个CUDA操作都分成若干等份，然后在每一个流中发布一个CUDA操作序列。
s t r e a m 0 : H 2 D − > K E R − > D 2 H \qquad \qquad \qquad \qquad stream0 : H2D -> KER -> D2H stream0:H2D−>KER−>D2H
s t r e a m 1 : H 2 D − > K E R − > D 2 H \qquad \qquad \qquad \qquad stream1 : \qquad \qquad H2D -> KER -> D2H stream1:H2D−>KER−>D2H
如上面使用两个流的情况，每个CUDA操作所处理的数据量只有使用一个CUDA流时的一半。其中H2D这里不能并发的执行，是受硬件资源的限制。如果H2D ，KER ，D2H这三个操作执行的时间都相同，那么可以有效的隐藏两个CUDA操作，使得总的执行效率相比使用单个CUDA流的情况提升 6 / 4 = 1.5 6/4=1.5 6/4=1.5倍。如下给出使用非默认CUDA流重叠核函数的执行与数据传递的示例：

// 部分代码
void __global__ add(const real *x, const real *y, real *z, int N)
{const int n = blockDim.x * blockIdx.x + threadIdx.x;if (n < N){// 和上面一样故意让核函数求和操作重复40次，让核函数和数据传输所用时间相当。for (int i = 0; i < 40; ++i){z[n] = x[n] + y[n];}}
}void timing
(const real *h_x, const real *h_y, real *h_z,real *d_x, real *d_y, real *d_z,const int num
)
{int N1 = N / num;int M1 = M / num;float t_sum = 0;float t2_sum = 0;for (int repeat = 0; repeat <= NUM_REPEATS; ++repeat){cudaEvent_t start, stop;CHECK(cudaEventCreate(&start));CHECK(cudaEventCreate(&stop));CHECK(cudaEventRecord(start));cudaEventQuery(start);// num个CUDA流for (int i = 0; i < num; i++){int offset = i * N1;// 使用异步数据传输CHECK(cudaMemcpyAsync(d_x + offset, h_x + offset, M1, cudaMemcpyHostToDevice, streams[i]));CHECK(cudaMemcpyAsync(d_y + offset, h_y + offset, M1, cudaMemcpyHostToDevice, streams[i]));int block_size = 128;int grid_size = (N1 - 1) / block_size + 1;// 非默认CUDA流add<<<grid_size, block_size, 0, streams[i]>>>(d_x + offset, d_y + offset, d_z + offset, N1);CHECK(cudaMemcpyAsync(h_z + offset, d_z + offset, M1, cudaMemcpyDeviceToHost, streams[i]));}

总结

CUDA流的基础知识总结

参考

如博客内容有侵权行为，可及时联系删除！
CUDA 编程：基础与实践
https://docs.nvidia.com/cuda/
https://docs.nvidia.com/cuda/cuda-runtime-api
https://github.com/brucefan1983/CUDA-Programming

CUDA编程之CUDA流相关推荐

java8函数式编程之Stream流处理的方法和案例讲解
函数式编程最早是数学家阿隆佐·邱奇研究的一套函数变换逻辑,又称Lambda Calculus(λ-Calculus),所以也经常把函数式编程称为Lambda计算. 为什么Java需要Lambda表达式 ...
CUDA编程之：Stream（流）
CUDA Stream(流):指在设备(Device)上按主机(Host)代码发出的顺序执行的一系列异步的CUDA操作.Stream封装这些操作,管理它们的顺序,允许在所有先前操作之后在流中排队执行操 ...
GPU编程之CUDA（三）——关于联想Y470
解决闪退问题-- 配置了几天,发现示例程序运行之后闪退以为是显卡驱动的问题,于是去NVIDIA官网下载了最新的驱动,但是之后还是闪退在这个过程中还尝试过禁用集成显卡,结果黑屏了...后来随意按了键 ...
CUDA编程之：cudaMemcpy()函数
cudaMemcpy用于在主机(Host)和设备(Device)之间往返的传递数据,用法如下: 主机到设备:cudaMemcpy(d_A,h_A,nBytes,cudaMemcpyHostToDevi ...
【Java八股文之基础篇（十九）】函数式编程之Stream流（上）
Stream流概述 Java8的Stream使用的是函数式编程模式,如同它的名字一样,它可以被用来对集合或数组进行链状流式的操作.可以更方便的让我们对集合或数组操作. 案例数据准备 <depe ...
java高级编程之IO流
1.流的分类 1.1按流向分类分为输入流和输出流 1.2按流角色分类分为节点流和处理流 1.3按数据单位分类分为字符流和字节流 2.流的四个抽象基类 InputStream.OutputStre ...
CUDA 网格级并发-流(2)
上节中描述到CUDA中可以通过流来实现网格级并发,按照流的划分主要分为两种: 空流(默认创建) 非空流(需要使用API创建) 空流为一个同步流,在空流中的大部分操作都会堵塞主机,kernel运行除外, ...
CUDA 网格级并发-流(1)
CUDA中的程序并发可以分为两种内核级并发网格级并发内核级并发是开发程序中经常使用到的,即通过划分block和thread实现同一个内核在GPU上同时并发,将同一个核分别部署到不同的SP上进行同 ...
java 网络编程简单聊天_网络编程之 TCP 实现简单聊天
网络编程之 TCP 实现简单聊天客户端 1.连接服务器 Socket 2.发送消息 package lesson02;import java.io.IOException;import java.i ...

CUDA编程之CUDA流

文章目录

前言