CN109542731A - 一种面向gpu的层次递进下钻的性能监测方法 - Google Patents

Abstract

本发明提供了一种面向GPU的层次递进下钻的性能监测方法，包括：(1)关注CPU程序与设备交互部分，识别CPU和GPU空闲时间、主机端CUDA函数、CPU与GPU数据传输开销以及热点kernel函数；(2)朝着细粒度方向监测，将CUDA模型中的网格grid级别作为一个监测层次，获得GPU设备级别的度量值；(3)向下进行SM级别监测，识别各SM上的负载平衡、执行的线程块数和SM级别的同步指令数；(4)监测其上运行的线程块block执行状况，包括线程块block的执行时间、分支数和同步指令数；(5)监测感兴趣的线程块中的warp行为，获得各warp执行时间、分歧数、同步指令数目以及存储器操作指令数信息；(6)监测CUDA线程的执行信息，包括各线程执行时间。

Description

一种面向GPU的层次递进下钻的性能监测方法

技术领域

本发明涉及一种信息技术的性能监测分析领域，特别是涉及一种面向GPU的层次递进下钻的性能监测方法。

背景技术

近年来，人工智能、大数据分析等技术的兴起，对高性能计算机性能提出新的要求。传统单一的计算结构逐渐被混合计算结构所替代，以期充分发挥各种组件的特长，为计算能力的提高发挥效力。在混合结构中，GPU作为当前混合结构的加速部分，发挥着越来越重要的作用。混合的计算模式增加了编程的复杂性。在实际开发中，存在很多语法正确但性能很差的程序，仅仅依靠调试环境很难满足实际需求，性能监测分析技术是GPU程序优化和应用性能进一步提高的重要途径。GPU作为一种外部设备，通过PCI-E总线与CPU连接。GPU与CPU地址空间相分离，它们存在异步执行的情形，指令集也不一致。这些硬件差异都给GPU的性能监测分析制造了障碍。目前，GPU的主流开发环境是基于CUDA模型展开的，CUDA程序的执行分为两个部分，一个是kernel函数级别的执行方式，属于粗粒度范畴，以任务并行为主，强调CPU与GPU之间的交互、并行以及CUDA流级别的并行性；而kernel函数内部的执行方式，属于细粒度范畴，以数据并行为主，它是采用了一种类似于早期向量机的并行执行模式：单指令多线程(Single Instruction Multiple Threads，SIMT)，程序尽量少用分支，线程规模庞大，通过快速切换线程来隐藏延迟。这种执行模式相对于传统CPU程序是非常复杂的。这两部分的并行执行模式不同，造成其监测和分析方法也存在较大差异。基于CUDA模型的GPU的性能监测分析技术也通常从这两个方面分别展开。

关于任务并行部分的性能监测，就是将GPU执行函数(即CUDA kernel函数)作为监测的基本单位来参与分析，主要通过设备端和主机端的交互、kernel函数的执行状况、kernel函数执行过程中的设备计数器度量值以及将设备kernel函数作为一个新的维度融入主机端编程模型的混合程序的性能分析，核心问题是如何能够在主机端准确地监测设备端kernel函数的执行信息(主要是时间度量)。通过主机端对设备kernel函数的监测，主要存在下面四种方法：

1)同步的方法：该方法适用于设备的同步操作，而不适于异步操作，比如数据传输与计算重叠的情况下，方法的主要思路是在kernel函数的开始和结束点通过主机端插桩操作进行测量，同步方法中，kernel函数就类似于传统CPU端程序的子函数调用，其优点是思路直观，无需特殊的性能监测机制；不足则表现为适用范围有限，无法支持高效灵活的设备异步执行模式。

2)基于设备事件的监测方法：该方法利用CUDA函数库中提供的API，在kernel函数(也支持其他设备操作，比如异步传输)前后加上设备端计时事件，本质上就是kernel启动函数的包装操作，优点是应用范围广，可以监测异步kernel函数。不足就是该方法依赖厂商提供的API，通常只能计时，度量比较单一。由于计时操作在设备端进行，CPU与GPU频率不同，因此还需要一个时间戳转换过程。另外，计时事件是一个全局的设备事件，可能会致使并行流串行化，从而影响程序的执行行为。

