Stream和event简介
Cuda stream是指一堆异步的cuda操作,他们按照host代码调用的顺序执行在device上。Stream维护了这些操作的顺序,并在所有预处理完成后允许这些操作进入工作队列,同时也可以对这些操作进行一些查询操作。这些操作包括host到device的数据传输,launch kernel以及其他的host发起由device执行的动作。这些操作的执行总是异步的,cuda runtime会决定这些操作合适的执行时机。我们则可以使用相应的cuda api来保证所取得结果是在所有操作完成后获得的。同一个stream里的操作有严格的执行顺序,不同的stream则没有此限制。
由于不同stream的操作是异步执行的,就可以利用相互之间的协调来充分发挥资源的利用率。典型的cuda编程模式我们已经熟知了:
将输入数据从host转移到device
在device上执行kernel
将结果从device上转移回host
在许多情况下,花费在执行kernel上的时间要比传输数据多得多,所以很容易想到将cpu和gpu之间的沟通时间隐藏在其他kernel执行过程中,我们可以将数据传输和kernel执行放在不同的stream中来实现此功能。Stream可以用来实现pipeline和双buffer(front-back)渲染。
Cuda API可分为同步和异步两类,同步函数会阻塞host端的线程执行,异步函数会立刻将控制权返还给host从而继续执行之后的动作。异步函数和stream是grid level并行的两个基石。
从软件角度来看,不同stream中的不同操作可以并行执行,但是硬件角度却不一定如此。这依赖于PCIe链接或者每个SM可获得的资源,不同的stream仍然需要等待别的stream来完成执行。下面会简单介绍在不同CC版本下,stream在device上的行为。
Cuda Streams
所有的cuda操作(包括kernel执行和数据传输)都显式或隐式的运行在stream中,stream也就两种类型,分别是:
cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority);
该函数创建一个stream,赋予priority的优先级,高优先级的grid可以抢占低优先级执行。不过优先级属性只对kernel有效,对数据传输无效。此外,如果设置的优先级超出了可设置范围,则会自动设置成最高或者最低。有效可设置范围可用下列函数查询:
cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);
顾名思义,leastPriority是下限,gretestPriority是上限。老规矩,数值较小则拥有较高优先级。如果device不支持优先级设置,则这两个值都返回0。
Cuda Events
Event是stream相关的一个重要概念,其用来标记strean执行过程的某个特定的点。其主要用途是:
同步stream执行
操控device运行步调
Cuda api提供了相关函数来插入event到stream中和查询该event是否完成(或者叫满足条件?)。只有当该event标记的stream位置的所有操作都被执行完毕,该event才算完成。关联到默认stream上的event则对所有的stream有效。
Recording Events and Mesuring Elapsed Time
Events标记了stream执行过程中的一个点,我们就可以检查正在执行的stream中的操作是否到达该点,我们可以把event当成一个操作插入到stream中的众多操作中,当执行到该操作时,所做工作就是设置CPU的一个flag来标记表示完成。下面函数将event关联到指定stream。
// 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操作可以被分为两类:
我们可以使用之前提到过的cudaDeviceSynchronize来同步该device上的所有操作。该函数会导致host等待所有device上的运算或者数据传输操作完成。显而易见,该函数是个heavyweight的函数,我们应该尽量减少这类函数的使用。
通过使用cudaStreamSynchronize可以使host等待特定stream中的操作全部完成或者使用非阻塞版本的cudaStreamQuery来测试是否完成。
Cuda event可以用来实现更细粒度的阻塞和同步,相关函数为cudaEventSynchronize和cudaEventSynchronize,用法类似stream相关的函数。此外,cudaStreamWaitEvent提供了一种灵活的方式来引入stream之间的依赖关系: