DAY71:阅读Device-side Launch from PTX

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第71天,我们正在讲解CUDA 动态并行,希望在接下来的30天里,您可以学习到原汁原味的CUDA,同时能养成英文阅读的习惯。

关注微信公众号,查看历史信息,可以看到之前的阅读

本文共计640字,阅读时间20分钟

前情回顾:

DAY61:关于Cooperative Groups

DAY62:阅读Glossary

DAY63:阅读Execution Environment

DAY64:阅读 Memory Model

DAY65:阅读Device-Side Kernel Launch

DAY66:阅读Streams

DAY67:阅读阅读Events、Synchronization和Device Management

DAY68:阅读 Memory Declarations

DAY69:阅读API Errors and Launch Failures

DAY70:阅读API Reference

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上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-08-29

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏技术博客

Ioc模式和MEF

  分离关注( Separation of Concerns : SOC)是Ioc模式和AOP产生最原始动力,通过功能分解可得到关注点,这些关注可以是 组件Co...

842
来自专栏Java成神之路

分布式_事务_02_2PC框架raincat源码解析

上一节已经将raincat demo工程运行起来了,这一节来分析下raincat的源码

2181
来自专栏ImportSource

JDK10要来了:下一代 Java 有哪些新特性?

JDK 10 目前正在Rampdown Phase One,开发正在努力的修复着bug。 排期 2017/12/14 Rampdown Phase One ...

5918
来自专栏老码农专栏

TodoBackend展示应用以及ActFramework的实现

1305
来自专栏Java后端技术栈

【面试题】2018年最全Java面试通关秘籍第三套!

注:本文是从众多面试者的面试经验中整理而来,其中不少是本人出的一些题目,网络资源众多,如有雷同,纯属巧合!禁止一切形式的碰瓷行为!未经允许禁止一切形式的转载和复...

1291
来自专栏互联网技术栈

Netflix Archaius 分布式配置管理依赖构件

archaius是Netflix公司开源项目之一,基于java的配置管理类库,主要用于多配置存储的动态获取。主要功能是对apache common config...

1482
来自专栏Kubernetes

深度解析Kubernetes Local Persistent Volume(二)

摘要:上一篇博客”深度解析Kubernetes Local Persistent Volume(一)“对local volume的基本原理和注意事项进行了分析,...

1.4K3
来自专栏编舟记

一步步编写SonarQube Plugin

插件确实不好写,因为插件是插入庞大的系统当中工作的,那也就意味着写插件需要具备一定的领域知识,包括系统架构、扩展点、业务共性及差异、API及其业务模型对应、安装...

7003
来自专栏微信公众号:Java团长

即将发布的 JDK 10 有 109 项新特性,你喜欢哪些?

按计划,JDK 10 将于 3 月 20 日正式发布。据前 Oracle 员工 Simon Ritter 的统计,JDK 10 总共包含 109 项新特性。当然...

662
来自专栏菩提树下的杨过

更好用的excel国际化多语言导出

不知道大家在开发中有没有遇到过『excel导出』的需求,反正我最近写了不少这种功能,刚开始利用poi,一行行的手动塞数据,生成excel,而且还有国际化需求,比...

1392

扫码关注云+社区

领取腾讯云代金券