CN103729180A - 一种快速开发cuda并行程序的方法 - Google Patents

Abstract

本发明提供一种快速开发CUDA并行程序的方法，涉及修改CPU串行程序和CUDA并行程序移植：按照CUDA并行程序格式要求修改原CPU串行程序，得到新版本的CPU串行程序；然后设计CUDA并行程序，把CPU串行程序移植到GPU平台上并进行进一步的优化。该一种快速开发CUDA并行程序的方法和现有技术相比，本发明利用先修改CPU串行程序后移植到GPU平台的方法，充分利用CPU和GPU上的资源，快速、有效地实现基于GPU平台进行程序并行化的方法，实用性强，易于推广。

Description

一种快速开发CUDA并行程序的方法

技术领域

本发明涉及计算机应用技术领域，具体的说是一种快速开发CUDA并行程序的方法。

背景技术

自从2006年英伟达（NVIDIA）公司推出图形处理器G80（包含了128个流式多处理器，最新的G200包含了240个多处理器）以来，图形处理器（GPU，Graphic Processing Unit）在某些大规模并行计算的应用上，相对于CPU来说性能提高可达100倍以上。尤其从2008年5月，NVIDIA公司推出用于GPU的开发平台统一计算设备架构（CUDA，Compute Unified Device Architecture）软件开发工具包（SDK，Software Development Kit）1.1以来，基于GPU平台的并行计算便得到了大规模的推广。CUDA为GPU计算提供了统一计算设备架构，使用户很容易地将GPU编程融于传统的编程工具（例如Visual Studio、Gcc等）和语言（例如C、C++及FORTRAN等）中。在短短的一年以来，CUDA被应用于加速大规模并行计算领域的许多方面，如在图像处理，物理模型模拟（如计算流体力学），工程和金融模拟与分析，生物医药工程，数据库及数据挖掘，搜索，排序等诸多方面都有很好的应用，在很多应用中取得了1至2个几何数量级的加速。

GPU拥有更多的晶体管，用于数据处理而不是像CPU那样去处理数据cache和指令控制，这意味着GPU具有巨大的并行计算能力。

CUDA C作为GPU的并行编程语言。CUDA C编程将CPU称之为主机，将GPU作为一个协处理器称为设备。在CUDA编程中，多个线程同时执行在一个GPU上，并由多个线程组成一个线程块（Block），多个线程块又组织成网格（Grid）；另外，每32个线程组成一个束（warp）。CUDA编程中常用到的优化技术有合理的网格配置，每个SM上有足够多的warp可隐藏访问延迟，并进行全局存储器的合并访问，共享存储器的使用，纹理存储器和常量存储器的使用，寄存器的合理使用等等。CUDA并行程序需要执行成千上万个线程，比CPU串行程序开发要复杂，设计不当会导致结果的错误和性能下降。

可见，为实现CUDA并行程序开发的需求，需要一种快速、有效地实现CUDA并行程序开发的方法。

发明内容

本发明的技术任务是解决现有技术的不足，提供一种快速开发CUDA并行程序的方法。

本发明的技术方案是按以下方式实现的，该一种快速开发CUDA并行程序的方法，其具体操作过程为：

步骤一、CPU串行程序修改，即按CUDA程序格式要求对CPU串行程序的分析和修改，其中

1.1）CPU串行程序分析具体包括：利用打印时间方式测试串行程序中的热点模块；根据算法特点和数据特点分析热点模块是否可以并行，是否可以采用CUDA细粒度并行；根据可并行的模块，找到CUDA内核将会使用到的数组，并对数组结构进行分析；

1.2）仿CUDA格式的CPU串行程序修改，具体包括：原程序算法修改，修改成可并行的代码；数组修改，修改成适合CUDA并行程序格式的数组形式；

步骤二、CUDA并行程序移植，即设计CUDA并行程序把CPU程序移植到GPU平台，以及CUDA并行程序的优化，其中

2.1）设计CUDA并行程序具体包括：线程块、网格的设计，用于数据的划分和内核的计算；通信函数的实现，用于进行CPU与GPU之间的数据传递；内核的设计，实现CUDA并行加速热点模块；

2.2）所述优化CUDA并行程序，具体包括：将上述CUDA并行程序利用优化技术进一步提高并行程序的性能，主要优化包括2个方面：通信优化和内核优化。

所述步骤1.1）的详细过程为：

热点测试是指根据时间的测试结果确定热点函数，作为后面移植的重点代码模块；

找出热点代码后，需要分析热点部分的算法、数据特点，根据算法和数据的特点分析其是否可以并行，是否可以采用CUDA细粒度并行，确定其并行性；

