一种基于CUDA标准的异构并行编程模型开发简介

最新推荐文章于 2022-11-10 18:45:09 发布

元夕0643

最新推荐文章于 2022-11-10 18:45:09 发布

阅读量625

点赞数 1

分类专栏：项目经验分享文章标签： cuda 并行计算编译器 opencl 编程语言

本文链接：https://blog.csdn.net/shuiliuyun0643/article/details/114434894

版权

项目经验分享专栏收录该内容

1 篇文章 0 订阅

订阅专栏

一种基于CUDA标准的异构并行编程模型开发简介

一、绪论

1.1研究背景及意义

程序的移植有源代码移植和二进制翻译等多种方法。由于现有的大多数串行机都遵从冯·诺依曼结构，因此串行程序移植的重点都放在机器指令翻译的二进制翻译方面，在源代码级不需要做过多的修改。而并行系统的体系结构差异很大，有多核、异构多核、众核、异构众核，这种较大的差异导致了并行程序的移植变得非常复杂。同时，并行编程模型滞后于并行系统的现状也促使研究人员研究相同的程序在不同系统中所获得的效能。因此，研究并行程序的移植变得很有意义[^1]。

NVIDIA公司率先提出了统一的计算设备体系结构（CUDA, Compute Unified Device Architecture）。这一CPU + GPU计算架构中，提供了类似C语言的开发环境，允许设计人员使用C语言和CUDA扩展库的形式编写程序。CUDA是一种CPU和GPU代码混合的显式异构并行编程模型，CPU代码和GPU代码相互分离，采用分层的线程和存储层次，利用线程的快速切换实现了大规模并行线程的快速执行。这种并行机制能够充分利用 GPU 硬件上的众多计算核心与存储结构，简化控制。从而降低了用户程序开发的复杂度，提高了开发效率[^1]。

CUDA 编程架构的出现加速了并行编程模型的发展，它的编程方式符合 GPU 的特点。国内外对它能否移植到已有的多核平台上并取得较好性能这一问题已有了初步的研究。

1.2目标平台体系结构简介

图1 M3A体系结构简图

如上图所示，目标平台M3A为ARM+DSP的异构平台，用户在使用M3A平台进行编程时，需要对device驱动api比较熟悉，在用户程序中显示的调用驱动api对DSP进行操作，这无疑提高了用户编程的难度。

为丰富我们的M3A软件生态，给用户提供更好的编程环境，基于CUDA标准，我为M3A开发了一套通用可移植的异构编程模型（HPPA，Heterogeneous Parallel Programming Architecture）。其简要设计实现如下文。

二、HPPA基本组成结构

图2 CUDA 软件体系结构图

如上图CUDA 软件体系结构图所示，用户开发CUDA应用程序时，可以基于CUDA函数库、CUDA runtime API、CUDA Driver API三种API来开发应用程序，用户为提供开发效率，自然是基于CUDA函数库、CUDA runtime API来开发应用程序最为高效。

HPPA基于CUDA标准，同样采用CUDA函数库、CUDA runtime API、CUDA Driver API三级开发流程来完善M3A软件生态。

图3 HPPA runtime API 软件层基本结构简图

如上图所示，HPPA 程序主要包括：
用关键字__global__标识的在host调用，在device执行的vecadd（）函数；
以及在host上执行的main函数（其省略了__host__关键字，无关键字函数默认在host上执行）。
程序执行需要调用host端编译器编译main（）函数部分，调用device端编译器编译vecadd（）函数部分。故HPPA 程序能够运行，首先是基于编译器，所以我们需要先开发HPPA 编译器 HPCC及其他相应的工具链。

三、编译工具链开发

HPPA 源程序是一个混合了host及device代码的.cu文件，以vecadd为例，输入vecadd.cu文件，我们首先需要将其拆分为vecadd_host.cpp及vecadd_device.cpp两个文件。

图4 HPPA程序编译流程

3.1 拆分工具HPCufe开发

