文章目录

前言
1. CUDA流的理解
2. C++ 实现CUDA流
3. 执行结果
总结
学习资料

VS2017 CUDA编程学习1：CUDA编程两变量加法运算
VS2017 CUDA编程学习2：在GPU上执行线程
VS2017 CUDA编程学习3：CUDA获取设备上属性信息
VS2017 CUDA编程学习4：CUDA并行处理初探 - 向量加法实现
VS2017 CUDA编程学习5：CUDA并行执行-线程
VS2017 CUDA编程学习6: GPU存储器架构
VS2017 CUDA编程学习7：线程同步-共享内存
VS2017 CUDA编程学习8：线程同步-原子操作
VS2017 CUDA编程学习9：常量内存
VS2017 CUDA编程学习10：纹理内存
VS2017 CUDA编程学习实例1：CUDA实现向量点乘
VS2017 CUDA编程学习11：CUDA性能测量

前言

今天跟大家分享CUDA编程中一个重要概念：CUDA流。

1. CUDA流的理解

之前介绍的GPU内核函数，通过多线程实现加速，可以理解为数据并行。而CUDA流可以理解为任务并行，即多个相互独立的内核函数同时执行。

CUDA流是GPU上的工作队列，队列中的工作将以特定的顺序执行。队列中的工作主要包括：内核函数的调用， cudaMemcpy系列的数据传输，以及CUDA事件的操作。CUDA流中工作的顺序以添加的顺序来执行，先加入的先执行，即每个CUDA流中的工作是串行执行的。

2. C++ 实现CUDA流

这里以GPU上加法操作为例实现CUDA流。

例子中主要调用一下API:
cudaHostAlloc/cudaFreeHost(): 使用页面锁定内存为CPU设备变量分配内存，页面锁定内存，简单理解就是内存数据永远存储在内存上，不会转储到硬盘空间来暂存，缺点是如果内存一直不释放，将造成内存高使用率，可能影响其他程序的运行；

cudaMemcpyAsync:用于异步传输数据，和cudaMemcpy函数相比，多了最后一个参数 - 指定一个特定的CUDA流，从而在这个特定的CUDA流中进行存储器传输操作。这里注意，异步操作是指当该函数返回时，实际的存储器传输操作可能尚未完成，甚至尚未开始。

这里要注意，每个CUDA流中的工作是串行的，而流和流之间则默认不保证顺序。每个流程都有传输和计算工作，这样不同的流间计算和传输工作可以并行执行，起到加速的作用，进而提升程序性能。

cudaStreamSynchronize(): 确保CUDA流中所有操作都结束后才返回控制权给主机（CPU）。

详细代码如下所示：

#include <stdio.h>
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
#include <device_launch_parameters.h>#define N 50000__global__ void gpuAdd(int* d_a, int* d_b, int* d_c)
{int tid = threadIdx.x + blockIdx.x * blockDim.x;while (tid < N){d_c[tid] = d_a[tid] + d_b[tid];tid += blockDim.x * gridDim.x;}
}int main()
{//定义CPU变量int* h_a, *h_b, *h_c;//为CUDA流 stream0定义GPU变量int* d_a0, *d_b0, *d_c0;//为CUDA流 stream1定义GPU变量int *d_a1, *d_b1, *d_c1;//定义CUDA流变量cudaStream_t stream0, stream1;//创建CUDA流cudaStreamCreate(&stream0);cudaStreamCreate(&stream1);//定义CUDA事件cudaEvent_t e_start, e_stop;//创建CUDA事件cudaEventCreate(&e_start);cudaEventCreate(&e_stop);//开始记录CUDA事件cudaEventRecord(e_start, 0);//分配CPU内存，使用页面锁定内存, 不会发生换页操作（将暂时不用的数据从内存中转存硬盘上），数据永远存在内存中cudaHostAlloc(&h_a, 2 * N * sizeof(int), cudaHostAllocDefault);cudaHostAlloc(&h_b, 2 * N * sizeof(int), cudaHostAllocDefault);cudaHostAlloc(&h_c, 2 * N * sizeof(int), cudaHostAllocDefault);//分配GPU内存cudaMalloc(&d_a0, N * sizeof(int));cudaMalloc(&d_b0, N * sizeof(int));cudaMalloc(&d_c0, N * sizeof(int));cudaMalloc(&d_a1, N * sizeof(int));cudaMalloc(&d_b1, N * sizeof(int));cudaMalloc(&d_c1, N * sizeof(int));//初始化CPU数据for (int i = 0; i < 2 * N; i++){h_a[i] = 2 * i * i;h_b[i] = i;}//为两个CUDA流执行异步内存拷贝操作cudaMemcpyAsync(d_a0, h_a, N * sizeof(int), cudaMemcpyHostToDevice, stream0);cudaMemcpyAsync(d_a1, h_a + N, N * sizeof(int), cudaMemcpyHostToDevice,stream1);cudaMemcpyAsync(d_b0, h_b, N * sizeof(int), cudaMemcpyHostToDevice,stream0);cudaMemcpyAsync(d_b1, h_b + N, N * sizeof(int), cudaMemcpyHostToDevice,stream1);//为两个CUDA流分别调用内核函数gpuAdd << <512, 512, 0, stream0 >> > (d_a0, d_b0, d_c0);gpuAdd << <512, 512, 0, stream1 >> > (d_a1, d_b1, d_c1);//从GPU拷贝数据到CPUcudaMemcpyAsync(h_c, d_c0, N * sizeof(int), cudaMemcpyDeviceToHost,stream0);cudaMemcpyAsync(h_c + N, d_c1, N * sizeof(int), cudaMemcpyDeviceToHost,stream1);//等待GPU所有线程运行结束cudaDeviceSynchronize();//等待所有CUDA流中工作完成cudaStreamSynchronize(stream0);cudaStreamSynchronize(stream1);//结束CUDA事件记录cudaEventRecord(e_stop, 0);//等待CUDA事件记录完成cudaEventSynchronize(e_stop);//GPU上运行时间统计float elapsedTime;cudaEventElapsedTime(&elapsedTime, e_start, e_stop);printf("GPU上加法操作 %d 个数耗时：%3.1f ms\n", 2 * N, elapsedTime);//验证GPU上运行结果的准确性int correct = 1;for (int i = 0; i < 2 * N; i++){if (h_a[i] + h_b[i] != h_c[i]){correct = 0;break;}}if (correct == 1){printf("GPU上加法操作结果准确！\n");}else{printf("GPU上加法操作结果不准确！\n");}//释放GPU内存cudaFree(d_a0);cudaFree(d_a1);cudaFree(d_b0);cudaFree(d_b1);cudaFree(d_c0);cudaFree(d_c1);//释放CPU页面锁定内存cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);system("pause");return 0;
}

