HGRTC API（v2.1）-真武 PPU 云服务(ppu)-阿里云帮助中心

1. 简介

HGRTC是一个运行时编译库，它接受字符串形式的CUDA/HGGC device源代码，产生AliPPU的二进制代码。AliPPU的二进制代码可以通过AliPPU Driver API hgModuleLoadData 、hgModuleLoadDataEx来加载，并在运行时执行。用户将通过以下方式从中受益：

使用release的SDK包简化部署
编译过程不会产生单独的进程，没有额外的磁盘I/O的开销

2. 开始

2.1 系统要求

HGRTC 运行所需要的系统要求：

Linux X86_64
PPU Card
PPU SDK ToolKit and Dirver

2.2 安装

HGRTC 包含在PPU SDK中，将随着PPU SDK一起被安装在用户环境中：

include/hgrtc.h
lib/libhgrtc.so

3. 用户接口

本章列举所有的HGRTC API。

3.1 ErrorHandling

3.1.1 Enumeration

enum hgrtcResult：hgrtc API调用的返回代码。

Values

 HGRTC_SUCCESS = 0,
 HGRTC_ERROR_OUT_OF_MEMORY = 1,
 HGRTC_ERROR_PROGRAM_CREATION_FAILURE = 2,
 HGRTC_ERROR_INVALID_INPUT = 3,
 HGRTC_ERROR_INVALID_PROGRAM = 4,
 HGRTC_ERROR_INVALID_OPTION = 5,
 HGRTC_ERROR_COMPILATION = 6,
 HGRTC_ERROR_BUILTIN_OPERATION_FAILURE = 7,
 HGRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 8,
 HGRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 9,
 HGRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 10,
 HGRTC_ERROR_INTERNAL_ERROR = 11

3.1.2 函数定义

const char* hgrtcGetErrorString(hgrtcResult result)
描述: 用来解释说明hgrtc返回代码的函数，该函数会根据输入的代码返回一串可读的字符串，例如HGRTC_SUCCESS to “HGRTC_SUCCESS”，如果遇到无法识别的代码，会返回“HGRTC_ERROR_unknown”。
参数: result — HGRTC API 返回代码
返回值: 用来解释返回代码的可读字符串

3.2 基础信息查询接口

hgrtcResult hgrtcVersion( int *major, int *minor)
描述: hgrtcVersion 用来设置HGGC实时编译的版本号.
参数: major — HGGC Runtime Compilation 主版本号
minor — HGGC Runtime Compilation 子版本号
返回值: HGRTC_SUCCESS、 HGRTC_ERROR_INVALID_INPUT
hgrtcResult hgrtcGetNumSupportedArchs(int* numArchs)
描述: 设置HGGC实时编译支持的架构的数量.
参数: numArchs — HGGC实时编译支持的架构的数量.
返回值: HGRTC_SUCCESS、 HGRTC_ERROR_INVALID_INPUT
hgrtcResult hgrtcGetSupportedArchs(int* supportedArchs)
描述: 设置HGGC实时编译支持的架构.
参数: supportedArchs — HGGC实时编译支持的架构.
返回值: HGRTC_SUCCESS、 HGRTC_ERROR_INVALID_INPUT

3.3 编译接口

HGRTC 为编译过程定义了如下的类型和函数：

类型

Typedefs typedef _hgrtcProgram* hgrtcProgram

函数

hgrtcResult hgrtcCreateProgram(hgrtcProgram* prog,
							   const char* src,
                               const char* name,
                               int numHeaders,
                               const char** headers,
                               const char** includeNames);
hgrtcResult hgrtcDestroyProgram(hgrtcProgram* prog);
hgrtcResult hgrtcAddNameExpression(hgrtcProgram prog,
							       const char* name_expression);
hgrtcResult hgrtcCompileProgram(hgrtcProgram prog,
								int numOptions,
								const char** options);
hgrtcResult hgrtcGetHGBIN(hgrtcProgram prog, char* hgbin);
hgrtcResult hgrtcGetHGBINSize(hgrtcProgram prog, size_t* hgbinSizeRet);
hgrtcResult hgrtcGetLTOIR(hgrtcProgram prog, char* LTOIR);
hgrtcResult hgrtcGetLTOIRSize(hgrtcProgram prog, size_t* LTOIRSizeRet);
hgrtcResult hgrtcGetLoweredName(hgrtcProgram prog,
								const char* name_expression,
								const char** lowered_name);
hgrtcResult hgrtcGetProgramLogSize(hgrtcProgram prog,
								   size_t* logSizeRet);
hgrtcResult hgrtcGetProgramLog(hgrtcProgram prog, char* log);
// Bellow APIs are deprecated and will be removed later.
hgrtcResult hgrtcGetCodeSize(hgrtcProgram prog, size_t* codeSizeRet);
hgrtcResult hgrtcGetCode(hgrtcProgram prog, char* code);

3.3.1 Typedefs

typedef _hgrtcProgram* hgrtcProgram
hgrtcProgram是控制编译的单元，需要在编译开始时创建。

3.3.2 函数

hgrtcResult hgrtcCreateProgram(hgrtcProgram* prog, const char* src, const char* name, int numHeaders, const char** headers, const char** includeNames)
hgrtcCreateProgram会根据输入的参数来创建一个hgrtcProgram实例，并通过参数返回这个实例。
参数
- prog: HGRTC 编译控制单元.
- src: 源文件buffer
- name: Input source name. Will use “default_program” when the name is NULL or empty.
- numHeaders: 参与编译头文件的数量，值大于或等于0。
- headers: 参与编译的头文件。当头文件数量为0时，该参数可以为NULL。
- includeNames: 参与编译的头文件的名称。当头文件数量为0时，该参数可以为NULL。
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_OUT_OF_MEMORY
- HGRTC_ERROR_PROGRAM_CREATION_FAILURE
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcDestroyProgram(hgrtcProgram* prog)
hgrtcDestroyProgram 负责销毁编译实例。
参数
- prog: HGRTC 编译控制单元
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcAddNameExpression(hgrtcProgram prog, const char* name_expression)
hgrtcAddNameExpression在编译过程中记录name_expression对应的global function或者是device/constant variable对应的地址。
参数
- prog: HGRTC 编译控制单元.
- name_expression: constant 表达式，用来表示 global function or device/constant variable 地址。
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION
Description 在后续调用API hgrtcGetLoweredName，并传递相同的name_expression时，可以得倒编译后的lowered name.
hgrtcResult hgrtcCompileProgram(hgrtcProgram prog, int numOptions, const char** options)
hgrtcCompileProgram 编译.
参数
- prog: HGRTC 编译控制单元.
- numOption: option数量.
- options: 接受C string数组类型的编译选项。当选项数量为0时，option可以是NULL。
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_OUT_OF_MEMORY
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
- HGRTC_ERROR_INVALID_OPTION
- HGRTC_ERROR_COMPILATION
- HGRTC_ERROR_BUILTIN_OPERATION_FAILURE
Description 在后续章节Supported Compile Options 中列出了编译支持的options。
hgrtcResult hgrtcGetHGBINSize(hgrtcProgram prog, size_t* hgbinSizeRet)
hgrtcGetHGBINSize 得到当前prog编译出的binary code的size，并通过设置参数hgbinSizeRet返回这一结果。
参数
- prog: HGRTC 编译控制单元.
- hgbinSizeRet: 生成的binary code的size.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcGetHGBIN(hgrtcProgram prog, char* hgbin)
hgrtcGetHGBIN 存储当前prog编译出的binary code。
参数
- prog: HGRTC 编译控制单元.
- hgbin: 编译出的binay code.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcGetLTOIRSize(hgrtcProgram prog, size_t* LTOIRSizeRet)
hgrtcGetLTOIRSize 得到当前prog编译出的LTOIR的size，并通过设置参数LTOIRSizeRet返回这一结果。
参数
- prog: HGRTC 编译控制单元.
- LTOIRSizeRet: 生成的 LTOIR的size.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcGetLTOIR(hgrtcProgram prog, char* LTOIR)
hgrtcGetLTOIR 存储当前prog编译出的LTO IR。
参数
- prog: HGRTC 编译控制单元.
- code: 编译结果.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcGetLoweredName(hgrtcProgram prog, const char* name_expression, const char** lowered_name)
hgrtcGetLoweredName 会抽取 global function or device/constant variable 的 lowered(mangled) name, 并且更新到对应的参数指针上. 当编译控制单元被销毁时，保存lowered(mangled) name的内存也会被释放。查询的name_expression必须在编译开始前通过API hgrtcAddNameExpression提供给编译控制单元。
参数
- prog: HGRTC 编译控制单元.
- name_expression: 用来表示a global function or device/constant variable 地址的常量string。
- lowered_name: 指向name_expression 对应的 lowered(mangled) name。
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION
- HGRTC_ERROR_NAME_EXPRESSION_NOT_VALID
hgrtcResult hgrtcGetProgramLogSize(hgrtcProgram prog, size_t* logSizeRet)
hgrtcGetProgramLogSize 得倒当前编译生成的log的size(包含了结尾处的NULL).
参数
- prog: HGRTC 编译控制单元.
- logSizeRet: 编译生成log的size (包含了结尾处的NULL).
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcGetProgramLog(hgrtcProgram prog, char* log)
hgrtcGetProgramLog 存储当前编译生成的log。
参数
- prog: HGRTC 编译控制单元.
- log: 编译log.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
Bellow APIs are deprecated and will be removed later.
hgrtcResult hgrtcGetCodeSize(hgrtcProgram prog, size_t* codeSizeRet)
hgrtcGetCodeSize 得到当前prog编译出的binary code的size，并通过设置参数hgbinSizeRet返回这一结果。
参数
- prog: HGRTC 编译控制单元.
- codeSizeRet: Size of generated binary code.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
hgrtcResult hgrtcGetCode(hgrtcProgram prog, char* code)
hgrtcGetCode 存储当前prog编译出的binary code。
参数
- prog: HGRTC 编译控制单元.
- code: 编译生成的code
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM

3.4 支持的编译选项

3.4.1 选项格式

HGRTC 编译选项需要符合以下规则：

长选项名由两个dash（--）开始，短选项名由一个dash（-）开始。
短选项名可以用来代替长选项名。
带参数的选项可以使用operator（=）来分割选项名和参数。
带参数的选项也可以使用空格来分隔。
单字符的短选项名，比如-D, -U, and -I, 不需要assignment operator, 编译选项可以和参数在同一个string中，可以使用或不使用空格隔开。比如, “-D=def”, “-Ddef”, and “-D def” 这些写法都支持。

3.4.2 HGRTC 编译选项

分离/整体编译 (Separate compilation/whole compilation)

HGRTC only support

--device-c(-dc)

该选项用来生成 relocatable code，可以link其他relocatable code，等效于 --relocatable-device-code=true.

--device-w(-dw)

该选项用来生成 non-relocatable code。等效于 --relocatable-device-code=false.

--relocatable-device-code={true|false} (-rdc)

设置是否生成 relocatable device code.

Default: false

Debugging support

--device-debug(-G)

生成 debug 信息. 如果 --dopt 没有设置, 该选项会关闭所有优化.

--generate-line-info(-lineinfo)

生成 line-number 信息

Code generation

--dopt on (-dopt) --dopt= on

当该选项和‘-G’联合使用时，会产生有限的debug信息（line number information），以便于生成优化过的device code。当没有设置‘-G’ 时, 默认‘-dopt=on’。

--ftz={true|false} (-ftz)

当执行单精度浮点计算时, 通过--ftz来控制flush denormal values to zero 或保留 denormal values。 --use_fast_math 表示 --ftz=true。

Default: false

--prec-sqrt={true|false} (-prec-sqrt)

当执行单精度浮点的square root计算时, 通过--prec-sqrt来控制使用 IEEE round-to-nearest mode 或使用更快速的近似计算。--use_fast_math 表示 --prec-sqrt=false.

Default: true

--prec-div={true|false} (-prec-div)

当执行单精度浮点的除法和倒数计算时, 通过--prec-div来控制使用IEEE round-to-nearest mode 或使用更快速的近似计算。--use_fast_math 表示 --prec-div=false.

Default: true

--fmad={true|false} (-fmad)

开启（关闭）浮点乘法和加法（减法）运算合并为乘加运算的操作。--use_fast_math 表示 --fmad=true.

Default: true

--use_fast_math (-use_fast_math)

使用 fast math operations. --use_fast_math 表示 --ftz=true --prec-div=false --prec-sqrt=false --fmad=true

--dlink-time-opt (-dlto)

生成中间code，用来做后续 link-time optimization，此时 -rdc=true.

Preprocessing

--define-macro= (-D)

def 可以为 name 或 name=definitions的格式。

name

预定义 name 为值为1 的宏。

name=definitions

The contents of definition are tokenized and preprocessed as if they appeared during translation phase three in a #define directive. In particular, the definition will be truncated by embedded new line characters.

--undefine-macro=<def> (-U)

删除当前的宏定义 def.

--include-path=<dir> (-I)

添加编译路径，用来查找编译头文件。

--pre-include=<header> (-include)

Preinclude header during preprocessing.

--no-source-include (-no-source-include)

预处理器会默认添加input sources的路径到编译include 路径中。这个编译选项会disable这种默认的行为。

Language Dialect

--std={c11, c14, c17} (-std={c11, c14, c17})

设置 language dialect to C11, C14 or C17. 默认的 languange dialect in HGRTC is c17.

--builtin-move-forward={true|false} (-builtin-move-forward)

当选择使用C++11时，提供内置的 definitions of std::move and std::forward。

Default: true

--builtin-initializer-list={true|false} (-builtin-initializer-list)

当选择使用C++11时，该选项可以提供内置的 std::initializer_list class and member functions。

Default: true

Misc

--disable-warnings (-w)

禁止waring messages。

--device-as-default-execution-space (-default-device)

默认source code中的entities 都是 device entities。

--device-int128 (-device-int128)

允许在device code中使用 __int128 type。

ppu arch 相关选项

--ppu-arch==

ppu arch，注意hgrtc不支持多arch的混合编译产物。

--ppu-arch==ppu001：编译ZW810、ZW610、ZW805、ZW810e、ZW610e系列编译产物；

--ppu-arch==ppu0015：编译ZW890系列编译产物；

3.4.3 Host help

HGRTC 定义了如下的辅助函数，方便用户同Host代码进行交互。

template

hgrtcResult hgrtcGetTypeName(std::string *result)[inline]

hgrtcGetTypeName stores the source level name of the template type argument T in the given std::string location.

参数

result: 指向用来表示类型名的std::string字符串指针

返回值

HGRTC_SUCCESS
HGRTC_ERROR_INTERNAL_ERROR

描述

当前函数只在宏 HGRTC_GET_TYPE_NAME 被定义为非0值时才会生效。在gcc/clang的编译器中，它通过使用 abi::__cxa_demangle 函数调用来获取类型名。

4. 语法

HGRTC 只支持device代码的编译，不接受host代码。

4.1 Include 语法

当 hgrtcCompileProgram() 被调用时, 当前的工作目录被添加进头文件搜索路径，通过#include "" 方式的头文件能够被找到(例如 #include “header.h”)。

4.2 预定义宏

NULL: null pointer constant

HGGCCC_RTC: 在用户代码中用来区分offline和runtime编译的部分。

HGGC_ARCH: 同offline编译一样的ARCH宏定义。

_cplusplus: c++ languange dialect.

4.3 Predefined Types

clock_t size_t ptrdiff_t

5. 基本用法

本章列举了一个使用HGRTC的简单例子。

HGGC 源代码 SAXPY:

const char *saxpy = "                                           \n\
extern \"C\" __global__                                         \n\
void saxpy(float a, float *x, float *y, float *out, size_t n)   \n\
{                                                               \n\
  size_t tid = blockIdx.x * blockDim.x + threadIdx.x;           \n\
  if (tid < n) {                                                \n\
    out[tid] = a * x[tid] + y[tid];                             \n\
  }                                                             \n\
}                                                               \n";

创建 hgrtcProgram 的实例

  // Create an instance of hgrtcProgram with the SAXPY code string.
  hgrtcProgram prog;
  hgrtcCreateProgram(&prog,         // prog
                     saxpy,         // buffer
                     "saxpy.cu",    // name
                     0,             // numHeaders
                     NULL,          // headers
                     NULL);         // includeNames

编译代码，选择--fmad=false的选项

  const char *opts[ ] = {"--fmad=false"};

  int numOpts = sizeof(opts)/sizeof(const char*);
  hgrtcResult compileResult = hgrtcCompileProgram(prog,    // prog
                                                  numOpts, // numOptions
                                                  opts);   // options

获取二进制代码和编译的log信息

  // Obtain compilation log from the program.
  size_t logSize;
  hgrtcGetProgramLogSize(prog, &logSize);
  char *log = new char[logSize];
  hgrtcGetProgramLog(prog, log);

  // Obtain PTX from the program.
  size_t ptxSize = 0;
  hgrtcGetCodeSize(prog, &ptxSize);
  char *ptx = new char[ptxSize];
  HGRTC_SAFE_CALL(hgrtcGetCode(prog, ptx));

销毁 hgrtcProgram 对象

  // Destroy the program.
  HGRTC_SAFE_CALL(hgrtcDestroyProgram(&prog));

执行编译产生的二进制kernel

  // Load the generated HGBIN and get a handle to the SAXPY kernel.
  HGdevice cuDevice;
  HGcontext context;
  HGmodule module;
  HGfunction kernel;
  HGGC_SAFE_CALL(hgInit(0));
  HGGC_SAFE_CALL(hgDeviceGet(&cuDevice, 0));
  HGGC_SAFE_CALL(hgCtxCreate(&context, 0, cuDevice));
  HGGC_SAFE_CALL(hgModuleLoadData(&module, ptx));
  HGGC_SAFE_CALL(hgModuleGetFunction(&kernel, module, "saxpy"));
  // Generate input for execution, and create output buffers.
  size_t n = NUM_THREADS * NUM_BLOCKS;
  size_t bufferSize = n * sizeof(float);
  float a = 5.1f;
  float *hX = new float[n], *hY = new float[n], *hOut = new float[n];
  for (size_t i = 0; i < n; ++i) {
    hX[i] = static_cast<float>(i);
    hY[i] = static_cast<float>(i * 2);
  }
  HGdeviceptr dX, dY, dOut;
  HGGC_SAFE_CALL(hgMemAlloc(&dX, bufferSize));
  HGGC_SAFE_CALL(hgMemAlloc(&dY, bufferSize));
  HGGC_SAFE_CALL(hgMemAlloc(&dOut, bufferSize));
  HGGC_SAFE_CALL(hgMemcpyHtoD(dX, hX, bufferSize));
  HGGC_SAFE_CALL(hgMemcpyHtoD(dY, hY, bufferSize));
  // Execute SAXPY.

  void *args[ ] = { &a, &dX, &dY, &dOut, &n };

  HGGC_SAFE_CALL(
    hgLaunchKernel(kernel,
                   NUM_BLOCKS, 1, 1,    // grid dim
                   NUM_THREADS, 1, 1,   // block dim
                   0, NULL,             // shared mem and stream
                   args, 0));           // arguments
  HGGC_SAFE_CALL(hgCtxSynchronize());
  // Retrieve and print output.
  HGGC_SAFE_CALL(hgMemcpyDtoH(hOut, dOut, bufferSize));
  for (size_t i = 0; i < n; ++i) {
    std::cout << a << " * " << hX[i] << " + " << hY[i]
              << " = " << hOut[i] << '\n';
  }
  // Release resources.
  HGGC_SAFE_CALL(hgMemFree(dX));
  HGGC_SAFE_CALL(hgMemFree(dY));
  HGGC_SAFE_CALL(hgMemFree(dOut));
  HGGC_SAFE_CALL(hgModuleUnload(module));
  HGGC_SAFE_CALL(hgCtxDestroy(context));

  delete[ ] hX;


  delete[ ] hY;


  delete[ ] hOut;

