20200307

@liwei 2020-03-07T00:00:00

Plan

CUDA 文档阅读 - 2hours
- graphs - 30mins 实际1hour
- events - 30mins 实际10mins
- synchronous calls - 30mins 实际10mins
LeetCode Daily - 30mins 实际1hour
C++ 30mins

Notes

CUDA Runtime

Graphs

图是一种新的CUDA执行流。图当中包含了一系列的操作，例如kernel launch，多个kernel之间的互相等待/同步。整个图是预先定义的，这种执行流，相比较原始的stream有很多的优势，例如：

cpu开销明显降低，因为大多数的任务已经预先完成设置
完整的计算拓扑，可以让cuda进行深入优化

整个执行的过程可以被分为三个过程：定义、实例化、执行。

图的结构

每个操作都是图中的一个node，node之间的依赖关系则是edge。但一个node的前序node节点都完成了，这一node就可以被调度执行。节点的类型包括了：

kernel
cpu function
memory copy
memset
empty node
子图

创建图

创建图的方式有两种，显式调用api或stream捕获。以下面计算图为例子：

/ -> b ->
a d \ -> c ->/

api表示的demo为：

// Create the graph - it starts out empty
cudaGraphCreate(&graph, 0);

// For the purpose of this example, we'll create
// the nodes separately from the dependencies to
// demonstrate that it can be done in two stages.
// Note that dependencies can also be specified 
// at node creation. 
cudaGraphAddKernelNode(&a, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&b, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&c, graph, NULL, 0, &nodeParams);
cudaGraphAddKernelNode(&d, graph, NULL, 0, &nodeParams);

// Now set up dependencies on each node
cudaGraphAddDependencies(graph, &a, &b, 1);     // A->B
cudaGraphAddDependencies(graph, &a, &c, 1);     // A->C
cudaGraphAddDependencies(graph, &b, &d, 1);     // B->D
cudaGraphAddDependencies(graph, &c, &d, 1);     // C->D

stream capture的demo为：

cudaGraph_t graph;

cudaStreamBeginCapture(stream);

kernel_A<<< ..., stream >>>(...);
kernel_B<<< ..., stream >>>(...);
libraryCall(stream);
kernel_C<<< ..., stream >>>(...);

cudaStreamEndCapture(stream, &graph);

stream capture 还通过cudaEventRecord() 和 cudaStreamWaitEvent()创建存在多个stream之间以来的graph。当我们在一个stream里面进行capture，相关联的stream中的任务都会进入capture模式。需要注意的是，begin 和 end 一定需要在一个stream中进行。在多个stream中创建上面的计算图，demo code为：

// stream1 is the origin stream
cudaStreamBeginCapture(stream1);

kernel_A<<< ..., stream1 >>>(...);

// Fork into stream2
cudaEventRecord(event1, stream1);
cudaStreamWaitEvent(stream2, event1);

kernel_B<<< ..., stream1 >>>(...);
kernel_C<<< ..., stream2 >>>(...);

// Join stream2 back to origin stream (stream1)
cudaEventRecord(event2, stream2);
cudaStreamWaitEvent(stream1, event2);

kernel_D<<< ..., stream1 >>>(...);

// End capture in the origin stream
cudaStreamEndCapture(stream1, &graph);

// stream1 and stream2 no longer in capture mode

需要注意的是，处于capture状态的stream，无法进行同步或者查询，更多的限制在使用这一api时，查询文档即可。

使用图

cudaGraph_t这一数据结构不是线程安全的。相同的cudaGraphExec_t无法同时执行。

Events

event机制用于提供进程监控功能，例如计时等。我们可以将一个event插入到stream开头，将一个event插入到stream末尾，以此获得设备运行时间。在我实际编程中，event最常见的两个用途为：

计时
进行stream/kernel之间的同步

Synchronous Calls

当一个同步函数被调用，在device完成task之间，控制权不会交还给host thread。在主机线程执行任何其他CUDA调用之前，可以通过调用带有某些特定标志的cudaSetDeviceFlags()来指定主机线程是yeild, block或者spin。

LeetCode Daily

今天的题目是 59. II. 队列的最大值。历程如下：

题目要求一个队列queue，能返回max_value，时间复杂度始终为O(1)。直觉想法是一个队列queue，每个value都加入queue，并且一个value记录最大值
如果max_value始终记录最大值，并且不断刷新，当最大值被pop之后，如何取第二个max_value，想了五分钟没思路
开始看答案，答案的核心思路其实也很简单：维护一个双端队列deque记录最大值候选值，每次有一个新值value，则将双端队列中小于value的值都出队列，之后再将value加入队尾；这一思路能保障正确的原因是deque中所有小于value的值都不会对max_value有任何影响，因为他们一定不是前一个最大值出deque后的最大值

下次要是这种问题想不明白，不妨模拟一下运算过程总结一下。今天的题目变成了medium，于是花了五分钟没得思路，答案看了20分钟...

明天有空闲继续阅读CUDA文档。