根据对串行程序的分析，确定哪些模块需要移植到GPU平台上运行，对于需要运行上GPU平台上的代码内的数据进行分析，确定数组在代码中的什么位置传递到CUDA内核中，传递的方向是CPUtoGPU还是GPUtoCPU，以及每次传递时的数据大小等信息。

所述步骤1.2）的详细过程为：

对于CPU串行程序，根据并行算法的要求修改原程序，改成可并行的模式；同时需要重新设计并行算法；

根据CUDA并行程序对数组格式的要求对串行程序中的数组进行修改，进而将对原CPU程序修改成一个仿CUDA格式的CPU串行程序。

所述步骤2.1）的详细过程为：

GPU数组设计、并行模型设计：设计GPU端数组，设计数组大小、类型；设计CUDA程序的线程并行模型：block和grid；

根据对原程序数组的分析，把CPU串行程序移植到GPU平台，根据热点模块的算法和代码实现CUDA内核代码。

所述步骤2.2）的详细过程为：

CPU与GPU通信优化，利用异步技术减少CPU和GPU之间通信的时间；CUDA内核优化，利用全局存储器合并访问、共享存储器、常量存储器、纹理存储器手段提高访存速度，利用指令流提高计算手段优化CUDA内核。

本发明与现有技术相比所产生的有益效果是：

本发明的一种快速开发CUDA并行程序的方法可以快速、有效地实现CPU串行程序移植到GPU平台上，降低CUDA并行程序开发的周期和难度，减少bug的调试时间，提高CUDA并行程序开发效率，降低CUDA并行程序开发周期，该方法按照逐步修改的原则，即先修改CPU串行程序后移植到GPU平台的原理，把需要在GPU上做的工作尽量先在CPU平台上修改，实用性强，易于推广。

附图说明

附图1是本发明的CUDA并行程序移植流程图。

具体实施方式

下面结合附图对本发明的一种快速开发CUDA并行程序的方法作以下详细说明。

如附图1所示，本发明提供的一种快速开发CUDA并行程序的方法是按照CUDA并行程序格式要求修改原CPU串行程序，得到新版本的CPU串行程序；然后设计CUDA并行程序，把CPU串行程序移植到GPU平台上并进行进一步的优化。本发明利用先修改CPU串行程序后移植到GPU平台的方法，充分利用CPU和GPU上的资源，快速、有效地实现基于GPU平台进行程序并行化的方法。其具体操作过程为：

1、CPU串行程序分析。

对于CPU串行程序，首先需要测试串行程序中的热点函数，以及分析热点函数的并行性。

热点测试：根据时间的测试结果确定热点函数，作为后面移植的重点代码模块。

并行性分析：找出热点代码后，需要分析热点部分的算法、数据特点，根据算法和数据的特点分析其是否可以并行，是否可以采用CUDA细粒度并行，确定其并行性。

确定CUDA内核使用的数组：根据对串行程序的分析，确定哪些模块需要移植到GPU平台上运行，对于需要运行上GPU平台上的代码内的数据进行分析，确定数组在代码中的什么位置传递到CUDA内核中，传递的方向是CPUtoGPU还是GPUtoCPU，以及每次传递时的数据大小等信息，然后设计这些数组的定义方式和大小。

2、仿CUDA格式的CPU串行程序

CUDA程序相对CPU程序比较复杂，当出现bug时，调试的难度也要比CPU程序大很多，为了降低CUDA程序开发难度和周期，可以把一些GPU上的移植工作提前在CPU平台上实现，具体涉及下面几个方面：

修改成可并行算法：对于CPU串行程序，有些代码理论上可以并行，但经过CPU版本的优化之后导致代码不能直接并行化，这时需要根据并行算法的要求修改原程序，改成可并行的模式；有些模块理论上是可以并行的，但串行算法无法直接并行化，需要重新设计并行算法。

数组修改：CPU串行程序中使用的数组定义的形式有可能无法在CUDA内核中直接使用，这时需要对数组的定义进行修改，如C语言程序，结构体中的指针需要改变成单独的指针/数组，才能进行CPU与GPU之间的数据传递。另外，考虑到全局存储器合并访问的问题，有时还需要对数组的访问方向进行修改，从而也需要改变数组的定义形式（如做行列变换）。总之，根据CUDA对数组使用和CPU串行程序之间的区别，提前把数组修改，方便程序的调试。

根据前面几条的修改方式，对原CPU程序修改成一个仿CUDA格式的CPU串行程序，为后面的移植工作做大量的准备，有利于后面CUDA程序的移植。

3、GPU数组设计、并行模型设计。

GPU数组设计：设计GPU数组大小、类型、维度等信息；

设计CPU与GPU之间的数组通信方式；

并行模型设计：block和grid的设计满足算法的数据特点。

4、CUDA并行程序基本版本。

根据对原程序数组的分析，把CPU串行程序移植到GPU平台，根据热点模块的算法和代码实现CUDA内核代码。