6. HGRTC CACHE

  <font style="color:#1d1d1d;">hgrtc cache在v1.6默认打开。在libhgrtc.so中实现对同一份源码在编译命令、相关环境变量、SDK版本等都相同的情况下，第一次调用hgrtcCompileProgram()时缓存输出文件(hgbin\lto ir)，在后续调用hgrtcCompileProgram()时，缓存文件将被复用，不再进行二次编译，节省时间。可以通过设置环境变量HGRTC_CACHE_DISABLE，或在命令行中添加 -no-cache 来关闭cache的使用。</font>

cache相关环境变量设置(可兼容支持CUDA的环境变量，但不要与HGRTC环境变量混用)：

HGRTC_CACHE_DISABLE 兼容CUDA_CACHE_DISABLE	默认=0，即cache打开，若设置为=1，cache关闭	当cache关闭，不会缓存文件到cache，也不会从cache中取文件
HGRTC_CACHE_MAXSIZE 兼容CUDA_CACHE_MAXSIZE	默认Cache容量为1G	以byte为单位设置cache容量，当缓存文件超过该大小时，将旧文件逐出，缓存新文件当HGRTC_CACHE_MAXSIZE=0时，关闭cache
HGRTC_CACHE_PRUN_INTERVAL_SECONDS	默认=0s	cache容量超限删除时间间隔，默认0s，可根据性能需求调整
HGRTC_CACHE_PATH 兼容CUDA_CACHE_PATH	缓存文件目录，默认为$HOME/.hg/ComputeCache/	当缓存位于共享或网络文件系统上时，如果文件系统竞争激烈或文件访问缓慢，则访问缓存的速度可能会很慢。可通过将 CUDA_CACHE_PATH 设置为本地文件系统中位置来加速
编译命令：-no-cache	设置该选项表示关闭cache	当cache关闭，不会缓存文件到cache，也不会从cache中取文件

当前支持的缓存文件：

nameExpressionsMap.txt	当需要获取lower name场景时缓存 hgrtcGetLoweredName()	相同编译环境： 1. SDK版本 2. hgrtc version/时间戳 3. 编译源码和命名 4. 前端编译选项 5. nameExpression输入 hgrtcAddNameExpression()
lto ir	当需要获取lto ir场景时缓存 hgrtcGetLTOIR()	相同编译环境同上5点
hgbin	当需要获取hgbin场景时缓存 hgrtcGetHGBIN()	相同编译环境为以上5点加以下： 6. HGGC_OPTIONS_FOR_BACKEND_TUNING环境变量 7. 后端编译选项

7. PCH

预编译头文件（PCH）是一个编译时优化特性，用于在连续的编译器调用中编译相同的一组“前缀”头文件场景。对两个不同的代码文件，在包含的头文件序列、编译命令、SDK版本等都相同的情况下，第一次调用hgrtcCompileProgram()时，hgrtc将会解析hggc的头文件序列并提取需要预编译的内容，根据该内容创建precompiled headers（pch），并缓存，在后续调用hgrtcCompileProgram()时，缓存文件将被复用，不再进行头文件解析AST构建，节省编译时间。

