CUDA by example Chapter10 流

为什么要使用CUDA流

前面章节实现的GPU并行性都是把同一个任务分成可以并行执行的子任务，利用GPU中数量众多的cuda核来实现单个任务的并行性。而如果要实现任务并行性，就需要使用CUDA流来实现。

页锁存主机内存

1. 页锁存内存的概念及用途

• 页所存内存也称为固定内存或者不可分页内存，它有一个重要属性：
  操作系统将不会对这块内存分页并交换到磁盘上，从而保证了该内存始终驻留在物理内存中。
• 由于GPU知道主机内存的物理地址，因此可以通过DMA(Direct Memory Access)不需要CPU参与地在GPU和主机之间复制数据。正是因为CPU不参与数据复制，CPU很可能在DMA执行过程中移动可分页的数据，这就可能对DMA操作造成延迟。
  事实上，当使用可分页内存进行复制时，复制操作将执行两遍：第一遍从可分页内存复制到一块"临时的"页锁定内存，然后再从这个页锁定内存复制到GPU上。
• 页锁定内存的使用将会在特定场合提高运行效率。
• 注意：当使用固定内存时，将会失去虚拟内存的所有功能，即系统将会更快地耗尽内存资源，因为没有办法将内存交换到磁盘中。

2. 页锁存内存的分配
通过cudaHostAlloc()分配页锁定的主机内存，通过cudaFreeHost()释放页锁定内存

//cudaHostAlloc()前两个参数是页锁定内存的地址和分配的大小
//最后一个参数的取值范围是一组标志，可以通过这些标志来修改cudaHostAlloc()的行为cudaHostAllocDefault将获得默认行为
cudaHostAlloc((void** )&a, size*sizeof(*a), cudaHostAllocDefault);

//...执行一些复制操作

//释放页锁定内存
cudaFreeHost(a);

3. 页锁定内存的测试

// 代码10.2页锁存主机内存
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

float cuda_malloc_test(int size, bool up)
{
	cudaEvent_t start, stop;
	int *a, *dev_a;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	a = (int *)malloc(size*sizeof(*a));
	cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));

	cudaEventRecord(start, 0);

	for (int i = 0; i < 100; i++)
	{
		if (up)
		{
			cudaMemcpy(dev_a, a, size*sizeof(*dev_a), cudaMemcpyHostToDevice);
		}
		else
		{
			cudaMemcpy(a, dev_a, size*sizeof(*a), cudaMemcpyDeviceToHost);
		}
	}

	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);

	free(a);
	cudaFree(dev_a);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	return elapsedTime;
}



float cuda_host_alloc_test(int size, bool up)
{
	cudaEvent_t start, stop;
	int *a, *dev_a;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	cudaHostAlloc((void **)&a,size*sizeof(*a),cudaHostAllocDefault);
	cudaMalloc((void **)&dev_a, size*sizeof(*dev_a));

	cudaEventRecord(start, 0);

	for (int i = 0; i < 100; i++)
	{
		if (up)
		{
			cudaMemcpy(dev_a, a, size*sizeof(*dev_a), cudaMemcpyHostToDevice);
		}
		else
		{
			cudaMemcpy(a, dev_a, size*sizeof(*a), cudaMemcpyDeviceToHost);
		}
	}

	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);

	cudaFreeHost(a);
	cudaFree(dev_a);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	return elapsedTime;
}


#define SIZE 10*1024*1024
int main()
{
	float elapsedTime;
	float MB = (float)100 * SIZE*sizeof(int) / 1024 / 1024;

	elapsedTime = cuda_malloc_test(SIZE, true);
	printf("Time using cudaMalloc:%3.1f ms\n", elapsedTime);
	printf("\t MB/s during copy up:%3.1f M/s\n", MB / (elapsedTime / 1000));

	elapsedTime = cuda_malloc_test(SIZE, false);
	printf("Time using cudaMalloc:%3.1f ms\n", elapsedTime);
	printf("\t MB/s during copy down:%3.1f M/s\n", MB / (elapsedTime / 1000));

	elapsedTime = cuda_host_alloc_test(SIZE, true);
	printf("Time using cudaHostAlloc:%3.1f ms\n", elapsedTime);
	printf("\t MB/s during copy up:%3.1f M/s\n", MB / (elapsedTime / 1000));

	elapsedTime = cuda_host_alloc_test(SIZE, false);
	printf("Time using cudaHostAlloc:%3.1f ms\n", elapsedTime);
	printf("\t MB/s during copy down:%3.1f M/s\n", MB / (elapsedTime / 1000));

	system("pause");

}