3)基于设备回调(callback)的监测方法：设备通过其驱动触发主机端先前已经注册的操作，从而完成监测的目的。监测库中要先注册感兴趣的设备API，而不是运行时API，才能在运行中回调，当然该方法的使用是需要驱动支持的。优点是可以减少开销，在理想情况下是最准确的测量方法。不足则是依赖于设备以及驱动，测量信息仍然是设备端的，需要进行转换操作。

4)设备信息采样的方法：主机端的线程查询设备端程序状态和计数器信息，是一种基于采样原理的监测方法。设备端kernel函数的测量精度极大依赖于探针的采样频率，通过识别设备状态的变化来监测kernel函数的运行。设备端采样可由特定事件触发，也可连续执行。该方法的优点是可监测丰富的性能信息，不足则是需要额外线程，从而增加开销，可能会扰动程序执行。

另外，NVIDIA公司还发布了一个第三方开发使用的性能监测工具接口CUPTI(CUDAProfiling Tools Interface)，以动态库的形式可供性能工具使用。该接口着力于实现上述多种主机端监测方法，为GPU性能监测的发展提供支持。CUPTI提供了四种模式：

1)Activity模式

这种模式是一种基于事件驱动的方法，需要先初始化，然后会将监测信息放入buffer存储空间，通常在全局同步点加以读取。这种模式可测量的事件相对较少，但信息量比较丰富。由于该模式与设备硬件调度器相关联，不会出现并行串行化的问题，结果更加准确。

2)Callback API模式

这种模式实现了回调功能，可对运行时函数、驱动函数、资源变化以及同步操作进行跟踪回调。这种模式可监测的事件类型较多，但每个事件获取的信息相对较简单，可用于解析一些函数的参数。

3)Event API模式

这种模式就是用来记录设备计数器信息，可通过回调机制触发，通常使用额外的线程采样。

4)Metric API模式，

这种模式用来处理和聚集特定的事件(计数器信息)来获取度量结果。GPU硬件结构的异步性、并发性、复杂性以及执行模型的新特征都致使传统的性能监测分析方法失效。目前GPU监测方法的主流研究，基本都是先使用基于设备事件的方法，后又加入CUPTI接口展开监测工作，受限于厂商的技术支持。这些研究只侧重于GPU与CPU的交互，以及kernel函数级别的执行情况(任务并行部分)。关于CUDA程序的性能监测分析方法主要存在的问题是:目前CUDA程序的监测方法，通常只是针对任务并行和数据并行部分分别展开工作，而以任务并行部分的研究为主，没有将两部分结合在一起共同研究，并未组织层次化的监测分析方法。

基于CUDA模型的GPU程序本身是有其运行特征的，作为一种外部设备，GPU设备本身无法自行执行程序，需要CPU主机驱动。最常见的流程是主机端准备数据并将其传输到设备端，启动设备端kernel函数执行计算程序，最后再将结果返回主机端，运算完成。CUDA程序在GPU设备端内部的并行执行与传统的CPU主机端的执行模式有着明显的区别，更类似向量机的数据并行方式，也就是所谓的SIMT模式。CUDA程序主要分为两个部分：分别是与主机交互的任务并行部分以及CUDA kernel函数在GPU设备上执行的数据并行部分。CUDA程序首先将数据由主机端传输到设备端，kernel函数级别的执行可使用流的方式实现并行(当然，多个GPU设备或上下文信息也可实现粗粒度的设备端并行)。CUDA kernel函数在抽象层上可分为网格(grid)，线程块(block)，CUDA线程(thread)三个层次，grid映射整个CUDA上下文，block映射于一个GPU SM(Streaming Multiprocessor)并执行且中途不会改变。block还会进一步划分为由一定CUDA线程组成的warp单元(目前warp是由32个线程组成)，warp是SM的调度单元，它在SM中被调度执行从而实现CUDA线程的并行执行。从上述过程的描述可知，CUDA程序的运行是一个复杂的过程。若要深入理解CUDA程序在GPU上执行的状况，发现性能瓶颈，需要解决下面几个方面的问题。GPU设备不是独立的执行单元，需被CPU主机端程序启动并进行传输、同步等交互操作。CUDA程序计算操作的关键部分就是GPU设备上执行的kernel函数，其执行方式与CPU程序截然不同。因此，将kernel函数级别(kernel函数的启动以及与CPU的交互)的任务并行部分以及kernel函数内部的数据并行部分相分离的监测分析都会限制程序全面深入的理解，一种层次递进下钻的性能监测方法是CUDA程序监测的重要需求。

