DAY62：阅读Glossary

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

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

本文共计456字，阅读时间15分钟

D. CUDA Dynamic Parallelism

D.1. Introduction

D.1.1. Overview

Dynamic Parallelism is an extension to the CUDA programming model enabling a CUDA kernel to create and synchronize with new work directly on the GPU. The creation of parallelism dynamically at whichever point in a program that it is needed offers exciting new capabilities.

The ability to create work directly from the GPU can reduce the need to transfer execution control and data between host and device, as launch configuration decisions can now be made at runtime by threads executing on the device. Additionally, data-dependent parallel work can be generated inline within a kernel at run-time, taking advantage of the GPU's hardware schedulers and load balancers dynamically and adapting in response to data-driven decisions or workloads. Algorithms and programming patterns that had previously required modifications to eliminate recursion, irregular loop structure, or other constructs that do not fit a flat, single-level of parallelism may more transparently be expressed.

This document describes the extended capabilities of CUDA which enable Dynamic Parallelism, including the modifications and additions to the CUDA programming model necessary to take advantage of these, as well as guidelines and best practices for exploiting this added capacity.

Dynamic Parallelism is only supported by devices of compute capability 3.5 and higher.

D.1.2. Glossary

Definitions for terms used in this guide.

• Grid
• A Grid is a collection of Threads. Threads in a Grid execute a Kernel Function and are divided into Thread Blocks.
• Thread Block
• A Thread Block is a group of threads which execute on the same multiprocessor (SMX). Threads within a Thread Block have access to shared memory and can be explicitly synchronized.
• Kernel Function
• A Kernel Function is an implicitly parallel subroutine that executes under the CUDA execution and memory model for every Thread in a Grid.
• Host
• The Host refers to the execution environment that initially invoked CUDA. Typically the thread running on a system's CPU processor.
• Parent
• A Parent Thread, Thread Block, or Grid is one that has launched new grid(s), the Child Grid(s). The Parent is not considered completed until all of its launched Child Grids have also completed.
• Child
• A Child thread, block, or grid is one that has been launched by a Parent grid. A Child grid must complete before the Parent Thread, Thread Block, or Grid are considered complete.
• Thread Block Scope
• Objects with Thread Block Scope have the lifetime of a single Thread Block. They only have defined behavior when operated on by Threads in the Thread Block that created the object and are destroyed when the Thread Block that created them is complete.
• Device Runtime
• The Device Runtime refers to the runtime system and APIs available to enable Kernel Functions to use Dynamic Parallelism.

本文备注/经验分享：