发现使用页锁存内存对性能有明显的提升！

CUDA流

CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。可以在流中添加一些操作，例如核函数的启动、内存复制，以及事件的启动和结束等。将这些操作添加到流的顺序也就是它们的执行顺序。可以将每个流视为GPU上的一个任务，并且这些任务可以并行执行。

注意到，在第六章介绍cuda事件时，cudaEventRecord()一共有两个参数，第一个参数是要记录的事件，第二个参数用于指定插入事件的流。

cudaEvent_t start;
cudaEventCreate(&start);
cudaEventRecord(start,0);

使用单个CUDA流

1. 设备硬件要求
GPU必须支持设备重叠功能才能使用CUDA流。支持设备重叠功能的GPU能够在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。我们最终将使用多个流来实现这种计算一数据传输的重叠，以提高运行效率。

	cudaDeviceProp prop;
	int whichDevice;
	cudaGetDevice(&whichDevice);
	cudaGetDeviceProperties(&prop, whichDevice);
	if (!prop.deviceOverlap)
	{
		printf("Device will not handle overlap,so no speed up from streams\n");
		return 0;
	}

2. 初始化流

	//初始化流
	cudaStream_t stream;
	cudaStreamCreate(&stream);//创建一个流

3. 用异步复制取代同步复制
cudaMencpy()将以同步的方式执行，当函数返回时，复制操作就已经完成。
而cudaMemcpyAsync()是CUDA提供的异步复制函数。在调用cudaMemcpyAsync()时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过参数stream来指定的。当函数返回时，我们无法确定复制操作是否已经启动，更无法保证它是否已经结束。任何传递给cudaMemcpyAsync()的主机内存指针都必须是通过cudaHostAlloc()分配的也锁定内存，即只能以异步方式对页锁定内存进行复制操作。

		cudaMemcpyAsync(dev_a, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream);
		cudaMemcpyAsync(dev_b, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream);

4. 用异步核函数调用取代同步调用

kernal << <N / 256, 256, 0, stream >> >(dev_a, dev_b, dev_c);

5.主机与GPU同步
主机在继续执行之前，要首先等待GPU执行完成。此时，可以调用cudaStreamSynchronize()并指定想要等待的流
6. 销毁对GPU操作进行排队的流
在退出应用程序之前，要调用cudaStreamDestroy()销毁对GPU操作进行排队的流。

cudaStreamDestroy(stream);

7. 完整代码
示例代码假设在这样的情形下运行：GPU内存远远少于主机内存，由于整个缓冲区无法一次性填充到GPU，因此需要分块进行填充，分块计算，分块取回。

// 代码10.4使用单个cuda流
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

#define N 1024*1024
#define FULL_DATA_SIZE 20*N

__global__ void kernal(int *a, int *b, int *c)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if (idx < N)
	{
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}


int main()
{
	cudaDeviceProp prop;
	int whichDevice;
	cudaGetDevice(&whichDevice);
	cudaGetDeviceProperties(&prop, whichDevice);
	if (!prop.deviceOverlap)
	{
		printf("Device will not handle overlap,so no speed up from streams\n");
		return 0;
	}

	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));
	//分配由流使用的页锁定内存
	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);
		//执行核函数
		kernal << <N / 256, 256, 0, stream >> >(dev_a, dev_b, dev_c);

		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost, stream);
	}

	//将GPU与主机同步
	cudaStreamSynchronize(stream);

	//停止计时器，收集性能数据
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("Time taken:%3.1f ms\n", elapsedTime);

	//释放流和内存
	cudaFreeHost(host_a);
	cudaFreeHost(host_b);
	cudaFreeHost(host_c);
	cudaFree(dev_a);
	cudaFree(dev_b);
	cudaFree(dev_c);
	cudaStreamDestroy(stream);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	system("pause");
	return 0;
}

