流

最新推荐文章于 2021-09-25 11:25:48 发布

scut许sir

最新推荐文章于 2021-09-25 11:25:48 发布

阅读量400

点赞数

分类专栏： CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/u010053524/article/details/50747692

版权

CUDA 专栏收录该内容

9 篇文章 1 订阅

订阅专栏

流

默认情况下所有cuda操作(kernel执行，数据传输)都运行载同一个流中。同一个流中的操作串行执行。通过手动创建多个流，可以实现：
主机-设备数据传输和设备计算的重叠
多个不同kernel的执行重叠
```
cudaError_t cudaStreamCreate(cudaStream_t* pStream);
cudaError_t cudaStreamDestroy(cudaStream_t stream);
```

设备和主机之间的数据传输可以是异步的，但必须指定所属流，并且主机端的目标内存时页锁定的

cudaError_t cudaMemcpyAsync(void* dst,
                            const void* src,
                            size_t count
                            cudaMemcpyKind kind
                            cudaStream_t stream = 0);

在调用kernel时可以指定所属流

kernelFunc <<<gridDim,blockDim,sharedMemSize,stream>>> (args)

查询流中的异步操作是否全部完成

cudaError_t cudaStreamQuery(cudaStream_t stream);

阻塞主机端直到流中异步操作完成

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
流分为null stream(默认)和non-null stream，其中non-null stream分为blocking stream和non-blocking stream。blocking stream和null stream之间的kernel执行会互相阻塞，blocking stream之间不会互相阻塞
```
cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);
//指定flag为cudaStreamDefault创建blocking stream
//指定flag为cudaStreamNonBlocking创建non-blocking stream
```
non-null stream中的操作必须是异步的

Event

Event可以理解为流中的一个标记点，当流在设备端执行到达标记点时，主机端对应的event中的flag会被修改

cudaError_t cudaEventCreate(cudaEvent_t* event);
cudaError_t cudaEventDestroy(cudaEvent_t event);

//在流中插入标记点
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

//在主机执行自旋锁直到对应流执行到标记点
cudaError_t cudaEventSynchronize(cudaEvent_t event);

//查询流是否执行到达标记点(不会阻塞主机)
cudaError_t cudaEventQuery(cudaEvent_t event);

Event可以用于不同流之间的同步，具体做法是在某一流中插入Event，在另一流中等待该Event被触发
```
cudaEventRecord(event,streamA);
cudaStreamWaitEvent(streamB,event,0);
```

Event是可配置的

cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);
//可以设置以下标志位
cudaEventDefault
cudaEventBlockingSync
cudaEventDisableTiming
cudaEventInterprocess

费米架构的假依赖

费米架构中，不同流中的操作最后还是被发送到同一工作队列中，队列中的操作必须按先入先出的顺序执行。如果连续两个操作属于不同的流，则后一个操作不必等待前一个操作完成才出列。假设有流A和B，A中有操作A1和A2，B中有B1操作，在主机端调用操作的顺序是A1，A2，B1，则队列的顺序为
```
B1      A2      A1
```
虽然B1不依赖与其他操作，但其在队列中的顺序决定了他必须在A2出列后才能出列，但由于A2与A1在同一流中，即A2依赖与A1，所以在A1执行完成前A2不能出列，导致B1也不能出列，这就是假依赖
开普勒架构的Grid Management Unit会为工作队列中的操作进行重排序和调度，能够有效避免假依赖
开普勒架构的Hyper-Q允许GPU最多维护32(默认为8)个操作队列，当stream数小于32时，cuda会为每一个stream申请单独一个队列，该方案也能够避免假依赖