qipccc的博客

0%

CUDA_03理解线程束的执行

发表于 更新于 分类于
本文字数： 2.6k 阅读时长 ≈ 4 分钟

CUDA执行的实质是线程束的执行。

来源：理解线程束执行的本质(Part I)

因为硬件根本不知道每个块谁是谁，也不知道先后顺序，硬件(SM)只知道按照机器码跑，而给他什么，先后顺序，这个就是硬件功能设计的直接体现了。
从外表来看，CUDA执行所有的线程，并行的，没有先后次序的，但实际上硬件资源是有限的，不可能同时执行百万个线程，所以从硬件角度来看，物理层面上执行的也只是线程的一部分，而每次执行的这一部分，就是我们前面提到的 线程束 。

理解CUDA中的线程束

来源

线程束是SM中基本的执行单元， 每个线程束包含了 32 个线程。所以，当我们定义的一个block中含有的线程数量不是 32 的整数倍时，那么该 block 会被分配给 n+1 个线程束， 其中 n 表示的是 线程数量与32 的整数倍，多出来的线程束会有部分线程不活跃，但因为线程束是基本执行单元，所以这个额外的线程束依旧会消耗和其他线程束相同的资源。

例如： 有一个一维的线程块包含有80个线程，那么映射到 硬件 上，系统就会为这个线程块分配 3个连续 的线程束，使总共96个硬件线程去支持80个软件线程。其中 第3个线程束 中有 16个线程不活跃，但是仍然 消耗 这系统资源，被浪费掉了。

每个线程束中的线程按照 单指令多线程 的方式执行，也就是说同属于一个线程束的线程均会执行 同一个指令。

线程束(warp)和线程块

执行过程：

  1. 一个网格被启动（网格被启动，等价于一个内核被启动，每个内核对应于自己的网格），网格中包含线程块；

  2. 线程块被分配到某一个SM上；

  3. SM上的线程块将分为多个线程束，每个线程束一般是 32 个线程；

  4. 在一个线程束中，所有线程按照单指令多线程SIMT的方式执行，每一步执行相同的指令，但是处理的数据是私有数据。

下图，对应 逻辑 、实际 和 硬件：

在块中，每个线程有唯一的编号， threadIdx;

在网格中，每个线程块也有唯一的编号，blockIdx;

所以每个线程就有在网格中的唯一编号。

因为线程块是逻辑层，被定义了一个三维结构，而线程束是硬件层，在硬件结构中是一维排列的，所以它们之间的对应关系就是：

  1. 对于一个二维的线程块，映射到一维的物理布局上是：
  1. 对于一个三维的线程块，映射到一维的物理布局上是：

线程束分化

线程束是硬件层面的线程集合，线程块是逻辑层面的线程集合。线程束被执行的时候会被分配相同的指令，但是是处理各自私有的数据。如果一个线程束中的不同线程包含不同的控制条件，那么当我们执行到这个控制条件时就会面临不同的选择，同一个线程束中的线程，执行不同的指令就叫做 线程束的分化。

GPU相较于CPU不时候大量计算逻辑复杂的任务，因为CPU在进行流水线作业时会采用 分支预测 的技术（之前有做个测试，如果对一个很大的随机数组在进行if…else逻辑判断前先进行排序会比未排序直接进行逻辑判断速度慢很多）,而GPU这方面就不擅长。

为了解决GPU中 线程束中的所有线程执行相同的指令 和 线程束又是分化的 之间的矛盾，只能让那些不成立的指令进行等待。

例如：

1
分水果，你不爱吃，那你就只能看着别人吃，等大家都吃完了，再进行下一轮（也就是下一个指令）

示意图：

线程束分化会产生严重的性能下降。条件分支越多，并行性削弱越严重。所以应该避免同一个线程束内的线程分化，控制线程块中线程分配到线程束是有规律的而不是随机的。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
//低效的分支：
__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;
}//奇数偶数分化，每次间隔都要分化一次，分化严重

//调整
__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;
}
//第一个线程束内的线程编号tid从0到31，tid/warpSize都等于0，那么就都执行if语句。
//第二个线程束内的线程编号tid从32到63，tid/warpSize都等于1，执行else线程束内没有分支，效率较高。

延迟隐藏-效率的最大化

效率的最大化就是要最大化硬件，尤其是计算部分的硬件跑满，都不闲着的情况下利用率是最高的，即最大化功能单元的利用率，而 利用率 和 常驻线程束 直接相关。硬件中的线程调度器负责线程束的调度，当每时每刻都有可用的线程束供其调度时，这时候可以达到计算资源的完全利用，依次来保证通过其他常驻线程束来发布其他指令，并可以隐藏每个指令的延迟。

对于指令的延迟，主要分为两种：

  1. 算数指令延迟：从一个算数操作开始，到产生结果之间的时间，这个时间段内只有某些计算单元处于工作状态，而其他逻辑计算单元处于空闲，10 ~ 20 个时钟周期。

  2. 内存指令延迟： 当产生内存访问的时候，计算单元要等数据从内存拿到寄存器，这个周期是非常长的， 400 ~ 800 个时钟周期。

提高并行的两种方式：

  1. 指令集并行（ILP）：一个线程中有很多独立的指令

  2. 线程级并行（TLP）：很多并发的符合条件的线程

指令隐藏的关键目的是使用全部的计算资源，而内存读取的延迟隐藏是为了使用全部的内存带宽

内核使用寄存器的数量会影响SM内线程束的数量，通过调整线程块内线程的多少来提高占用率

  • 小的线程块： 每个线程块中线程太少，会在所有资源没用完就达到了线程束的最大要求；

  • 大的线程块： 每个线程块中太多线程，会导致每个SM中每个线程可用的硬件资源较少。