发明内容

为了解决这些问题，本发明提出了一种面向GPU的层次递进下钻的性能监测方法，包括如下步骤：

(1)首先关注CPU程序与设备交互的部分，从而识别CPU和GPU的空闲时间、主机端CUDA函数、CPU与GPU数据传输开销以及热点kernel函数；

(2)在识别所述热点kernel函数之后，进一步朝着细粒度的方向进行监测，将CUDA模型中的网格grid级别作为一个监测层次，获得一些设备级别的IPC(Instruction PerClock)度量值，包括kernel占有率和周期指令数；

(3)向下进行SM级别的监测，识别各个SM上的负载平衡、执行的线程块数和SM级别的同步指令数；

(4)进一步监测其上运行的线程块block的执行状况，从而探索造成热点SM的根源，这里监测的是同一kernel函数内部的线程块block，识别各个所述线程块block的执行时间、分支数以及同步指令数；

(5)线程块block的基本调度单位是warp单元，进一步监测感兴趣的线程块block中的warp行为，可以获得各warp执行时间、各warp的分歧数以及其他warp级别的度量值，所述其他warp级别的度量值包括同步指令数目、存储器操作指令数的性能相关的信息；

(6)最后监测CUDA线程的执行信息，包括各线程执行时间。

优选的，所述层次递进下钻的性能监测方法是基于层次递进监测框架的，所述层次递进监测框架分为两个部分，分别是任务并行部分粗粒度的监测以及数据并行部分细粒度的监测。所述任务并行部分粗粒度的监测基于CUPTI接口，根据规则选择用户感兴趣的部分结果，即kernel函数集作为后一部分监测的输入条件，所述数据并行部分细粒度的监测是基于Lynx开源项目展开的。

优选的，所述任务并行部分粗粒度的监测流程为：监测系统充分利用CallbackAPI模式的资源域跟踪功能，在每次创建资源的时刻都通过回调函数来触发Activity API监测kernel函数的执行和数据传输信息，同时通过Callback API模式运行时函数域的跟踪功能，通过回调cudalaunch函数来触发底层信息的监测，在监测驱动API时，调用相应驱动函数的回调功能来触发系统监测；当系统中没有专门从事监测计数器信息的进程或线程时，则创建一个监测进程或线程，利用Event API和Metric API专门用来收集硬件计数器的信息；当先前的监测进程/线程处于休眠状态时，则重新启动；利用Callback API的同步域和运行时域的数据传输函数cudaMemcpy()的跟踪来触发两个功能：一是将Activity API采集的信息从buffer之中dump出来，另一个则是终止线程/进程的数据采集，使它们处于休眠状态；主机端CUDA函数的监测，包括运行时和驱动API函数，使用传统的包装库的方法或CUPTICallback API本身提供的对运行时域和驱动域的监测接口完成监测。

优选的，所述数据并行部分细粒度的监测在设备端进行，基于GPU Lynx系统展开，Lynx在Ocelot框架下提供了类C语言的插桩规范说明，所述插桩过程和方法为：依据事先设定的插桩规范说明按照监测需求描述插桩层次以及所需度量，所述插桩规范说明通过C-on-demand，即CoD编译器编译成一种RISC的IR，即Internal Representation，再使用C-to-PTX翻译器将其翻译为PTX指令，此指令可用于实现插桩功能；CUDA应用程序通过nvcc将主机端代码和PTX kernel函数分离开，PTX kernel函数再链接GPU ocelot运行时系统，GPUocelot的PTX-PTX路径转换管理器实现了将PTX桩指令插入到原始PTX kernel函数的功能；将插桩后的kernel函数生成可执行文件在GPU设备上运行；与此同时，GPU ocelot框架在设备上开辟了部分存储空间用来保存程序运行中产生的性能信息；程序运行完毕，除了计算结果，监测数据也同样会返回主机端，监测系统会对监测数据进行规范化处理并存储于性能数据库，在此基础上，可进行性能分析和可视化展示，从而发现性能瓶颈进行kernel函数优化。