新增PCH相关环境变量：

HGRTC_DISABLE_PCH = 1/0 兼容支持NVRTC_DISABLE_PCH	默认=0，即支持pch功能，根据pch相关编译选项支持各功能，若设置为=1，关闭pch功能，忽略其相关的编译选项。	关闭pch功能支持
兼容支持NVRTC_PCH_HEAP_SIZE	默认值=4096，若设置为=0，pch关闭	兼容cuda pch heap size设置的环境变量，cuda pch 复用要求存储空间虚拟地址一致，故开辟了一段固定虚拟空间的heap，但在hgrtc pch方案下，pch文件和虚拟地址无关，故不提供heap size的接口

新增PCH相关编译命令：

-pch		自动pch模式，该模式下编译器会自动创建和使用pch文件，若发现无可用的pch，则创建pch文件，并完成编译，若发现可复用的pch，直接复用文件。pch文件的管理由编译器负责，当前放在cache中管理。
--create-pch=filename		用户指定生成pch文件模式，该模式下用户负责管理pch文件，编译器生成路径为filename的文件，并完成编译。
--use-pch=filename		用户指定使用pch文件模式，该模式下用户负责管理pch文件，编译器使用路径为filename的文件完成编译。
--pch-dir=		当在-pch 模式下时，dir内容会影响计算文件hash值，用来限定pch查找范围在(`-create-pch` or `-use-pch`)模式下，当filename非绝对路径时，作为filename的路径前缀，输出用户可见的pch文件
--pch-verbose={true\|false}	默认true	在 -pch模式下有效，若复用的pch文件不兼容，输出编译详细信息
--pch-messages={true\|false}	默认true	输出pchfile使用或创建信息
--instantiate-templates-in-pch={true\|false}	默认true	在创建PCH启用或禁用模板的实例化。实例化模板可能会增加PCH文件的大小

新增PCH相关pragma：

#pragma hg_hdrstop 兼容支持#pragma nv_hdrstop	只在编译源文件生效，使用示例： pch_test.cu: #include "pch_common.h" #pragma nv_hdrstop #include "pch_util.h" global void another(double a, double b) { a = b + doit(); }	代表生成pch的代码截止点
#pragma hg_no_pch 兼容支持#pragma nv_no_pch	只在编译源文件生效，使用示例： pch_test.cu: #include "pch_common.h" #pragma nv_no_pch global void other(double a, double b) { a = b + doit(); }	关闭pch功能（一般用于在编译选项添加了-pch时仍关闭pch功能）

新增接口函数：

hgrtcResult hgrtcGetPCHCreateStatus
hgrtcGetPCHCreateStatus 获取pch文件创建状态
参数
- prog: HGRTC 编译控制单元.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
- HGRTC_ERROR_INVALID_PROGRAM
- HGRTC_ERROR_NO_PCH_CREATE_ATTEMPTED
- HGRTC_ERROR_PCH_CREATE
- HGRTC_ERROR_PCH_CREATE_HEAP_EXHAUSTED
hgrtcResult hgrtcGetPCHHeapSize(size_t *ret)
hgrtcGetPCHHeapSize 获取pch heapsize，默认4096
参数
- ret: 返回size指针
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
hgrtcResult hgrtcGetPCHHeapSizeRequired(hgrtcProgram prog, size_t *size)
hgrtcGetPCHHeapSizeRequired 获取需要的pch heapsize
参数
- size: 返回size指针
- prog: HGRTC 编译控制单元.
返回值
- HGRTC_SUCCESS
- HGRTC_ERROR_INVALID_INPUT
hgrtcResult hgrtcSetPCHHeapSize(size_t size)
hgrtcSetPCHHeapSize 设置pch heapsize，当set size为0时，关闭pch 功能
参数
- size: pch heap size
返回值
- HGRTC_SUCCESS