DAY71：阅读Device-side Launch from PTX
DAY71：阅读Device-side Launch from PTX
我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第71天，我们正在讲解CUDA 动态并行，希望在接下来的30天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。
关注微信公众号，查看历史信息，可以看到之前的阅读
本文共计640字，阅读时间20分钟
前情回顾：
DAY61：关于Cooperative Groups
DAY65:阅读Device-Side Kernel Launch
DAY67：阅读阅读Events、Synchronization和Device Management
DAY69：阅读API Errors and Launch Failures
D.3.2. Device-side Launch from PTX
This section is for the programming language and compiler implementers who target Parallel Thread Execution (PTX) and plan to support Dynamic Parallelism in their language. It provides the low-level details related to supporting kernel launches at the PTX level.
D.3.2.1. Kernel Launch APIs
Device-side kernel launches can be implemented using the following two APIs accessible from PTX: cudaLaunchDevice() and cudaGetParameterBuffer(). cudaLaunchDevice() launches the specified kernel with the parameter buffer that is obtained by calling cudaGetParameterBuffer() and filled with the parameters to the launched kernel. The parameter buffer can be NULL, i.e., no need to invoke cudaGetParameterBuffer(), if the launched kernel does not take any parameters.
D.3.2.1.1. cudaLaunchDevice
At the PTX level, cudaLaunchDevice()needs to be declared in one of the two forms shown below before it is used.
The CUDA-level declaration below is mapped to one of the aforementioned PTX-level declarations and is found in the system header file cuda_device_runtime_api.h. The function is defined in the cudadevrt system library, which must be linked with a program in order to use device-side kernel launch functionality.
The first parameter is a pointer to the kernel to be is launched, and the second parameter is the parameter buffer that holds the actual parameters to the launched kernel. The layout of the parameter buffer is explained in Parameter Buffer Layout, below. Other parameters specify the launch configuration, i.e., as grid dimension, block dimension, shared memory size, and the stream associated with the launch (please refer to Execution Configuration for the detailed description of launch configuration.
D.3.2.1.2. cudaGetParameterBuffer
cudaGetParameterBuffer() needs to be declared at the PTX level before it's used. The PTX-level declaration must be in one of the two forms given below, depending on address size:
The following CUDA-level declaration of cudaGetParameterBuffer() is mapped to the aforementioned PTX-level declaration:
The first parameter specifies the alignment requirement of the parameter buffer and the second parameter the size requirement in bytes. In the current implementation, the parameter buffer returned by cudaGetParameterBuffer() is always guaranteed to be 64- byte aligned, and the alignment requirement parameter is ignored. However, it is recommended to pass the correct alignment requirement value - which is the largest alignment of any parameter to be placed in the parameter buffer - to cudaGetParameterBuffer() to ensure portability in the future.
D.3.2.2. Parameter Buffer Layout
Parameter reordering in the parameter buffer is prohibited, and each individual parameter placed in the parameter buffer is required to be aligned. That is, each parameter must be placed at the nth byte in the parameter buffer, where n is the smallest multiple of the parameter size that is greater than the offset of the last byte taken by the preceding parameter. The maximum size of the parameter buffer is 4KB.
For a more detailed description of PTX code generated by the CUDA compiler, please refer to the PTX-3.5 specification.
本文备注/经验分享：
今天的内容主要说, 涉及到PTX的场合, 结合动态并行的处理.
我们都知道, CUDA做为一种私有的开发接口, 却实际上比公开的API具有更广泛的应用.这种广泛的应用在于私有能更加贴近自家的卡本身,而无需像是公开的API(OpenCL)那样, 为了适合不同的硬件(GPU, CPU, 或者手机GPU以及加速器), 必须求同存异,掩盖掉自家的产品特色.N卡的这种自家的特色发挥, 很多是依靠PTX实现的.一般的用户可以通用的适用CUDA. 更贴近硬件本身的能力, 则可以使用PTX.例如carry bit(整数加法)的时候, 可以很方便的PTX来处理长进位链.PTX这里也不例外,在较大篇幅的使用了PTX的优化程度较深的代码,临时从PTX状态切换到CUDA C状态, 就为了使用后者的<<<>>>语法, 然后再切换回PTX,实际上将会很折腾人, 而本章节中的讲述的做法, 以及, 较多的范例代码,可以让你保持在PTX状态, 直接利用动态并行能力启动kernel. 本章节的内容实际上很少, 主要都是范例的代码, 我稍微说一下要点: 参数的设定方式不同.CUDA中实际上长期总有3种参数的设定方式的,最简单的是<<<>>>()语法, 这个语法能自动推导参数, 放入合适的位置,然后同时启动kernel.另外一种方式则是driver api, cuda runtime api的隐藏内容, 以及, 9.2新加的cooperative group的启动方式.最后一个新的能在特定的情况下, 方便完成全局同步的新内容, 我们之前越过了.而这三种方式, 都是通过直接将每个参数的首地址罗列好, 然后一并传输完成的.这是一种常见的除了<<<>>>的传输方式, 很多用户迟早会需要使用它,例如使用NVENC的用户. 因为NV只提供了在cuda driver api + nvenc的时候, 的使用说法.这样没有直接在runtime api下使用nvenc的用户, 必须学习这种非<<<>>>的参数设定方式，不过这个是后话, 我们就先不详细说了. 实际上, OpenCL也是这种方式, 只是分步进行的, 所以有机会我们会讲一下这个方式了(本手册没说)，而本章节的ptx调用动态并行时候, 的参数设定方式, 和其他的都不同，PTX的这种方式更加复杂. 比另外两种。具体说是你需要在特定的一个缓冲区中, 在特定的位置上, 讲参数放置在上面.这里面的主要容易出错的点在于放置的位置. 这里实际重点问题只有一个, 每个参数必须放置到自己的大小的整数倍.例如一个16B的参数(double2类似)，必须放置到16B的, 在参数缓冲区中的边界.(类似的, 32B的参数需要在32B边界. 这点和CUDA的对齐类似, 但CUDA的对齐最多到16B就了了,这个可能会更多) 这样导致了2个问题: (1)不正确的参数位置放置, 例如你连续放置了多个参数, 会挂. (2)参数空间较小, 不恰当的参数放置, 会导致你连4KB都用不了. 例如这种: __global__ void your_kernel(uint8_t a, double b) 那么实际上在参数缓冲区中的布局是: 第1个字节: a 第8个字节: b 中间的7个字节都是空位. 你不能使用它们. 这个特性有的时候和本章节说的不同,在PTX中可能不容易看出来, 但是直接编译成目标代码(例如你的exe)， 然后直接用cuobjdump看的时候, 会看到类似这种: st [R0 + 0], ... //保存参数a st [R0 + 8], ... //保存参数b 注意中间的7个字节没有使用. 不小心处理这点, 会导致PTX启动的kernel无法正常运行. 实际上, 手册是这样说的: Each parameter must be placed at the nth byte in the parameter buffer, where n is the smallest multiple of the parameter size that is greater than the offset of the last byte taken by the preceding parameter. 这是关于布局的原始说法.注意手册这里说的比较晦涩.正常的解说是刚才的做法. 这个做法是安全的. 原文字面: 每个参数必须在(参数缓冲区中)放置到N号字节位置. N这里是, 大于之前的最后一个参数所占据的字节的最小整数, 该最小整数满足是当前参数的大小的倍数.实际上这等于刚才的: 必须对齐到每个参数自己大小那么大的边界, 同时参数不可以重叠(不能一个参数覆盖另外一个参数---这显然的). 这是唯一的坑. 此外, 请主要, 这里的launch device函数, 所启动的kernel的参数缓冲区的参数位置(有点绕口, 就是将参数缓冲区本身做为一个参数),和其他的刚才说过的, 另外的driver api / runtime api 没说的启动方式 / 以及cooperative group启动的方式, 有参数位置的略微不同. 这点的时候需要注意了.知道了如何在参数缓冲区中放入参数, 然后即用launch device来启动kernel了.这样就完成了全然不退出PTX的情况下,例如很多时候, 在CUDA C里面的嵌入PTX,一旦要退出, 临时一下, 像是.pred这种数据类型, 如何有效的临时保存起来, 是个问题(CUDA C没又直接的1-bit数据类型),而通过本章节的PTX就地动态并行启动kernel的方式, 不仅仅减少了反复进入离开PTX状态的烦恼.还为保存一些不方面的数据类型提供了可能(你先在不需要保存它了)
也为纯PTX代码, 完全不使用CUDA C的代码生成(例如不使用NVVM IR而是选择PTX, 做为二次代码生成选择)的平台,提供了可能.还是很有意义的. 请注意! 本章节的一些代码行较多,建议认真看一下.以及, 如同之前说过的, 和本章节提供的, 一些CUDA C调用动态并行时候的另外的一些手工操作函数(缓冲区设定+kernel启动),建议用户在切换到PTX里调用动态并行之前(本章节内容), 先尝试在CUDA C里进行手工启动(不使用<<<>>>语法),确定自己明确了具体参数在缓冲区中的放置位置后, 能从CUDA C中启动kernel了,再切换到PTX.这样有时候可以少走一些弯路.
有不明白的地方，请在本文后留言
或者在我们的技术论坛bbs.gpuworld.cn上发帖