优选的，各类性能数据采集完毕之后，系统进行信息汇总、处理和存储的工作，将性能信息放入数据库，为可视化展示和程序行为分析做好信息源的准备。

优选的，所述面向GPU的层次递进下钻的性能监测方法也可由上至下跨级进行。

根据下文结合附图对本发明具体实施例的详细描述，本领域技术人员将会更加明了本发明的上述以及其他目的、优点和特征。

附图说明

后文将参照附图以示例性而非限制性的方式详细描述本发明的一些具体实施例。附图中相同的附图标记标示了相同或类似的部件或部分。本领域技术人员应该理解，这些附图未必是按比例绘制的。本发明的目标及特征考虑到如下结合附图的描述将更加明显，附图中：

图1为根据本发明实施例的层次递进监测架构LPMF(Level ProgressiveMonitoring Framework)性能监测方法示意图。

图2为根据本发明实施例的任务并行部分性能监测技术实现原理图。

图3为根据本发明实施例的设备端数据并行部分监测技术实现原理图。

具体实施方式

为了使得本发明能够针对其发明要点更加明显易懂，下面将结合附图和实例对本发明作进一步的说明。在下面的描述中阐述了很多细节和具体实例，提供这些实例是为了能够更透彻地理解本发明，并且能够将本发明完整形象地传达给本领域的技术人员。虽然本发明能够以很多不同于此描述的其它方式实施，但是本领域技术人员可以在不违背本发明内涵的情况下做相应的推广，因此本发明不受下面公开的具体实例及具体附图所限制。

CUDA程序中运行于设备的kernel函数，内部结构复杂，相关联的性能度量数目繁多，如何能够快速识别最重要的度量，从而有效地定位关键问题，是提高优化效率的重要步骤。

参见图1，一种面向GPU的层次递进下钻的性能监测方法，包括如下步骤：

(1)首先关注CPU程序与设备交互的部分，从而识别CPU和GPU的空闲时间、主机端CUDA函数、CPU与GPU数据传输开销以及热点kernel函数；

(2)在识别了热点kernel函数之后，进一步朝着细粒度的方向进行监测，将CUDA模型中的网格grid级别作为一个监测层次，获得一些设备级别的IPC度量值，包括kernel占有率和周期指令数，当然根据本领域技术人员的知识，也可以获得一些其他设备级别IPC度量值，均在本发明保护范围内；

(3)向下进行SM级别的监测，识别各个SM上的负载平衡、执行的线程块数和SM级别的同步指令数，当然根据本领域技术人员的知识，也可以获得其他一些SM级别的度量值，均在本发明保护范围内；

(4)进一步监测其上运行的线程块block的执行状况，从而探索造成热点SM的根源，这里监测的是同一kernel函数内部的线程块block，识别各个线程块block的执行时间、分支数以及同步指令数，当然根据本领域技术人员的知识，也可以获得一些其他block级别的度量值，均在本发明保护范围内；

(5)线程块的基本调度单位是warp单元，进一步监测感兴趣的线程块block中的warp行为，可以获得各warp执行时间、各warp的分歧数以及其他warp级别的度量值，所述其他warp级别的度量值包括但不限于同步指令数目、存储器操作指令数的性能相关的信息，本领域技术人员可以选择其他一些warp级别的度量值，同样也属于本发明的保护范围；

(6)最后监测CUDA线程的执行信息，包括各线程执行时间，根据本领域技术人员的知识，也可以获得一些其他线程级别的度量值，同样也属于本发明的保护范围。

层次递进下钻监测方法是基于层次递进监测框架的，层次递进监测框架分为两个部分，分别是任务并行部分粗粒度的监测以及数据并行部分细粒度的监测。任务并行部分粗粒度的监测基于CUPTI接口，根据规则选择用户感兴趣的部分结果，即kernel函数集作为后一部分监测的输入条件，所述数据并行部分细粒度的监测是基于Lynx开源项目展开的。

