CUDA高性能并行计算(11)--CUDA流

CUDA流

CUDA程序的并行层次主要有两个，一个是在核函数内部的并行，另一个是在核函数外部的并行。之前的线程块、计算网格之类的都是内部的并行，而外部的并行主要包括核函数计算与数据传输之间的并行、主机计算与数据传输之间的并行、不同数据传输(方向)之间的并行、核函数计算与主机计算之间的并行以及不同核函数之间并行。

不过核函数外部的并行并不是编程时考虑的重点内容。但是如果需要考虑核函数外的并行，就需要涉及到流的概念了。

流概述

一个CUDA流是指由主机发出的在一个设备中执行的CUDA操作序列。除了主机端发送的流外，设备端也可以发送流，不过此处我们只考虑前者。

一个流的哥哥操作的次序是由主机控制的，按照主机发布的次序执行，但是来自于两个不同的流的操作不一定按照某个次序执行，而是按照并发或交错地执行。

任何CUDA操作都存在于某个流中，要么是默认流，也称空流，要么是明确的非默认流。如果没有显式地指定一个流，那么所有的CUDA操作都是在空流中进行的。

非默认的流，需要在主机端产生与销毁。一个CUDA流可以由以下的CUDA运行时API产生:cudaError_t cudaStreamCreate(cudaStream_t *)，其输入参数是cudaStream_t类型的一个指针。


cudaStream_t stream_1;
cudaStreamCreate(&stream_1); // 创建,这个需要传递流地址
cudaStreamDestroy(stream_1); // 销毁,这个不需要流地址

为了实现不同的CUDA流之间的并发，主机在向某个CUDA流中发布一系列命令后必须马上获得程序的控制权，不用等待CUDA流在GPU中执行完毕。这样就可以通过主机产生多个并行的CUDA流。

可以用以下函数检查一个CUDA流的所有操作是否在设备中执行完。

cudaError_t cudaStreamSynchronize(cudaStream_t stream); 强制阻塞主机，直到所有CUDA流中的操作执行完毕。
cudaError_t cudaStreamQuery(cudaStream_t stream); 检查CUDA流中操作是否执行完毕，但不阻塞主机。

在默认流中重叠主机和设备计算

同一个CUDA流中的所有CUDA操作都是顺序执行的，但是依然可以重叠默认流和主机之间的计算。

由于一个CUDA流内是顺序执行的，而我们前面也说过，主机执行完核函数后会立刻执行下一条指令，而不等待设备。因此如果主机执行完核函数后的下一条指令仍然是CUDA相关的(比如从设备中复制计算结果数据到主机)，就会导致这条指令被阻塞(因为设备会等待核函数结束后再执行这条复制相关的指令)。

而如果主机在执行核函数后执行的指令与设备无关，那么就能实现主机和设备同时计算。这样就让设备和主机之间的计算重叠，这个技巧有时候会有用，前提是设备和主机执行的计算量差不多。

用非默认的CUDA流来重叠多个核函数的执行

要实现多个核函数之间并行必须需要使用多个流。

核函数中流参数

其实一个核函数在尖括号内可以传递4个参数，分别是网格尺寸，线程块尺寸，共享内存大小，流。因此要配置不同的流，就可以按照前述方法先创建出流，然后传递进核函数就可以了。其中共享内存大小可以为0。

xxx_kernel<<<gridDim, blockDim, sharedMem, streamId>>>(...)

重叠多个核函数的例子

下面给出一个简单的例子。


#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdlib.h>
#include <stdio.h>


void __global__ stream_kernel(int streamid)
{
	printf("%d thread in %d block in %d\n", threadIdx.x, blockIdx.x, streamid);
}

int main()
{
	cudaStream_t streams[5];
	for (auto &x : streams)
	{
		cudaStreamCreate(&x); // 创建流
	}

	for (int i = 0; i < 3;i++)
	{
		stream_kernel << <2, 64, 0, streams[i] >> > (i); // 调用核函数
	}

	for (auto &x : streams)
	{
		cudaStreamDestroy(x); // 销毁流
	}
}

不同流之间核函数的执行与数据传递

要实现核函数执行与数据传输的并发，必须让这两个操作处于不同的非默认流中，且数据传输必须使用cudaMemcpy()的异步版本cudaMemcpyAsync()(这个由GPU中DMA来实现的)；另外，由于我们需要将数据传输实现成异步形式，因此需要使用不可分页内存而不能使用常规的可分页内存(也就是默认的分配内存，因为这些内存在传输时需要和主机进行同步，比如缺页，达不到异步的效果)，可以用cudaMallocHost()函数来实现，不可分页内存需要用cudaFreeHost()来进行释放，否则会出现错误。下面来具体介绍一下这些API。

cudaError_t cudaMallocHost(void **ptr, size_t size) : 将*ptr指向的地址分配成不可分页地址。
cudaError_t cudaFreeHost(void *ptr) : 释放掉不可分页地址
cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t size, enum cudaMemcpyKind kind, cudaStream_t stream) : 前面几个参数与cudaMemcpy()一致，最后一个参数是需要复制到的流中。

总结

cudaStreamCreate()和cudaStreamDestroy()可以用于创建和销毁流。
cudaStreamSynchronize()强制阻塞主机，直到所有CUDA流中的操作执行完毕。cudaStreamQuery()，检查CUDA流中操作是否执行完毕，但不阻塞主机。
核函数调用中尖括号内第四个参数是流。
cudaMallocHost()和cudaFreeHost()可以用于分配和销毁不可分页内存。
cudaMemcpyAsync()可以在流内异步传输数据，且数据类型必须为不可分页内存。