GPU 中分支指令的处理

阅读本文,你将知道:

  • 什么是 Warp
  • 什么是线程发散(Thread Diverge)
  • 什么是线程合并(Thread Reconverge)
  • 硬件如何实现分支发散和分支合并的处理
  • 降低线程发散带来的性能降低的方法有哪些

Warp

Warp 是 NVIDIA CUDA 中的一个术语,代表一组同时执行相同指令的线程。一个 Warp 包含一定数量的线程(NVIDIA GPU 上通常是 32 个线程)。Warp 与 AMD GPU 中的 Wavefront 是一个概念。

需要强调的是:Warp 中的线程以一个单一的指令流同时执行相同的指令序列。这种方式,即单指令流多数据流(SIMD)并行性,是建立在所有线程可以在任何时候执行相同操作的假设之上的。

线程发散

分支指令会非常影响 GPU 的性能。

依赖于数据的分支指令(例如if-else语句)可能导致线程之间的执行路径分岔。当一个 Warp 中的线程执行分支指令且分支判断的结果在不同的线程之间有所不同时,GPU 必须处理每个分支路径。这就是所谓的线程发散(Thread Diverge)

例如,假设一个 Warp 是 4 个线程,Kernel 程序如下:

int i = 0;
int thread_id = (int)(get_global_id(0)); // 获取当前线程的 ID
if (thread_id < 2) { // 当前线程的 ID 如果小于 2,则给 i 变量加一
    i += 1;
}
i -= 1;

第一个 Warp 中的四个线程 ID 分别为:0、1、2和3。2、3 号线程逻辑上不用执行加一操作,但由于 0、1 号线程需要执行且一个 Warp 中的线程永远执行的是相同的指令,因此 2、3 号线程也需要陪跑。这种陪跑的现象可以看成执行过程中的“气泡”,降低了硬件的利用率。

硬件实现

硬件如何处理线程发散。

每个 Warp 设置一个 PC,代表 Warp 中的所有线程当前要执行的指令的 PC。

Warp 中的每个线程再设置一个 NEXT_PC,代表当前线程接下来要跳转到的 PC。在没有分支指令的情况下, NEXT_PC 等于 PC + 1。

假设一个 Warp 包含 4 个线程。

当 4 个线程的 NEXT_PC 不一样时,代表发生了线程发散。线程发散发生之后,选择 NEXT_PC 最小的线程来执行,其他线程陪跑。陪跑线程的 NEXT_PC 在陪跑过程中不变。

退出分支执行时,同样的,选择 NEXT_PC 最小的线程来执行。

举个例子,

__kernel void
test(__global int* OUT) {
  int tid = (int)(get_global_id(0)); // Inst 0
  if (tid < 2) {                     // Inst 1
    OUT[tid] += 10;                  // Inst 2
  }
  OUT[tid] += 100;                   // Inst 3
}

执行过程如下,

图中,

  • 每一行的最前面的 PC 为 Warp 正在执行的指令 PC,紧接着的带方框的 4 个 PC 为四个线程的 NEXT_PC
  • Warp PC 为 Inst 1 时,发生了分支发散(Diverge)
  • 灰色方框代表此时线程处于陪跑状态
  • EOS(End of Shader) 代表线程结束。

线程合并

我们称退出分支执行并更新线程的执行状态的操作为线程合并(Thread Reconverge)。线程发散只有在线程的实际执行过程中才能发现,因此需要硬件自行判断,但是线程合并一般不让硬件判断,因为编译器可以提前知道。编译器在需要做线程合并的指令做好标记,硬件识别出该标记后,做线程合并操作,即选择 NEXT_PC 最小的线程来执行。

举例:嵌套分支

嵌套分支的处理逻辑同上,为了让读者更好的理解,给出一个更加复杂的情形的执行过程,加深理解。

__kernel void
test(__global int* OUT) {
/*Inst  0*/ int tid = (int)(get_global_id(0));
/*Inst  1*/ if (tid < 2) {
/*Inst  2*/   OUT[tid] += 10;
/*Inst  3*/   if (tid < 1) {
/*Inst  4*/     OUT[tid] -= 10;
/*Inst  -*/   } else {
/*Inst  5*/     OUT[tid] -= 5;
/*Inst  -*/   }
/*Inst  -*/ } else {
/*Inst  6*/   OUT[tid] += 20;
/*Inst  7*/   if (tid < 3) {
/*Inst  8*/     OUT[tid] -= 10;
/*Inst  -*/   } else {
/*Inst  9*/     OUT[tid] -= 5;
/*Inst  -*/   }
/*Inst  -*/ }
/*Inst 10*/ OUT[tid] += 100;
}

降低线程发散带来的性能降低的方法

