CUDA编程：笔记2

最新推荐文章于 2024-03-17 19:09:41 发布

jinchengwu3344

最新推荐文章于 2024-03-17 19:09:41 发布

阅读量936

点赞数 3

分类专栏： CUDA编程文章标签： cuda

本文链接：https://blog.csdn.net/longlongqin/article/details/115140407

版权

CUDA编程专栏收录该内容

3 篇文章 2 订阅

订阅专栏

本笔记主要是阅读：谭升的博客的 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位于最外层。

计算出三维对应的线性地址是： $threadIdx.x+threadIdx.y\times blockDim.x+threadIdx.z\times blockDim.x \times blockDim.y$

想象下c语言的数组，如果把上面这句话写成c语言，假设三维数组t保存了所有的线程，那么(threadIdx.x,threadIdx.y,threadIdx.z)表示为**：t[z][y][x];**

一个线程块包含多少个线程束：

$\text{WarpsPerBlock}=\text{ceil}\begin{pmatrix}\frac{\text{ThreadsPerBlock}}{\text{warpSize}}\end{pmatrix}$ 。注：ceil函数是向正无穷取整的函数，比如 $ceil(\frac{9}{8})=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个任务。

那么平局一个任务用时 $\frac{20}{10}=2$ 分钟/任务。
但是我们还有一种情况，因为任务还有很多，第十分钟你的第一个任务结束的时候你继续向你的计算机添加任务，那么这个循环将继续进行，那么第二十分钟你停止添加任务，等待第三十分钟所有任务执行完，那么平均每个任务的时间是： $\frac{30}{20}=1.5$ 分钟/任务。
如果一直添加下去，理想情况下， $lim_{n\to\infty}\frac{n+10}{n}=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 ，转化成时钟周期 $\frac{144GB/s}{1.566GHz}=92 B/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

可见，当多层调用子网格的时候，同一家的（就是用相同祖宗线程的子网）是隐式同步的，而不同宗的则是各跑各的。

jinchengwu3344

关注

3
点赞
踩
11

收藏

觉得还不错? 一键收藏
1
评论
CUDA编程：笔记2

本笔记主要是阅读：谭升的博客的 GPU编程（CUDA）3_2 理解线程束执行的本质(PartⅠ)从外表来看，CUDA执行所有的线程，并行的，没有先后次序的；但实际上硬件资源是有限的，不可能同时执行百万个线程，所以从硬件角度来看，物理层面上执行的也只是线程的一部分，而每次执行的这一部分，就是我们前面提到的线程束。1.1 warp与blockwarp是SM中基本的执行单位。是硬件层面的线程集合。当一个网格被启动（网格被启动，等价于一个内核被启动。每个内核对应于自己的网格），线程块被分配到某一个S.
复制链接

扫一扫