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