线程发散会带来 GPU 执行性能的降低。减少分支发散的方法一般有:

  1. 减少一个 Warp 包含的线程的数量。例如,假设程序逻辑为:当线程 ID 小于 2 时执行 if 语句内容,且一个 Warp 是 4 个线程。那么此时会发生线程分散。如果设置一个 Warp 只包含 2 个线程,那么就不会发生线程分散了。该方法降低了线程发散的概率。
    • 英伟达一个 Warp 包含 32 个线程
    • ARM 一个 Warp 包含 4/8/16 个线程
  2. 将一个 Warp 中的线程分多次执行。例如,还是上面所述的程序,假设一个 Warp 包含 4 个线程,第一个时钟执行前两个线程,第二个时钟执行后两个线程。这也会降低线程分散的概率。
    • 细心的读者会发现,这样做硬件资源减半,程序执行时间也加倍了。解决办法是:在保持硬件资源不变的情况下,该方案可以一次调度两个 Warp 执行。
    • 英伟达的 Fermi 架构就是这样做的
  3. 分支合并
    参看:fetch.php (cmu.edu)
最后编辑于
©著作权归作者所有,转载或内容合作请联系作者
  • 序言:七十年代末,一起剥皮案震惊了整个滨河市,随后出现的几起案子,更是在滨河造成了极大的恐慌,老刑警刘岩,带你破解...
    沈念sama阅读 159,835评论 4 364
  • 序言:滨河连续发生了三起死亡事件,死亡现场离奇诡异,居然都是意外死亡,警方通过查阅死者的电脑和手机,发现死者居然都...
    沈念sama阅读 67,598评论 1 295
  • 文/潘晓璐 我一进店门,熙熙楼的掌柜王于贵愁眉苦脸地迎上来,“玉大人,你说我怎么就摊上这事。” “怎么了?”我有些...
    开封第一讲书人阅读 109,569评论 0 244
  • 文/不坏的土叔 我叫张陵,是天一观的道长。 经常有香客问我,道长,这世上最难降的妖魔是什么? 我笑而不...
    开封第一讲书人阅读 44,159评论 0 213
  • 正文 为了忘掉前任,我火速办了婚礼,结果婚礼上,老公的妹妹穿的比我还像新娘。我一直安慰自己,他们只是感情好,可当我...
    茶点故事阅读 52,533评论 3 287
  • 文/花漫 我一把揭开白布。 她就那样静静地躺着,像睡着了一般。 火红的嫁衣衬着肌肤如雪。 梳的纹丝不乱的头发上,一...
    开封第一讲书人阅读 40,710评论 1 222
  • 那天,我揣着相机与录音,去河边找鬼。 笑死,一个胖子当着我的面吹牛,可吹牛的内容都是我干的。 我是一名探鬼主播,决...
    沈念sama阅读 31,923评论 2 313
  • 文/苍兰香墨 我猛地睁开眼,长吁一口气:“原来是场噩梦啊……” “哼!你这毒妇竟也来了?” 一声冷哼从身侧响起,我...
    开封第一讲书人阅读 30,674评论 0 203
  • 序言:老挝万荣一对情侣失踪,失踪者是张志新(化名)和其女友刘颖,没想到半个月后,有当地人在树林里发现了一具尸体,经...
    沈念sama阅读 34,421评论 1 246
  • 正文 独居荒郊野岭守林人离奇死亡,尸身上长有42处带血的脓包…… 初始之章·张勋 以下内容为张勋视角 年9月15日...
    茶点故事阅读 30,622评论 2 245
  • 正文 我和宋清朗相恋三年,在试婚纱的时候发现自己被绿了。 大学时的朋友给我发了我未婚夫和他白月光在一起吃饭的照片。...
    茶点故事阅读 32,115评论 1 260
  • 序言:一个原本活蹦乱跳的男人离奇死亡,死状恐怖,灵堂内的尸体忽然破棺而出,到底是诈尸还是另有隐情,我是刑警宁泽,带...
    沈念sama阅读 28,428评论 2 254
  • 正文 年R本政府宣布,位于F岛的核电站,受9级特大地震影响,放射性物质发生泄漏。R本人自食恶果不足惜,却给世界环境...
    茶点故事阅读 33,114评论 3 238
  • 文/蒙蒙 一、第九天 我趴在偏房一处隐蔽的房顶上张望。 院中可真热闹,春花似锦、人声如沸。这庄子的主人今日做“春日...
    开封第一讲书人阅读 26,097评论 0 8
  • 文/苍兰香墨 我抬头看了看天上的太阳。三九已至,却和暖如春,着一层夹袄步出监牢的瞬间,已是汗流浃背。 一阵脚步声响...
    开封第一讲书人阅读 26,875评论 0 197
  • 我被黑心中介骗来泰国打工, 没想到刚下飞机就差点儿被人妖公主榨干…… 1. 我叫王不留,地道东北人。 一个月前我还...
    沈念sama阅读 35,753评论 2 276
  • 正文 我出身青楼,却偏偏与公主长得像,于是被迫代替她去往敌国和亲。 传闻我的和亲对象是个残疾皇子,可洞房花烛夜当晚...
    茶点故事阅读 35,649评论 2 271

推荐阅读更多精彩内容