LLVM工程提供了丰富的集成库用于各种功能：
1 libsupport 来自于 LLVM 的基本支撑库
2 libsystem 来自于 LLVM 的系统抽象库
3 libbasic 用于输入资源文件的诊断、资源定位、资源缓冲抽象、文件系统快存
4 libast 表示 C 抽象语法树、C 类型系统、内置函数等的类库，也提供多种分析和操作抽象语法树的对象
5 liblex 词法分析与预处理、标识符哈希表、语用处理、记号以及宏指令扩展
6 libparse 语法分析
7 libsema 语义分析
8 libcodegen 从抽象语法树转换到 LLVM 中间表示，用以优化和产生代码
9 librewrite 文本缓冲的编辑
10 libanalysis 静态分析支撑[^2].

由于我一直是开发llvm后端，对前端的文件处理分析并不了解，故在此处使用了一种更加简单粗暴的拆分实现思路。
1）读取vecadd.cu每一行，创建文本数据包基本块BB，根据关键字为每个BB添加属
性，例如空行设置属性为Empty ，“//”注释设置属性为DNote，其他例如host头文件、device头文件，kernel函数声明，kernel函数体，kernel函数launch部分等等，具体关键字、类别及实现细节不在细说。
2）根据每个属性，寻找BB结束位置，完成一个BB创建。
3）在BB创建过程中，对kernel函数声明及函数实现进行检查，并获取kernel函数名及参数类型列表vv_typelist（vector<vector>类型，用于存放多个kernel信息）。
4）在kernel函数launch部分的BB中，获取kernel函数名及参数列表vv_paralist。
5）重写kernel函数launch部分，例如，将字符串：

vecadd<<<blocksPerGrid, threadsPerBlock>>>(d_x, d_y, d_z, N);

重写为：

hpSetupArgument(d_x, 0);
hpSetupArgument(d_y, sizeof(d_x));
hpSetupArgument(d_z, sizeof(d_x)+sizeof(d_y));
hpSetupArgument(N, sizeof(d_x)+sizeof(d_y)+sizeof(d_z));
hpLaunch("vecadd");

6)根据每个属性，将完成的BB输入到host端文件vecadd_host.cpp.i或者device端文件vecadd_device.cpp文件中。
7)维护一个全局的vector<kernel_msg_t> kernel_list，用于存放kernel信息，将
vv_paralist及vv_typelist中的信息提取到kernel_list。
8)拆分工作基本完成。

3.2 HPfrontend

这个工具做一些简单的宏替换，死代码消除之类的优化工作，并将HPPA头文件定义的一些数据类型，builtins函数等放到vecadd_host.cpp.i文件开头，并生成可以供给host编译器编译的vecadd_host.cpp文件。

3.3 device端工具链开发

我们Device端（DSP）的编译工具链有2套，一是基于GCC开发的，还有基于LLVM开发的编译工具链。完整的device工具链包括编译器gcc/clang，汇编器as,连接器ld，其他的ar、nm、objdump、readelf等等工具。同时还包括llvm工程的一些独有的工具链工具如opt。

在编译vecadd_device.cpp时，可以一步直接生成最后的设备代码描述符文件vecadd_device.fatbin.c，也可以先使用clang++将其编译成vecadd_device.bc，再使用opt工具，对其执行特定的优化pass，得到优化后的vecadd_device.bc，最后再调用其他工具，最终生成vecadd_device.fatbin.c。

3.4 start.asm

vecadd函数在device上执行时，它不像使用编译器编译生成的普通程序，存在主调函数，所以它执行时需要的参数并没有存放布置好在相应的寄存器以及堆栈，如此vecadd函数执行时，从相应的寄存器与堆栈取不到正确的值。所以，在vecadd函数执行之前，需要先执行一个入口程序，这个程序模拟编译器布置函数参数的过程，将参数值从与host约定好的存放args的shared mem 中取出参数，存放到相应的寄存器与堆栈中，而后再修改pc指针与dp指针，让vecadd能够获得正确的参数并执行完成。

将该程序写成一个start.asm，编译成start.o，在编译vecadd_device.cpp生成vecadd_device.o之后，与vecadd_device.o，需要的库文件等一起编译生成vecadd_device.cubin文件。

start.asm代码由同事提供。

