`

CUDA warp divergence与bank conflict

    博客分类:
  • CUDA
阅读更多

1.Warp Divergence

   warp是SM的基本执行单元。一个warp包含32个并行thread,这32个thread执行于SMIT模式。也就是说所有thread 执行同一条指令,并且每个thread会使用各自的data执行该指令。

   因为所有同一个warp中的thread必须执行相同的指令,那么如果这些线程在遇到控制流语句时,如果进入不同的分支,那么同一时刻除了正在执行的分之外,其余分支都被阻塞了,十分影响性能。这类问题就是warp divergence(warp分歧)。

  warp divergence问题只会发生在同一个warp中。为了避免warp分歧就要避免同一个warp存在不同的执行路径。

参考 
http://www.cnblogs.com/1024incn/p/4541313.html

 

2.Bank Conflict

  对于同一个wrap中的线程(一个wrap内包含了32个线程),访问共享存储器时,以half-wrap的形式分两次访问。同一half-wrap内的线程同时可以访问不同的bank,而不同线程对同一个bank 的访问只能顺序进行。

   所谓的bank-conflict,就是同一half-wrap内的线程,访问了同一bank里的共享内存。bank-conflict会让原本并行的对共享内存的访存操作变成串行从而极大的降低程序效率。 特殊情况是:half-wrap内所有的线程访问同一个共享内存中的同一地址,会产生一次广播,在这种情况下不会发生bank conflict。

  下面有一些小技巧可以避免bank conflict 或者提高global存储器的访问速度

       1. 尽量按行操作,需要按列操作时可以先对矩阵进行转置

       2. 划分子问题时,使每个block处理的问题宽度恰好为16的整数倍,使得访存可以按照 s_data[tid]=i_data[tid]的形式进行

       3. 使用对齐的数据格式,尽量使用nvidia定义的格式如float3,int2等,这些格式本身已经对齐。

       4. 当要处理的矩阵宽度不是16的整数倍时,将其补为16的整数倍,或者用malloctopitch而不是malloc

        5. 利用广播,例如s_odata[tid] = tid%16 < 8 ? s_idata[tid] : s_idata[15];会产生8路的块访问冲突而用:s_odata[tid]=s_idata[15];s_odata[tid]= tid%16 < 8 ? s_idata[tid] : s_odata[tid]; 则不会产生块访问冲突

参考 
http://blog.csdn.net/endlch/article/details/47043069

 

分享到:
评论

相关推荐

    Inter-warp divergence aware execution on GPUs

    当WARP中的线程需要执行不同分支的代码时,就会出现分支分化(divergence)现象。当这种情况发生时,WARP的执行效率会下降,因为硬件需要串行处理每个分支,这会导致部分线程处于空闲状态,影响整体性能。 在CUDA...

    CUDA学习----sp, sm, thread, block, grid, warp概念 .docx

    在深入理解CUDA编程之前,我们需要先了解几个关键的概念:SP、SM、thread、block、grid和warp。 1. SP(Streaming Processor):SP,即流处理器,是GPU硬件中的基本执行单元,通常被称为CUDA核心。每个SP都能执行...

    cuda合并访问共享存储器bank冲突实用PPT学习教案.pptx

    CUDA 合并访问和共享存储器 Bank 冲突是GPU编程中的关键概念,尤其是在优化计算效率时。CUDA(Compute Unified Device Architecture)是NVIDIA提供的一个并行计算平台,允许程序员利用GPU的强大处理能力来执行计算...

    cuda.zip_CUDA ustc_cuda_ustc cuda

    5. **同步与通信**:CUDA提供了多种同步机制,如`__syncthreads()`用于线程块内的同步,`cudaStream_t`用于管理异步执行。同时,主机和设备之间的数据传输也是CUDA编程的重要部分。 6. **错误处理**:CUDA编程中...

    warp-rnnt:CUDA-Warp RNN-换能器

    CUDA-Warp RNN-换能器 一个GPU实现RNN换能器(格雷夫斯 , )。 该代码是从移植的(由Awni Hannun编写),并且充分利用了CUDA warp机制。 损失的主要瓶颈是基于动态规划算法的前进/后退。 特别是,存在一个嵌套循环...

    cuda性能指南--用与提升gpu编程cuda的性能提升

    例如,若条件只依赖于(threadIdx / WSIZE),其中WSIZE是warp大小,那么可以避免warp的分散,因为这样的条件与warp的排列是同步的。 CUDA编译器有时能够通过分支预测、循环展开等技术优化控制流,以避免warp分散。...

    CUDA C编程权威指南.pdf

    - 线程束(warp)的概念及其在CUDA执行模型中的重要性。 - 线程束分化的影响因素及解决方案。 - **3.3 并行性的表现** - 使用nvprof工具检测线程束活跃程度和内存操作情况。 - 增大并行性的方法。 - **3.4 避免...

    《CUDA C编程》CodeSamples源代码

    - warp:CUDA的最小调度单位是warp,每个warp包含32个线程,这些线程会以锁步方式执行。 - 异步执行:CUDA支持多个流(Stream),使得计算和数据传输可以重叠,提高资源利用率。 4. CUDA编程技巧与优化: - 内存...

    cuda_scan2_cudascan_cuda_scan_

    在实际应用中,CUDA Scan可以通过多种方法实现,如Tree-Based Scan、Warp-Based Scan、Block Reduction等。每种方法都有其优缺点,适用于不同的数据规模和场景。 为了提高效率,还可以采用分治策略,将大数组拆分成...

    CUDA编程指南,教您如何写CUDA程序

    CUDA中的线程按照层级结构组织,包括网格(Grid)、区块(Block)和线程束(Warp)。网格包含多个区块,每个区块又包含多个线程。这种分层结构使得开发者能够更灵活地控制并行度和线程间的协作。 #### 2.2 存储器...

    CUDA架构 .pdf

    在CUDA中,GPU的执行单元以32个线程为一组进行调度,这被称为一个warp(线程束)。在内存读取等操作中,所有32个线程必须同步执行,如果遇到数据依赖等问题,它们会被挂起。当CPU上的CUDA程序调用内核函数时,这些...

    cuda检测工具 devicequery.exe

    Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a ...

    Mac Warp IP 优选.zip.cab

    Mac Warp IP 优选.zip.cab

    CUDA_C_Programming_Guide v10.0.pdf

    此外,文档还涵盖了矩阵函数(Warp matrix functions)支持的新矩阵乘法操作,包括m=32,n=8,k=16以及m=8,n=32,k=16的矩阵乘法操作,这一变化扩展了CUDA在机器学习和深度学习领域的应用能力,因为矩阵乘法是这类应用...

    CUDA 的经典入门

    CUDA编程模型的核心是线程层次结构,包括线程块(Thread Block)、线程束(Warp)和网格(Grid)。线程块是由一组线程组成的逻辑单位,它们可以在硬件上并行执行。线程束是线程块内的基本执行单元,通常包含32个线程...

    Warp_Dash_v1.5.1.ipa iphone

    Warp_Dash ipa iphone

    warp 一种快速远程登录软件

    warp 一种快速远程登录软件 warp 一种快速远程登录软件

    CUDA编程学习资料

    #### 一、CUDA简介与编程模型 **1.1 CUDA:可伸缩并行编程模型** CUDA(Compute Unified Device Architecture)是英伟达推出的一种并行计算平台和编程模型,它使开发人员能够利用GPU(图形处理单元)的强大计算能力...

    relax and timewarp_RELAX_osu_timewarp_

    在"relax and timewarp"这个压缩包中,很可能包含了与这两个功能相关的资源、教程或者玩家的自定义设置。可能包括了设置文件、自定义地图、教学视频或者是一些关于如何有效利用relax和timewarp提高游戏技能的文章。...

Global site tag (gtag.js) - Google Analytics