CUDA流

我们已经看到了通过单指令的数据流（Single Instruction Multiple Data，SIMD)的方式进行数据并行，GPU性能取得了巨大的提升。但我们还没有看到任务并行的效果，后者是指多个互相独立的内核函数同时执行。例如，CPU上的一个函数可能正在计算像素的值,而另外一个函数则可能正在从Internet上下载东西。GPU也提供这种能力，但不如CPU灵活。在GPU上是通过使用CUDA流来实现任务并行的，具体怎么做我们将在本小节说明。
CUDA流是GPU上的工作队列，队列里的工作将以特定的顺序执行。这些工作可以包括:内核函数的调用，cudaMemcpy系列传输，以及对CUDA事件的操作。它们添加到队列的顺序将决定它们的执行顺序。每个CUDA流可以被视为单个任务，因此我们可以启动多个流来并行执行多个任务。(译者注:很多情况下，来自多个流中的工作可能同时执行。因此通过同时使用多个流，就有可能在GPU上取得作者说的任务并行的效果。

#include "stdio.h"
#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>

//Defining number of elements in Array
#define N	50000

//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c)
{
	//Getting Thread index of current kernel
	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()
{
	//Defining host arrays
	int *h_a, *h_b, *h_c;
	//Defining device pointers for stream 0
	int *d_a0, *d_b0, *d_c0;
	//Defining device pointers for stream 1
	int *d_a1, *d_b1, *d_c1;
	cudaStream_t stream0, stream1;
	cudaStreamCreate(&stream0);
	cudaStreamCreate(&stream1);
	cudaEvent_t e_start, e_stop;
	cudaEventCreate(&e_start);
	cudaEventCreate(&e_stop);
	cudaEventRecord(e_start, 0);

	cudaHostAlloc((void**)&h_a, N * 2 * sizeof(int), cudaHostAllocDefault);
	cudaHostAlloc((void**)&h_b, N * 2 * sizeof(int), cudaHostAllocDefault);
	cudaHostAlloc((void**)&h_c, N * 2 * sizeof(int), cudaHostAllocDefault);

	// allocate the memory
	cudaMalloc((void**)&d_a0, N * sizeof(int));
	cudaMalloc((void**)&d_b0, N * sizeof(int));
	cudaMalloc((void**)&d_c0, N * sizeof(int));
	cudaMalloc((void**)&d_a1, N * sizeof(int));
	cudaMalloc((void**)&d_b1, N * sizeof(int));
	cudaMalloc((void**)&d_c1, N * sizeof(int));
	//Initializing Arrays
	for (int i = 0; i < N * 2; i++) 
	{
		h_a[i] = 2 * i * i;
		h_b[i] = i;
	}
	
	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);
	//Calling kernels passing device pointers as parameters
	gpuAdd << <512, 512, 0, stream0 >> > (d_a0, d_b0, d_c0);
	gpuAdd << <512, 512, 0, stream1 >> > (d_a1, d_b1, d_c1);
	//Copy result back to host memory from device memory
	cudaMemcpyAsync(h_c , d_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
	cudaMemcpyAsync(h_c + N, d_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
	
	cudaDeviceSynchronize();
	cudaStreamSynchronize(stream0);
	cudaStreamSynchronize(stream1);
	cudaEventRecord(e_stop, 0);
	cudaEventSynchronize(e_stop);
	float elapsedTime;
	cudaEventElapsedTime(&elapsedTime, e_start, e_stop);
	printf("Time to add %d numbers: %3.1f ms\n",2* N, elapsedTime);

	int Correct = 1;
	printf("Vector addition on GPU \n");
	//Printing result on console
	for (int i = 0; i < 2 * N; i++) 
	{
		if ((h_a[i] + h_b[i] != h_c[i]))
		{
			Correct = 0;
		}
	}
	if (Correct == 1)
	{
		printf("GPU has computed Sum Correctly\n");
	}
	else
	{
		printf("There is an Error in GPU Computation\n");
	}

	//Free up memory
	cudaFree(d_a0);
	cudaFree(d_b0);
	cudaFree(d_c0);
	cudaFree(d_a0);
	cudaFree(d_b0);
	cudaFree(d_c0);
	cudaFreeHost(h_a);
	cudaFreeHost(h_b);
	cudaFreeHost(h_c);
	return 0;
}

我们定义了两个cudaStream_t类型的CUDA流对象，stream 0和 stream 1，然后通过cudaStreamCreate函数创建了流。我们也定义了两套相关的主机和设备指针，将分别会在2个流中使用。我们定义并创建了2个CUDA 事件对象，用来进行该程序的性能测量。
CUDA流（中的传输）需要使用页面锁定内存，所以我们这里使用了cudaHostAlloc函数，而不是常规的malloc进行内存分配。在上一节，我们已经了解过页面锁定内存的优点。我们用cudaMalloc分配了两套设备指针。需要注意的是，主机指针容纳了整体的数据，所以它(（的分配大小)是2*N * sizeof(int);而设备指针则因为每次只需要计算一半的数据，所以它(的分配大小)是N * sizeof(int)。我们还用一些用来做加法的随机值初始化了主机数组。
不同于之前调用简单的cudaMemcpy 函数，我们现在改用cudaMemcpyAsync，它用于进行异步传输。它额外可以在函数的最后一个参数中指定一个特定的流，从而在这个特定的流中进行存储器传输操作。当该函数返回的时候，(实际)的存储器传输操作可能尚未完成，甚至尚未开始，所以它叫作异步操作。它只是在特定的流中添加了一个传输工作。具体的我们能看到,stream0中进行了从0到N的数据传输工作，而 stream1中则进行了从N+1到2N的数据传输工作。
需要说明的是，每个流中的工作是串行的，而流和流之间则默认不保证顺序。因此,当我们想让(一个流中的)传输工作和(另外一个流中的)内核计算操作并行执行的话，需要特别地注重流中操作的顺序问题。你不能将所有的传输工作都放到一个流中，而将所有的计算工作都放到另外一个流中，相反，我们在每个流中都放入了(依次的)传输和计算工作。这将会保证每个流中的工作顺序正确，同时在硬件的辅助下，两个流之间的计算和传输工作并行执行。这样如果传输和计算工作都原本正好占据一半的时间，我们能取得将近2倍的性能提升(因为随着系统、卡、问题规模的不同，不一定是精确的时间2倍提升)。
时间线是从上往下展开的。我们能够看到两个传输过程和内核计算是在同一个时间段内执行的。这样会加速你的程序。我们也可以看到，调用cudaMemcpyAsync进行的几次传输操作是异步的。所以，当一个流(中的传输调用)返回的时候，实质的传输操作可能还没有开始。如果我们想用最后一次回传cudaMemcpyAsync操作的结果，必须等待这两个流中的操作完成。
cudaStreamSynchronize()确保了流中的所有操作都结束后才返回控制权给主机。记录结束时刻，这样，基于结束时刻和开始时刻的时间差值，能计算出程序的整体执行时间，并在控制台上输出。在程序的最后，通过使用cudaFree释放掉了设备上分配的显存，并通过cudaFreeHost释放掉了cudaHostAlloc分配的(页面锁定)内存。