Cuda Stream流分析
Stream
一般来说，cuda c并行性表现在下面两个层面上：
- Kernel level
- Grid level
Stream和event简介
Cuda stream是指一堆异步的cuda操作，他们按照host代码调用的顺序执行在device上。
典型的cuda编程模式我们已经熟知了：
- 将输入数据从host转移到device
- 在device上执行kernel
- 将结果从device上转移回host
Cuda Streams
所有的cuda操作（包括kernel执行和数据传输）都显式或隐式的运行在stream中，stream也就两种类型，分别是：
- 隐式声明stream（NULL stream）
- 显示声明stream（non-NULL stream）
异步且基于stream的kernel执行和数据传输能够实现以下几种类型的并行：
- Host运算操作和device运算操作并行
- Host运算操作和host到device的数据传输并行
- Host到device的数据传输和device运算操作并行
- Device内的运算并行
下面代码是常见的使用形式，默认使用NULL stream:
cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaMemcpy(..., cudaMemcpyDeviceToHost);
下面版本是异步版本的cudaMemcpy：
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,cudaMemcpyKind kind, cudaStream_t stream = 0);
上面代码使用了默认stream，如果要声明一个新的stream则使用下面的API定义一个：
cudaError_t cudaStreamCreate(cudaStream_t* pStream);
Pinned memory的分配如下：
cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);
在执行kernel时要想设置stream的话，只要加一个stream参数就好：
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);
// 非默认的stream声明
cudaStream_t stream;
// 初始化
cudaStreamCreate(&stream);
// 资源释放
cudaError_t cudaStreamDestroy(cudaStream_t stream);
所有stram的执行都是异步的，需要一些API在必要的时候做同步操作：
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);
看一下代码片段：
for (int i = 0; i < nStreams; i++) {
int offset = i * bytesPerStream;
cudaMemcpyAsync(&d_a[offset], &a[offset], bytePerStream, streams[i]);
kernel<<grid, block, 0, streams[i]>>(&d_a[offset]);
cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream, streams[i]);
}
for (int i = 0; i < nStreams; i++) {
cudaStreamSynchronize(streams[i]);
}
使用了三个stream，数据传输和kernel运算都被分配在了这几个并发的stream中。
kernel数目是依赖于device本身的，Fermi支持16路并行，Kepler是32。并行数是受限于shared memory，寄存器等device资源。
Stream Scheduling
C和P以及R和X是可以并行的，因为他们在不同的stream中，但是ABC，PQR以及XYZ却不行，比如，在B没完成之前，C和P都在等待。
Hyper-Q
Hyper-Q的技术， Kepler上出现了32个工作队列。实现了TPC上可以同时运行compute和graphic的应用。当然，如果超过32个stream被创建了，依然会出现伪依赖的情况。
Stream Priorities
对于CC3.5及以上版本，stream可以有优先级的属性：
cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority);
该函数创建一个stream，赋予priority的优先级，高优先级的grid可以抢占低优先级执行。
cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);
leastPriority是下限，gretestPriority是上限。数值较小则拥有较高优先级。如
Cuda Events
Event是stream用来标记strean执行过程的某个特定的点。其主要用途是：
- 同步stream执行
- 操控device运行步调
Creation and Destruction
// 声明
cudaEvent_t event;
// 创建
cudaError_t cudaEventCreate(cudaEvent_t* event);
// 销毁
cudaError_t cudaEventDestroy(cudaEvent_t event);
streeam的释放，在操作完成后自动释放资源。
Recording Events and Mesuring Elapsed Time
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);
等待event会阻塞调用host线程，同步操作调用下面的函数：
cudaError_t cudaEventSynchronize(cudaEvent_t event);
类似于cudaStreamSynchronize，等待event而不是整个stream执行完毕。使用API来测试event是否完成，该函数不会阻塞host：
cudaError_t cudaEventQuery(cudaEvent_t event);
该函数类似cudaStreamQuery。此外，还有专门的API可以度量两个event之间的时间间隔：
cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop);
返回start和stop之间的时间间隔，单位是毫秒。Start和stop不必关联到同一个stream上。
下面代码简单展示了如何使用event来度量时间：
// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// record start event on the default stream
cudaEventRecord(start);
// execute kernel
kernel<<<grid, block>>>(arguments);
// record stop event on the default stream
cudaEventRecord(stop);
// wait until the stop event completes
cudaEventSynchronize(stop);
// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);
// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);
Stream Synchronization
由于所有non-default stream的操作对于host来说都是非阻塞的，就需要相应的同步操作。
从host的角度来看，cuda操作可以被分为两类：
- Memory相关的操作
- Kernel launch
Kernel launch对于host来说都是异步的，许多memory操作则是同步的，比如cudaMemcpy，cuda runtime也会提供异步函数来执行memory操作。
阻塞和非阻塞stream
使用cudaStreamCreate创建的是阻塞stream，也就是说，该stream中执行的操作会被早先执行的同步stream阻塞。
例如：
kernel_1<<<1, 1, 0, stream_1>>>();
kernel_2<<<1, 1>>>();
kernel_3<<<1, 1, 0, stream_2>>>();
可以通过下面的API配置生成非阻塞stream：
cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);
// flag为以下两种，默认为第一种，非阻塞便是第二种。
cudaStreamDefault: default stream creation flag (blocking)
cudaStreamNonBlocking: asynchronous stream creation flag (non-blocking)
Implicit Synchronization
Cuda有两种类型的host和device之间同步：显式和隐式。已经了解到显式同步API有：
- cudaDeviceSynchronize
- cudaStreamSynchronize
- cudaEventSynchronize
这三个函数由host显式的调用，在device上执行。
许多memory相关的操作都会影响当前device的操作，比如：
- A page-locked host memory allocation
- A device memory allocation
- A device memset
- A memory copy between two addresses on the same device
- A modification to the L1/shared memory confi guration
Explicit Synchronization
从grid level来看显式同步方式，有如下几种：
- Synchronizing the device
- Synchronizing a stream
- Synchronizing an event in a stream
- Synchronizing across streams using an event
可以使用cudaDeviceSynchronize来同步该device上的所有操作。通过使用cudaStreamSynchronize可以使host等待特定stream中的操作全部完成或者使用非阻塞版本的cudaStreamQuery来测试是否完成。
Cuda event可以用来实现更细粒度的阻塞和同步，相关函数为cudaEventSynchronize和cudaEventSynchronize，用法类似stream相关的函数。此外，cudaStreamWaitEvent提供了一种灵活的方式来引入stream之间的依赖关系：
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);
该函数会指定该stream等待特定的event，该event可以关联到相同或者不同的stream，对于不同stream的情况，如下图所示：
Stream2会等待stream1中的event完成后继续执行。
Configurable Events
Event的配置可用下面函数：
cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);
cudaEventDefault
cudaEventBlockingSync
cudaEventDisableTiming
cudaEventInterprocess
