`
ld_hust
  • 浏览: 171624 次
  • 性别: Icon_minigender_1
  • 来自: 武汉
社区版块
存档分类
最新评论

Tesla架构下的CUDA程序优化

阅读更多

CUDA优化的最终目的是:在最短的时间内,在允许的误差范围内完成给定的计算任务。在这里,“最短的时间”是指整个程序运行的时间,更侧重于计算的吞吐量,而不是单个数据的延迟。在开始考虑使用GPUCPU协同计算之前,应该先粗略的评估使用CUDA是否能达到预想的效果,包括以下几个方面:

精度:目前GPU的单精度性能要远远超过双精度性能,整数乘法、求模、求余等运算的指令吞吐量也较为有限。在科学计算中,由于需要处理的数据量巨大,往往采用双精度或者四精度才能获得可靠的结果,目前的Tesla架构还不能很好的满足高精度计算的需要。如果你的计算需要很高的精度,或者需要进行很多轮的迭代,最好考虑在关键的步骤中使用双精度,而在其他部分仍然使用单精度浮点以获得指令吞吐量和精度的平衡。而如果你对精度有更高的要求,那么现在的架构还不能太高的加速比。不过,在2010年将会普及的下一代架构中,双精度浮点和整数处理能力将有很大的提升,这种情况会有根本性的改变。

延迟:目前CUDA还不能单独为某个处理核心分配任务,因此必须先缓冲一定量的数据,再交给GPU进行计算。这样的方式可以获得很高的数据吞吐量,不过单个数据经过缓冲、传输到GPU计算、再拷贝回内存的延迟就比直接由CPU进行串行处理要长很多。如果对应用实时性要求很高,比如必须在数毫秒内完成对一个输入的处理,那么使用CUDA可能会影响系统的整体性能。对于要求人机能够实时交互的系统,应该将延迟控制在数十毫秒,以响应用户的输入。通过减小缓冲,可以减小延迟,但至少要保证每个内核程序处理的一批数据能够让GPU满负荷工作。不过在大多数情况下,在计算吞吐量较大,需要由GPU才能实时实现的系统,投入相同成本使用CPU很难做到接近实时。如果确实对实时性和吞吐量都有很高要求,应该考虑ASICFPGA或者DSP实现,这需要更多的投入,更长的开发时间和硬件开发经验。

计算量:如果计算量太小,那么使用CUDA是不合算的。衡量计算量有绝对和相对两种方式。

从绝对量来说,如果要优化的程序使用频率比较低,并且每次调用需要的时间也可以接受,那么使用CUDA优化并不会显著改善使用体验。对于一些计算量非常小(整个程序在CPU上可以在几十毫秒内完成)的应用来说,使用CUDA计算时在GPU上的执行时间无法隐藏访存和数据传输的延迟,此时整个应用程序需要的时间反而会比CPU更长。此外,虽然GPU的单精度浮点处理能力和显存带宽都远远超过了CPU,但由于GPU使用PCI-E总线与主机连接,因此它的输入和输出的吞吐量受到了IO带宽的限制。当要计算的问题的计算密集度很低时,执行计算的时间远远比IO花费的时间短,那么整个程序的瓶颈就会出现在PCI-E带宽上。此时无论如何提高浮点处理能力和显存带宽,都无法提高系统性能。

相对的计算量指可以并行的部分在整个应用中所占的时间。如果整个应用中串行部分占用时间较长,而并行部分较短,那么也需要考虑是否值得使用GPU进行并行计算。例如,假设一个程序总的执行时间为1.0,其中串行部分占0.8,而并行部分只占0.2,那么使用GPU将并行部分加速10倍,总的执行时间也只能从1.0降低到0.82。即使是在CPUGPU可以同时并行计算的应用中,执行时间也至少是CPU串行计算需要的0.8。只有在并行计算占用了绝大多数计算时间的应用中,使用CUDA加速才能获得很高的加速比。不过,随着GPU+CPU并行计算的普及和GPU架构的进一步改进,即使获得的加速比较小,也可能会由GPU执行。

完成对GPU加速能够达到的效果的粗略评估后,就可以开始着手编写程序了。为了在最短的时间内完成计算,需要考虑算法、并行划分、指令流吞吐量、存储器带宽等多方面问题。总的来说,优秀的CUDA程序应该同时具有以下几个特征:

在给定的数据规模下,选用算法的计算复杂度不明显高于最优的算法;

active warp的数量能够让SM满载,并且active block数量大于2,能够有效的隐藏访存延迟;

当瓶颈出现在指令流时,指令流效率已经经过了充分优化;

当瓶颈出现在访存或者IO时,已经选用了恰当的存储器和和存储器访问方式以获得最大带宽。

按照开发流程的先后顺序,CUDA程序的编写与优化需要解决以下问题:

1.确定任务中的串行部分和并行部分,选择合适的算法。首先,需要将问题分为几个步骤,并确定哪些步骤可以用并行算法实现,并确定要使用的算法。

2.按照算法确定数据和任务的划分方式,将每个需要并行实现的步骤映射为一个满足CUDA两层并行模型的内核函数。在这里就要尽量让每个SM上拥有至少6个活动warp和至少2个活动线程块。

3.编写一个能够正确运行的程序,作为优化的起点。程序必须能够稳定运行,不能发生存储器泄漏的情况。为了保证结果正确,在必要的时候必须使用memory fence、同步、原子操作等功能。在精度不足或者发生溢出时必须使用双精度浮点或者更长的整数类型。

3.优化显存访问,避免显存带宽成为瓶颈。在显存带宽得到完全优化前,其他优化不会产生明显结果。显存访问优化中可以使用的技术包括:

将可以采用相拓扑实现的几个kernel合并为一个,减少对显存访问;

除非非常必要,应该尽力避免将线程私有变量分配到local memory

为满足合并访问,采用cudaMallocPitch()或者cudaMalloc3D()分配显存;

为满足合并访问,对数据类型进行对齐(使用__align);

为满足合并访问,保证访问的首地址从16的整数倍开始,如果可能,尽量让每个线程一次读32bit

在数据只会被访问一次,并且满足合并访问的情况下可以考虑使用zerocopy

在某些情况下,考虑存储器控制器负载不均衡造成分区冲突的影响。

使用用有缓存的常数存储器和纹理存储器提高某些应用的实际带宽。

4.优化指令流。由于编译器会进行一些优化,而编译过程基本无法控制,所以指令流优化不一定能获得立竿见影的效果。但是,仍然有一些准则可以参考,包括:

如果只需要少量线程进行操作,那么一定记得要使用类似(if threadID < N)的方式,避免多个线程同时运行占用更长时间或者产生错误结果;

在不会出现不可接受的误差的前提下采用CUDA算术指令集中的快速指令;

使用#unroll,让编译器能够有效的展开循环;

采用原子函数实现更加复杂的算法,并保证结果的正确性;

避免多余的同步;

如果不产生bank conflict的算法不会造成算法效率的下降或者非合并访问,那么就应该避免bank conflict

5.资源均衡。调整shared memroyregister的使用量。为了使程序能够获得更高的SM占用率,应该调整每个线程处理的数据数量、shared memoryregister的使用量。这需要在三者间进行调整。当每个线程处理的子任务间有一定的公用部分时,可以考虑让一个线程处理更多的数据来提高指令流和访存效率。为了获得更高的SM占用率,必须控制每个线程的shared memoryregister的使用量。通过调整block大小,修改算法和指令,以及动态分配shared memory,都可以提高shared的使用效率。而减小register的使用则相对困难,因为register的使用量并不是由内核程序中声明的变量大小决定,而是由内核程序中使用寄存器最多的时刻的用量决定的。由于编译器会尽量减小寄存器的用量,因此实际使用的寄存器有可能会小于在程序中声明的量。但是在通常情况下,由于需要暂存中间结果并且一些指令也需要更多的寄存器,一般寄存器用量都大于内核程序中声明的私有变量的总数量。使用以下方法可能可以节约一些寄存器的使用:使用shared memory存储变量;使用括号更加明确的表示每个变量的生存周期;用对[u]long型或的处理代替对两个相邻的[u]short型或者四个相邻的[u]char型的处理;使用占用寄存器较小的等效指令代替原有指令,如用__sin函数代替sin函数。不过,由于不能对编译器的优化过程进行控制,即使使用了这些手段也不一定能减小寄存器的用量。值得注意的是,采用--maxrregcount编译选项只是让编译器将超出限制的私有寄存器分配在local memory中,造成较大的访存延迟。

6. 与主机通信优化。由于PCI-E带宽相对较小,应该尽量减少CPUGPU间传输的数据量,并通过一些手段提高可用带宽。可用的技术包括:

使用cudaMallocHost分配主机端存储器,可以获得更大的带宽;

一次缓存较多的数据,再一并传输,可以获得较高的实际带宽;

需要将结果显示到屏幕时,直接使用与图形学API互操作功能完成,避免将数据传回;

       使用流式处理隐藏与主机的通信时间。

       对于用CUDA C语言编写的程序,按照上述流程进行优化,是比较适合的。不过在优化中,各种因素往往相互制约,很难同时达到最优。读者需要按照要处理问题的类型,瓶颈出现的部位和原因具体分析。按照预想进行优化也不是总能达到预想中的效果,有时优化手段反而会降低性能。在实践中,仍然需要不断实验各种优化方法,在不断试验与迭代中一步步排除不可行的方案,最后得到一个比较理想的方案。

使用CUDA C并不总是能够编译到最优的指令。如果确实必要,可以用PTX优化程序中最关键的步骤。

除此以外,还要灵活采用宏和模版,动态分配内存和显存以及动态划分数据等手段提高程序的通用性,并在处理不同规模、不同数据类型的问题时选用不同的优化策略。

2010年将要推出的新处理器中,各种存储器的带宽和延迟会有一定的调整,大部分指令的吞吐量也会有非常显著的提升。随着GPU架构的进一步改进和编译器性能的不断提高,下一代GPU上的CUDA程序优化工作会变得更加简单。

分享到:
评论

相关推荐

    cuda优化文档

    - 随着技术进步,新一代GPU架构(如Tesla架构)在双精度浮点和整数处理能力上有了显著提升,有助于改善高精度计算场景下的性能。 2. **延迟** - 由于CUDA不能为每个处理核心单独分配任务,因此需要先将数据缓存...

    Tesla GPU架构分析2

    了解这些硬件特性对于优化CUDA程序至关重要。理解线程如何在SM中分配,以及如何有效地利用SIMT架构,可以显著提升代码的性能。例如,减少线程间的通信,合理分配线程块大小,以及充分利用共享内存,都是提升GPU计算...

    2-Tesla GPU架构分析1

    NVIDIA Parallel Nsight是专为Visual Studio开发人员设计的工具,便于CUDA应用程序的调试和优化。 8. **硬件规格**:Tesla C2050和C2070 GPU拥有448个CUDA核心,工作频率1.15 GHz,配备3GB或6GB的GDDR5存储器,...

    2.Tesla GPU架构分析1

    在Tesla架构的SM设计中,每个SM通常包括32个运算核心、16个LD/ST模块、4个SFU、一定大小的寄存器和L1缓存。此外,还有纹理读取单元、纹理缓存和PolyMorph Engine,负责处理顶点和曲面细分。Warp Schedulers负责调度...

    胡文美 CUDA_PDF教材

    - **Tesla架构概述**:Tesla是NVIDIA针对高性能计算市场推出的一系列GPU产品线。该架构专为科学计算和数据中心应用设计,具备高精度浮点运算能力和大量内存带宽。 - **Tesla与CUDA的关系**:Tesla GPU集成了对CUDA...

    NVIDIA-CUDA相关资料

    理解如何有效地使用这些内存层次对于优化CUDA程序至关重要。 4. **CUDA库和工具**:NVIDIA提供了许多CUDA库,如cuBLAS(用于线性代数)、cuFFT(快速傅里叶变换)、cuDNN(深度神经网络加速)等,这些库极大地简化...

    cuda_7.5.18_windows.zip

    2. 支持更多硬件:CUDA 7.5支持更多的NVIDIA GPU型号,包括Kepler、Maxwell和早期的Tesla架构GPU,使得更广泛的硬件可以利用CUDA编程的优势。 3. 改进的库:CUDA 7.5包含了更新和优化的数学库,如cuBLAS(线性代数)...

    cuda8.0,下载包

    - 编写CUDA程序时,需要包含头文件`#include &lt;cuda_runtime.h&gt;`和其他CUDA库头文件,使用特殊的`__global__`和`__device__`关键字定义GPU上的函数和变量。 6. **CUDA优化**: - 合理分配CUDA线程块和网格大小,...

    CUDA教程,开发者必备

    CUDA支持异构串行-并行编程模式,其中NVIDIA的Tesla GPU架构能够加速CUDA程序的运行。CUDA不仅能够充分利用NVIDIA GPU的强大计算能力,还能够适应多核CPU架构。 CUDA编程模型的核心在于其异构编程特性。在CUDA程序...

    电子-NVIDIACUDA开发者系列培训之三CUDA演化.pdf

    总之,CUDA的发展历程体现了NVIDIA为GPU计算所做的不断努力,从最初的基础编程支持,到后续深度优化以及最后的CUDA演化,不仅提供了更加丰富的库和工具,还优化了软件平台的安装和配置流程,使得开发者能够更加专注...

    cuda toolkit 工具包 3.1 win7 64版

    3. **Nsight Eclipse Edition**:一个集成开发环境(IDE),专门为CUDA编程提供了调试和性能分析工具,使得开发者可以更方便地构建、调试和优化CUDA程序。 4. **CUDA可视化工具**:如Nsight Visual Studio Edition...

    并行编程cuda

    NVIDIA的GPU配备了新的Tesla统一图形和计算架构,这些GPU能够运行CUDA C程序,并且在笔记本电脑、个人电脑、工作站和服务器上广泛可用。 CUDA提供三个关键的抽象概念:线程层次结构的层次、共享内存以及屏障同步。...

    《GPU高性能计算之CUDA》实例

    本书除了详细介绍了CUDA的软硬件架构以及C for CUDA程序开发和优化的策略外,还包含有大量的实例供读者学习参考用。 下表是各个实例的介绍列表。 文件夹 对应书中章节 备注 ACsearch_DPPcompact_with_driver ...

    中科曙光HPC培训教程汇总:D31-并行编程—CUDA程序设计简介.ppt

    【CUDA 程序设计简介】 CUDA(Compute Unified Device Architecture)是由NVIDIA推出的并行计算平台和编程模型,旨在利用GPU(图形处理器)的强大计算能力解决复杂计算问题。CUDA允许开发者使用C、C++、Fortran等...

    cudnn-for-cuda9.1-win10

    5. **验证安装**: 运行简单的CUDA程序或深度学习框架,如TensorFlow、PyTorch,以确保CUDNN正确安装并工作。 **总结** CUDNN 7.0.5与CUDA 9.1的结合为Windows 10用户提供了强大的深度学习计算能力,加速了图像处理...

    CUDA 编程指南4.0中文版

    介绍如何在CUDA程序中混合使用运行时API和驱动API,以实现更灵活和高级的编程功能。 **3.5 版本和互操作性** 说明不同版本的CUDA之间的兼容性和互操作性问题,以及如何处理跨版本编程的问题。 **3.6 计算模式** ...

    cudnn_6 /CUDA8适用

    1. 安装CUDA Toolkit:这是运行TensorFlow所必需的,包含了开发和运行CUDA程序所需的所有工具。 2. 下载CUDNN:根据你的CUDA版本和操作系统选择正确的CUDNN版本。 3. 解压CUDNN并复制到CUDA安装目录:通常需要将解压...

    CUDA4.0编程指导(中文版)

    这部分介绍了如何设计和优化CUDA程序,以获得最佳性能。 **5.2 最大化利用率** - **5.2.1 应用层次**:这部分介绍了如何在整体应用程序层面优化性能。 - **5.2.2 设备层次**:这部分说明了如何在GPU层面优化性能。...

    CUDA PART1 姓名

    特别是到了2009年,NVIDIA推出了针对高性能计算优化的产品,如Tesla C1060,该产品拥有240个核心,峰值性能达到了936 GFLOPS,远远超过了同期的多核CPU(如Intel Core i7 965,仅4个核心,峰值性能为102 GFLOPS)。...

Global site tag (gtag.js) - Google Analytics