流(stream)是一种基于context之上的任务管道抽象，一个context可以创建n个流；如果想要在cuda上实现高性能异步多并发，可以通过流来实现任务的异步控制。nullptr是默认流每个线程都有自己的默认流。

1. 同步与异步

1.1 同步任务

• 以女朋友让男朋友买东西为例，同步任务下女朋友发消息让男朋友买东西，在发消息到拿到东西之间的时间，女朋友什么也不能做，只能在那边等待，程序卡在那里。
• 同步模式：当前面一个任务完成，才能继续去执行下一个任务，这是典型的串行同步的方式，执行任务效率比较低下。

1.2 异步任务

• 异步任务：女朋友想吃苹果发一条消息给男朋友，发完消息之后还可以继续做其他事情比如写作业，这个时间是解放出来的想干嘛干嘛。发完买苹果的任务，此时又想吃西瓜，给男朋友发买西瓜的任务。此时男朋友的任务队列新增了买西瓜任务。。。

总结

• 上面的例子中，男朋友的微信消息，就是任务队列，流的一种抽象
• 女朋友发出指令后，她可以做任何事情，无需等待指令执行完毕才去做其他事情，也就是执行异步操作.
• 女朋友发的指令被送到流中排队，男朋友根据流的队列，顺序执行
• 女朋友选择性的在需要的时候等待所有执行结果
• 新建一个流，就是新建一个男朋友，给他发指令就是给他发微信，你可以新建很多个男朋友
• 通过cudaEvent可以选择性等待任务队列中的部分任务是否就绪。

2. stream流代码实现

// CUDA 运行时头文件
#include <cuda_runtime.h>

#include <stdio.h>
#include <string.h>

#define checkRuntime(op)  __check_cuda_runtime((op),#op,__FILE__,__LINE__)

bool __check_cuda_runtime(cudaError_t code,const char* op,const char* file,int line)
{
	if(code != cudaSuccess){
		const char* err_name = cudaGetErrorName(code);
		const char* err_message = cudaGetErrorString(code);
		print("runtime_error %s%d %s failed.\n code=%s,message = %s\n",file,line,op,*err_name,*err_message)
		return false;
	}
	return true;
}

int main(){
	int device_id=0;
	checkRuntime(cudaSetDevice(device_id));
	
	// 通过cudaStreamCreate创建流
	cudaStream_t stream = nullptr;
	checkRuntime(cudaStreamCreate(&stream));
	
	//在GPU上开辟空间
	float* memory_device=nullptr;
	checkRuntime(cudaMalloc(&memory_device,100*sizeof(float)));
	
	// 在CPU上开辟空间并且放数据进去，将数据复制到GPU
	float* memory_host=new float[100];
	memory_host[2]=520.5;
	checkRuntime(cudaMemcpyAsyc(memory_device,memory_host,sizeof(float)*100,cudaMemcpyHostToDevice,stream))  // 异步
	
	// 在CPU上开辟pinned memory，并将GPU数据复制回来
	float* memory_page_locked = nullptr;
	checkRuntime(cudaMallocHost(&memory_page_locked,100*sizeof(float)));
	checkRuntime(cudaMemcpyAsyc(memory_page_locked,memory_device,sizeof(float)*100,cudaMemcpyDeviceToHost,stream))  // 异步	
	checkRuntime(cudaStreamSynchronize(stream));
	
	print("%f\n",memory_page_locked[2]);
	
	//释放内存
	checkRumtime(cudaFreeHost(memory_page_locked));
	checkRuntime(cudaFree(memory_device));
	checkRuntime(cudaStreamDestory(stream));
	
	delete[] memory_host;
	return 0;
}

• 通过cudaStreamCreate创建流
• cudaMemcpyAsync,其中cudaMemcpy它是同步的赋值方式，加上了Async就变成了异步的拷贝方式，异步的拷贝方式多了一个参数stream, 把流stream作为参数传入，函数就知道是通过异步方式拷贝。执行这条指令是立即返回的，也就是此时的拷贝操作被没有完成
• 最后通过cudaStreamSynchronize来执行stream等待，来同步等待所有的异步任务全部结束,也就是真正的耗时是在等待所有异步任务执行完成，也就是在这条命令上。这条命令之后，所有异步任务就全部完成了。

我们把打印命令print("%f\n",memory_page_locked[2]);放到cudaMemcpyAsyc之后，同步命令之前，按道理打印输出命令，拷贝的任务被没有结束,打印应该无法得到正确结果

checkRuntime(cudaMemcpyAsyc(memory_page_locked,memory_device,sizeof(float)*100,cudaMemcpyDeviceToHost,stream))  // 异步	
print("%f\n",memory_page_locked[2]);
checkRuntime(cudaStreamSynchronize(stream));

打印出来的是0.00，打印结果不对。这是因为异步赋值任务还没有结束，你就去拿它的结果了，肯定得不到预期结果

总结

1. 要十分注意，指令发出后，流队列中存储的是指令参数，不能加入队列后立即释放指针，这会导致流队列执行该指令时指针失效而出错,如下代码所示

checkRuntime(cudaMemcpyAsyc(memory_device,memory_host,sizeof(float)*100,cudaMemcpyHostToDevice,stream))  // 异步
delete[] memory_host;

异步拷贝任务还没结束就释放，会造成未知的结果，比如程序奔溃等不明确结果，正确释放时机要确保cudaMemcpyAsyc这条指令已经全部执行完成。已经不依赖于memory_host指针的时候，你才可以去释放它。
3. 应当在十分肯定流已经不需要这个指针后，才进行修改或者释放，否则会出现非预期结果出现
4. 举个例子：你给钱让男朋友买西瓜，他刚到店拿好西瓜，你把转的钱撤回去了。此时你无法预知他是否会跟店家闹矛盾，还是屁颠回去。如果想得到预期的结果，必须得让卖西瓜结束再处理钱的事情

流的代码实现是挺简单的，就是创建流，然后利用带Async这种异步的函数，把创建好的流指针作为参数传入异步函数，那他就可以执行异步任务了。另外创建多个流，并且多个流同时并发也是允许的。