Dou*_*oug 5 cuda warp-scheduler
有关如何选择#blocks和blockSize的讨论很多,但我仍然遗漏了一些东西.我的许多问题都解决了这个问题:CUDA阻止/变形/线程如何映射到CUDA核心？ (为了简化讨论,有足够的perThread和perBlock内存.内存限制在这里不是问题.)
kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal);
1)为了让SM尽可能忙,我应该设置nThreads为倍数warpSize.真正？
2)SM一次只能执行一个内核.这就是SM的所有HWcores只执行kernelA.(不是某些HWcores运行kernelA,而其他运行kernelB.)因此,如果我只有一个线程可以运行,那么我就"浪费"了其他的HWcores.真正？
3)如果warp-scheduler以warpSize(32个线程)为单位发出工作,并且每个SM有32个HWcore,则SM将被充分利用.SM有48个HWcores时会发生什么？当调度程序以32块为单位发布工作时,如何充分利用所有48个核心？(如果上一段是真的,如果调度程序以HWcore大小为单位发布工作会不会更好？)
4)看起来warp-scheduler一次排队2个任务.因此当当前正在执行的内核停止或阻塞时,第二个内核被交换.(目前尚不清楚,但我猜这里的队列深度超过2个内核.)这是正确的吗？
5)如果我的HW的上限为每块512个线程(nThreadsMax),这并不意味着具有512个线程的内核将在一个块上运行得最快.(同样,mem也不是问题.)如果我将512线程内核分布在多个块中,而不仅仅是一个块,那么我很有可能获得更好的性能.该块在一个或多个SM上执行.真正？
5a)我认为越小越好,但是我做得多小也一样重要nBlocks吗？问题是,如何选择nBlocks那个体面的价值？(不一定是最优的.)是否有选择的数学方法nBlocks,或者它只是试验性的.
1)是的.
2)CC 2.0 - 3.0设备可同时执行多达16个网格.每个SM限制为8个块,因此为了达到完全并发,设备必须至少具有2个SM.
3)是的warp调度程序在时间选择并发出warp.忘掉他们无关的CUDA核心概念.为了隐藏延迟,您需要具有高指令级并行性或高占用率.对于CC 1.x,建议> 25%,对于CC> = 2.0,建议> 50%.通常,由于调度程序加倍,CC 3.0需要比2.0设备更高的占用率,但每SM的warp仅增加33%.Nsight VSE问题效率实验是确定您是否有足够的扭曲来隐藏指令和内存延迟的最佳方法.不幸的是,Visual Profiler没有此指标.
4)没有记录warp调度程序算法; 但是,它不考虑线程块发起的网格.对于CC 2.x和3.0设备,CUDA工作分配器将在从下一个网格分配块之前分配网格中的所有块; 但是,编程模型无法保证这一点.
5)为了保持SM忙,你必须有足够的块来填充设备.在此之后,您需要确保有足够的经线以达到合理的占用率.使用大型线程块有利有弊.大螺纹块通常使用较少的指令缓存并且在缓存上具有较小的占用空间; 然而,大型线程块在syncthreads处停滞(SM可能变得效率较低,因为可以选择较少的warp)并且倾向于使指令在类似的执行单元上执行.我建议每个线程块尝试128或256个线程来启动.较大和较小的螺纹块都有充分的理由.5a)使用占用率计算器.挑选太大的线程块大小通常会导致您受到寄存器的限制.挑选太小的线程块大小可能会发现您受限于共享内存或每个SM限制8个块.
让我尝试一一回答你的问题。
\n\n根据NVIDIA Fermi 计算架构白皮书：“SM 以 32 个并行线程为一组进行调度，称为 warp。每个 SM 具有两个 warp 调度程序和两个指令调度单元，允许同时发出和执行两个 warp。Fermi\xe2\x80 \x99s 双 warp 调度器选择两个 warp，并从每个 warp 向一组 16 个核心、16 个加载/存储单元或 4 个 SFU 发出一条指令。因为 warp 独立执行，所以 Fermi\xe2\x80\x99s 调度器不需要检查指令流中的依赖性”。
\n\n此外，NVIDIA Kepler 架构白皮书指出：“Kepler\xe2\x80\x99s 四核扭曲调度程序选择四个扭曲，并且每个扭曲可以在每个周期调度两个独立指令。”
\n\n因此，通过一次调度多个扭曲来使用“多余”核心。
warp 调度程序调度同一内核的 warp，而不是不同内核的 warp。