3.5 M3fatbinary生成与设备代码描述符文件生成

NVPTX架构下有许多不同计算能力的架构比如sm_20,sm_30…，使用cuda的nvcc生成的kernel程序二进制文件会有不同计算力的多个版本的.cubin文件，之后，CUDA会把这些.cubin文件打包成一个.fatbin文件，这样的好处是，若最开始设置的GPU架构为sm_60，但是在实际执行时，硬件架构达不到sm_60这个版本，如此，显卡驱动可以从fatbin中选取符合硬件版本的.cubin程序，而不需要再重新编译整个CUDA程序。

我们的异构平台目前并没有多个算力多个版本，但为了统一标准以及为可能存在的后续版本预留位置，我保留了fatbin的名字。

生成fatbin文件之后，就是生成设备代码描述符文件，也就是一个文本文件vecadd_device.fatbin.c。

生成.fatbin.c文件主要需要做的如下：

1）打开vecadd_device.fatbin二进制文件，维护一个全局变量vector<section_msg_t> sec_list，用来存放从vecadd_device.fatbin二进制文件的文件头中解析获得的section信息。
section_msg_t数据类型定义如下：

typedef struct 
{
  string   secName;
  uint64_t secAddr;
  int      secSize;
  int      secOffset;
}section_msg_t;

注：sec_list是提供给runtime api用来load kernel程序时使用的，所以secOffset指的是该section在vecadd_device.fatbin.c中相对与第一个section的偏移。而不是vecadd_device.fatbin二进制文件中获取的section偏移。

2）遍历每一个section，根据需要，将相应的段内容取出，以ansic码形式输出到vecadd_device.fatbin.c文本文件中。其中，.text、.data及section flags为SHF_ALLOC和section type 为SHT_PROGBITS的段为程序执行所必须，同时，还需将与host约定好，存放kernel函数参数的section，例如“_cl_args_addr”取出，输出到文本文件中。同时，若需要调试kernel程序，还需将debug相关的段存入文本文件。
输出的vecadd_device.fatbin.c文本文件头及尾应该如下：

文件头：

