CUDA系列-Kernel Launch-8

本章主要追踪一下kernel launch的流程，会不断完善。

---

kernel launch

先抛出一个问题，如果在一个循环中不断的发送kernel（kernel 内部while死循环），会是什么结果。

// kernel 函数
__global__ void kernel(float *a, int n) {int id = threadIdx.x + blockIdx.x * blockDim.x;while(1) {//a[id] = sqrt(a[id] + 1);//这句注释掉对结果没有影响}
}// 持续不断的把kernelfun送入某一个具体stream
int main() {
//1. 声明变量（略）//2. 设置cudaLimitDevRuntimePendingLaunchCount为128/1000等cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 128)；//3. 创建stream
StreamCreate(&stream);//4. launch kernel ctrl+C 退出while (1) {
//grid_dim, block_dim一次性占满所有资源或者<<<1,1>>>kernel<<<grid_dim, block_dim, 0, stream>>>(buffer, size);}
...
//5. 销毁资源sync();StreamDestroy(&stream1);
}上面3，4，5可以改为多线程，一个线程一个stream.
其中还有一个简单的办法，首先在stream中发射一个阻塞的hostfun，然后发送空kernel也能计算到其大小，参考部分有相关代码

结果：

持续的发送一个小kernel到1个stream中，在1022次kernal launch 后，host出现block。3个stream中现象和1个stream一样，也是在1022次后被阻塞住。

详细参数如下(无论cudaLimitDevRuntimePendingLaunchCount设置为多少，下面结果没有变化)

index	griddim	blockdim	strem	result
1	1	1	1	从1022次开始阻塞
2	1	1	3	从1022次开始阻塞
3	1728	128	1	从1022次开始阻塞
4	1728	128	3	从1022次开始阻塞
6	1728	128	8	从1022次开始阻塞
7	1/1728	1/128	12	约~743次开始阻塞
8	1/1728	1/128	16	约~550次开始阻塞
9	1/1728	1/128	48	约~224次开始阻塞
10	1/1728	1/128	128	约~12-33次开始阻塞

1. cudaLimitDevRuntimePendingLaunchCount 的设置对结果没有影响，上面表格中，cudaLimitDevRuntimePendingLaunchCount无论设置为128，256，1000等，最后结果都是一样的。因为它是CUDA Dynamic Parallelism 嵌套launch的一个控制参数（后面会有证明）。

2）grid_dim, block_dim大小对结果也没有影响，因为此刻限制issue item的数量的是cuda runtimes 中的stream queue（实际上该queue 被map到另外一个叫做channel的对象上），对于正在执行或者pending状态的item，都会在queue中占用资源，直到其完成。上面都说明限制在某个stream上launch kernel的瓶颈点是一个host端侧的资源，所以排除cudaLimitDevRuntimePendingLaunchCount，因为它是一个device侧的资源控制量。

3）stream用户可以创建很多个，但是stream queue最后都是被map到channel上，channel的数量是有限的，并且channel又分为很多类型,不同类型其capacity也不一样，其中看实验结果，其中CU_CHANNEL_COMPUTE类型的只有8个，并且每个channel容纳1022个kernel func。

4. 如果stream数量少于channel的数量，那么每个stream对应一个channel，如果stream的数量大于channel，distributes work evenly across all channels。

cuiLaunchstreamBeginPushWithFlagsstreamBeginPushWithDescstreamBeginPushOnChannelWithFlagschannelBeginPushInternalchannelBeginPushInternal_UnderLockchannelMustAdvance_UnderlockchannelMustAdvance_WaitForGPFIFOchannelCanAdvanceGPFIFO（在这里判断pushbuf和fifo entry）gpfifoHasPushbufferSpacegpfifoAdvanceGpuGetpushbufferHasSpace

整个调用链是这样的：
当向stream中下发kerneL的时候，stream会找到一个channel，该channel中有一个gpfifo的queue，其内部有一个ring_buffer(4M)，另外还维护着一个semaphore queue(Max：1024)，我们下发的每个kernel都会写道对应的ring_buffer中，并且每个kernel对专门对应一个semaphore entry放在semaphore queue中，当GPU开始执行kernel的时候，ring_buffer中的kernel data是不能删除的，只有当kernel执行完后，GPU 发送一个semaphore signal给CPU，CPU收到后会找对应的semaphore entry，让其释放资源，因为fifo有顺序要求，所以如果前面的kernel没有执行完，后面的kernel执行完，那么依然会block.也就是说只要第一个kernel在执行，即使后面全部是empty kernel 那么依然会block.

(继续完善)