近年来,随着人工智能、高性能数据分析和金融分析等计算密集型领域的兴起,传统通用计算已经无法满足对计算性能的需求,异构计算越来越引起学术界和产业界的重视。
异构计算是指采用不同类型的指令集和体系架构的计算单元组成系统的计算方式。相比传统CPU,异构计算可以实现更高的效率和更低的延迟。目前的异构计算引擎主要有图形处理器(GPU,GraphicsProcessingUnit)、现场可编程门阵列(FPGA,FieldProgrammingGateArray)、专用集成电路(ASIC)等。
当前的通用CPU设计得已经很复杂,配有几十个核心,运行频率高达几GHz,每个核心有自己的独立缓存。通常CPU已具备一级、二级、三级缓存。而GPU是目前科研领域比较常用的硬件计算工具。GPU的计算核心数通常是CPU的上百倍,运行频率尽管比CPU的低,但是核心数量多,整体性能好。所以,GPU比较适合计算密集型应用,比如视频处理、人工智能等,现在传统的科学计算、工程计算等也开始越来越适合在GPU上运行。相比来说,CPU的缺点就是太通用了,数据读写、计算、逻辑等各种功能都得照顾,反而影响了计算性能。
通用图形处理器(GPGPU,GeneralPurposeGraphicsProcessingUnit)最早由NVIDIA公司的MarkJ.Harris于2002年提出。基于图形渲染管线的流水线特征,GPU本质上是一个可同时处理多个计算任务的硬件加速器。由于GPU中包含了大量的计算资源,MarkJ.Harris自2002年就开始尝试在GPU上做通用并行计算方面的研究。在此阶段,由于架构及编程平台的限制,研究人员采用将目标计算算法转换为图形运算算法的方式,使用GPU来实现通用并行计算需求。
NVIDIA公司提出Tesla统一渲染架构以及CUDA(ComputeUnifiedDeviceArchitecture,计算统一设备架构)编程模型后,NVIDIA公司的GPU开始了对通用并行计算的全面支持。在CUDA提出近两年之后,开放计算语言标准OpenCL1.0发布,这标志着利用GPU进行通用并行计算已基本成熟。目前市场上应用甚广的GPU芯片除了完成高质量的图形渲染之外,通用并行计算也已经成为一个主流应用。GPGPU在各个方面得到了不同GPU厂家为GPU通用计算提供的编程模型与平台,如CUDA和OpenCL,这些编程模型在C/C++基础之上做了面向大规模通用并行计算的语法扩展,为程序员提供了更好的、面向GPU的编程接口。
GPGPU通常由成百上千个架构相对简易的基本运算单元组成。在这些基本运算单元中,一般不提供复杂的诸如分支预测、寄存器重命名、乱序执行等处理器设计技术来提高单个处理单元性能,而是采用极简的流水线进行设计。每个基本运算单元可同时执行一至多个线程,并由GPGPU中相应的调度器控制。GPGPU作为一个通用的众核处理器,凭借着丰富的高性能计算资源以及高带宽的数据传输能力在通用计算领域占据了重要的席位。虽然各个GPGPU厂商的芯片架构各不相同,但几乎都是采用众核处理器阵列架构,在一个GPU芯片中包含成百上千个处理核心,以获得更高的计算性能和更大的数据带宽。
GPU中执行的线程对应的程序通常成为内核(kernel),这与操作系统中的内核是完全不同的两个概念。除此之外,GPU中执行的线程与CPU或者操作系统中定义的线程也有所区别,GPU中的线程相对而言更为简单,所包含的内容也更为简洁。在GPU众核架构中,多个处理核心通常被组织成一个线程组调度执行单位,线程以组的方式被调度在执行单元中执行,如NVIDIA的流多处理器、AMD的SIMD执行单元。同一个线程组中的线程执行相同的程序指令,并以同步的方式执行,每个线程处理不同的数据,实现数据级并行处理。不同GPU架构对线程组的命名各不一样,如NVIDIA将线程组称为warp,AMD将线程组称为wavefront。线程组中包含的线程数量各不相同,从4个到128个不等。除此以为,线程组的组织执行模式也各不相同,常见的执行模式有SIMT(SingleInstructionMultipleThreads,单指令多线程)执行模式和SIMD(SingleInstructionMultipleData,单指令流多数据流)执行模式两种。
了解CPU和GPU之间差异的一种简单方法是比较它们处理任务的方式。CPU由几个针对顺序串行处理优化的内核组成,而GPU具有大规模并行架构,由数千个更小、更高效的内核组成,旨在同时处理多个任务。
在GPU上解决计算问题原则上类似于使用多个CPU解决问题。手头的任务必须拆分为小任务,其中每个任务由单个GPU内核执行。GPU内核之间的通信由GPU芯片上的内部寄存器和内存处理。CUDA或OpenCL等特殊编程语言不是使用消息传递进行编程,而是提供主机CPU之间的数据交换和同步GPU内核的机制。
一个现代超级计算系统实际上可能由大量节点组成,每个节点包含2到32颗常规CPU以及1到16个GPU。通常还会有一个高速网络和一个数据存储系统。该系统的软件可以使用传统编程语言(如C/C++、Fortran等)的组合编写,结合用于CPU并行化的消息传递系统以及用于GPU的CUDA或OpenCL。所有这些组件都必须进行调整和优化,以实现整个系统的最佳性能。
CUDA是由NVIDIA公司推行的一套并行编程框架,目前只有NVIDIA的GPU支持该框架,其开发语言主要为CUDAC。作为一种GPU的并行开发语言,CUDA的API涉及设备管理、存储管理、数据传输、线程管理、事件管理等功能。
CUDA的存储模型主要分为全局存储(globalmemory)、局部存储(localmemory)、共享存储(sharedmemory)、常量存储(constantmemory)和纹理存储(texturememory)等存储类型。不同的存储类型,其存储容量、可见程度、读写速度差异巨大,需要在程序设计中根据各自特点和应用问题的需求合理调配。
CUDA的编程模型和执行模型按照层次结构分层设计。CUDA的执行模型由3个层级组成,最基础的执行单位是线程(thread),多个线程组成一个线程块(block),多个线程块形成线程网格(grid)。
CUDA编程模型作为一个异构模型,其中使用了CPU和GPU。在CUDA中,主机(host)指的是CPU及其存储器,设备(device)是指GPU及其存储器。在主机上运行的代码可以管理主机和设备上的内存,还可以启动在设备上执行的内核函数(kernel)。这些内核由许多GPU线程并行执行。
鉴于CUDA编程模型的异构性,CUDAC程序的典型操作序列是:
CUDA应用程序的源代码由传统的C++主机代码和GPU设备函数混合组成。CUDA编译过程将设备函数与主机代码分开,使用专有的NVIDIA编译器和汇编器编译设备函数,使用可用的C++主机编译器编译主机代码,之后将编译过的GPU函数嵌入到主机对象(object)文件。在链接阶段,向最终生成的可运行二进制文件添加特定的CUDA运行时库函数,如支持远程SPMD过程调用的运行时库函数、显式GPU操作的运行时库函数(如分配GPU内存缓冲区,主机与设备之间的数据传输等)。
编译过程涉及每个CUDA源文件的拆分、编译、预处理和合并步骤。为了将上述复杂的编译过程向开发人员隐藏,NVIDIA公司设计了CUDA编译器引擎程序nvcc。nvcc接受一系列常规编译器选项,例如宏定义和头文件、函数库路径设置,支持编译过程的组合。所有非CUDA编译步骤都被转发到nvcc支持的主机C++编译器,编译选项到主机C++编译选项的转换也由nvcc自动完成。
预定义宏
含义
__NVCC__
编译C/C++/CUDA源文件时预定义
__CUDACC__
编译CUDA源文件时预定义
__CUDACC_VER_MAJOR__
NVCC主版本号
__CUDACC_VER_MINOR__
NVCC次版本号
__CUDACC_VER_BUILD__
NVCC编译版本号
输入文件后缀
描述
.cu
CUDA源文件,包含主机代码和设备函数
.c
C源文件
.cc,.cxx,.cpp
C++源文件
.o,.obj
目标文件(objectfile)
.a,.lib
库文件(libraryfile)
.res
资源文件(resourcefile)
.so
共享目标文件(sharedobjectfile)
选项
-ofile
配置输出文件名和路径
-llibrary,…
配置链接阶段链接的库文件
-Ddef,…
定义预处理阶段使用的宏
-Udef,…
取消宏定义
-Ipath,…
配置头文件搜索路径
-Lpath,…
配置库文件搜索路径
-cudart{none|shared|static}
配置CUDA运行时库的类型,默认使用静态(static)
-cudadevrt{none|static}
配置CUDA设备运行时库的类型,默认使用静态(static)
-pg
生成供gprof使用的可执行代码
-g
编译带调试信息的主机代码
-G
编译带调试信息的设备代码
-Olevel
指定主机代码的优化级别
-doptkind
允许设备端代码优化。当不指定-G选项时,设备端代码优化是默认的行为
-shared
生成共享库
-x{c|c++|cu}
显式指定待编译的输入文件的编程语言,而不是由编译器根据文件后缀自动判断
-std{c++03|c++11|c++14|c++17}
指定c++标准的版本
下表列举了可以直接传递给nvcc封装的内部编译工具的编译选项。通过这些选项的应用,nvcc不需要具备对内部编译工具的过多细节的了解。
-Xcompileroptions,…
指定直接传递给编译器/预处理器的编译选项
-Xlinkeroptions,…
指定直接传递给链接器的编译选项
-Xarchiveoptions,…
指定直接传递给库管理器的编译选项
-Xptxasoptions,…
指定直接传递给ptxas(PTX优化汇编器)的编译选项
-Xnvlinkoptions,…
指定直接传递给nvlink(设备链接器)的编译选项
-arch{arch|native|all|allmajor}
指定编译阶段使用的虚拟GPU类型
-codecode,…
指定汇编优化使用的具体GPU类型
-use_fast_math
使用快速数学计算库
为了实现架构演进,NVIDIA的GPU以不同的世代(generation)发布。新一代产品在功能和/或芯片架构方面进行重大改进,同时同一代产品中的GPU型号仅在配置方面存在次要差别,对功能、性能的影响适中。不同代的GPU其应用程序的二进制兼容性是没有保证的。例如,为FermiGPU编译的CUDA应用程序很可能无法在KeplerGPU上运行(反之亦然)。这是因为每一代的指令集和指令编码与其他世代的指令编码都不相同。同一代的GPU由于共享相同的指令集,在满足特定条件下其二进制兼容性可以得到保证。特定条件通常是指两个没有功能差异的GPU版本之间的情况(例如,当一个版本是另一个版本的缩减版),或者当一个版本在功能上完全包含在另一个版本中。后者的一个例子是基础Maxwell版本sm_52,其功能是所有其他Maxwell版本的一个子集:任何为sm_52编译的代码将可以在所有MaxwellGPU上运行。
nvcc编译命令总是使用两个架构:一个虚拟的中间架构,加上一个真实的GPU架构(指定代码将运行的平台)。要使nvcc命令有效,真实架构必须是虚拟架构的实现。
虚拟GPU完全由提供给应用程序的能力和特征定义。虚拟架构提供了一个通用的指令集合,并且不涉及二进制编码格式。虚拟架构列表如下:
虚拟架构(-arch参数)
特征描述
真实架构(-code参数)
下面介绍一个简单的CUDAC程序例子,演示如何在瀚海22上编译运行CUDA代码。
例子展示的是两个向量相加的CUDA代码add.cu。
#include float*x,*y,*z,*d_x,*d_y,*d_z;x=(float*)malloc(N*sizeof(float));y=(float*)malloc(N*sizeof(float));z=(float*)malloc(N*sizeof(float));cudaMalloc(&d_x,N*sizeof(float));cudaMalloc(&d_y,N*sizeof(float));cudaMalloc(&d_z,N*sizeof(float));指针x、y和z分别指向使用malloc分配的主机内存空间,d_x、d_y和d_z指针分别指向使用CUDA运行时APIcudaMalloc函数分配的设备存储空间。CUDA中的主机和设备有独立的内存空间,这两个空间都可以从主机代码进行管理。 为了初始化设备数组,使用cudaMemcpy将数据从x和y复制到相应的设备数组d_x和d_y,它的工作方式与标准的Cmemcpy函数一样,只是增加了第四个参数,指定拷贝的方向。在这里,我们使用cudaMemcpyHostToDevice指定第一个(目标)参数是设备指针,第二个(源)参数是主机指针。 cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice);cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice);add内核由以下语句启动: add<<<(N+256-1)/256,256>>>(N,d_x,d_y,d_z);<<<和>>>符号之间的信息是执行配置,指示有多少设备线程并行执行内核。在CUDA中,软件中有一个线程层次结构,它模仿线程处理器在GPU上的分组方式。执行配置中的第一个参数指定网格中线程块的数量,第二个参数指定线程块中的线程数。线程块和网格可以通过为这些参数传递dim3(一个由CUDA用x、y和z成员定义的简单结构)值来生成一维、二维或三维的线程块和网格。对于add这个示例,只需要一维线程组,所以我们只传递整数。在本例中,我们使用包含256个线程的线程块启动内核,并使用“上整计算”来确定处理数组全部N个元素所需的线程块数((N+256-1)/256)。 由于数组的元素数有不能被线程块大小整除的可能,内核代码必须检查内存访问是否越界。 在运行内核之后,使用cudaMemcpy(拷贝方向:cudaMemcpyDeviceToHost),从d_z指向的设备数组复制到z指向的主机数组,将结果返回给主机。 cudaMemcpy(z,d_z,N*sizeof(float),cudaMemcpyDeviceToHost);程序的最后,使用cudaFree()和free()分别清理设备端和主机端申请的内存。 备注 本小节的简单例子用于展示瀚海22超级计算系统上CUDA程序的编译运行(交互式)的一般流程。 瀚海22超级计算系统系统预装了NVIDIAHPCSDK(nvhpc),方便用户充分利用GPU计算资源,实现科学计算应用和性能优化等工作。 NVIDIAHPCSDK是NVIDIA公司提供的一个包含编译器、函数库和软件工具的软件包,包含了用于方便用户开发、增强程序性能和可移植性的一系列工具。NVIDIAHPCSDKC、C++和Fortran编译器支持使用标准C++和Fortran、OpenACC指令和CUDA对HPC建模和模拟应用程序进行GPU加速。GPU加速的数学库使普通HPC算法的性能最大化,优化的通信库使基于标准的多GPU和可扩展系统编程成为可能。性能剖析和调试工具简化了HPC应用程序的移植和优化,而容器化工具则使得在企业内部或在云端的部署变得容易。HPCSDK支持NVIDIAGPU和运行Linux的Arm、OpenPOWER或x86-64CPU,为用户提供了构建NVIDIAGPU加速的HPC应用程序所需的工具。 陈国良,吴俊敏.并行计算机体系结构(第2版)[M].北京:高等教育出版社,2021.