CN114116208A - 一种基于gpu的短波辐射传输模式三维加速方法 - Google Patents

Abstract

本发明公开了一种基于GPU的短波辐射传输模式三维加速方法，通过对大气短波辐射传输模式RRTMG_SW进行基于GPU的三维并行计算以及性能优化，提升了RRTMG_SW的计算效率。该方法包括以下步骤：①三维并行计算：在CPU端为短波辐射过程定义、初始化数据；定义三维向量dim3类型的变量；为GPU端分配合适的内存空间来存储计算数据以及最终结果；将CPU内存中参与计算的数据复制到GPU的全局内存中；进行主机和设备端的数据传输并启动kernel函数开始三维并行计算；将GPU端的计算结果复制回CPU端。②算法性能优化：使用GPU寄存器降低计算过程中的全局内存访问时间；使用CUDA流覆盖数据传输与kernel计算过程。本发明有益效果：大大提升了RRTMG_SW的计算效率，使其更加适合使用GPU进行并行计算。

Description

一种基于GPU的短波辐射传输模式三维加速方法

技术领域

本发明涉及高性能计算技术领域，具体涉及一种基于GPU的短波辐射传输模式三维加速方法。

背景技术

太阳辐射是大气物理中的关键因素之一。能够合理准确地模拟来自地球表面的大气辐射分布及变化情况，对预测未来天气和气候变化有很大影响。大气辐射传输消耗了整个大气物理过程大量的计算资源，因此，开发出高精度和高速度的辐射传输模型对于全球大气建模非常有意义。快速辐射传输模型(RRTMG)是大气环流中计算长波和短波辐射通量的一个相关K分布模型。由于RRTMG的高准确性，已将其大量应用于天气预报和多种气候模型中。虽然RRTMG的计算速度相当快，但它仍然占据整个大气物理过程25％到35％的计算时间，所以非常有必要对其进一步加速。目前，已经开发出了基于CUDA C的短波辐射模式加速模型CC-RRTMG_SW，但是其仅使用了一维并行计算且没有针对并行算法进行性能优化，未充分发挥出使用GPU计算的优势，仍然有很大的优化空间。

RRTMG_SW由两个子例程mcica_subcol_sw和rrtmg_sw组成，其中rrtmg_sw是RRTMG_SW的驱动程序，计算耗时最多，所以主要对rrtmg_sw进行基于GPU的三维加速计算。在rrtmg_sw中，各个子例程的计算时间各异，其中spcvmc_sw耗时最长，占rrtmg_sw总计算时间的71.4％。对于这四个子例程，它们在水平、垂直、g-point以及jp-band维度的数据计算具有部分独立性，可以使用一维、二维和三维并行相结合的方式对其进行加速计算。

表1 RRTMG_SW各子程序计算时间(s)

发明内容

本发明的目的在于提供一种基于GPU的短波辐射传输模式三维加速方法，用以提高短波辐射传输模式RRTMG_SW的计算效率。

为实现上述目的，本发明的技术方案为：

步骤1.对RRTMG_SW中inatm_sw、clddprmc_sw、setcoef_sw以及spcvmc_sw四个子例程进行二维和三维并行化计算，并行后的kernel分别为inatm_d、cldprmc_d、setcoef_d、taumol_d以及spcvmc_d。

inatm_d:由于数据依赖性以及数据同步需求，需要将inatm_sw划分为五个kernel进行计算，分别为inatm_d1、inatm_d2、inatm_d3、inatm_d4以及inatm_d5。其中，inatm_d1、inatm_d2和inatm_d4在水平和垂直维度的计算无依赖性，可以进行二维并行计算；inatm_d3在水平、垂直以及g-point维度的计算无依赖性，可以进行三维并行计算；inatm_d5仅在水平维度无计算依赖性，可以进行一维并行计算。

cldprmc_d:cldprmc_sw用于传递三维辐射云属性参数，其计算部分在水平、垂直以及g-point维度均无依赖性，因此cldprmc_sw使用三维并行计算。

setcoef_d:setcoef_sw被划分为了两个kernel，setcoef_d1和setcoef_d2。其中，setcoef_d1的计算在水平和垂直维度无依赖性，因此使用二维并行计算；setcoef_d2中的计算部分为累加操作，仅在水平维度计算无依赖性，因此使用一维并行计算。

spcmvc_d:spcvmc_sw调用三个子例程taumol、reftra以及vrtqdr。在spcvmc_d的实现过程中，将taumol单独作为一个kernel即taumol_d，由于taumol_d在水平和垂直维度的计算无依赖性，因此采用二维并行计算。spcvmc_sw剩余的计算部分划分为两个kernel，分别为spcvmc_d1以及spcvmc_d2。spcvmc_d1的计算在水平和jp-band维度(jp-band维度用于区分短波十四个波段的计算)无依赖性，因此采用二维并行计算；spcvmc_d2是对短波十四个波段辐射通量的积分操作，仅在水平维度计算无依赖性，因此使用一维并行计算。

