CUDA系列-Kernel Launch-8

泰勒朗斯

已于 2024-06-26 10:18:21 修改

阅读量660

点赞数 25

分类专栏： GPU 文章标签： linux 机器学习 cuda

于 2024-06-19 16:34:13 首次发布

本文链接：https://blog.csdn.net/weixin_43360707/article/details/139806834

版权

GPU 专栏收录该内容

9 篇文章 1 订阅

订阅专栏

这里写目录标题

kernel launch

本章主要追踪一下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次开始阻塞

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

2）grid_dim, block_dim大小对结果也没有影响，因为此刻限制issue item的数量的是cuda runtimes 中的stream对应channel中的gpfifo->entries和gpfifo→Pushbuffer两个变量。

其原理为：

当app向stream中下发kerneL的时候，stream会找到一个CU_CHANNEL_COMPUTE类型的channel（该channel是CPU和GPU沟通通道，默认8个，但是可以通过环境变量CUDA_DEVICE_MAX_CONNECTIONS来修改，最大不超过CU_CHANNEL_MAX_COMPUTE（32）个）。

该channel中有一个gpfifo对象，该对象中有一个pushbuffer(ring_buffer) default 4M + 一个gpfifoEntry数组default 1024个，它们两个一一对应。

当我们向gpfifo中的pushbuffer写入一个kerne（l代码中称之为method）的时候，首先会检查：

a）ring_buffer中的space是否足够；

b）有没有free的gpfifoEntry。

当上述两个条件满足的时候，先在ring_buffer中写入method，然后在对应位置的gpfifoEntry中记录该method的相关信息（trackSemValEnd/trackSemValStart）这些信息用来sync以及记录从GPU返回该kernel完成的信息。

如果上述两个条件不满足，就会busy waiting，这就是我们上述代码中看到的阻塞现象。

因为gpfifoEntry默认总大小是1024个，在实际使用的时候会预留2个不用（具体原因位置，看代码是作为padding），那么如果pushbuffer的space足够的情况下，那么最多能使用的gpfifoEntry就1022个。

这里要注意，因为channel是通过fifo来管理下发的kernel的，所以如果我们第一个kernel（假设while循环）一直执行，那么即使后面的kernel为空kernel，那么也还是会产生阻塞。

目前最新版本ring_buffer和gpfifoEntry数量无法调整，旧版本是有两个宏定义可以调整（CUDA_GPFIFO_ENTRY_COUNT和CUDA_COMPUTE_PUSHBUFFER_SIZE已经被舍弃了），新版本目前不清楚具体哪个参数来设置。看文档提供了CUDA_SCALE_LAUNCH_QUEUES这个环境变量，但是设置了也不起作用。

3）stream和channel的对应关系：

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

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

  c）channel的数量可以根据环境变量CUDA_DEVICE_MAX_CONNECTIONS来修改，其最大为 不能超过CU_CHANNEL_MAX_COMPUTE（32)。

关于Execution环境变量的设置可以参考：1. Introduction — CUDA C Programming Guide (nvidia.com)

其中CUDA_LAUNCH_BLOCKING设置为1后，会依阻塞方式运行kernel，比如第一个kernel执行完成后，第二个kernel才能发送，不然会在CPU端block。默认是0.

CUDA_SCALE_LAUNCH_QUEUES设置了没有起作用。

cuiLaunch
    streamBeginPushWithFlags
        streamBeginPushWithDesc
            streamBeginPushOnChannelWithFlags
                channelBeginPushInternal
                    channelBeginPushInternal_UnderLock
                        channelMustAdvance_Underlock
                              channelMustAdvance_WaitForGPFIFO
                                  channelCanAdvanceGPFIFO（在这里判断pushbuf和fifo entry）
                                       gpfifoHasPushbufferSpace
                                           gpfifoAdvanceGpuGet
                                           pushbufferHasSpace

整个调用链是这样的：
当向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.

(继续完善)