参见图2，任务并行部分粗粒度的监测流程为：监测系统充分利用Callback API模式的资源域跟踪功能，在每次创建资源的时刻都通过回调函数来触发Activity API监测kernel函数的执行和数据传输信息，同时通过Callback API模式运行时函数域的跟踪功能，通过回调cudalaunch函数来触发底层信息的监测，在监测驱动API时，调用相应驱动函数的回调功能来触发系统监测；当系统中没有专门从事监测计数器信息的进程或线程时，则创建一个监测进程或线程，利用Event API和Metric API专门用来收集硬件计数器的信息；当先前的监测进程/线程处于休眠状态时，则重新启动；利用Callback API的同步域和运行时域的数据传输函数cudaMemcpy()的跟踪来触发两个功能：一是将Activity API采集的信息从buffer之中dump出来，另一个则是终止线程/进程的数据采集(这是由于在通常情况下，CUDA程序会在kernel函数运行完毕之后，将数据传回主机端，在这之前常常需要使用同步操作，完成GPU与CPU的同步)，使它们处于休眠状态；主机端CUDA函数的监测，包括运行时和驱动API函数，使用传统的包装库的方法或CUPTICallback API本身提供的对运行时域和驱动域的监测接口完成监测(在此进行说明：图2中仅示出运行时函数部分作为示例，本发明针对运行时函数的监测方法同样可以推广到驱动API函数，只是由于驱动API函数属于底层函数，且函数关注程度小于运行时函数，因此没有示出)，本实施例选择包装库的方法，是为了兼容传统的性能监测工具。各类性能数据采集完毕之后，系统还需要进行信息汇总、处理和存储的工作，将性能信息放入数据库，为可视化展示和程序行为分析做好信息源的准备。除了可以识别CPU与GPU交互过程中的性能问题，还可以识别热点kernel函数，为kernel函数内部的性能监测做出前期准备。

参见图3，数据并行部分细粒度的监测在设备端进行，基于GPU Lynx系统展开，Lynx在Ocelot框架下提供了类C语言的插桩规范说明，所述插桩过程和方法为：依据事先设定的插桩规范说明按照监测需求描述插桩层次以及所需度量，所述插桩规范说明通过C-on-demand，即CoD编译器编译成一种RISC的IR，即Internal Representation，再使用C-to-PTX翻译器将其翻译为PTX指令，此指令可用于实现插桩功能；CUDA应用程序通过nvcc将主机端代码和PTX kernel函数分离开，PTX kernel函数再链接GPU ocelot运行时系统，GPUocelot的PTX-PTX路径转换管理器实现了将PTX桩指令插入到原始PTX kernel函数的功能；将插桩后的kernel函数生成可执行文件在GPU设备上运行；与此同时，GPU ocelot框架在设备上开辟了部分存储空间用来保存程序运行中产生的性能信息；程序运行完毕，除了计算结果，监测数据也同样会返回主机端，监测系统会对监测数据进行规范化处理并存储于性能数据库，在此基础上，可进行性能分析和可视化展示，从而发现性能瓶颈进行kernel函数优化。由于Lynx系统原始功能仅提供了部分度量的监测，这里通过修改插桩规范说明以及改进系统数据输出部分，以实现和完善上述层次化的监测功能，满足实际分析需求。

当然，面向GPU的层次递进下钻的性能监测方法也可由上至下跨级进行。本实施例采用跨越线程块block，关注SM下的所有warp执行信息或CUDA线程状态，本领域技术人员也可以进行其他跨级操作，均在本发明保护范围内。

虽然本发明已经参考特定的说明性实施例进行了描述，但是不会受到这些实施例的限定而仅仅受到附加权利要求的限定。本领域技术人员应当理解可以在不偏离本发明的保护范围和精神的情况下对本发明的实施例能够进行改动和修改。

Claims (6)

1.一种面向GPU的层次递进下钻的性能监测方法，其特征在于包括如下步骤：

(1)首先关注CPU程序与GPU设备交互的部分，从而识别CPU和GPU的空闲时间、主机端CUDA函数、CPU与GPU数据传输开销以及热点kernel函数；

(2)在识别所述热点kernel函数之后，进一步朝着细粒度的方向进行监测，将CUDA模型中的网格grid级别作为一个监测层次，获得GPU设备级别的度量值，包括kernel占有率、周期指令数IPC(Instruction Per Clock)；