使用多个CUDA流

我们用两个流**“分块”计算并且实现内存复制与核函数的重叠**。

在一些新的NVIDIA GPU中同时支持核函数和两次内存复制操作，一次是从主机到设备，另一次是从设备到主机。这样能使应用程序的整体性能获得再一步提升。

// 代码10.6使用多个cuda流
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

#define N 1024*1024
#define FULL_DATA_SIZE 20*N

__global__ void kernal(int *a, int *b, int *c)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if (idx < N)
	{
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}


int main()
{
	cudaDeviceProp prop;
	int whichDevice;
	cudaGetDevice(&whichDevice);
	cudaGetDeviceProperties(&prop, whichDevice);
	if (!prop.deviceOverlap)
	{
		printf("Device will not handle overlap,so no speed up from streams\n");
		return 0;
	}

	cudaEvent_t start, stop;
	float elapsedTime;

	//启动计时器
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);

	//初始化流
	cudaStream_t stream0,stream1;
	cudaStreamCreate(&stream0);
	cudaStreamCreate(&stream1);

	int *host_a, *host_b, *host_c;
	int *dev_a0, *dev_b0, *dev_c0;
	int *dev_a1, *dev_b1, *dev_c1;
	//在GPU上分配内存
	cudaMalloc((void **)&dev_a0, N*sizeof(int));
	cudaMalloc((void **)&dev_b0, N*sizeof(int));
	cudaMalloc((void **)&dev_c0, N*sizeof(int));
	cudaMalloc((void **)&dev_a1, N*sizeof(int));
	cudaMalloc((void **)&dev_b1, N*sizeof(int));
	cudaMalloc((void **)&dev_c1, N*sizeof(int));
	//分配由流使用的页锁定内存
	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 += 2*N)
	{
		//第一个流
		cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		//执行核函数
		kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);


		//第二个流
		cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		//执行核函数
		kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
	}

	//将GPU与主机同步
	cudaStreamSynchronize(stream0);
	cudaStreamSynchronize(stream1);

	//停止计时器，收集性能数据
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("Time taken:%3.1f ms\n", elapsedTime);

	//释放流和内存
	cudaFreeHost(host_a);
	cudaFreeHost(host_b);
	cudaFreeHost(host_c);
	cudaFree(dev_a0);
	cudaFree(dev_b0);
	cudaFree(dev_c0);
	cudaFree(dev_a1);
	cudaFree(dev_b1);
	cudaFree(dev_c1);
	cudaStreamDestroy(stream0);
	cudaStreamDestroy(stream1);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	system("pause");
	return 0;
}

实践表明，使用两个流确实比使用一个流效率更高。

GPU工作调度机制

程序员编程时可以将流视为有序的操作序列，其中既包含内存复制操作，又包含核函数调用。然而在硬件中并没有流的概念，而是包含一个或多个引擎来执行内存复制操作，以及一个引擎来执行核函数。这些引擎彼此独立地对操作进行排队。
将上一节两个流映射到GPU引擎：

注意到，在两个流的示例中，流队列的加入顺序为：

		//第一个流
		cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		//执行核函数
		kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);


		//第二个流
		cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		//执行核函数
		kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);

应用程序首先将第0个流的所有操作放入队列，然后是第1个流的所有操作。下图说明了流内部的依赖性，其中从复制操作到核函数的箭头表示：复制操作要等核函数执行完成之后才能开始。

在硬件上执行的时间线如下：

由于第0个流中将c复制回主机的操作要等待核函数执行完成，因此第一个流中将a和b复制到GPU的操作虽然是完全独立的，但却被阻塞了。

高效地使用多个CUDA流