#include "fatBinaryCtl.h"
#define __CUDAFATBINSECTION  ".nvFatBinSegment"
#define __CUDAFATBINDATASECTION  ".nv_fatbin"
asm(
".section .nv_fatbin, \"a\"\n"
".align 8\n"
"fatbinData:\n"
".quad 0x00100001ba55ed50,0x00000000000008a0,0x0000000000000618\n"
......

文件尾：

extern const unsigned long long fatbinData[278];
static const __fatBinC_Wrapper_t __fatDeviceText __attribute__ ((aligned (8))) __attribute__ ((section (__CUDAFATBINSECTION)))= 
	{ 0x466243b1, 1, fatbinData, 0 };
......

如此，只需在fatBinaryCtl.h头文件中将__fatDeviceText变量声明为extern，则runtime就可以通过__fatDeviceText获取fatbinData地址，进而得到kernel程序执行需要的代码。同时得到程序入口地址0x466243b1。

注：此处也有另外一种做法，即并不将vecadd_device.fatbin转换为vecadd_device.fatbin.c文本文件，或者转换为vecadd_device.fatbin.c之后，接下来3.7节编译host代码时，并不将vecadd_device.fatbin.c与vecadd_host.cpp一起编译，提供fatbinData给runtime，而是在runtime中，launch kernel，执行kernel之前，从vecadd_device.fatbin或vecadd_device.fatbin.c文件中临时提取kernel执行需要的信息及代码。但这种方式明显会增加用户程序运行时间，所以我尽可能的将能够做的工作放到编译时，从而提升运行时的效率。

注：此工具参考同事的工作完成。

3.6 关键信息传递

runtime在执行kernel程序时，需要一些关键信息，这些关键信息在编译vecadd_device.cpp以及将其转换为vecadd_device.fatbin.c时，已经存入全局变量kernel_list及sec_list，我们需要编译时将其传递给runtime。

我采取文件传递的方式。在vecadd_device.fatbin.c生成完成之后，再生成一个kernel_msg_init.cpp的文件，用来将kernel_list及sec_list传递给runtime。

首先，需要将#include “vecadd_device.fatbin.c” 输入到kernel_msg_init.cpp文件中。
而后，还需要这些字符串输入到kernel_msg_init.cpp文件中：

	"#include \"kernel_msg.h\" \n\n"
    "vector<kernel_msg_t> kernel_list;\n"
    "fatbin_msg_t fatbin_info;\n"
    "void init_kernel_msg(){ \n\n"
    "\targs_t arg;\n"
    "\tsection_msg_t section;\n"
    "\tkernel_msg_t kernel;\n";
......

之后，就是将kernel_list及sec_list的值显示打印输入到kernel_msg_init.cpp，将信息传递给runtime的kernel_list及sec_list对象。runtime就获得了在执行kernel程序时，需要的这些关键信息。

3.7 可执行代码生成

生成可执行代码没什么好说的，直接调用host编译器，将需要的头文件，库，路径设置好，编译链接即可。

3.8 编译引导程序hpcc

之前的3.1-3.7已经将生成HPPA程序可执行文件的过程分步列出，编译引导程序hpcc就是分别调用之前的工具以及实现每个过程，中间使用一些类似opt_keep、opt_alias、opt_tmpdir、opt_builtin、opt_llvm、opt_gnu的opt选项控制每个过程，若用户不输入参数，则直接生成可执行程序：

hpcc  vecadd.cu  -->  vecadd.out

若用户输入各种参数，-k，-g，–TOOLS=GNU --MODE=FATBIN，等等，分别控制相应的过程，选择不同分支，产生相应的输出文件。

同时，hpcc基于host与device的编译工具链，编译host与device程序时，还需要相应的头文件，链接相应的库，故需要先安装好host的编译工具链与device的交叉编译工具链，并将其路径添加到环境变量，再添加一个runtime库路径的环境变量：

export HOST_HPPA_CGT_INSTALL=/usr/share/.../arm_cgt
export DEVICE_HPPA_CGT_INSTALL=/usr/share/.../m3_cgt
export M3A_HPPA_INSTALL=/usr/share/.../hppa

同时，在各个过程中间，可以添加一些编译器未做的优化，或者针对特定的数据结构，特定的硬件功能单元，做一些特殊的优化。例如，设备上存在向量处理单元VPU，其由8个同构VPE构成，那么在编译device程序时，可以对数据进行自动向量化优化，将数据组织成如 int vec[8]，float vec[8]等等类似的向量数据类型，调用相应的向量算法、函数库，执行程序，提高程序性能。

3.9 编译hpcc工程

编译hpcc工程由分布在hpcc工程各个目录下的CMakeLists.txt文件指定，最后将编译好的工具链install到环境变量M3A_HPPA_INSTALL指定的目录下的bin文件夹下即可。为了hpcc使用方便，也可将其路径设置为全局PATH。

如何编写一个工程的CMakeLists.txt文件在此不做细说。

四、hppa-rt库开发

Hpcc编译工具链开发完成之后，整个异构程序执行框架也就大致建立起来了，runtime库开发也就很简单了。

Runtime库开发主要是包括以下这些内容：

host端上：一个平台M3A（platform object）、多个设备（ARM+DSP，ARM在操作系统看来也是一个device，device object）、多级存储（memory objects）、一个设备程序（program object）、一或多个kernel函数（kernel object）、命令队列（commandQueue object）。

Device端上：移植到嵌入式dsp平台上的标准库libc.a，数学库libm.a，提高device程序性能的builtins库，其他一些高性能库，如多核间通信，算法库之类的，如果有的话。（我是之前将MPI通信库移植到了dsp平台上，让其作为一个用在多核间通信，算法优化的高性能库提供给dsp用户使用）。Device端这些库最终整合为一个libdev.a静态库提供给hpcc编译器链接device程序使用，这些库的开发在此不多赘言。

4.1 Platform 对象

Platform主要包含2个成员对象 p_devices、p_shmFactory，以及相应的方法来获取device informations和Shared Memory informations。

4.2 device 对象

Device 对象包括2个部分：CPUDevice及DSPDevice，（将host也作为一个device对象来管理。）CPU主要由操作系统来管理，所以runtime中CPUDevice主要用作获取一些informations的作用，重要的是DSPDevice对象。
DSPDevice对象主要在驱动模块devdrv.ko的基础上，对dsp设备进行管理。

其主要包含这些对象：

	SharedMemory*                   p_shmHandler;
	profiling_t                     p_profiling;
	std::set<uint8_t>               p_compute_units;
	......

由其子类DSPDeviceManager实现如下功能：

virtual bool DeviceInit()   const  = 0;
virtual bool DeviceReset()  const  = 0;
virtual bool DeviceLoad()  const  = 0;
virtual bool DeviceConfig()  const = 0;
virtual bool DeviceRun()   const  = 0;
virtual bool DeviceStop()   const  = 0;
......

DSPDevice对象还需实现如下功能：

std::set<uint8_t> const          GetComputeUnits()    const {...}
    	virtual void             pushEvent(Event* event) = 0;
......
    	virtual bool             mail_query()            = 0;
    	virtual Msg_t*           mail_from(const uint8_t dev_id,
                                       int* retcode = nullptr) = 0;
    	virtual void             mail_to(Msg_t& msg,
                                     const uint8_t dev_id) = 0;
......

在DSPDevice对象中使用虚函数virtual 是为了使用不同版本的驱动，利用C++的多态性，我们可以使用基于PCIE、JTAG、SOCKET或基于操作系统ioctl等方式的驱动，在DSPDevice的派生类中override该虚函数，以实现不同版本的DSPDevice方法。

4.3 memory 对象

1）首先对异构平台所有的内存进行分类：

    enum class Kind { 
MEM_LOCAL, /*  */
MEM_SHARED, 
MEM_CONST, 
MEM_GLOBAL };
    enum class Location {ONCHIP, OFFCHIP};

2）而后对内存布局进行一个划分，不同的内存用以不同的数据：

class MemoryRange
{
    DSPDevicePtr64 GetBase()     const { return start; }
    uint64_t       GetSize()     const { return size; }
    Kind           GetKind()     const { return kind; }
    Location       GetLocation() const { return loc; }
}

......

3）对Shared Memory及Global Memory 做内存管理，实现如下功能：

class SharedMemory
{
    // Allocate/Free on chip shared memory
    virtual uint64_t  AllocateMSMC(size_t size) =0;
    virtual void      FreeMSMC(uint64_t addr) =0;

    // Allocate/Free off chip shared memory
    virtual uint64_t  AllocateGlobal(size_t size, bool prefer_32bit) =0;
    virtual void      FreeGlobal(uint64_t addr) =0;

    // Determine which memory the pointer was allocated from and free it
    virtual void      FreeMSMCorGlobal(uint64_t addr)  =0;

    // Allocate/Free functions with ability to query whether an address
    // has been allocated by the clMalloc method
    virtual void* clMalloc     (size_t size, MemoryRange::Location l) =0;
    virtual void  clFree       (void* ptr) =0;
    virtual bool  clMallocQuery(void* ptr, uint64_t* p_addr, size_t* p_size) =0;
......
}

类似DSPDevice对象，使用虚函数virtual ，基于不同版本的驱动实现以上功能。

4）创建class MemObject基类：

class MemObject 
{
 int32_t init();
 bool allocate(DeviceInterface *device); /*!< \brief Allocate this memory object on the given \p device */
 size_t size() const = 0;                /*!< \brief Device-independent size of the memory object */
  Type type() const = 0;                  /*!< \brief Type of the memory object */
  void *host_ptr() const;                         /*!< \brief Host pointer */
  DeviceBuffer *deviceBuffer(DeviceInterface *device) const; /*!< \brief \c Coal::DeviceBuffer for the given \p device */
  DeviceBuffer *deviceBuffer(SharedMemory *shm) const;
......
}

5）基于class MemObject，再派生出不同数据结构的memory子类：

class Buffer : public MemObject
{...}
class SubBuffer : public MemObject
{...}
class Image2D : public MemObject
{...}
class Image3D : public Image2D
{...}

4.4 program对象

program对象用来标记程序状态State，类型Type，并控制程序编译build，加载load，运行run。

因为采用fatbin模式，程序已经离线编译好，并作为可执行程序的一个.nv_fatbin 数据段存在。HPPA并没有在线模式，所以build在此处并没用。

Load则为从fatbinData中取出程序，按照sec_list中的section informations，将程序加载到指定位置。
Run则是控制device执行kernel程序。

主要为调用device对象的方法来实现以上功能。

4.5 kernel对象

Kernel对象主要包含arg对象，arg对象标记kernel函数的args位置，类型，并提供对arg的一些操作方法，比如:refineKind()。

由于host与device体系结构差异，host上的数据结构与device上可能并不一致，例如，host上的long类型可能是32位，而device上的long类型可能是64位，故需要对args的type进行host到device的转换。

Kernel对象还实现了一系列对args的操作方法，如setArg()，实现将args部署到与dsp约定好的args_addr。

Kernel对象还需对任务空间进行布置，例如执行如下kernel：

vecadd<<<100, 256>>>(d_x,d_y,d_z,N)；

一个kernel在一个device上执行，针对M3A体系结构，将block映射为超节点，thread映射为核，则需按用户程序指定的blocksPerGrid = 100，threadsPerBlock = 256，将kernel任务均匀的布置到计算节点core上。

4.6 command queue对象

CommandQueue对象包含Event成员，主要是将各种操作转换为event对象，再将event对象放入CommandQueue，之后按in order 或 out of order模式来执行CommandQueue。

在class Event中，主要实现标记Event对象的状态State，类型Type，以及实现一些对Event操作的方法，比如插入计时函数timeing，用以做profiling。
CommandQueue对象则主要实现对Event的管理。

4.7 context对象

Context是上下文对象，类似通信域，主要实现一些关键信息在各个对象之间传递。例如kernel_list及sec_list就是Context的成员。

4.8 API接口

以上各部分完成之后，就是实现api接口，提供给用户使用。

使用cmake功能，完成hppa-runtime库编译生成。最后将生成的libhppart.so及libdev.a库install到环境变量M3A_HPPA_INSTALL指定路径下的lib64目录，将include中的头文件install到环境变量M3A_HPPA_INSTALL指定路径下的include目录，如此，hpcc在编译程序时，就可以链接相应的库。完成可执行程序的生成。

到此，整个异构并行编程模型开发完成，接下来就是测试及调优了。

总结

以上为我最近1年半的项目开发经验，在此做一个总结，一个记录，也算是重新学习的过程。由于个人能力有限，疏漏错误在所难免，希望各位指正。

在此，要感谢我的领导，给我充分的信任，让我从C++一知半解开始开发这个项目，也给我充足的时间，一点点查资料，按自己思路完成这个项目的开发。同时，也要感谢同事，给予技术上的一些帮助。

后续有时间会慢慢增加一些内容，或做一些修订，并且陆续将以前的一些项目开发经验分享出来，不一直做伸手党。

2021/3/6 元夕

参考文献

[1] 岳峰，庞建民，张一弛，余勇; 《CUDA 程序到 Cell平台的源代码移植》(计算机工程，2012）。
[2] 龚丹，苏小红，王甜甜；《Clang 编译平台优势分析》（智能计算机与应用
，2017）。

元夕0643

关注

1
点赞
踩
1

收藏

觉得还不错? 一键收藏
0
评论
一种基于CUDA标准的异构并行编程模型开发简介

一种基于CUDA标准的异构并行编程模型开发简介目录一、绪论1.1研究背景及意义1.2目标平台体系结构简介二、HPPA基本组成结构三、编译工具链开发3.1 拆分工具HPCufe开发3.2 HPfrontend3.3 device端工具链开发3.4 start.asm3.5 M3fatbinary生成与设备代码描述符文件生成3.6 关键信息传递3.7 可执行代码生成3.8 编译引导程序hpcc3.9 编译hpcc工程四、hppa-rt库开发4.1 Platform 对象4.2 device 对象4.3 mem
复制链接

扫一扫