分支分歧(branch divergence)造成SIMT性能降低的原因

最新推荐文章于 2023-10-22 17:30:06 发布
最新推荐文章于 2023-10-22 17:30:06 发布
阅读量2.3k 收藏 5
点赞数 2
分类专栏： GPU 文章标签： GPU

分支分歧(branch divergence)造成SIMT性能降低的原因

warp中的线程按照SIMD(Single Instruction Multiple Data)模式执行——它们执行相同的代码，但处理不同的数据。当遇到分支指令时，执行不同分支路径的线程只能串行执行。这种因不一致分支行为引起的串行执行现象被称作Branch Divergence，它的出现严重影响了GPU的执行效率，降低了程序性能。针对这个问题，人们提出了线程交换的优化方法，其主要思想为：将具有相同分支行为(即执行相同分支路径)的线程放在同一个warp中，这样warp中的线程就可以并行执行，从而达到减少Branch Divergence、提高性能的目的。

于齐, 王博千, 沈立, 王志英, & 陈微. (2016). Gpu平台上面向性能和功耗的分支优化. 计算机科学, 43(5), 22-26.

yuyuanhang14
关注
  • 2
    点赞
  • 踩
  • 5
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
7. warp divergence
WuYuFffan的博客
09-18 362
7. warp divergence 因为CUDA是SIMD架构，所以当一个cuda核执行选择分支时，其他非该分支的核会强制进入等待状态。 int tid = threadIdx.x; if (tid % 2 == 0) { //do something } else { //do something else } tid为奇数的设备执行if时，tid为偶数的设备拥塞。反之亦然。 可见如果同一个warp中的Thread有很多分支,会导致warp divergence,这会严重降低程序的运
关于warp中遇到Branch Divergence的执行情况
文森的学术之路
03-18 2179
CUDA官方的文档[1]里原文是： If threads of a warp diverge via a data-dependent conditional branch, the warp serially executes each branch path taken, disabling threads that are not on that path, and when all pa
CUDA编程——Warp Divergence
Zhang Junior 的博客
01-19 8855
Warp　　回顾一下CUDA的线程层次 　　CUDA编程中，warp是调度和运行的基本单元，目前，每个warp包含32个threads。软件逻辑上，程序员的所有thread是并行的，但是，从硬件的角度来说，实际上并不是所有的thread能够在同一时刻执行。　　例如：一个block 有128个thread，那么可以分为warp0~warp3，SM同时只执行一个warp，其他warp以就绪，挂起等状态
CUDA学习笔记（八）Branch Divergence and Unrolling Loop
qq_45788429的博客
10-22 347
和之前的图示相比，工作的thread数目没有变化，但是，每个thread的load/store global memory的位置是不同的。从高级语言层面是无法看出性能提升的原因的，需要从low-level instruction层面去分析，第二段代码循环次数减少了一半，而循环体两句语句的读写操作的执行在CPU上是可以同时执行互相独立的，所以相对第一段，第二段性能要好。global array index也要相应的调整，因为，相对之前的版本，同样的数据，我们只需要原来一半的thread就能解决问题。
CUDA之Branch/Divergent branches详解
Bruce_0712的博客
03-25 1990
https://devtalk.nvidia.com/default/topic/463316/branch-divergent-branches/
[system-track][computing][cuda] branch divergence and __syncthreads
giantpoplar的专栏
07-05 155
#include<stdio.h> #include<cuda_runtime.h> __global__ void branch_divergence_0(){ __shared__ int counter; if(threadIdx.x == 0 ) counter = 0; int c = 0; whi...
Warp divergence & Bank conflict
GaryW666的博客
10-08 728
Warp divergence & Bank conflict warp、bank概念         warp是并行计算时最小的并发架构，通常由32个thread         bank是对SM中共享内存的划分，划分个数与对应硬件warp中所含thread数一致。         对应使用的计算能力3.x版本的显卡，一个warp含有32个thread，因此划分的bank数也为3
divergence.zip_divergence_矩形网格matlab
07-13
用于简单的解偏微分方程的学习，使用最简单的矩形网格
GPU平台上面向性能和功耗的分支优化.pdf
09-25
在SIMT架构下，当线程组内的线程执行到分支指令时，如果线程间分支选择不一致，GPU会执行所有可能的分支路径，以满足所有线程的需求，这就会产生分支分歧。这种分歧不仅降低了执行效率，还增加了功耗，因为GPU需要...
Divergence Dashboard_forexindicator_indicator_divergence_zip_MET
09-30
indicator metatrader4.........
Inter-warp divergence aware execution on GPUs
06-15
GPU的编程细节专著：“分支问题”。CUDA里，天生带来的就有WARP的分支执行问题。这里，作者专门研究了如何优化，非常有收获！
Phan ky_rsi_TheHidden_divergence_
09-30
DivergencesLookback windows to identify pivot points .-The signal will be delayed by the number of bars used in the right lookback window.Maximum range to detect divergence is the largest distance ...
PZ_DivergenceLite_PZ_DivergenceLite_divergence_
10-03
Indicator of Pz Divergence Lite version
CUDA ---- Branch Divergence and Unrolling Loop
weixin_30583563的博客
06-02 279
Avoiding Branch Divergence 有时，控制流依赖于thread索引。同一个warp中，一个条件分支可能导致很差的性能。通过重新组织数据获取模式可以减少或避免warp divergence（该问题的解释请查看warp解析篇）。 The Parallel Reduction Problem 我们现在要计算一个数组N个元素的和。这个过程用CPU编程很容易实现： int ...
深入理解git分支分叉(branch diverged)的问题
热门推荐
ikscher的专栏
02-27 7万+
问： I googled and read many posts, but none could make me understand the branch divergence problem yet. If I've remote tracking branch, I often get into the following: $ git status # On branch
CUDA学习笔记（七）分支发散和展开循环
云中寻雾的博客
07-20 3664
这节是关于Branch Divergence and Unrolling Loop。（分支发散和展开循环） Avoiding Branch Divergence（避免分支发散） 有时，控制流依赖于thread索引。同一个warp中，一个条件分支可能导致很差的性能。通过重新组织数据获取模式可以减少或避免warp divergence（该问题的解释请查看warp解析篇）。 The Paralle...
idia连不上连接池_AI芯片：具有Nv​​idia Ampere架构的A100 GPU
weixin_26746861的博客
07-17 740
idia连不上连接池Nvidia GPUs have dominated the AI chip market for most of the decade. However, when Google’s TPU (Tensor Processor Unit) was publicly available in 2018 through the Google Cloud, it changed p...
SIMT与SIMD的区别？
bdfpxsfh的专栏
03-12 1335
来源： SIMD指令和SIMT指令有什么区别？在出现SIMD和SIMT之前，占主导地位的是SISD（Single instruction single data）。但很快人们发现，在一些逻辑简单但计算量大的场景下，这些计算可以并行完成。因此，直接想到的解决方案就是多块芯片并行计算，即多核系统，与此对应的编程模型是多线程编程。多核系统的产生使得编程变得轻松，只需解决多线程互斥、同步等一些小问题即可。但很快有人意识到，从硬件的角度看，这么做很不划算。我们只是为了并行计算，按理说只需要多块计算元件和寄存器就行了，
SecondactivityMainActivity.java
最新发布
06-24
SecondactivityMainActivity.java
warp divergence
12-14
warp divergence（指变形发散）是计算机图形学中的概念，用于描述在并行计算中的处理器线程之间的差异性。在图形渲染过程中，为了提高性能，常常采用并行计算的方式来对多个像素或顶点进行处理。 然而，不同处理器线程之间的运算速度、内存访问和负载的差异性会导致某些线程完成任务的时间比其他线程更快或更慢。这种差异性被称为warp divergence。 当一个warp（一组连续的线程）中的线程执行不同的代码路径时，就会发生warp divergence。这种情况下，线程无法同时进行计算，需要等待其他线程完成执行，从而影响了整个warp的性能。 常见引起warp divergence的情况包括分支语句（如if/else、switch等），当分支条件不同的时候，不同线程会进入不同的代码分支，从而产生warp divergence。此外，也会发生在访问不连续内存、常量内存的操作中，或者在执行不同的指令集等。 为了避免warp divergence对并行计算的性能造成负面影响，可以采用一些优化技术来减少warp divergence的发生。例如，可以将分支条件相同的线程放在一个warp中，这样它们可以同时执行相同的代码分支。此外，通过合并指令、重新组织数据、考虑内存访问模式等方式，也可以降低warp divergence的风险。 综上所述，warp divergence是计算机图形学中一个重要的概念，用于描述并行计算中处理器线程之间的差异性。理解和优化warp divergence对于提高并行计算性能具有重要意义。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值