(3)向下进行SM(Streaming Multiprocessor)级别的监测，识别各个SM上的负载平衡、执行的线程块数和SM级别的同步指令数；

(4)进一步监测其上运行的线程块block的执行状况，从而探索造成热点SM的根源，这里监测的是同一kernel函数内部的线程块block，识别各个所述线程块block的执行时间、分支数和同步指令数；

(5)线程块block的基本调度单位是warp单元，进一步监测感兴趣的线程块block中的warp行为，可以获得各warp执行时间、各warp的分歧数以及其他warp级别的度量值，所述其他warp级别的度量值包括同步指令数目、存储器操作指令数的性能相关的信息；

(6)最后监测CUDA线程的执行信息，包括各线程执行时间。

2.根据权利要求1所述的一种面向GPU的层次递进下钻的性能监测方法，其特征在于：所述层次递进下钻的性能监测方法是基于层次递进监测框架的，所述层次递进监测框架分为两个部分，分别是任务并行部分粗粒度的监测以及数据并行部分细粒度的监测。所述任务并行部分粗粒度的监测基于CUPTI接口，根据规则选择用户感兴趣的部分结果，即kernel函数集作为后一部分监测的输入条件，所述数据并行部分细粒度的监测是基于Lynx开源项目展开的。

3.根据权利要求2所述的一种面向GPU的层次递进下钻的性能监测方法，其特征在于：所述任务并行部分粗粒度的监测流程为：监测系统充分利用Callback API模式的资源域跟踪功能，在每次创建资源的时刻都通过回调函数来触发Activity API监测kernel函数的执行和数据传输信息，同时通过Callback API模式运行时函数域的跟踪功能，通过回调cudalaunch函数来触发底层信息的监测，在监测驱动API时，调用相应驱动函数的回调功能来触发系统监测；当系统中没有专门从事监测计数器信息的进程或线程时，则创建一个监测进程或线程，利用Event API和Metric API专门用来收集硬件计数器的信息；当先前的监测进程/线程处于休眠状态时，则重新启动；利用Callback API的同步域和运行时域的数据传输函数cudaMemcpy()的跟踪来触发两个功能：一是将Activity API采集的信息从buffer之中dump出来，另一个则是终止线程/进程的数据采集，使它们处于休眠状态；主机端CUDA函数的监测，包括运行时和驱动API函数，使用传统的包装库的方法或CUPTICallback API本身提供的对运行时域和驱动域的监测接口完成监测。

4.根据权利要求2所述的一种面向GPU的层次递进下钻的性能监测方法，其特征在于：所述数据并行部分细粒度的监测在设备端进行，基于GPU Lynx系统展开，Lynx在Ocelot框架下提供了类C语言的插桩规范说明，所述插桩过程和方法为：事先按照监测需求描述插桩层次以及所需度量设定的插桩规范说明，将所述插桩规范说明通过C-on-demand，即CoD编译器编译成一种RISC的IR，即Internal Representation，再使用C-to-PTX翻译器将其翻译为PTX指令，此指令可用于实现插桩功能；CUDA应用程序通过nvcc将主机端代码和PTXkernel函数分离开，PTX kernel函数再链接GPU ocelot运行时系统，GPU ocelot的PTX-PTX路径转换管理器实现了将PTX桩指令插入到原始PTX kernel函数的功能；将插桩后的kernel函数生成可执行文件在GPU设备上运行；与此同时，GPU ocelot框架在设备上开辟了部分存储空间用来保存程序运行中产生的性能信息；程序运行完毕，除了计算结果，监测数据也同样会返回主机端，监测系统会对监测数据进行规范化处理并存储于性能数据库，在此基础上，可进行性能分析和可视化展示，从而发现性能瓶颈进行kernel函数优化。

5.根据权利要求1-4任一所述的一种面向GPU的层次递进下钻的性能监测方法，其特征在于：各类性能数据采集完毕之后，系统进行信息汇总、处理和存储的工作，将性能信息放入数据库，为可视化展示和程序行为分析做好信息源的准备。

6.根据权利要求1-5任一所述的一种面向GPU的层次递进下钻的性能监测方法，其特征在于：所述面向GPU的层次递进下钻的性能监测方法也可由上至下跨级进行。