今天是6月10号，记得自己是去年这个时候开始接触CUDA的，目前只会写一点简单的程序，对顺程序的优化设计，还是一个小萌新，我在实验室的项目里面，看到的CUDA 程序使用了多个流的设计，但是它的时间线却让我匪夷所思，为什么不是并行的呢？带着这个疑惑，我决定在最近这段时间，通过Google，搞明白这个问题，然后分享给大家，希望有朝一日，碰巧会有一个少年，碰巧也是通信CS 领域，碰巧也要用CUDA做加速计算，碰巧看到了我这篇文章，那么，我的所作所为就有意义了，帮助他节省了一段时间，最重要的，在他刚开始接触这个领域的时候，给了他一股学下去的兴趣~

1.cuda单个流的应用

话不多说，先看一个Nsight里面的Timeline

图中Streams里面表示了一个流里面的执行顺序，基本上就是Memory和Compute的结合，毕竟，拷贝数据和计算，就是我们想要完成的事；
这是一个完成数组相加的cuda程序，基本的思路是“拷贝a到GPU,拷贝b到GPU，计算，拷贝c到CPU”

核函数如下，每次可以算1024个点，采用4个block，每个block256个thread，一个thread算一个点；

__global__ void kernel(int *a, int *b, int *c)
{
   
 int idx = threadIdx.x + blockDim.x * blockIdx.x;
 if (idx < N)
 {
   
  c[idx] = a[idx] + b[idx];
 }
}

然后，现在我把数组的大小设置为1024*4，就需要核函数执行四次，从而算出结果；这是一个有着先后顺序的任所以，将这些任务顺序的丢到流里面，就会按照顺序进行下去；
（本人使用的是cuda10.2+vs2015）
第一部分完整代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

#define N 1024*1024
__global__ void kernel(int *a, int *b, int *c)
{
   
	int idx = threadIdx.x + blockDim.x * blockIdx.x;
	if (idx < N)
	{
   
		c[idx] = a[idx] + b[idx];
	}
}
int main()
{
   
	//查询设备是否支持overlap
	cudaDeviceProp prop;
	int whichDevice;
	cudaGetDevice(&whichDevice);
	cudaGetDeviceProperties(&prop, whichDevice);
	if (!prop.deviceOverlap)
	{
   
		printf("Devices will not handle overlaps");
		return 0;
	}

	int FULL_DATA_SIZE = N * 4;

	cudaEvent_t start, stop;
	float elapsedTime;
	//启动计时器
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);
	//初始化流
	cudaStream_t stream;
	cudaStreamCreate(&stream);
	//数据分配操作
	int *host_a, *host_b, *host_c;
	int *dev_a, *dev_b, *dev_c;
	//在 GPU 上分配内存
	cudaMalloc((void**)&dev_a, N * sizeof(int));
	cudaMalloc((void**)&dev_b, N * sizeof(int));
	cudaMalloc((void**)&dev_c, N * sizeof(int));
	//在 CPU 上分配内存，是流使用的页锁定内存
	cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
	cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
	cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault);
	//随机数填充主机内存
	for (int i = 0; i < FULL_DATA_SIZE; i++)
	{
   
		host_a[i] = rand();
		host_b[i] = rand();
	}
	//在整体数据上循环，每个数据块的大小是 N 
	for (int i = 0; i < FULL_DATA_SIZE; i += N)
	{
   
		//将锁定内存以异步方式复制到设备上
		cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);
		cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream);
		kernel << <N / 256, 256, 0, stream >> > (dev_a, dev_b, dev_c);
		//将数据复制到锁定内存中
		cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyDeviceToHost, stream);
	}
	//同步
	cudaStreamSynchronize(stream);
	//
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
	//释放流和内存
	cudaFreeHost(host_a);
	cudaFreeHost(host_b);
	cudaFreeHost(host_c);
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	//销毁流
	cudaStreamDestroy(stream);
	return 0;
}

2020.6.17更新

2.cuda多个流的应用

首先，可以通过cudaSample查询一下GPU支持几个流；
打开这个目录下的vs，
“…\ProgramData\NVIDIA Corporation\CUDA Samples\v10.2\1_Utilities\deviceQuery”

图中的MultiProcessor代表实际可以使用的流的个数

突然接到了项目任务，这一块以后再更新吧！