今天这章节开始, 我们开始步入了一个重要的特性--动态并行. 传统的CUDA代码, 总是通过GPU上的kernel来进行计算, 而通过Host(CPU)上的代码来控制GPU. 在今天之前我们从手册上看到的所有例子都是这种情况.你可以直接理解成GPU是干重活的员工, 而CPU则是负责调度的领导.很多情况下, 这需要一定的CPU上的性能, 才能较好的调度. 我们之前ASUS机器, 统一都升级到了skylake平台, 更好的CPU性能, 往往会带来更好的GPU表现.而本章节的特性则是从第二代Kepler开始才能具有的(计算能力3.5+), 能让GPU给GPU发布任务的新方式,所谓的动态并行.好在CUDA 9.2已经放弃了计算能力2.X和之前的所有的卡, 这样目前CUDA Toolkit所支持的版本中,只有计算能力3.0的Kepler不被支持了.其他所有的卡(Kepler二代, Maxwell, Pascal, Volta)都可以使用它.因此这将变成一个普遍的特性了. 在当年可是只有K40, K80之类的专业卡, 外加一款特别的家用卡(GT730, 只有2个SM, 专门为调试动态并行而生)才能使用. 回到动态并行本身来. 我们都知道, 被移植到GPU上的代码, 往往需要经过很多改写,这种改写, 一是为了能充分利用GPU的并行特性. 二则是因为GPU不支持常规的kernel递归, CPU上的很多递归算法只能进行改写, 不能直接实现.而动态并行的出现, 因为GPU现在能直接自己给自己启动kernel了, 因此一些传统的算法可以更好的映射到GPU上, 扩大了GPU的应用面,也减轻了用户的使用成本.需要说明一点的是, 到目前为止, 所有的GPU上的代码对动态并行的应用, 都可以改成相应的没有动态并行的版本.也就是说, 这个特性并不是必须的.例如之前举例的kernel自身递归, 可以通过笨办法反复从Host上传回数据, 然后根据数据进行分析, 选择性的继续启动kernel.但是这种实现起来很繁琐, 同时性能也容易不好.因此虽然很多特性不总是必须有的(例如之前的warp系列相关函数也是如此), 但是如果有它们你可以活的更好一点.也能充分的发挥卡的能力.但用户如果只需要能用CUDA, 也可以只学习一个最小的特性子集, 就如同使用CPU, 理论上说, 只需要一条能任意访存的加法指令即可(感兴趣的可以参考最小图灵完备的指令集) 但是实际上往往只这样会带来诸多不变, 所以该看到的还是应该看的.在计算能力3.5引入动态并行的时候, 当时同时引入了Hyper-Q, 这是一个很有意思的现象,Hyper-Q带来了32个硬件队列. 就如同AMD在引入了OpenCL版本的动态并行实现的时候(AMD那边叫"设备端kernel启动"), 同样增强了ACE到64个队列一样.目前尚未知道Hyper-Q特性和动态并行的关系(但是它们总是同时出现的), 不清楚前者是否是后者实现的必要条件, 但是在具有动态并行能力的卡上,使用多个kernel同时执行的时候(包括通过动态并行启动的, 也包括从Host上通过多个streams启动的), 往往具有更好的性能.用户在考虑性能是否是因为动态并行而提升的时候, 需要注意这点(很可能性能并非来自动态并行, 而是因为多个队列)..因为本章节是一个介绍性质的章节, 我们继续说一下动态并行带来的主要好处, 或者说应用上的优势,然后再对本章后部分的一些动态并行引入的新词汇, 和涉及到的重点旧词汇做一下解说. 主要的好处有3点: (1)点是, 增加了kernel调用kernel的能力. 刚才说过, 将一些算法改写到GPU上的时候, 能直接具有此能力, 将方便不少. 这里的kernel是指的你的一次__global__函数启动, 它将具有调用另外一个__global__函数的能力,还记得之前我们章节的<<<>>>钻石形状的语法吗? (启动配置语法),之前的章节只能从host上这样用, 现在可以从GPU上了. 需要注意的是, 一些书说的, 计算能力2.0开始就可以开始递归了, 为何我们这里却说动态并行需要3.5+?请注意以前的2.0+开始的递归的确可以, 但这个是__global__调用__device__, 并不能__global__调用__global__,前者只能在1个线程的上下文中进行调用, 而动态并行的这种调用可以直接启动一堆线程(一个grid, grid是blocks的集合, 而block是threads的集合).这大大增加了适用面. 第(2)点好处是则是, 可以直接更好的均衡任务了. 以前曾经有过很多论文,头疼一种叫block间工作量不均衡的问题. 我们都知道, 从之前的<<<>>>的语法章节, 你知道启动的是N个blocks, 每个blocks里面都有M个线程, 或者看到的很多资料中的图片, CUDA的模型是一个2维的方形, 或者规则的3维的格子, 对于一些不规则的问题, 或者说, 具体形状不能用规则的方式映射，往往需要用户的很精心的单独处理，例如说, kernel只有在运行到一半的时候, 才能根据数据, 知道本block的计算量，而host是在运行前就指定kernel的启动配置的,这个时候就很尴尬了.很容易出现, 1个kernel, 具有100个blocks, 其中99个都很快就执行完毕了,然后剩下最后1个block, 用了3倍长的时间才执行完毕.因为一次kernel启动必须所有里面的blocks和线程们都执行结束后才能完成,此时, 这kernel的执行时间被延长了3倍, 就因为某个拖累的block迟迟没有完成，很多文章指出了很多技巧试图解决这个问题, 其中有一些也很有成效(之前说过, 动态并行不是必须的)，例如一些来自arxiv.org上的文章，但是这些都需要用户去额外学习, 去研究它们. 特别是公司的开发, 等于需要额外的研发成本。而有了动态并行后, 可以直接将原本host上的经验(启动kernel)就地在GPU端适用，例如本例中的某个block发现自己的任务过多, 可以直接启动一次子kernel.子kernel的执行不同, 它可以由多个新blocks构成, 重新铺满整个GPU上多个SM(例如GTX1080的20个SM)，而无需像原本的父block那样, 只能慢吞吞的在一个SM(CUDA要求每个block不能跨越多个SM, 之前的章节说过这个问题)上执行.从而有效的提升了速度. 例如从原本的3X时间, 瞬间缩小了到了1/20.这样整体kernel从原本的3X(假设该拖累的block从时刻0开始执行), 变成了只增加0.05X的时间, 有效的提升了速度.这就如同一个公司, 原本大领导给手下的8个员工指派活, 因为总有无法在干活前就能确定的因素的影响,其中7个人半天就干完了, 而剩下的1个人需要3天,此时动态并行入场, 直接这1个人对另外7个人说, 我的活分的太多了, 同事们一起上啊.于是这样直接的就平衡了任务, 用了半天多一天, 而不是3天的时间, 大家就完成了总任务.而无需反复的领导进行情况判断, 反复用各种技巧(例如频繁的沟通下属员工的进度情况)进行指派或者重新安排.减轻了不少麻烦. 这就是动态并行的最经典应用的场景几乎. 此外, 第三点需要说的则是,刚才你看到我们说过, 完全可以不用动态并行也能完成任务的,例如在执行一个图像搜索的时候, 可能会具有2个搜索的kernel,其中第一个kernel进行快速大致范围的搜索,然后第二个后续kernel对第一个kernel的大概结果范围内部的精细定位搜索.这种情况的应用, 虽然不是递归(2个独立的kernel, 前后执行), 但是也可以应用动态并行.考虑我们常规的做法是什么? 第一个kernel先将结果做成JobDescriptor_t之类的东西(一个结构体, 名字随意),里面标注了范围的4个点的位置和其他辅助信息.然后将这个信息保存到global memory,然后继续从host上在同1个stream(保持前后顺序么)启动第二个kernel,启动的规模则是刚才判断出来的有多少个JobDescriptor_t区域等等,然后第二个kernel再继续读取每个descriptor_t,判断任务信息, 然后得到后继续进行精细任务.这个做法是再没有动态并行时候的经典做法.之前的章节中其实还有提到过(原子操作章节)，需要维护一个任务列表信息, 头部还需要用原子操作保护,非常繁琐.这还没完, 将前一个任务的数据保存成显存中的二进制数据, 然后下一个任务在读取, 需要进行某种类似序列化/反序列化的操作,虽说是同一个架构(CUDA)上的应用, 不需要像ASN.1之类的那样麻烦,但是依然够折腾的. 这两点带来了巨大的工作量.而现在动态并行极大的降低了这个压力.直接第一个kernel发现了可疑的位置后(例如有20个), 可以就地要求启动精细kernel处理:

__global__ void kernel1(....) { .... calculate possible x0,x1,y0,y1; if (condition_satisfied) kernel2<<<...>>>(x0, x1, y0, y1, ...); ... }

你看到这里直接就可以在2个kernel间传递参数数据. 不需要考虑维护任务描述符列表之类的东西,也不需要考虑这个信息如何保存和读取.方便的很. 你看到, 虽然有些东西, 没有动态并行也能用(刚才说的传统做法), 但是有了后, 直接简化成了1行.提高了很大的工作效率.降低了公司的运营成本.大致这三点常见的好处. 然后我说一下本章节末尾的一些词汇列表： Grid, 这个是一次kernel启动(用你的<<<>>>语法或者其他方式)得到的blocks的集合.本手册的封面的那个图就是. 非常直观.也经常被翻译成"网格". 但早期这样翻译, 后面的CUDA书往往都保持英文了.这是一个. 然后block和thread这两个词. 在前面的章节又过具体描述, 这里就不说了. 请注意这里强调了1个block只能在1个SM上运行的概念. 这个之前也说过.这点需要额外注意.也是为何刚才的(2)点优势的原因. 注意这里将Multi-processor简称了SMX. 这实际上有点小问题.一般我们都叫它是SM(Stream Multi-processor, 流多处理器). NV后期因为Kepler出现, 改称了SMX, 用来彰显Kepler的伟大(然而. 这代显卡实际上比较悲剧)，Maxwell出现后, 改称了SMM(M=Maxwell)，Pascal出现后, 改称了SMP(P=Pascal)，不过一般的, 总是称为SM(无任何一代的后缀). 这里为何叫SMX, 是因为手册直接抄了以前的版本, 动态并行(DP)是出现在Kepler的概念.当年能运行动态并行代码的, 只有Kepler的SMX.CUDA 9.2的手册这里忘记去掉后缀了.用户知道有这么回事即可.Host这个就不说了. 注意两个新词语. Parent(父) thread/block/grid, 和Child(子) thread/block/grid,这描述的是适用动态并行的时候的启动关系,被启动的叫子. 启动的一方叫父.(可以直接理解成有父才能有子, 这样就容易记住了)。 注意这里还提到的, 父block/grid(kernel的一次启动)必须等待它所启动的子kernels们都结束, 才能结束. 这点需要注意, 后续的章节应当会有更详细的说明.动态并行实际上为了维持易用(和以前的CPU上启动kernel几乎一样的方式), 还是有一些小坑的. 例如这里的同步问题.后续的章节应当会说.用户心里可以先注意一下. 然后还提到了使用动态并行的时候, 创立的stream/event之类的对像只在父kernel的1个block中有效的问题.这个也需要注意.这两点都是和CPU上的启动略微不同的.本章节无其他需要注意的事项了

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