如果同时调度某个流的所有操作(即按照流深度优先的方式调用)，那么很容易在无意中阻塞另一个流的复制操作或者核函数执行。要解决这个问题，在将操作放入流的队列时应该采用宽度优先方式，而非深度优先方式。
流的深度优先是指：先添加第0个流的所有操作(a复制，b复制，核函数，c复制)，再添加第1个流的所有操作。
流的广度优先是指：将a复制添加到第0个流，将a复制添加到第1个流；将b复制添加到第0个流，将b复制添加到第1个流；将核函数添加到第0个流，将核函数添加到第1个流；将c复制添加到第0个流，将c复制添加到第1个流。
代码修改为：

		cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		//执行核函数
		kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
		kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
		cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);

新的硬件上执行时间线如下：

完整代码：

// 代码10.7高效地使用多个cuda流
//时间：2019.07.30
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

#define N 1024*1024
#define FULL_DATA_SIZE 20*N

__global__ void kernal(int *a, int *b, int *c)
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if (idx < N)
	{
		int idx1 = (idx + 1) % 256;
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		c[idx] = (as + bs) / 2;
	}
}


int main()
{
	cudaDeviceProp prop;
	int whichDevice;
	cudaGetDevice(&whichDevice);
	cudaGetDeviceProperties(&prop, whichDevice);
	if (!prop.deviceOverlap)
	{
		printf("Device will not handle overlap,so no speed up from streams\n");
		return 0;
	}

	cudaEvent_t start, stop;
	float elapsedTime;

	//启动计时器
	cudaEventCreate(&start);
	cudaEventCreate(&stop);
	cudaEventRecord(start, 0);

	//初始化流
	cudaStream_t stream0, stream1;
	cudaStreamCreate(&stream0);
	cudaStreamCreate(&stream1);

	int *host_a, *host_b, *host_c;
	int *dev_a0, *dev_b0, *dev_c0;
	int *dev_a1, *dev_b1, *dev_c1;
	//在GPU上分配内存
	cudaMalloc((void **)&dev_a0, N*sizeof(int));
	cudaMalloc((void **)&dev_b0, N*sizeof(int));
	cudaMalloc((void **)&dev_c0, N*sizeof(int));
	cudaMalloc((void **)&dev_a1, N*sizeof(int));
	cudaMalloc((void **)&dev_b1, N*sizeof(int));
	cudaMalloc((void **)&dev_c1, N*sizeof(int));
	//分配由流使用的页锁定内存
	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 += 2 * N)
	{
		
		cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		cudaMemcpyAsync(dev_b0, host_b + i, N*sizeof(int), cudaMemcpyHostToDevice, stream0);
		cudaMemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int), cudaMemcpyHostToDevice, stream1);
		//执行核函数
		kernal << <N / 256, 256, 0, stream0 >> >(dev_a0, dev_b0, dev_c0);
		kernal << <N / 256, 256, 0, stream1 >> >(dev_a1, dev_b1, dev_c1);
		//将数据从设备复制到锁定内存
		cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int), cudaMemcpyDeviceToHost, stream0);
		cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int), cudaMemcpyDeviceToHost, stream1);
		
	}

	//将GPU与主机同步
	cudaStreamSynchronize(stream0);
	cudaStreamSynchronize(stream1);

	//停止计时器，收集性能数据
	cudaEventRecord(stop, 0);
	cudaEventSynchronize(stop);
	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("Time taken:%3.1f ms\n", elapsedTime);

	//释放流和内存
	cudaFreeHost(host_a);
	cudaFreeHost(host_b);
	cudaFreeHost(host_c);
	cudaFree(dev_a0);
	cudaFree(dev_b0);
	cudaFree(dev_c0);
	cudaFree(dev_a1);
	cudaFree(dev_b1);
	cudaFree(dev_c1);
	cudaStreamDestroy(stream0);
	cudaStreamDestroy(stream1);
	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	system("pause");
	return 0;
}

实际代码测试注意到，宽度优先的效率不及深度优先，个人猜想这可能是新版本的CUDA对其做了相关优化。