步骤2.对kernel spcvmc_d2使用GPU寄存器进行访存优化。线程访问寄存器的速度远快于访问全局内存global memory，但是寄存器是GPU每个核中的有限资源，如果一个线程使用了过多的寄存器，就会使多核处理器的占用率降低，因此并不是使用越多的寄存器，程序的效率就会越高。在spcvmc_d2中，考虑寄存器的使用率和占用率的最佳平衡，将六个存储辐射通量的二维数组临时存储于寄存器中，在计算过程中将其降维为一维数组，计算完成之后再通过增加循环过程将其重新赋值到global memory的二维数组中。

步骤3.使用CUDA流对数据传输进行优化。创建四个CUDA流，将spcvmc_d的数据传输、kernel计算划分为四部分，每部分运行于一个CUDA流中，单个CUDA流分别依次完成CUDAMemcpy HtoD、kernel execution、CUDA Memcpy DtoH三个过程，这样四个CUDA流的计算和数据传输可以互相重叠，降低数据传输的时间。

本发明具有如下优点：

本发明将基于GPU的加速技术和CUDA并行计算架构应用到了短波辐射传输模式RRTMG_SW中，在不考虑I/O传输下，在一个NVIDIA Tesla V100GPU上实现了99.09倍的加速，在考虑I/O传输的情况下依然有27.19倍的加速。大幅提升了短波辐射传输模式RRTMG_SW的计算效率。

附图说明

图1本发明的三维加速算法流程示意图。

图2RRTMG短波辐射程序结构。

图3当水平列数量(ncol值)不同时，RRTMG_SW的三维加速版本在一个V100GPU上相较于串行版本的运行时间加速比。

图4当ncol＝2048时，RRTMG_SW的三维加速版本在不同类型GPU上相较于串行版本运行时间的加速比。

具体实施方式

下面将结合本发明实施例中的附图和表格，对本发明实施例中的技术方案进行清楚、完整地描述。

RRTMG_SW基于GPU的三维加速算法如下：

以下算法2～6分别为核函数inatm_d、cldprmc_d、setcoef_d、taumol_d和spcvmc_d二维和三维并行算法实现的伪代码。

RRTMG_SW在Tesla K20，K40和V100GPU上的运行时间和相较于串行版本的加速比如表2所示。表中串行时间为RRTMG_SW使用CPU串行计算的运行时间，GPU时间为RRTMG_SW在三种GPU上三维并行计算的运行时间。实验数据规模为1024×768，为了使RRTMG_SW在不同GPU上并行计算效率能达到最高，表中分别在不同的GPU上取所能达到的最大ncol值(ncol为短波辐射水平列数量)。实验结果表明在一个Tesla V100 GPU上RRTMG_SW的加速比最高能够达到99.09×，证明了该发明对提升短波辐射传输模式RRTMG_SW计算效率的有效性。

表2不同GPU上的计算时间(s)及加速比(数据规模1024×768，block size＝128)

虽然，上文中已经用一般性说明及具体实施例对本发明作了详尽的描述，但在本发明基础上，可以对之作一些修改或改进，这对本领域技术人员而言是显而易见的。因此，在不偏离本发明精神的基础上所做的这些修改或改进，均属于本发明要求保护的范围。

Claims (5)

1.一种基于GPU的短波辐射传输模式三维加速方法，其特征在于：通过GPU计算及CUDA架构对短波辐射传输模式RRTMG_SW进行三维并行计算及算法性能优化，以此来提升大气短波辐射传输模式RRTMG_SW的计算效率，并使得模式本身更加适合使用GPU进行计算。

2.根据权利要求1所述的基于GPU的短波辐射模式三维加速方法，其特征在于，该方法包括以下步骤：

数据初始化与内存分配。在CPU端进行RRTMG_SW所需数据的定义和初始化；在CPU端定义启动核函数所需要的dim3类型的三维向量，三维向量具有x、y和z三个维度，用于三维并行计算的线程组织。在GPU端，采用指针数组，并通过cudaMalloc()对GPU端的数组进行内存分配。

3.根据权利要求2所述的基于GPU的短波辐射模式三维加速方法，其特征在于：三维并行。通过对短波辐射中水平、垂直和g-point维度数据的分析，使用三维CUDA线程对核函数的三重循环进行三维并行计算。启动核函数的格式为kernel<<<Grid,tBlock>>>(parameters)，Grid和tBlock分别用于指定网格中块和线程的组织形式，kernel中获取CUDA线程号的格式为blockDim*blockIdx+threadIdx。

4.根据权利要求3所述的基于GPU的短波辐射模式三维加速方法，其特征在于：使用寄存器进行访存优化。在CUDA程序中，如果不指定数据在GPU端的存储位置，数据默认存储于global memory中，而线程对global memory的访问速度远低于GPU寄存器。在kernel中，通过设置特定数量的临时变量存储于寄存器中，加快了线程计算时对数据的访存时间，得到计算结果后，再将寄存器中的临时变量重新传回global memory。

5.根据权利要求4所述的基于GPU的短波辐射模式三维加速方法，其特征在于：使用CUDA流进行数据传输优化。该方法包括以下步骤：

创建多个CUDA流，将核函数数据传输与kernel计算根据创建的CUDA流数量划分为多个部分来完成；使用异步拷贝函数cudaMemcpyAsync通过不同的流将数据传到GPU；在不同的流上启动核函数进行计算；使用cudaMemcpyAsync将不同流上的计算结果传回CPU；销毁CUDA流。

Images