本笔记主要是阅读:谭升的博客的 GPU编程(CUDA)
3_2 理解线程束执行的本质(PartⅠ)
从外表来看,CUDA执行所有的线程,并行的,没有先后次序的;
但实际上硬件资源是有限的,不可能同时执行百万个线程,所以从硬件角度来看,物理层面上执行的也只是线程的一部分,而每次执行的这一部分,就是我们前面提到的线程束。
1.1 warp与block
warp是SM中基本的执行单位。是硬件层面的线程集合。
当一个网格被启动(网格被启动,等价于一个内核被启动。每个内核对应于自己的网格),线程块被分配到某一个SM上以后,将分为多个warp(每个线程束一般是32个线程),在一个线程束中,所有线程按照单指令多线程SIMT的方式执行,每一步执行相同的指令,但是处理的数据为私有的数据。
[外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传(img-1ZMLktvm-1616500678241)(https://face2ai.com/CUDA-F-3-2-理解线程束执行的本质-P1/3_10.png)]
线程块是个逻辑产物。因为在计算机里,内存总是一维线性存在的,所以执行起来也是一维的访问线程块中的线程。需要计算出线程的唯一ID。
当一个线程块中有128个线程的时候,其分配到SM上执行时,会分成4个块:
warp0: thread 0,........thread31
warp1: thread 32,........thread63
warp2: thread 64,........thread95
warp3: thread 96,........thread127
注意:
当编号使用三维编号时,x位于最内层,y位于中层,z位于最外层。
计算出三维对应的线性地址是: t i d = t h r e a d I d x . x + t h r e a d I d x . y × b l o c k D i m . x + t h r e a d I d x . z × b l o c k D i m . x × b l o c k D i m . y tid = threadIdx.x+threadIdx.y\times blockDim.x+threadIdx.z\times blockDim.x \times blockDim.y tid=threadIdx.x+threadIdx.y×blockDim.x+threadIdx.z×blockDim.x×blockDim.y
想象下c语言的数组,如果把上面这句话写成c语言,假设三维数组t保存了所有的线程,那么(threadIdx.x,threadIdx.y,threadIdx.z)表示为**:
t[z][y][x];
**
一个线程块包含多少个线程束:
WarpsPerBlock = ceil ( ThreadsPerBlock warpSize ) \text{WarpsPerBlock}=\text{ceil}\begin{pmatrix}\frac{\text{ThreadsPerBlock}}{\text{warpSize}}\end{pmatrix} WarpsPerBlock=ceil(warpSizeThreadsPerBlock)。注:ceil函数是向正无穷取整的函数,比如 c e i l ( 9 8 ) = 2 ceil(\frac{9}{8})=2 ceil(89)=2
1.2 线程束分化
warp被执行的时候会被分配给相同的指令,处理各自私有的数据。当遇到分支时,warp内的线程有的会执行分支内容,有的不执行。
分支:在CUDA中支持C语言的控制流,比如if…else, for ,while 等。
题外话:当CPU在处理分支判断时,处理器都采用分支预测技术,所以,处理速度会提升很多。即:CPU适合逻辑复杂计算量不大的程序,比如操作系统,控制系统,GPU适合大量计算简单逻辑的任务,所以被用来算数。
如:
if (con){
//do something
}
假设这段代码是核函数的一部分,那么当一个线程束的32个线程执行这段代码的时候,如果其中16个执行if中的代码段,而另外16个执行else中的代码块,同一个线程束中的线程,执行不同的指令,这叫做线程束的分化。
在每个指令周期,warp中的所有线程执行相同的指令,但是线程束又是分化的,所以这似乎是相悖的,但是事实上这两个可以不矛盾。
- 解决矛盾的办法:就是每个线程都执行所有的if和else部分.
- 当一部分con成立的时候,执行if块内的代码,
- 有一部分线程con不成立,那这些线程只能等待其他线程执行完if内的代码,才能执行下面的指令。线程束分化会产生严重的性能下降。条件分支越多,并行性削弱越严重。
注意线程束分化研究的是一个线程束中的线程,不同线程束中的分支互不影响。
减少线程束分化的方法:线程束内的线程是可以被我们控制的,那么我们就把都执行if的线程塞到一个线程束中,或者让一个线程束中的线程都执行if,另外线程都执行else的这种方式可以将效率提高很多。
/******* 假设只配置一个x=64的一维线程块,那么只有两个个线程束 *****/
// 1. 这个kernel可以产生一个比较低效的分支
__global__ void mathKernel1(float *c)
{
int tid = blockIdx.x* blockDim.x + threadIdx.x;
float a = 0.0;
float b = 0.0;
if (tid % 2 == 0)
{
a = 100.0f;
}
else
{
b = 200.0f;
}
c[tid] = a + b;
}
// 2. 进行优化:
// 第一个线程束内的线程编号tid从0到31,tid/warpSize都等于0,那么就都执行if语句。
// 第二个线程束内的线程编号tid从32到63,tid/warpSize都等于1,执行else
// 线程束内没有分支,效率较高。
__global__ void mathKernel2(float *c)
{
int tid = blockIdx.x* blockDim.x + threadIdx.x;
float a = 0.0;
float b = 0.0;
if ((tid/warpSize) % 2 == 0)
{
a = 100.0f;
}
else
{
b = 200.0f;
}
c[tid] = a + b;
}
1.3 事件和指标
事件:是可计算的活动,比如这个分支就是一个可以计算的活动,对应一个在内核执行期间被搜集的硬件计数器。
指标:是内核的特征,由一个或多个事件计算得到。
3_2 理解线程束执行的本质(Part II)
最近这几篇应该是CUDA最核心的部分,并不是编程模型,而是执行模型,通过执行模型我们去了解GPU硬件的具体运行方式,这样才能保证我们写出更快更好的程序。
1.1 资源分配
每个SM上执行的基本单位是warp。即:
- 单指令通过指令调度器广播给某线程束的全部线程,这些线程同一时刻执行同一命令;
- 当然,也有分支情况。此时应注意 线程束分化 。
还有很多线程束是没有执行程序的,那么这些没执行的线程束情况又如何呢?
答:(原作者分了两类),我们离开 线程束 内的角度(线程束内是观察线程行为,离开线程束我们就能观察线程束的行为了):
- 一类是已经激活的,也就是说这类线程束其实已经在SM上准备就绪了,只是没轮到他执行,这时候他的状态叫做阻塞;
- 还有一类可能分配到SM了,但是还没上到片上,这类称之为未激活线程束;
而每个SM上有多少个线程束处于激活状态,取决于以下资源:
- 程序计数器
- 寄存器
- 共享内存
线程束一旦被激活来到片上,那么他在执行结束之前就不会再离开SM。
一个SM上被分配多少个线程块和线程束取决于:
- SM中可用的寄存器和共享内存,以及内核需要的寄存器和共享内存大小。
每个SM都有32位的寄存器组,每个架构寄存器的数量不一样,其存储于寄存器文件中,为每个线程进行分配.
同时,固定数量的共享内存,在线程块之间分配。
当SM内的资源没办法处理一个完整块,那么程序将无法启动。
当寄存器和共享内存分配给了线程块,这个线程块处于活跃状态,所包含的线程束称为活跃线程束。
活跃的线程束又分为三类:
-
选定的线程束:SM要执行的某个线程束;
-
阻塞的线程束:线程束不符合条件还没准备好;
-
符合条件的线程束:准备要执行;
满足下面的要求,线程束才算是符合条件的:
- 32个CUDA核心可以用于执行
- 执行所需要的资源全部就位
由于计算资源是在线程束之间分配的,且线程束的整个生命周期都在片上,所以线程束的上下文切换是非常快速的。
下面我们介绍如何通过大量的活跃的线程束切换来隐藏延迟
1.2 延迟隐藏
延迟是什么?
答:
比如一个算法验证,你交给计算机,计算机会让某个特定的计算单元完成这个任务,共需要十分钟。而这十分钟内计算机利用率可能没有达到100%,也就是机器的某些功能是空闲的。
-
而如果,在这10分钟之内再跑一个同样的程序不同的数据,这时候你发现还没有完全利用完资源,于是有继续加任务给计算机,结果加到第十分钟了,已经加了十个了。你还没加完,但是第一个任务已经跑完了,如果你这时候停止加任务,等陆陆续续的你后面加的任务都跑完了共用时20分钟,共执行了10个任务。
那么平局一个任务用时 20 10 = 2 \frac{20}{10}=2 1020=2 分钟/任务 。
-
但是我们还有一种情况,因为任务还有很多,第十分钟你的第一个任务结束的时候你继续向你的计算机添加任务,那么这个循环将继续进行,那么第二十分钟你停止添加任务,等待第三十分钟所有任务执行完,那么平均每个任务的时间是: 30 20 = 1.5 \frac{30}{20}=1.5 2030=1.5分钟/任务。
-
如果一直添加下去,理想情况下, l i m n → ∞ n + 10 n = 1 lim_{n\to\infty}\frac{n+10}{n}=1 limn→∞nn+10=1也就是极限速度,一分钟一个,隐藏了9分钟的延迟。
这是理想情况,有一个必须考虑的就是虽然你十分钟添加了10个任务,可是没准添加5个计算机就满载了,这样的话 极限速度只能是:$lim_{n\to\infty}\frac{n+10}{n\times 5}=0.2 $分钟/任务 了。
所以最大化是要最大化硬件,尤其是计算部分的硬件满跑,都不闲着的情况下利用率是最高的,利用率与常驻线程束直接相关。
硬件中 线程调度器 负责调度 线程束调度,当每时每刻都有可用的线程束供其调度,这时候可以达到计算资源的完全利用,以此来保证通过其他常驻线程束中发布其他指令的,可以隐藏每个指令的延迟。
1.2.1 指令隐藏延迟
GPU的延迟隐藏及其重要。对于指令的延迟,通常分为两种:
-
算术指令
算数指令延迟:是一个算术操作从开始,到产生结果之间的时间内,只有某些计算单元处于工作状态,而其他逻辑计算单元处于空闲。
延迟:算术延迟 10~20 个时钟周期
-
内存指令
内存指令延迟:当产生内存访问的时候,计算单元要等数据从内存拿到寄存器,这个周期是非常长的。
延迟:内存延迟400~800个时钟周期
1.2.1.1 计算所需线程束
下图就是阻塞线程束到可选线程束的过程逻辑图:
其中线程束0在阻塞两段时间后恢复可选模式,但是在这段等待时间中,SM没有闲置。
那么至少需要多少线程,线程束来保证最小化延迟呢?
答:little法则给出了下面的计算公式" 所需线程束 = 延迟 × 吞吐量 \text{所需线程束} = \text{延迟} \times \text{吞吐量} 所需线程束=延迟×吞吐量"
带宽:一般指的是理论峰值,最大每个时钟周期能执行多少个指令;
吞吐量:是指实际操作过程中每分钟处理多少个指令。
另外有两种方法可以提高并行:
- 指令级并行(ILP):一个线程中有很多独立的指令
- 线程级并行(TLP):很多并发地符合条件的线程
1.2.2 内存隐藏延迟
与指令周期隐藏延迟类似,内存隐藏延迟是靠内存读取的并发操作来完成的。
-
指令隐藏的关键目的是使用全部的计算资源;
-
内存读取的延迟隐藏是为了使用全部的内存带宽;内存延迟的时候,计算资源正在被别的线程束使用,所以我们不考虑内存读取延迟的时候计算资源在做了什么。
我们的根本目的是把计算资源,内存读取的带宽资源全部使用满,这样就能达到理论的最大效率。
同样下表根据Little 法则给出了需要多少线程束来最小化内存读取延迟:
这里有个单位换算过程,机器的性能指标内存读取速度给出的是GB/s 的单位。
我们需要的是每个时钟周期读取字节数,所以要用这个速度除以频率。例如C 2070 的内存带宽是144 GB/s ,转化成时钟周期 144 G B / s 1.566 G H z = 92 B / t \frac{144GB/s}{1.566GHz}=92 B/t 1.566GHz144GB/s=92B/t,这样就能得到单位时间周期的内存带宽了。
注意:需要说明的是这个速度不是单个SM的而是整个GPU设备的,我么用的内存带宽是GPU设备的而不是针对一个SM的。
如:Fermi 需要并行的读取74的数据才能让GPU带宽满载,如果每个线程读取4个字节,我们大约需要18500个线程,大约579个线程束才能达到这个峰值。
所以,延迟的隐藏取决于活动的线程束的数量,数量越多,隐藏的越好,但是线程束的数量又受到上面的说的资源影响。所以这里就需要寻找最优的执行配置来达到最优的延迟隐藏。
问:那么我们怎么样确定一个线程束的下界呢,使得当高于这个数字时SM的延迟能充分的隐藏?
答:这个公式很简单,也很好理解,就是SM的计算核心数乘以单条指令的延迟。
比如32个单精度浮点计算器,每次计算延迟20个时钟周期,那么我需要最少 32x20 =640 个线程使设备处于忙碌状态。
1.3 占用率
占用率:是一个SM种活跃的线程束的数量,占SM最大支持线程束数量的比,
CUDA工具包中,提供一个叫做CUDA占用率计算器的电子表格,填上相关数据可以帮你自动计算网格参数。
上面我们已经明确内核使用寄存器的数量会影响SM内线程束的数量,nvcc的编译选项也有手动控制寄存器的使用。
也可以通过调整block内线程的多少来提高占用率,当然要合理不能太极端:
- 小的线程块:每个线程块中线程太少,会在所有资源没用完就达到了线程束的最大要求;
- 大的线程块:每个线程块中太多线程,会导致每个SM中每个线程可用的硬件资源较少。
1.4 同步__syncthread()
并发程序对同步非常有用,比如pthread中的锁,openmp中的同步机制,这样做的主要目的是避免内存竞争
CUDA同步这里只讲两种:
- block内同步
- 系统级别
块级别的就是同一个块内的线程会同时停止在某个设定的位置,用:
__syncthread();
//注意:
该函数只能同步同一个块内的线程,不能同步不同块内的线程。
问:如何同步不同块内的线程?
答:就只能让核函数执行完成,控制程序交换主机,这种方式来同步所有线程。
1.5 可扩展性
可扩展性其实是相对于不同硬件的:
- 当某个程序在设备1上执行的时候时间消耗是T。
- 当我们使用设备2时,其资源是设备1的两倍,我们希望得到T/2的运行速度,
这种性质是CUDA驱动部分提供的特性,目前来说 Nvidia正在致力于这方面的优化,如下图:
3_4 避免线程束的分化
https://www.yuque.com/longlongqin/qalbrf/fbogu6
归约的方式基本包括如下几个步骤:
- 将输入向量划分到更小的数据块中:数据分块保证我们可以用一个线程块来处理一个数据块。
- 用一个线程计算一个数据块的部分和;
- 对每个数据块的部分和再求和得到最终的结果。
归约问题最常见的加法计算是把向量的数据分成对,然后用不同线程计算每一对元素,得到的结果作为输入继续分成对,迭代的进行,直到最后一个元素。
成对的划分常见的方法有以下两种:
-
相邻配对:元素与他们相邻的元素配对
-
交错配对:元素与一定距离的元素配对
[外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传(img-er5tRMOT-1616500678254)(https://face2ai.com/CUDA-F-3-4-避免分支分化/jiaocuo.png)]
3_6 动态并行
本节作为CUDA执行模型的最后一篇介绍动态并行。
到目前为止,我们所有的内核都是在主机线程中调用的,那么我们肯定会想,是否我们可以在内核中调用内核,这个内核可以是别的内核,也可以是自己,那么我们就需要动态并行了,这个功能在早期的设备上是不支持的。
动态并行的好处:
-
能让复杂的内核变得有层次,坏处就是写出来的程序更复杂,因为并行行为本来就不好控制。
-
另一个好处是等到执行的时候再配置创建多少个网格,多少个块,这样就可以动态的利用GPU硬件调度器和加载平衡器了,通过动态调整,来适应负载。并且在内核中启动内核可以减少一部分数据传输消耗。
1.1 嵌套执行
前面我们大费周章的其实也就只学了,网格,块,和启动配置,以及一些线程束的知识,现在我们要做的是从内核中启动内核。
内核中启动内核,和cpu并行中有一个相似的概念,就是父线程和子线程。子线程由父线程启动,但是到了GPU,这类名词相对多了些,比如父网格,父线程块,父线程,对应的子网格,子线程块,子线程。子网格被父线程启动,且必须在对应的父线程,父线程块,父网格结束之前结束。所有的子网格结束后,父线程,父线程块,父网格才会结束。
上图清晰地表明了父网格和子网格的使用情况,一种典型的执行方式:
主机启动一个网格(也就是一个内核)-> 此网格(父网格)在执行的过程中启动新的网格(子网格们)->所有子网格们都运行结束后-> 父网格才能结束,否则要等待。
1.1.1 同步
如果调用的线程没有显示同步启动子网格,那么运行时保证,父网格和子网格隐式同步。
父网格中的不同线程启动的不同子网格,这些子网格拥有相同的父线程块,他们之间是可以同步的。
线程块中所有的线程创建的所有子网格完成之后,线程块执行才会完成。如果块中的所有线程在子网格完成前退出,那么子网格隐式同步会被触发。隐式同步就是虽然没用同步指令,但是父线程块中虽然所有线程都执行完毕,但是依旧要等待对应的所有子网格执行完毕,然后才能退出。
前面我们讲过隐式同步,比如
cudaMemcpy
就能起到隐式同步的作用,但是主机内启动的网格,如果没有显式同步,也没有隐式同步指令,那么cpu线程很有可能就真的退出了,而你的gpu程序可能还在运行,这样就非常尴尬了。
父线程块启动子网格需要显示的同步,也就是说不同的线程束需要都执行到子网格调用那一句,这个线程块内的所有子网格才能依据所在线程束的执行,一次执行。
1.1.2 内存竞争
内存竞争对于普通并行就很麻烦了,现在对于动态并行,更麻烦,主要的有下面几点:
- 父网格和子网格共享相同的全局和常量内存;
- 父网格子网格有不同的局部内存;
- 有了子网格和父网格间的弱一致性作为保证,父网格和子网格可以对全局内存并发存取;
- 有两个时刻父网格和子网格所见内存一致:子网格启动的时候,子网格结束的时候;
- 共享内存和局部内存分别对于线程块和线程来说是私有的;
- 局部内存对线程私有,对外不可见;
1.2 例子:在GPU上嵌套执行
为了研究初步动态并行,我们先来写个简单的程序进行操作,代码如下:
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void nestHelloWorld(int iSize, int iDepth)
{
unsigned int tid = threadIdx.x;
printf("depth : %d blockIdx: %d, threadIdx, %d\n",iDepth, blockIdx.x,threadIdx.x); // 从输出结果可以看出,grid中的block并不是按照其ID进行执行的,哪一个block先执行是随机的。
if(iSize == 1) return;
int nthread = (iSize >> 1);
if(tid == 0 && nthread > 0) // 从输出结果得出,不同block中的线程执行的速度也是不一样的,有的先执行到这一步,有的晚一些才执行到这一步。
{
nestHelloWorld<<<1, nthread>>>(nthread, ++iDepth);
printf("-----------> nested execution depth: %d\n",iDepth);
}
}
int main(int argc,char* argv[])
{
int size=16;
int block_x=2;
dim3 block(block_x,1);
dim3 grid((size-1)/block.x+1,1);
nestHelloWorld<<<grid,block>>>(size,0);
cudaGetLastError();
cudaDeviceReset();
return 0;
}
执行:
nvcc -o 3_6_nestedHelloWorld.out 3_6_nestedHelloWorld.cu -Wno-deprecated-gpu-targets -lcudadevrt --relocatable-device-code true
这个程序的功能如下
第一层: 有多个线程块,执行输出,然后在tid==0
的线程,启动子网格,子网格的配置是当前的一半,包括线程数量,和输入参数 iSize。
第二层: 有很多不同的子网格,因为我们上面多个不同的线程块都启动了子网格,我们这里只分析一个子网格,执行输出,然后在tid==0
的子线程,启动子网格,子网格的配置是当前的一半,包括线程数量,和输入参数 iSize。
第三层: 继续递归下去,直到iSize==0
结束。
结果:
depth : 0 blockIdx: 6, threadIdx, 0
depth : 0 blockIdx: 6, threadIdx, 1
depth : 0 blockIdx: 0, threadIdx, 0
depth : 0 blockIdx: 0, threadIdx, 1
depth : 0 blockIdx: 7, threadIdx, 0
depth : 0 blockIdx: 7, threadIdx, 1
depth : 0 blockIdx: 3, threadIdx, 0
depth : 0 blockIdx: 3, threadIdx, 1
depth : 0 blockIdx: 2, threadIdx, 0
depth : 0 blockIdx: 2, threadIdx, 1
depth : 0 blockIdx: 5, threadIdx, 0
depth : 0 blockIdx: 5, threadIdx, 1
depth : 0 blockIdx: 1, threadIdx, 0
depth : 0 blockIdx: 1, threadIdx, 1
depth : 0 blockIdx: 4, threadIdx, 0
depth : 0 blockIdx: 4, threadIdx, 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
-----------> nested execution depth: 1
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
depth : 1 blockIdx: 0, threadIdx, 0
depth : 1 blockIdx: 0, threadIdx, 1
depth : 1 blockIdx: 0, threadIdx, 2
depth : 1 blockIdx: 0, threadIdx, 3
depth : 1 blockIdx: 0, threadIdx, 4
depth : 1 blockIdx: 0, threadIdx, 5
depth : 1 blockIdx: 0, threadIdx, 6
depth : 1 blockIdx: 0, threadIdx, 7
-----------> nested execution depth: 2
-----------> nested execution depth: 2
-----------> nested execution depth: 2
-----------> nested execution depth: 2
-----------> nested execution depth: 2
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
-----------> nested execution depth: 2
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
-----------> nested execution depth: 2
-----------> nested execution depth: 2
depth : 2 blockIdx: 0, threadIdx, 0
depth : 2 blockIdx: 0, threadIdx, 1
depth : 2 blockIdx: 0, threadIdx, 2
depth : 2 blockIdx: 0, threadIdx, 3
-----------> nested execution depth: 3
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
-----------> nested execution depth: 3
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
-----------> nested execution depth: 3
-----------> nested execution depth: 3
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
-----------> nested execution depth: 3
-----------> nested execution depth: 3
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
-----------> nested execution depth: 3
depth : 3 blockIdx: 0, threadIdx, 0
depth : 3 blockIdx: 0, threadIdx, 1
-----------> nested execution depth: 3
-----------> nested execution depth: 4
-----------> nested execution depth: 4
depth : 4 blockIdx: 0, threadIdx, 0
depth : 4 blockIdx: 0, threadIdx, 0
-----------> nested execution depth: 4
depth : 4 blockIdx: 0, threadIdx, 0
depth : 4 blockIdx: 0, threadIdx, 0
-----------> nested execution depth: 4
depth : 4 blockIdx: 0, threadIdx, 0
-----------> nested execution depth: 4
depth : 4 blockIdx: 0, threadIdx, 0
-----------> nested execution depth: 4
depth : 4 blockIdx: 0, threadIdx, 0
depth : 4 blockIdx: 0, threadIdx, 0
-----------> nested execution depth: 4
-----------> nested execution depth: 4
可见,当多层调用子网格的时候,同一家的(就是用相同祖宗线程的子网)是隐式同步的,而不同宗的则是各跑各的。