设计调用语句：

Kernel<<<grid, block, …>>>(…);

设计CUDA内核。

根据算法的并行性分析，设计内核，划分每个线程的计算任务，利用同步语句满足内核程序的逻辑正确性。

5、CUDA并行程序优化版本。

根据步骤4实现的基本版本的CUDA并行程序，利用CUDA的优化技术进一步提高并行程序的性能，主要优化包括2个方面：通信优化和内核优化。

GPU通信优化：GPU计算需要CPU与GPU之间进行数据的传递，合理的利用通信优化技术有利用提高CUDA并行程序的性能，如异步。

CUDA内核优化：CUDA内核的优化对其性能更为重要，主要涉及存储器访问优化和指令流优化，存储器访问优化包括：全局存储器合并访问，利用共享存储器、常量存储器、纹理存储器替换全局存储器的访问，提高访问速度；指令流优化是指利用高效的指令代替低效的指令，如CUDA中的快速函数。

由本发明的技术方案可见，本发明该方法按照先修改CPU串行程序后移植到GPU平台的原理，把需要在GPU上做的工作尽量先在CPU平台上修改，降低了程序的开发难度，同时有利用bug的调试。通过实现一种快速、有效地CUDA并行程序开发的方法，提高CUDA并行程序开发效率，降低CUDA并行程序开发周期和难度。

以上所述仅为本发明的实施例而已，凡在本发明的精神和原则之内，所作的任何修改、等同替换、改进等，均应包含在本发明的保护范围之内。

Claims (5)

1.一种快速开发CUDA并行程序的方法，其特征在于其具体操作过程为：

步骤一、CPU串行程序修改，即按CUDA程序格式要求对CPU串行程序的分析和修改，其中

1.1）CPU串行程序分析具体包括：利用打印时间方式测试串行程序中的热点模块；根据算法特点和数据特点分析热点模块是否可以并行，是否可以采用CUDA细粒度并行；根据可并行的模块，找到CUDA内核将会使用到的数组，并对数组结构进行分析；

1.2）仿CUDA格式的CPU串行程序修改，具体包括：原程序算法修改，修改成可并行的代码；数组修改，修改成适合CUDA并行程序格式的数组形式；

步骤二、CUDA并行程序移植，即设计CUDA并行程序把CPU程序移植到GPU平台，以及CUDA并行程序的优化，其中

2.1）设计CUDA并行程序具体包括：线程块、网格的设计，用于数据的划分和内核的计算；通信函数的实现，用于进行CPU与GPU之间的数据传递；内核的设计，实现CUDA并行加速热点模块；

2.2）所述优化CUDA并行程序，具体包括：将上述CUDA并行程序利用优化技术进一步提高并行程序的性能，主要优化包括2个方面：通信优化和内核优化。

2.根据权利要求1所述的一种快速开发CUDA并行程序的方法，其特征在于：所述步骤1.1）的详细过程为：

热点测试是指根据时间的测试结果确定热点函数，作为后面移植的重点代码模块；

找出热点代码后，需要分析热点部分的算法、数据特点，根据算法和数据的特点分析其是否可以并行，是否可以采用CUDA细粒度并行，确定其并行性；

根据对串行程序的分析，确定哪些模块需要移植到GPU平台上运行，对于需要运行上GPU平台上的代码内的数据进行分析，确定数组在代码中的什么位置传递到CUDA内核中，传递的方向是CPUtoGPU还是GPUtoCPU，以及每次传递时的数据大小等信息。

3.根据权利要求1或2所述的一种快速开发CUDA并行程序的方法，其特征在于：所述步骤1.2）的详细过程为：

对于CPU串行程序，根据并行算法的要求修改原程序，改成可并行的模式；同时需要重新设计并行算法；

根据CUDA并行程序对数组格式的要求对串行程序中的数组进行修改，进而将对原CPU程序修改成一个仿CUDA格式的CPU串行程序。

4.根据权利要求3所述的一种快速开发CUDA并行程序的方法，其特征在于：所述步骤2.1）的详细过程为：

GPU数组设计、并行模型设计：设计GPU端数组，设计数组大小、类型；设计CUDA程序的线程并行模型：block和grid；

根据对原程序数组的分析，把CPU串行程序移植到GPU平台，根据热点模块的算法和代码实现CUDA内核代码。

5.根据权利要求4所述的一种快速开发CUDA并行程序的方法，其特征在于：所述步骤2.2）的详细过程为：

CPU与GPU通信优化，利用异步技术减少CPU和GPU之间通信的时间；CUDA内核优化，利用全局存储器合并访问、共享存储器、常量存储器、纹理存储器手段提高访存速度，利用指令流提高计算手段优化CUDA内核。

Images