DAY 75:阅读Configuration Options
DAY 75:阅读Configuration Options
我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第75天，我们正在讲解CUDA 动态并行，希望在接下来的25天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。
关注微信公众号，查看历史信息，可以看到之前的阅读
本文共计415字，阅读时间15分钟
前情回顾：
DAY61：关于Cooperative Groups
DAY65:阅读Device-Side Kernel Launch
DAY67：阅读阅读Events、Synchronization和Device Management
DAY69：阅读API Errors and Launch Failures
DAY71：阅读Device-side Launch from PTX
DAY72：阅读Toolkit Support for Dynamic Parallelism
DAY73：阅读Programming Guidelines
D.4.3.1.4. Configuration Options
Resource allocation for the device runtime system software is controlled via the cudaDeviceSetLimit() API from the host program. Limits must be set before any kernel is launched, and may not be changed while the GPU is actively running programs.
The following named limits may be set:
D.4.3.1.5. Memory Allocation and Lifetime
cudaMalloc() and cudaFree() have distinct semantics between the host and device environments. When invoked from the host, cudaMalloc() allocates a new region from unused device memory. When invoked from the device runtime these functions map to device-side malloc() and free(). This implies that within the device environment the total allocatable memory is limited to the device malloc() heap size, which may be smaller than the available unused device memory. Also, it is an error to invoke cudaFree() from the host program on a pointer which was allocated by cudaMalloc() on the device or vice-versa.
D.4.3.1.6. SM Id and Warp Id
Note that in PTX %smid and %warpid are defined as volatile values. The device runtime may reschedule thread blocks onto different SMs in order to more efficiently manage resources. As such, it is unsafe to rely upon %smid or %warpid remaining unchanged across the lifetime of a thread or thread block.
D.4.3.1.7. ECC Errors
No notification of ECC errors is available to code within a CUDA kernel. ECC errors are reported at the host side once the entire launch tree has completed. Any ECC errors which arise during execution of a nested program will either generate an exception or continue execution (depending upon error and configuration).
本文备注/经验分享：
本章节最后说明了动态并行时候的剩余事项.
前者对应了昨日章节里面的, 所说过的嵌套深度, 和同步深度的说法. 并告诉了如何能修改最大同步深度的方式. 注意这里的任何修改(超过默认两层, 这两层包含从Host启动的那一层), 可能会导致显存的使用量增加.过大的保留的同步层数可能会导致原本的程序无法正常工作.第二点则说明了, 如何设定Pending的Launch Count, 这里明确了两点(相比以前的版本的说法)：
(1) 子kernel暂时无法启动, 出于等待状态. 可能是由于卡上的资源不足(例如SM都被父kernel占满了), 或者是因为逻辑执行上的依赖(例如同一个流中的先后关系, 或者多个设备端的流之间的event等待).
(2) 该设定大小的等待启动的kernel总数满了后, 会以慢速的方式继续允许保存更多需要启动的子kernel---这点就和GTC的说法一致了. 请参考昨日的说法看一下.
这两者的设定方式都是通过cudaDeviceSetLimit来进行的，这样, 结合我们的, 动态并行之前的章节, CUDA目前一共有5个限制可以用户设定:
(1)每个线程的堆栈大小(cudaLimitStackSize). 这个一般不会出问题.
(2)设备端的printf缓冲区大小(cudaLimitPrintfFifoSize), 这个少量会出问题. 可以有选择的扩大一些. 特别是很多使用printf进行大量调式输出的人. 论坛经常会出现相关问题.
(3)以及, 设备端的动态分配的显存大小(cudaLimitMallocHeapSize), 这个还是需要注意的. 特别是动态并行的时候, 很多时候需要使用设备端的malloc甚至cudaMalloc来分配显存. 而这个默认的大小很小, 往往需要扩大.
注意第(3)点无论是否使用动态并行都会遭遇到. 论坛上也经常出现这点。
(4)和(5)则是今日章节里面动态并行最大同步深度, 和动态并行最大等待启动的子kernel数量.注意最后这两点只会在动态并行的时候遭遇到.
大致这5点很重要.因为第(3)点是在设备端可以在动态并行和非动态并行的时候都可以使用(计算能力2.0+),
本章节的第二段, 还强调了, 发配和释放时候的方式问题.这点和之前的非动态并行的时候的设备端分配是一致的:设备端的分配不能从Host上释放.Host上的分配也不能从设备端方式.每个上的分配只能由自己的上面的释放函数来释放.原因实际上在动态并行之前的章节我们曾经提到过, 因为这可以看成是2个独立的heap. 每个heap都有自己的对应的堆管理函数, 不能互换.(他们的最小分配粒度和对齐大小也不同的)。
第二段落给出了一个3x3的表格, 穷尽了所有的情况. 可以看一下.注意这里的动态并行时候的cudaMalloc()依然受到cudaLimitMallocHeapSize的限制(上面说的5点之一)。注意动态并行时候的limit都含有cudaLimitDevRuntime字样开头,而普通的另外3个limit都是cudaLimit开头,但本处的MallocHeapSize的限制虽然没有DevRuntime字样, 但依然会影响到你的设备端的cuda runtime的cudaMalloc(), 或者普通malloc()的使用的. 需要稍微注意一下.
注意倒数第二段说明了, SMID(表明当前执行的SM的虚拟ID),和WARPID(用来识别当前的warp id)的两个值,在使用了动态并行后,因为你知道动态并行会有父kernel的block被动态的切出SM来冻结, 再在时机成熟后, 再切换回去SM继续执行.这样原本一个有20个SM的卡, 原始1个block在编号为3的SM上执行,那么通过动态并行调用了子kernel, 并同步等待, 然后继续等待完成后继续执行的时候, 可能已经被切换到了4号SM上执行了.很多算法依赖于SM位置进行的优化:例如你知道某kernel的block最多某卡能1个SM上执行一个,而一共某卡有20个SM,这样可以一共准备20个保存结果的位置(例如20个队列之类的东西), 每个位置通过SM的ID确定和kernels里面的blocks的对应关系,(例如一个10000个blocks的kernel启动) ，这样有时候可以用这种优化, 不使用原子操作来访问每个保存结果的位置. 从而能提升性能. (因为最大可能的并发冲突数量是20, 而你又通过SM ID解决了这20个冲突)，但是根据本章节的提示, 在使用了动态并行后, 那么Block被切换出去和切换出来后, 可能执行位置发生了变化, 自己原本的SM位置可能被其他伙伴blocks占据了, 从而不再安全. 用户需要特别注意这点.
此外, 根据NV近期的暗示, 1个block在执行到一半的时候被切换出去(例如做为一个大的context switch的一部分), 然后重新再切换回来,这种依然哪怕不使用动态并行, 未来也可能不安全的.用户需要注意一定要自己的特定kernel, 在特定的卡上 + 特定的驱动下, 这种行为经过测试没有问题, 才要这样用(大部分自家购买的集群上面的这种行为在目前的卡上一般是安全的，自己的一个独立的kernel + 不会更新的系统环境--集群么. 从来不打补丁的).
以及, 倒数第二段这里的SM和warp id, 都需要你通过PTX来访问,CUDA C里面没有直接导出他们.这段落等于手册是对ptx用户提醒的---但用CUDA C的用户难免总会或多或少的需要使用PTX.当然, 深度学习用户除外, 他们号称是连计算都不需要懂的. 只要会写文本文档即可.
注意最后一段提示了ECC错误的反馈. 这点和之前的普通卡上的访存错误有点类似.都是从host启动的最初的父kernel, 和它的所有子kernel做为一个整体反馈的.但是和普通卡上的访存错误有区别的是:普通卡上的访存错误是用户自己造成的(例如有BUG的代码), 而ECC报告和错误则是硬件环境造成的(例如显存颗粒不稳定, 或者某时某刻太阳黑子爆发之类的导致的) ，注意ECC有些错误能修复, 有些不能了.我使用过的带有ECC的卡, 没有遭遇过ECC报告和错误, 无论是能修复的还是不能的(用nvidia-smi可以看到和重置报告).但是Host上的ECC的内存的报告, 还是见过的.概率较低, 大约一年一台机器24x7的运行, 能遇到1-2次,但这么多年来, 遇到的这些错误都被修复了. 尚未遭遇到不能修复的错误.可以仅供参考. 也欢迎用户提供一下自己的使用期间的遭遇ECC的情况.
有不明白的地方，请在本文后留言
或者在我们的技术论坛bbs.gpuworld.cn上发帖