3. 执行结果

总结

CUDA流，说白了就是实现了多任务并行的功能。今天就分享到这里，本人还是小菜鸟，如果上面描述有问题，还请大神不吝指正！

学习资料

《基于GPU加速的计算机视觉编程》

VS2017 CUDA编程学习12：CUDA流相关推荐

VS2017 CUDA编程学习11：CUDA性能测量
文章目录前言 1. CUDA事件API 2. C++ 实现CUDA事件例子 3. CUDA事件例子的执行结果 4. NVIDIA Visual Profiler工具总结学习资料 VS2017 C ...
VS2017 CUDA编程学习4：CUDA并行处理初探 - 向量加法实现
文章目录前言 CUDA实现向量加法 2. 结果展示总结学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2:在GPU上执行线程 VS2017 ...
VS2017 CUDA编程学习8：线程同步-原子操作
文章目录前言 1. 原子操作的理解 2. C++ CUDA实现原子操作 3. 执行结果总结学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2 ...
VS2017 CUDA编程学习5：CUDA并行执行-线程
文章目录前言 1. 线程 2. 代码实现总结学习资料 VS2017 CUDA编程学习1:CUDA编程两变量加法运算 VS2017 CUDA编程学习2:在GPU上执行线程 VS2017 CUDA编 ...
【Java】函数式编程学习笔记——Stream流
学习视频:https://www.bilibili.com/video/BV1Gh41187uR?p=1 (1)[Java]函数式编程学习笔记--Lambda表达式 (2)[Java]函数式编程学习笔 ...
CUDA 编程学习
0 简单的CUDA简介 1.简单教程 CUDA C ++只是使用CUDA创建大规模并行应用程序的方法之一.它允许您使用功能强大的C ++编程语言来开发由GPU上运行的数千个并行线程加速的高性能算法.许 ...
CUDA编程学习3——并行计算初窥CUDA的软硬件架构
目录并行计算概述查看GPU相关信息软硬件架构基础物理层(物理结构) 逻辑层(kernel组织) 物理层和逻辑层的总结参考并行计算概述所谓并行计算的概念定义同时多个计算资源一起工作(逻辑 ...
CUDA编程第六章: 流和并发
本章内容: 理解流和事件的本质理解网格级并发重叠内核执行和数据传输重叠CPU和GPU执行理解同步机制避免不必要的同步调整流的优先级注册设备回调函数通过NVIDIA可视化性能分析器显示应 ...
CUDA编程学习0——环境搭建环境详解
目录环境配置软件安装 1.支持最高的cuda版本查询,下载cuda开发软件: 3.配置环境(~/.bashrc添加环境变量) 4.后续维护查询补:关于windows下的cuda环境配置一.Vi ...
Cuda编程学习（一）
使用Cuda进行GPU编程 --GPU高性能编程Cuda实战第四章C并行编程总结开始学习Cuda和GPU编程的相关知识啦. 感觉GPU编程会越来越重视,尤其是在移动端的计算复杂度优化方面.掌 ...

VS2017 CUDA编程学习12：CUDA流