.la域名做的网站,行业数据可以从哪里获取,中国最近的军事新闻大事,济南城市建设集团网站这里写目录标题 kernel launch 本章主要追踪一下kernel launch的流程#xff0c;会不断完善。 kernel launch
先抛出一个问题#xff0c;如果在一个循环中不断的发送kernel#xff08;kernel 内部while死循环#xff09;#xff0c;会是什么结果。
// kernel 函数
__glo… 这里写目录标题 kernel launch 本章主要追踪一下kernel launch的流程会不断完善。 kernel launch
先抛出一个问题如果在一个循环中不断的发送kernelkernel 内部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 ctrlC 退出while (1) {
//grid_dim, block_dim一次性占满所有资源或者1,1kernelgrid_dim, block_dim, 0, stream(buffer, size);}
...
//5. 销毁资源sync();StreamDestroy(stream1);
}上面345可以改为多线程一个线程一个stream.
其中还有一个简单的办法首先在stream中发射一个阻塞的hostfun然后发送空kernel也能计算到其大小参考部分有相关代码结果
持续的发送一个小kernel到1个stream中在1022次kernal launch 后host出现block。3个stream中现象和1个stream一样也是在1022次后被阻塞住。
详细参数如下(无论cudaLimitDevRuntimePendingLaunchCount设置为多少下面结果没有变化)
indexgriddimblockdimstremresult1111从1022次开始阻塞2113从1022次开始阻塞317281281从1022次开始阻塞417281283从1022次开始阻塞617281288从1022次开始阻塞71/17281/12812约~743次开始阻塞81/17281/12816约~550次开始阻塞91/17281/12848约~224次开始阻塞101/17281/128128约~12-33次开始阻塞
cudaLimitDevRuntimePendingLaunchCount 的设置对结果没有影响上面表格中cudaLimitDevRuntimePendingLaunchCount无论设置为1282561000等最后结果都是一样的。因为它是CUDA Dynamic Parallelism 嵌套launch的一个控制参数后面会有证明。
2grid_dim, block_dim大小对结果也没有影响因为此刻限制issue item的数量的是cuda runtimes 中的stream queue实际上该queue 被map到另外一个叫做channel的对象上对于正在执行或者pending状态的item都会在queue中占用资源直到其完成。上面都说明限制在某个stream上launch kernel的瓶颈点是一个host端侧的资源所以排除cudaLimitDevRuntimePendingLaunchCount因为它是一个device侧的资源控制量。
3stream用户可以创建很多个但是stream queue最后都是被map到channel上channel的数量是有限的并且channel又分为很多类型,不同类型其capacity也不一样其中看实验结果其中CU_CHANNEL_COMPUTE类型的只有8个并且每个channel容纳1022个kernel func。
如果stream数量少于channel的数量那么每个stream对应一个channel如果stream的数量大于channeldistributes work evenly across all channels。
cuiLaunchstreamBeginPushWithFlagsstreamBeginPushWithDescstreamBeginPushOnChannelWithFlagschannelBeginPushInternalchannelBeginPushInternal_UnderLockchannelMustAdvance_UnderlockchannelMustAdvance_WaitForGPFIFOchannelCanAdvanceGPFIFO在这里判断pushbuf和fifo entrygpfifoHasPushbufferSpacegpfifoAdvanceGpuGetpushbufferHasSpace整个调用链是这样的 当向stream中下发kerneL的时候stream会找到一个channel该channel中有一个gpfifo的queue其内部有一个ring_buffer(4M)另外还维护着一个semaphore queue(Max1024)我们下发的每个kernel都会写道对应的ring_buffer中并且每个kernel对专门对应一个semaphore entry放在semaphore queue中当GPU开始执行kernel的时候ring_buffer中的kernel data是不能删除的只有当kernel执行完后GPU 发送一个semaphore signal给CPUCPU收到后会找对应的semaphore entry让其释放资源因为fifo有顺序要求所以如果前面的kernel没有执行完后面的kernel执行完那么依然会block.也就是说只要第一个kernel在执行即使后面全部是empty kernel 那么依然会block.
(继续完善)
