2. 希姆计算异构编程手册
2.1. 版本历史
| 文档版本 | 对应产品版本 | 作者 | 日期 | 描述 |
|---|---|---|---|---|
| V1.6.0 | STCRP V1.7.0 | 希姆计算 | 2024-01-15 | 更新支持PCIe P2P和设置32核的HS group。 |
| V1.5.0 | STCRP V1.4.0 | 希姆计算 | 2023-04-07 | 更新stc-smi回显示例。 |
| V1.4.0 | STCRP V1.2.0 | 希姆计算 | 2022-11-30 | - 更新stc-smi回显示例。 - 编辑优化。 |
| V1.3.0 | HPE V1.4.0 | 希姆计算 | 2022-09-09 | 更新stc-smi回显示例。 |
| V1.2.0 | HPE V1.3.0 | 希姆计算 | 2022-07-07 | 更新stc-smi回显示例。 |
| V1.1.0 | HPE V1.2.0 | 希姆计算 | 2022-04-11 | - 文档版本号对齐HPE发版。 - 整篇编辑优化，确认模糊的描述。 |
| V1.0.0 | Unknown | 希姆计算 | 2021-09-01 | 初始版本。 |
2.2. 概述
希姆计算推出了异构编程环境HPE（Heterogeneous Programming Engine），安装HPE后您可以使用C/C++语言开发异构程序，方便地使用NPU进行并行计算。
异构程序包括在CPU上运行的主机端程序和在NPU上运行的设备端程序，希姆计算的异构编程环境为开发、编译、运行主机端程序和设备端程序提供了完整的工具链。主要包括以下模块：
| 模块 | 全称 | 说明 |
|---|---|---|
| SHC | Stream Computing Heterogeneous C++ | 异构编程语言，完整兼容C++17标准，并针对异构程序扩展了语法和函数库，例如支持调用核函数的<<< >>>符号。 |
| stc-dkms | Stream Computing Dynamic Kernel Module Support | 异构驱动模块，对应有异构驱动stc.ko文件。主机端与设备端使用PCIe总线连接，异构驱动将主机端的操作请求通过PCIe总线传递到设备端。 stc-dkms方便异构驱动适应不同的Linux内核版本，不用逐版本手动编译也能保证异构驱动始终可用，主机端和设备端的交互不会因Linux内核版本变化而产生异常。 |
| hpert | Heterogeneous Programming Engine Runtime | 主机端运行时库，为主机端程序提供控制设备端内存访问、核函数执行等操作的接口。 |
| npurt | NPU Runtime | 设备端运行时库，为设备端程序提供信息打印、内存拷贝等操作的接口。 |
| stcc | Stream Computing Heterogeneous C++ Compiler | 异构程序编译器，统一编译异构程序中的主机端程序和设备端程序，生成可执行文件。 |
| stc-smi | Stream Computing System Management Interface | 设备管理工具，管理和监控NPU设备的状态。 |
| stc-prof | Stream Computing Profiler | 性能调优工具，采集异构程序的性能数据，分析异构程序的性能。 说明：stc-prof是命令行工具，希姆计算还提供了接口形式的性能调优工具STCPTI（ Stream Computing Profiling Tool Interface ）。 |
| stc-vprof | Stream Computing Visual Profiler | 可视化性能调优工具，主要用于异构程序的可视化性能分析。 |
| stc-gdb | Stream Computing Debugger | 异构程序调试工具，监视程序运行状态，获取和修改程序的中间运行结果。stc-gdb完全兼容GDB原生命令，并添加了希姆计算的扩展命令。 |
说明：更多希姆计算产品相关的概念，请参见希姆计算基本概念。
2.3. 异构编程模型
在异构编程环境中，设备端与主机端在计算和存储结构上存在差别，因此需要通过不同的方式执行函数和访问内存。异构编程模型中需要考虑核函数、内存管理、异步函数等因素。
2.3.1. 核函数
2.3.1.1. 定义核函数
异构程序包括主机端程序和设备端程序，并引入核函数的概念来关联两类程序，核函数是两类代码的桥梁。核函数在主机端调用，由主机端部署到设备端，并在设备端的NPC上并行执行。
核函数使用__global__作为函数修饰符。在调用核函数时，需要通过<<< N >>>指定执行核函数时所用NPC的个数。__global__的详细用法，请参见C++语言扩展接口章节。
在下方的异构程序用例中，核函数的作用为打印hello word信息，然后在主机端指定8个NPC同时执行该核函数，执行结果为8个NPC分别输出一条hello word信息并在主机端显示。
说明：异构程序的源文件必须以.hc作为扩展名。
#include <hpe.h>
#include <npurt.h>
__global__ void kernel(void) {
printf("hello world in NPC %d/%d.\n", CoreID, CoreNum);
}
int main(void) {
kernel<<<8>>>();
stcDeviceSynchronize();
return 0;
}
2.3.1.2. 编译核函数
使用stcc统一编译异构程序，包括核函数。示例如下：
编译单个源文件：编译hello_world.hc，输出名为hello_world的二进制文件。
$ stcc hello_world.hc -o hello_world
编译多个源文件：编译hello_world_1.hc和hello_world_2.hc，输出名为hello_world的二进制文件。
说明：编译多个源文件时需要添加
--shc-combine-device选项。$ stcc hello_world_1.hc hello_world_2.hc --shc-combine-device -o hello_world
2.3.1.3. 执行核函数
运行编译生成的二进制文件，在执行到核函数时，会自动转到设备端。
$ ./hello_world
默认在NPC Cluster 0上执行核函数，您也可以自行指定NPC Cluster。支持以下方式：
通过环境变量STC_SET_DEVICES修改默认的起始NPC Cluster。例如
STC_SET_DEVICES为3指定在NPC Cluster 3上执行核函数。通过
stcSetDevice单次指定NPC Cluster，stcSetDevice的效果为STC_SET_DEVICES基础上的增量。例如：STC_SET_DEVICES为0，则stcSetDevice(1)指定在NPC Cluster 1上执行核函数。STC_SET_DEVICES为2，则stcSetDevice(1)指定在NPC Cluster 3上执行核函数。
说明：一个NPC Cluster只能同时运行一个核函数，即使在执行当前核函数时只使用了部分NPC，也不能在空闲的NPC上立即执行下一个核函数。
2.3.2. 内存管理
2.3.2.1. 内存布局
希姆计算自研NPU中的内存分为本地内存（L1）、共享内存（LLB）、全局内存（DDR），各类型内存的大小和访问速度存在差别。
全局内存：每个NPC Cluster私有的内存，由NPC Cluster内的NPC共享，访问速度最慢。STCP920中，每个NPC Cluster的全局内存大小为4GiB。
共享内存：每个NPC Cluster私有的内存，由NPC Cluster内的NPC共享，访问速度较快。STCP920中，每个NPC Cluster的共享内存大小为8MiB。
本地内存：每个NPC私有的内存，访问速度最快。STCP920中，每个NPC的本地内存大小为1.25MiB。
NPU中各NPC Cluster的内存布局完全相同，如下所示：
2.3.2.2. 访问内存
在STCP920中，每个NPC Cluster的4GiB全局内存分为以下类型：
NPC可以直接访问的3GiB内存（0GiB ~ 3GiB）。
只能通过sysDMA访问的1GiB内存（3GiB ~ 4GiB），也称为高端内存。
主机端支持动态分配内存：
在主机端调用
stcMalloc()动态分配设备端全局内存（0GiB ~ 3GiB）。在主机端调用
stcMallocHigh()动态分配设备端全局内存（3GiB ~ 4GiB）。
设备端不支持动态分配内存，希姆计算提供了变量修饰符在指定的地址空间定义局部变量：
__device__：定义位于全局内存的局部变量和全局变量。如果定义局部变量时不添加变量修饰符，则默认位于全局内存。__shared__：定义位于共享内存的局部变量。__local__：定义位于本地内存的局部变量。
说明：本地内存（1.25MiB）、共享内存（8MiB）以及全局内存设备端运行栈（64KiB）的大小有限，请避免定义过大的局部变量。
各类型内存支持的访问来源和访问方式的差别如下：
支持从设备端访问本地内存、共享内存，支持从主机端和设备端访问全局内存。
在设备端调用
memcpy读写设备端本地内存、共享内存、全局内存。在主机端调用
stcMemcpy、stcMemcpyAsync读写设备端全局内存。
说明：在设备端访问全局内存时（访问局部变量除外），请确保访问区域不会被其他NPC缓存。因为NPC的cache line大小为32字节，如果多个NPC访问同一个32字节地址对齐且大小不超过32字节的区域，NPC刷新缓存会导致无法保障全局内存中数据的正确性。
2.3.2.3. 访问内存用例
下方的用例中，在主机端分配全局内存并写入数据，然后在设备端使用8个NPC分别将全局内存中的数据拷贝到共享内存和本地内存，并打印拷贝结果。
说明：各变量修饰符和接口的详细说明，请参见变量修饰符章节和运行时接口章节。
#include <hpe.h>
#include <npurt.h>
#define NCORE 8
__global__ void kernel(int *global_data) {
__shared__ int shared_data[NCORE];
__local__ int local_data[NCORE];
memcpy(shared_data, global_data, sizeof(shared_data));
memcpy(local_data, global_data, sizeof(local_data));
printf("core %d read shared data %d local data %d\n", CoreID,
shared_data[CoreID], local_data[CoreID]);
}
int main(void) {
int host_data[] = {1, 2, 3, 4, 5, 6, 7, 8};
int *dev_data;
stcMalloc((void **)&dev_data, sizeof(host_data));
stcMemcpy(dev_data, host_data, sizeof(host_data), stcMemcpyHostToDevice);
kernel<<<NCORE>>>(dev_data);
stcDeviceSynchronize();
stcFree(dev_data);
return 0;
}
2.3.3. 异步函数
部分运行时接口的函数类型设计为异步函数，在主机端调用异步函数后会立即返回，不用等待完成异步函数规定的所有操作，有利于主机端和设备端并行处理任务。例如，调用stcLaunchKernel在主机端启动核函数后，主机端无需等待完成核函数的所有操作即可开始处理下一个任务；调用stcMemcpyAsync在主机端和设备端之间拷贝数据后，主机端无需等待拷贝完所有数据即可开始处理下一个任务。
异步函数的返回值不是异步操作的结果。如果您调用了异步函数，但仍然需要等待所有操作完成，可以调用stcDeviceSynchronize、stcStreamSynchronize或stcStreamSynchronizeUnified等待设备或流上的所有操作完成后再处理下一个任务。stcDeviceSynchronize、stcStreamSynchronize和stcStreamSynchronizeUnified的返回值是异步操作的结果，但仅返回最后一个异步操作的结果。
说明：核函数没有返回值，如果核函数在真正开始执行前出错退出，只能通过
stcGetLastError获取错误信息。例如指定的NPC数量超过实际的NPC数量，导致启动核函数后还未实际执行就出错退出。
2.4. 异构编程典型操作
希姆计算提供了丰富的运行时接口，方便您从主机端控制在设备端执行任务。
2.4.1. 指定运行设备
主机端访问设备端时需要指定运行设备。在主机端程序中启动核函数后，默认在NPC Cluster 0上执行核函数，您也可以提前通过调用stcSetDevice或修改STC_SET_DEVICES指定其他NPC Cluster。示例如下：
推导NPC Cluster ID。
以STCP920为例，执行
stc-smi -q命令获取NPU的信息，包括NPU设备标识符、NPC Cluster设备标识符等。代码中使用的NPC Cluster ID由NPU设备标识符、NPC Cluster设备标识符推导得出。假设需要指定一个NPC Cluster，查看得知其对应的NPU设备标识符为x、NPC Cluster设备标识符为y，则推导方式为NPC Cluster ID = 4 * x + y。例如，NPU设备标识符为为0，NPC Cluster设备标识符为1，则NPC Cluster ID为1。
在代码中调用
stcSetDevice指定NPC Cluster ID。下方的用例中，指定使用NPC Cluster 1，然后8个NPC分别输出一条hello word信息并在主机端显示。
说明：一个进程中各个线程调用
stcSetDevice互不影响，如果需要修改使用的NPC Cluster，请分别设置。#include <hpe.h> #include <npurt.h> __global__ void kernel(void) { printf("hello world in NPC %d/%d.\n", CoreID, CoreNum); } int main(void) { stcSetDevice(1); kernel<<<8>>>(); stcDeviceSynchronize(); return 0; }
2.4.2. 访问设备端全局内存
访问设备端全局内存的运行时接口如下：
在主机端调用
stcMalloc动态分配设备端全局内存（0GiB ~ 3GiB）。在主机端调用
stcMallocHigh动态分配设备端全局内存（3GiB ~ 4GiB）。在主机端调用
stcMemcpy、stcMemcpyAsync读写设备端全局内存。
在STCP920中，每个NPC Cluster的4GiB全局内存分为以下类型：
NPC可以直接访问的3GiB内存（0GiB ~ 3GiB）。
只能通过sysDMA访问的1GiB内存（3GiB ~ 4GiB），也称为高端内存。
下方的用例中，分别在不同的内存范围中分配了内存：
在主机端分配0GiB ~ 3GiB范围的全局内存并写入数据，然后在设备端打印数据。
#include <hpe.h> #include <npurt.h> __global__ void kernel(int *data) { printf("core %d read data %d\n", CoreID, data[CoreID]); } int main(void) { int host_data[] = {1, 2, 3, 4, 5, 6, 7, 8}; int *dev_data; stcMalloc((void **)&dev_data, sizeof(host_data)); stcMemcpy(dev_data, host_data, sizeof(host_data), stcMemcpyHostToDevice); kernel<<<8>>>(dev_data); stcDeviceSynchronize(); stcFree(dev_data); return 0; }
在主机端分配3GiB ~ 4GiB范围的全局内存并写入数据，然后在设备端将数据从高端内存拷贝到NPC可以直接访问的范围，并打印拷贝结果。
说明：NPC不能直接访问高端内存，因此本用例中，核函数内的
memcpy不能替换为data = high[CoreID]，否则会触发NPC读访问异常。#include <npurt.h> #include <hpe.h> __global__ void kernel(int *high) { int data; memcpy(&data, &high[CoreID], sizeof(data)); printf("core %d read data %d\n", CoreID, data); } int main(void) { int host_data[] = {1, 2, 3, 4, 5, 6, 7, 8}; int *dev_data; stcMallocHigh((void **)&dev_data, sizeof(host_data)); stcMemcpy(dev_data, host_data, sizeof(host_data), stcMemcpyHostToDevice); kernel<<<8>>>(dev_data); stcDeviceSynchronize(); stcFree(dev_data); return 0; }
2.4.3. 并行执行
主机端发送给设备端的请求类型包括执行核函数（Kernel）、从主机端向设备端拷贝数据（H2D）、从设备端向主机端拷贝数据（D2H）。主机端可以并行发起不同类型的请求，但一个或多个NPC Cluster内是顺序处理请求的。希姆计算提供了流的运行时接口，您可以基于流实现一个或多个NPC Cluster内并行处理请求。
在异构编程环境中，流分为UnifiedHS模式的流和DividedHS模式的流。在DividedHS模式上流在一个Cluster上运行，每一个同步域中包含8个NPC。在UnifiedHS模式下可以将一个核函数放在多个NPC Cluster乃至多张板卡上运行，适用于需要更多NPC并行以及更大片上内存的场景，例如运行大模型。在UnifiedHS模式上流在四个Cluster上运行，每一个同步域中包含32个NPC。
说明： UnifiedHS流上的核函数在执行时，因为其占用了所有32个NPC，所以其它普通流的核函数均无法被调度。
2.4.3.1. 流定义
流（stream）是由主机端发起、设备端处理的一系列请求。同一个流内的请求顺序处理，不同流间的不同类型请求可以并行处理。 流包括以下类型：
隐式声明流：默认创建的流，只有一个，可以包括同步、异步请求。执行核函数、拷贝数据时默认使用隐式声明流。
显示声明流：您自行创建的流，只能包括异步请求。执行核函数、拷贝数据（异步）时可以使用显示声明流。
说明：在一个NPC Cluster上创建的DividedHS模式的流只能在该NPC Cluster上使用。
下方的用例中，创建了一个显示声明流来执行核函数。如果需要等待流上的所有操作完成后再处理下一个请求，调用stcStreamSynchronize或stcDeviceSynchronize即可。
#include <npurt.h>
#include <hpe.h>
__global__ void kernel(void) {
printf("hello world in NPC %d/%d.\n", CoreID, CoreNum);
}
int main(void) {
stcStream_t stream;
stcStreamCreate(&stream);
kernel<<<8, stream>>>();
stcStreamSynchronize(stream);
stcStreamDestroy(stream);
return 0;
}
2.4.3.2. 流调度
同一流内的请求只能顺序调度，不同流间的不同类型请求可以并行调度。假定有三组请求，均包括H2D、Kernel、D2H请求，且不同组的请求之间没有依赖关系，采取并行调度可以提高性能。顺序调度和并行调度的性能对比如下：
说明：仅当所有cluster都空闲时才能执行UnifiedHS模式流中的kernel任务。在执行UnifiedHS模式下的kernel任务时，其它DividedHS流中的kernel任务不能被调度。
顺序调度和并行调度的代码示例如下：
顺序调度，顺序在NPC Cluster上处理每组请求。
#include <hpe.h> #include <npurt.h> #include <stdio.h> __global__ void kernel(int *in, int *out) { *out = *in; } #define NJOB 3 int main(void) { int host_in[NJOB] = {1, 2, 3}; int host_out[NJOB]; int *dev_in[NJOB], *dev_out[NJOB]; for (int i = 0; i < NJOB; i++) { stcMalloc((void **)&dev_in[i], sizeof(int)); stcMalloc((void **)&dev_out[i], sizeof(int)); } for (int i = 0; i < NJOB; i++) { stcMemcpyAsync(dev_in[i], &host_in[i], sizeof(int), stcMemcpyHostToDevice); kernel<<<1>>>(dev_in[i], dev_out[i]); stcMemcpyAsync(&host_out[i], dev_out[i], sizeof(int), stcMemcpyDeviceToHost); } stcDeviceSynchronize(); for (int i = 0; i < NJOB; i++) { printf("%d, ", host_out[i]); stcFree(dev_in[i]); stcFree(dev_out[i]); } printf("\n"); return 0; }
并行调度，为每组请求创建一个显示声明流，并在NPC Cluster上并行处理三个流中的请求。
#include <hpe.h> #include <npurt.h> #include <stdio.h> __global__ void kernel(int *in, int *out) { *out = *in; } #define NJOB 3 int main(void) { int host_in[NJOB] = {1, 2, 3}; int host_out[NJOB]; int *dev_in[NJOB], *dev_out[NJOB]; stcStream_t stream[NJOB]; for (int i = 0; i < NJOB; i++) { stcMalloc((void **)&dev_in[i], sizeof(int)); stcMalloc((void **)&dev_out[i], sizeof(int)); stcStreamCreate(&stream[i]); } for (int i = 0; i < NJOB; i++) { stcMemcpyAsync(dev_in[i], &host_in[i], sizeof(int), stcMemcpyHostToDevice, stream[i]); kernel<<<1, stream[i]>>>(dev_in[i], dev_out[i]); stcMemcpyAsync(&host_out[i], dev_out[i], sizeof(int), stcMemcpyDeviceToHost, stream[i]); } stcDeviceSynchronize(); for (int i = 0; i < NJOB; i++) { printf("%d, ", host_out[i]); stcFree(dev_in[i]); stcFree(dev_out[i]); stcStreamDestroy(stream[i]); } printf("\n"); return 0; }
2.4.3.3. 流同步
事件（event）用于在流中插入标记，当流中该标记前的请求处理完毕后，会将事件置为完成状态。事件具有以下用途：
监控流的进展：调用
stcEventElapsedTime获取处理两个事件间请求所消耗的时间。同步流的执行：在多流场景中，如果不同流中的请求之间有依赖关系，可以调用
stcStreamWaitEvent进行同步，在事件被置为完成状态后再开始处理其他流中的请求。
下方的用例中，定义了两个核函数kernel1、kernel2，kernel1循环执行共20次，在kernel1执行10次后添加事件，触发执行1次kernel2。
说明：调用
stcEventRecord添加事件时，如果不指定流，则为所有流添加事件，在所有流中标记前的请求都处理完毕后，才会将事件置为完成状态。
#include <hpe.h>
#include <npurt.h>
__global__ void kernel1(int index) {
printf("%s index %d come in\n", __func__, index);
}
__global__ void kernel2(void) {
printf("%s come in\n", __func__);
}
int main(void) {
stcStream_t stream1, stream2;
stcEvent_t event1;
stcStreamCreate(&stream1);
stcStreamCreate(&stream2);
stcEventCreate(&event1);
for (int i = 1; i <= 20; i++) {
kernel1<<<1, stream1>>>(i);
if (i == 10)
stcEventRecord(event1, stream1);
}
stcStreamWaitEvent(stream2, event1);
kernel2<<<1, stream2>>>();
stcDeviceSynchronize();
stcStreamDestroy(stream1);
stcStreamDestroy(stream2);
stcEventDestroy(event1);
return 0;
}
2.5. C++语言扩展接口
SHC完整兼容C++17标准，并针对异构程序扩展了语法和函数库。基于SHC提供的运行时接口编写代码时，您可以使用C++扩展语言接口方便地控制执行代码的逻辑。
2.5.1. 调用核函数
在SHC中，调用核函数的方式如下：
kernel_function<<<NCORE, stream, flags>>>(arg0, ...)
命令中配置部分和参数部分的含义如下所示：
| 命令内容 | 说明 |
|---|---|
| << |
指定设备端的配置，配置项含义如下： - NCORE：执行核函数所使用NPC的个数。 - stream：指定执行核函数时所在的流，默认为0代表使用隐式声明流。详细的流使用说明，请参见并行执行章节。 - flags：指定核函数的运行标志，默认为0（stcKernelFlagNone）代表无运行标志。详细的运行标志含义，请参见stcKernelFlag_t章节。 |
| (arg0, ...) | 指定核函数的参数，参数需要满足以下条件： - 参数列表中的变量类型和个数必须和核函数的定义相匹配。 - 每个参数的大小不能超过4字节。 |
2.5.2. 函数修饰符
SHC支持函数修饰符__host__、__global__、__device__，用于区分不同用途的函数。函数修饰符以及函数用途说明如下所示：
| 函数用途 | 修饰符 | 说明 |
|---|---|---|
| 主机端函数 | __host__ |
具有以下特点： - 在主机端执行，对设备端程序不可见。 - 主机端函数中可以调用主机端函数、核函数、双边函数，可以使用STL、libc、libc++库，支持递归调用自身。 - 参数来自于用户输入。 - 与普通C++函数没有任何区别，可以作为函数模板、类方法或匿名函数。 说明：如果函数没有添加任何函数修饰符，默认是主机端函数，编译器会自动添加函数修饰符。 |
| 核函数 | __global__ |
具有以下特点： - 能且只能由主机端函数调用，并在设备端执行。 - 核函数中可以调用设备端函数、双边函数，可以使用libnpurt库。 - 从栈上取参数。 - 能且只能是普通函数。 - 返回类型必须是void。 具有以下限制： - 不支持递归调用自身。 - 不能包含 long、longlong或double类型的参数。- 不支持使用其他变量修饰符修饰核函数的参数。 |
| 设备端函数 | __device__ |
具有以下特点： - 在设备端执行，对主机端程序不可见。 - 设备端函数中可以调用设备端函数、双边函数，可以使用libnpurt库，支持递归调用自身。 - 从寄存器取参数。 - 与普通C++函数没有任何区别，可以作为函数模板、类方法或匿名函数。 说明：如果需要函数需要在设备端执行，则不可省略函数修饰符。 |
| 双边函数 | __host__和__device__ |
具有以下特点： - 可以在主机端、设备端执行。 - 能且只能调用双边函数，支持递归调用自身。 - 可以被主机端函数、核函数、双边函数调用，参数来自于调用者。 - 双边函数一般是一些主机端和设备端都会用到的小型辅助函数，例如求数组最大值。 |
说明：不建议使用
__device__和__host__各自修饰类方法，会导致在主机端和设备端看到的类定义不同。
2.5.3. 变量修饰符
设备端不支持动态分配内存，但SHC支持通过变量修饰符__device__、 __local__、 __shared__、 __mutable__静态分配内存。每种变量修饰符对应不同的分配规则，各变量类型的基本属性如下：
| 变量修饰符 | 存放位置 | 有效范围 | 读写属性 | 变量初始化 | 多核访问 |
|---|---|---|---|---|---|
__device__ |
全局内存 | 全局变量、全局静态变量、局部变量、静态局部变量 | - 单独使用：只读 - 和 __thread联用：可读写 |
只读时，必须显式初始化。 | - 单独使用：可以被多个核共同访问。 - 和 __thread联用：变量在每个核都保存一份副本。对单核上的副本的修改，对其它核不可见。 |
__local__ |
本地内存 | 局部变量 | 可读写 | 推荐定义__local__数组作为buffer使用。说明： __local__变量允许初始化，但是初始化会降低执行效率，因此不推荐初始化。 |
定义的__local__变量在每个NPC都保存一份副本。对单核上的副本的修改，对其它核不可见。 |
__shared__ |
共享内存 | 局部变量 | 可读写 | 推荐定义__shared__数组作为buffer使用。说明： __shared__变量允许初始化，但是初始化会降低执行效率，因此不推荐初始化。另外，如果需要从多个NPC访问__shared__变量，读写变量前注意执行sync指令，保证在多个NPC中完成并发初始化。 |
定义的__shared__变量可以被多个NPC共同访问。如果希望使用__shared__变量在NPC间传递信息，需要使用volatile修饰。 |
__mutable__说明：不能单独使用，需要和 __device__一起使用，且不可以修饰指针。 |
全局内存 | 全局变量、全局静态变量、局部变量、局部静态变量 说明：全局变量地址默认非32字节对齐，您需要自行添加对齐属性 __attribute__((aligned(32)))。 |
可读写 | 可读写，因此无需显式初始化。 | 定义的__mutable__变量可以被多个核共同访问。如果希望使用__mutable__变量在核间传递信息，需要使用volatile修饰。 |
__imb__ |
本地内存 | 局部变量 | 可读写 | __imb__变量不允许初始化，只能作为高速buffer使用。 |
定义的__imb__变量在每个NPC都保存一份副本。对单核上的副本的修改，对其它核不可见。 |
变量类型会影响访问速度和内存占用情况：
访问速度
由变量存放位置决定，本地内存最快（
__local__、__imb__），共享内存次之（__shared__），全局内存最慢（__device__、__mutable__）。
内存占用
函数内所有的
__imb__变量和__local__变量的总大小不能超出本地内存空间（1.25MiB）。函数内所有的
__shared__变量的总大小不能超出共享内存空间（8MiB）。__local__、__shared__、__imb__所需内存在调用函数时动态分配，因此未调用函数时不占用空间，递归调用会占用多倍空间。全局内存空间（4GiB）充足，不太可能出现超出问题。
使用本章节所述的变量修饰符时，请注意以下限制：
仅支持在核函数、设备端函数中使用，不支持在主机端函数、双边函数中使用。
变量修饰符都不支持修饰函数参数。
变量修饰符都不支持修饰C++对象，仅支持修饰POD类型（例如普通变量、数组、结构体）。
2.6. 主机端编程
2.6.1. 接口调用要求
调用主机端运行时接口时，需要包含对应的头文件：
#include <hpe.h>
接口使用的限制：
因为PCIe BAR0空间增大为16GB，不能在大部分桌面型CPU平台上运行（例如物理地址寻址位数小于40的CPU）。
在没有PCIe Switch的intel平台上存在P2P性能很差的问题（约为0.85GB/s）。
在带PCIe Switch的Intel平台上，跨NUMA的P2P存在性能很差的问题（约为0.85GB/s）
主机端运行时接口提供以下功能：
设备管理：提供操作设备（NPC Cluster）相关的功能，例如指定待使用的设备、获取设备信息等。
内存管理：提供操作内存相关的功能，例如分配/释放内存、拷贝内存数据等。
执行控制：提供执行目标程序相关的功能，例如注册/释放目标程序、指定运行配置、启动核函数、加载/卸载目标程序等。
流管理：提供操作流相关的功能，例如创建/销毁流、创建/销毁事件、添加事件等。
错误处理：提供获取错误信息相关的功能，例如获取错误码、获取错误详情等。
说明：主机端运行时接口涉及的数据类型和环境变量，请参见数据类型和环境变量章节。
2.6.2. 运行时接口
2.6.2.1. 设备管理
2.6.2.1.1. stcSetDevice
函数描述：调用stcSetDevice设置用于执行设备端程序的NPC Cluster。
函数类型：同步函数
函数定义：
stcError_t stcSetDevice (int device)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| device | 输入参数 | int | NPC Cluster ID，由NPU设备标识符、NPC Cluster设备标识符推导得出。详细的推导说明，请参见指定NPC Cluster ID章节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.1.2. stcGetDevice
函数描述：调用stcGetDevice获取用于执行设备端程序的NPC Cluster。
函数类型：同步函数
函数定义：
stcError_t stcGetDevice (int *device)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| device | 输出参数 | int* | 指向所获取NPC Cluster ID的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.1.3. stcGetDeviceAttribute
函数描述：调用stcGetDeviceAttribute获取一个NPC Cluster所在NPU的属性，例如NPU硬件版本、NPC Cluster数量等。
函数类型：同步函数
函数定义：
stcError_t stcGetDeviceAttribute (int *value, stcDeviceAttr attr, int device)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| value | 输出参数 | int* | 指向所获取属性数据的指针。 |
| attr | 输入参数 | stcDeviceAttr | 属性名称，例如代表NPU硬件版本的stcDevAttrChipHWVersion。详细的属性名称含义，请参见stcDeviceAttr章节。 |
| device | 输入参数 | int | NPC Cluster ID。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.1.4. stcGetDeviceName
函数描述：调用stcGetDeviceName获取一个NPC Cluster所在NPU的设备名称。
函数类型：同步函数
函数定义：
stcError_t stcGetDeviceName (const char **name, int device)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| name | 输出参数 | const char ** | 指向所获取设备名称的二级指针。 |
| device | 输入参数 | int | NPC Cluster ID。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.1.5. stcGetDeviceCount
函数描述：调用stcGetDeviceCount获取所有NPU上NPC Cluster的数量。
函数类型：同步函数
函数定义：
stcError_t stcGetDeviceCount (int *count)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| count | 输出参数 | int* | 指向所获取NPC Cluster数量的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.1.6. stcDeviceSynchronize
函数描述：调用stcDeviceSynchronize等待当前进程的所有设备端操作执行结束。如果核函数执行异常退出，则输出触发异常时核函数的调用栈。
说明：只在对应cluster上有效，不能用来推测其他cluster的状态。
函数类型：同步函数
函数定义：
stcError_t stcDeviceSynchronize (void)
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2. 内存管理
2.6.2.2.1. stcMalloc
函数描述：调用stcMalloc在设备端全局内存的0GiB ~ 3GiB范围动态分配内存。
函数类型：同步函数
函数定义：
stcError_t stcMalloc (void **devPtr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| devPtr | 输出参数 | void** | 指向所分配内存地址的指针。 |
| size | 输入参数 | size_t | 所需分配的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.2. stcMallocHigh
函数描述：调用stcMallocHigh在设备端全局内存的3GiB ~ 4GiB范围（也称为高端内存）动态分配内存。
函数类型：同步函数
函数定义：
stcError_t stcMallocHigh(void **devPtr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| devPtr | 输出参数 | void** | 指向所分配内存地址的指针。 |
| size | 输入参数 | size_t | 所需分配的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.3. stcFree
函数描述：调用stcFree释放在设备端动态分配的全局内存。
函数类型：同步函数
函数定义：
stcError_t stcFree (void *devPtr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| devPtr | 输入参数 | void* | 指向待释放内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.4. stcMallocHost
函数描述：调用stcMallocHost在主机端分配内存，并设置为不会被换出的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcMallocHost (void **ptr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输出参数 | void** | 指向所分配内存地址的指针。 |
| size | 输入参数 | size_t | 所需分配的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.5. stcFreeHost
函数描述：调用stcFreeHost释放在主机端分配的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcFreeHost (void *ptr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输入参数 | void* | 指向待释放内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.6. stcHostRegister
函数描述：调用stcHostRegister将主机端的内存设置为不会被换出的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcHostRegister (void *ptr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输入参数 | void* | 指向待锁定内存的指针。 |
| size | 输入参数 | size_t | 所需锁定的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.7. stcHostUnregister
函数描述：调用stcHostUnregister解锁主机端的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcHostUnregister (void *ptr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输入参数 | void* | 指向待解锁内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.8. stcMemcpy
函数描述：调用stcMemcpy拷贝内存，该函数支持主机间拷贝、主机端向设备端拷贝、设备端向主机端拷贝。
函数类型：同步函数
函数定义：
stcError_t stcMemcpy (void *dst, const void *src, size_t count, stcMemcpyKind kind)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | void* | 拷贝操作的目的地址，向该指针指向的内存写入数据。 |
| src | 输入参数 | const void* | 拷贝操作的源地址，从该指针指向的内存读取数据。 |
| count | 输入参数 | size_t | 待拷贝内存的大小，单位为字节。 |
| kind | 输入参数 | stcMemcpyKind | 拷贝的方向，支持主机间拷贝、主机端向设备端拷贝、设备端向主机端。详细的类型说明，请参见stcMemcpyKind章节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.9. stcMemcpyAsync
函数描述：调用stcMemcpyAsync异步拷贝内存，调用后立即返回，不用等待拷贝完成。该函数支持主机间拷贝、主机端向设备端全局内存拷贝、设备端全局内存向主机端拷贝。
函数类型：异步函数
函数定义：
stcError_t stcMemcpyAsync (void *dst, const void *src, size_t count, stcMemcpyKind kind, stcStream_t stream=0)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | void* | 拷贝操作的目的地址，向该指针指向的内存写入数据。 |
| src | 输入参数 | const void* | 拷贝操作的源地址，从该指针指向的内存读取数据。 |
| count | 输入参数 | size_t | 待拷贝内存的大小，单位为字节。 |
| kind | 输入参数 | stcMemcpyKind | 拷贝的方向，支持主机间拷贝、主机端向设备端拷贝、设备端向主机端。详细的类型说明，请参见stcMemcpyKind章节。 |
| stream | 输入参数 | stcStream_t | 流标识符，默认为0，代表隐式声明流。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.10. stcMemcpyPeer
函数描述：调用stcMemcpyPeer在NPC Cluster间拷贝全局内存。
函数类型：同步函数
函数定义：
stcError_t stcMemcpyPeer (void *dst, int dstDevice, const void *src, int srcDevice, size_t count)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | void* | 拷贝操作的目的地址，向该指针指向的内存写入数据。 |
| dstDevice | 输入参数 | int | 目的NPC Cluster的ID。 |
| src | 输入参数 | const void* | 拷贝操作的源地址，从该指针指向的内存读取数据。 |
| srcDevice | 输入参数 | int | 源NPC Cluster的ID。 |
| count | 输入参数 | size_t | 待拷贝内存的大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.2.11. stcMemcpyPeerAsync
函数描述：调用stcMemcpyPeerAsync在NPC Cluster间异步拷贝全局内存，调用后立即返回，不用等待拷贝完成。
函数类型：异步函数
函数定义：
stcError_t stcMemcpyPeerAsync (void *dst, int dstDevice, const void *src, int srcDevice, size_t count, stcStream_t stream=0)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | void* | 拷贝操作的目的地址，向该指针指向的内存写入数据。 |
| dstDevice | 输入参数 | int | 目的NPC Cluster的ID。 |
| src | 输入参数 | const void* | 拷贝操作的源地址，从该指针指向的内存读取数据。 |
| srcDevice | 输入参数 | int | 源NPC Cluster的ID。 |
| count | 输入参数 | size_t | 待拷贝内存的大小，单位为字节。 |
| stream | 输入参数 | stcStream_t | 流标识符，默认为0，代表隐式声明流。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.3. 执行控制
2.6.2.3.1. stcRegisterFatBinary
函数描述：调用stcRegisterFatBinary注册在设备端执行的目标程序。
函数类型：同步函数
函数定义：
stcError_t stcRegisterFatBinary (const void *data)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| data | 输入参数 | const void* | 指向目标程序所占内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.3.2. stcUnregisterFatBinary
函数描述：调用stcUnregisterFatBinary释放所有已注册的目标程序。
函数类型：同步函数
函数定义：
stcError_t stcUnregisterFatBinary (void)
函数参数：
无
函数返回值：
stcSuccess
2.6.2.3.3. stcConfigureCall
函数描述：调用stcConfigureCall指定执行核函数的配置。
函数类型：同步函数
函数定义：
stcError_t stcConfigureCall(int core_num, stcStream_t stream=0, unsigned int flags=stcKernelFlagNone)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| core_num | 输入参数 | int | 在一个NPC Cluster上并行执行核函数所使用的NPC个数。 |
| stream | 输入参数 | stcStream_t | 流标识符，默认为0，代表将配置应用于隐式声明流。 |
| flags | 输入参数 | unsigned int | 指定核函数的运行标志，默认无运行标志（stcKernelFlagNone）。详细的运行标志含义，请参见stcKernelFlag_t章节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.3.4. stcLaunchKernel
函数描述：调用stcLaunchKernel启动核函数。
函数类型：异步函数
函数定义：
stcError_t stcLaunchKernel（stcModule_t *module, const char *kname, stcKernelParams_t kernelParams)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输出参数 | stcModule_t* | 指向核函数所在目标程序模块的指针。 |
| kname | 输入参数 | const char* | 核函数的名称。 |
| kernelParams | 输入参数 | stcKernelParams_t | 核函数的参数信息。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.3.5. stcModuleLoadData
函数描述：调用stcModuleLoadData将目标程序加载到设备端。
函数类型：同步函数
函数定义：
stcError_t stcModuleLoadData (stcModule_t *module, const void *data, size_t size, stcStream_t stream=0)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输出参数 | stcModule_t* | 指向加载到设备端的目标程序模块的指针。 |
| data | 输入参数 | const void* | 待加载目标程序的内存地址。 |
| size | 输入参数 | size_t | 待加载目标程序占用的内存大小，单位为字节。 |
| stream | 输入参数 | stcStream_t | 流标识符，默认为0，代表将配置应用于隐式声明流。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.3.6. stcModuleUnload
函数描述：调用stcModuleUnload从设备端卸载目标程序模块。
函数类型：同步函数
函数定义：
stcError_t stcModuleUnload (stcModule_t module)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输入参数 | stcModule_t | 待释放的目标程序模块。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4. 流管理
2.6.2.4.1. stcStreamCreate
函数描述：调用stcStreamCreate在cluster范围内创建一个具有DividedHS模式的流。
函数类型：同步函数
函数定义：
stcError_t stcStreamCreate (stcStream_t *pStream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pStream | 输出参数 | stcStream_t* | 指向创建的流标识符的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 返回值为0：成功； 返回值为其他值：错误码，其中详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.2. stcStreamCreateUnified
函数描述：使用stcStreamCreateUnified在NPU范围内创建的一个具有UnifiedHS模式的流。
函数类型：同步函数
函数定义：
stcError_t stcStreamCreateUnified(stcStream_t *pStream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pStream | 输出参数 | stcStream_t* | 指向创建的流标识符的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 返回值为0：成功； 返回值为其他值：错误码，其中详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.3. stcStreamDestroy
函数描述：调用stcStreamDestroy销毁指定流。如果流中有未完成的操作，会终止执行并释放相关资源。
函数类型：同步函数
函数定义：
stcError_t stcStreamDestroy (stcStream_t stream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 待销毁流的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 返回值为0：成功； 返回值为其他值：错误码，其中详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.4. stcStreamDestroyUnified
函数描述：调用stcStreamDestroyUnified销毁指定的UnifiedHS流。在销毁前会清除正在运行和待运行的任务。
函数类型：同步函数
函数定义：
stcError_t stcStreamDestroyUnified(stcStream_t pStream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pstream | 输入参数 | stcStream_t | 待销毁流的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 返回值为0：成功；返回值为其他值：错误码，其中详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.5. stcStreamSynchronize
函数描述：调用stcStreamSynchronize等待指定流上的所有操作执行结束。如果核函数执行异常退出，则输出触发异常时核函数的调用栈。
函数类型：同步函数
函数定义：
stcError_t stcStreamSynchronize (stcStream_t stream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 流标识符，等待该流上的所有操作执行结束。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 返回值为0：成功；返回值为其他值：错误码，其中详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.6. stcStreamSynchronizeUnified
函数描述：调用stcStreamSynchronizeUnified等待指定流上的所有操作执行结束。如果核函数执行异常退出，则输出触发异常时核函数的调用栈。
说明：在UnifiedHS模式下，核函数需要加载四次到所有四个Cluster上，不同Cluster开始执行核函数的时间不是精确同步的。如果需要保证核函数在同一时间开始执行有效代码，用户需要自行在核函数最开始处显示地执行同步指令。同样的，如果需要保证核函数在所有Cluster同时完成，您需要自行在核函数结束前显示地执行同步指令。
函数类型：同步函数
函数定义：
stcError_t stcStreamSynchronizeUnified(stcStream_t pStream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pstream | 输入参数 | stcStream_t | 流标识符，等待该流上的所有操作执行结束。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 返回值为0：成功；返回值为其他值：错误码，其中详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.7. stcStreamClean
函数描述：调用stcStreamClean停止处理流上的请求并销毁所有请求，但不会销毁流本身。
函数类型：同步函数
函数定义：
stcError_t stcStreamClean(stcStream_t stream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 流标识符，停止处理该流的请求，并销毁所有请求。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.8. stcEventCreate
函数描述：调用stcEventCreate创建一个事件。
函数类型：同步函数
函数定义：
stcError_t stcEventCreate(stcEvent_t *pEvent)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pEvent | 输出参数 | stcEvent_t* | 指向创建的事件标识符的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.9. stcEventDestroy
函数描述：调用stcEventDestroy销毁指定事件。
函数类型：同步函数
函数定义：
stcError_t stcError_t stcEventDestroy(stcEvent_t event)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| event | 输入参数 | stcEvent_t | 待销毁事件的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.10. stcEventSynchronize
函数描述：调用stcEventSynchronize等待指定事件进入完成状态。
函数类型：同步函数
函数定义：
stcError_t stcEventSynchronize(stcEvent_t event)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| event | 输入参数 | stcEvent_t | 事件标识符，等待流中该事件前的所有操作执行结束后，才会将事件置为完成状态。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.11. stcEventRecord
函数描述：调用stcEventRecord在指定流的当前运行点添加事件。
函数类型：同步函数
函数定义：
stcEventRecord(stcEvent_t event, stcStream_t stream=0)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| event | 输入参数 | stcEvent_t | 待添加事件的标识符。 |
| stream | 输入参数 | stcStream_t | 待添加事件的流的标识符，默认为0，代表在所有流的当前运行点添加事件。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.12. stcEventElapsedTime
函数描述：调用stcEventElapsedTime获取处理两个事件间请求所消耗的时间。
函数类型：同步函数
函数定义：
stcError_t stcEventElapsedTime(float *ms, stcEvent_t start, stcEvent_t end)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ms | 输出参数 | float* | 指向所获取消耗时间的指针，消耗时间的单位为ms。 |
| start | 输入参数 | stcEvent_t | 开始事件的标识符。 |
| end | 输入参数 | stcEvent_t | 结束事件的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.4.13. stcStreamWaitEvent
函数描述：调用stcStreamWaitEvent在多流场景中建立流之间的同步关系，指定某个流需要等指定事件被置为完成状态后再开始处理请求。
函数类型：同步函数
函数定义：
stcError_t stcStreamWaitEvent(stcStream_t stream, stcEvent_t event)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 流标识符，该流在对应事件被置为完成状态后再开始处理请求。 |
| event | 输入参数 | stcEvent_t | 事件标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.5. 卡卡互联
2.6.2.5.1. stcSetP2PMap
函数描述：设置跨卡程序需要使用的板卡信息，包含板卡的数量，以及板卡的编号列表。
说明：所有的板卡均需要设置，并且在使用前请使用stcSetDevice接口指定板卡。
函数类型：同步函数
函数定义：
stcError_t stcSetP2PMap(int mapcount, int *pmaparray)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| mapcount | 输入参数 | int | 跨卡程序使用的卡数量。 |
| *pmaparray | 输入参数 | int | 跨卡程序使用的卡的物理ID集合。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.5.2. stcClearP2PMap
函数描述：清除跨卡程序的板卡信息。
说明：所有的板卡均需要设置，并且在使用前请使用stcSetDevice接口指定板卡。
函数类型：同步函数
函数定义：
stcError_t stcClearP2PMap()
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.6. 错误处理
2.6.2.6.1. stcGetLastError
函数描述：调用stcGetLastError返回一个线程调用主机端运行时接口时产生的最后一个错误，然后将结果重置为stcSuccess。如果没有错误，则返回stcSuccess。
函数类型：同步函数
函数定义：
stcError_t stcGetLastError(void)
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
2.6.2.6.2. stcGetErrorName
函数描述：调用stcGetErrorName从获取的错误码得到错误名称。
函数类型：同步函数
函数定义：
const char* stcGetErrorName (stcError_t error)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| error | 输入参数 | stcError_t | 调用stcGetLastError获取到的错误码。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| const char* | 指向错误名称字符串的指针，字符串以NULL结尾。 |
2.6.2.6.3. stcGetErrorString
函数描述：调用stcGetErrorString从获取的错误码得到错误详情。
函数类型：同步函数
函数定义：
const char* stcGetErrorString (stcError_t error)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| error | 输入参数 | stcError_t | 调用stcGetLastError获取到的错误码。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| const char* | 指向NULL结尾的字符串。 |
2.6.2.7. 数据类型
2.6.2.7.1. stcKernelParams_t
数据描述：记录一个核函数的所有参数信息。
2.6.2.7.2. stcModule_t
数据描述：记录一个NPC Cluster设备目标程序模块的相关信息。
2.6.2.7.3. stcStream_t
数据描述：记录一个流的相关信息。
2.6.2.7.4. stcEvent_t
数据描述：记录一个事件的相关信息。
2.6.2.7.5. stcError_t
数据描述：记录调用主机端运行时接口时返回的错误。支持返回的错误类型如下所示：
| 枚举成员 | 枚举值 | 说明 |
|---|---|---|
| stcSuccess | 0 | 函数调用成功，未返回错误。 |
| stcErrorInvalidValue | 1 | 一个或多个参数的取值超出了有效值范围。 |
| stcErrorInvalidDevice | 2 | 使用了无效的NPC Cluster ID。 |
| stcErrorHostMemoryAllocation | 3 | 在主机端分配内存失败。 |
| stcErrorDeviceMemoryAllocation | 4 | 在设备端分配内存失败。 |
| stcErrorInvalidDevicePointer | 5 | 使用了无效的设备端内存地址。 |
| stcErrorLinkFailure | 6 | 设备端目标程序链接失败。 |
| stcErrorInvalidKernel | 7 | 使用了无效的核函数名称。 |
| stcErrorInvalidImage | 8 | 设备端目标程序不可用。 说明：设备端目标程序对应fat binary，而设备端目标模块则对应具体的binary。 |
| stcErrorNoImage | 9 | 设备端目标程序不存在。 |
| stcErrorInvalidModule | 10 | 设备端目标模块不可用。 说明：设备端目标程序对应fat binary，而设备端目标模块则对应具体的binary。 |
| stcErrorNoModule | 11 | 设备端目标模块不存在。 |
| stcErrorInvalidStream | 12 | 流不可用。 |
| stcErrorInvalidEvent | 13 | 事件不可用。 |
| stcErrorDeviceImageException | 15 | 设备端目标程序运行时出现异常。 |
| stcErrorSyscallFailure | 16 | 主机端的系统调用失败。 |
| stcErrorForkForbidden | 17 | 父进程已调用过运行时接口，禁止子进程再次调用。 |
| stcErrorInvalidCoreNum | 18 | 使用了无效的核数量。 |
| stcErrorGDBFailure | 19 | stc-gdb跟踪失败，无法获取调试信息。 |
| stcErrorDriverMismatch | 20 | NPU驱动版本和NPU设备不匹配。 |
| stcErrorDeviceBreakdown | 21 | NPU设备出现故障，无法继续使用。 |
| stcErrorInvalidNpuTask | 24 | NPU任务不可用。 |
| stcErrorInvalidImageDataSection | 25 | 设备端目标程序包含了太多可写数据段，导致程序无法正常运行。 |
| stcErrorOutOfHostMemory | 27 | 主机端内存不足。 |
| stcErrorTooManyOpenStreams | 28 | 打开流的数量过多，已超过上限。 |
| stcErrorTooManyOpenEvents | 29 | 打开事件的数量过多，已超过上限。 |
| stcErrorDriverFailure | 99 | NPU驱动错误。 |
2.6.2.7.6. stcMemcpyKind_t
数据描述：设置拷贝内存操作的类型。支持设置的操作类型如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| stcMemcpyHostToHost | 0 | 在同一台主机的不同内存段间拷贝数据。 |
| stcMemcpyHostToDevice | 1 | 从主机端向设备端拷贝数据。 |
| stcMemcpyDeviceToHost | 2 | 从设备端向主机端拷贝数据。 |
2.6.2.7.7. stcDeviceAttr_t
数据描述：记录NPU设备的属性。支持查看的属性如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| stcDevAttrChipHWVersion | 0 | NPU硬件版本。 |
| stcDevAttrBoardHWVersion | 1 | 板卡硬件版本。 |
| stcDevAttrClusterCount | 2 | NPC Cluster的数量。 |
| stcDevAttrNPCPerCluster | 3 | 每个NPC Cluster包含的NPC数量。 |
| stcDevAttrSharedmemPerCluster | 4 | 每个NPC Cluster中共享内存的大小。 |
| stcDevAttrGlobalmemPerCluster | 5 | 每个NPC Cluster中全局内存的大小。 |
| stcDevAttrConcurrentKernels | 6 | 每个NPC Cluster可以并行执行的核函数的数量。 |
| stcDevAttrPciBusId | 7 | PCIe总线的ID。 |
| stcDevAttrPciDeviceId | 8 | PCIe设备的ID，即希姆计算板卡的ID，例如STCP920的为0100。 |
| stcDevAttrFirmwareVersion | 9 | 设备端固件的版本。 |
| stcDevAttrDriverVersion | 10 | 主机端安装的设备驱动版本。 |
| stcDevAttrCount | 11 | 支持查看的属性数量，即本表格中除stcDevAttrCount外枚举定义的数量。 |
2.6.2.7.8. stcKernelFlag_t
数据描述：设置核函数的运行标志。支持设置的运行标志如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| stcKernelFlagNone | 0 | NPC执行核函数前后均处理DCache，可以视为无运行标志。 |
| stcKernelFlagInputDataBypassDcache | 1 | NPC执行核函数前，不需要处理DCache。 |
| stcKernelFlagOutputDataBypassDache | 2 | NPC执行核函数后，不需要处理DCache。 |
2.6.2.7.9. stcKernelData_t
数据描述：核函数输入输出对应的主机端数据区的属性。包含的成员变量如下所示：
| 成员变量 | 说明 |
|---|---|
| data | 数据区的地址。 |
| size | 数据区的大小。 |
2.6.2.7.10. stcP2PMap_t
数据描述：板卡信息，包含板卡的数量，以及板卡的编号列表。包含的成员变量如下所示：
| 成员变量 | 说明 |
|---|---|
| npu_nr | 板卡的数量。 |
| npu_ids | 板卡的编号列表。 |
2.6.2.8. 环境变量
2.6.2.8.1. STC_SET_DEVICES
数据描述：运行时修改stcSetDevice的执行结果，支持同时设置多个NPC Cluster ID，以半角逗号分隔即可。每个设置的索引为原NPC Cluster ID，设置的值为新NPC Cluster ID。示例如下：
| 设置 | 说明 |
|---|---|
| export STC_SET_DEVICES=2 | 设置STC_SET_DEVICES为2后：- 默认使用NPC Cluster 2。 - 执行 stcSetDevice的起始NPC Cluster ID为2。例如，stcSetDevice(0)代表使用NPC Cluster ID为2，stcSetDevice(1)代表使用NPC Cluster ID为3。 |
| export STC_SET_DEVICES=2, 3 | 设置STC_SET_DEVICES为2和3后：- 默认使用NPC Cluster 2，可以使用NPC Cluster 2和NPC Cluster 3。 - 执行 stcSetDevice的起始NPC Cluster ID为2。例如，stcSetDevice(0)代表使用NPC Cluster 2，stcSetDevice(1)代表使用NPC Cluster 3。 |
2.7. 设备端编程
2.7.1. 接口调用要求
调用设备端运行时接口时，需要包含对应的头文件：
#include <npurt.h>
说明：设备端运行时接口涉及的数据类型和变量，请参见数据类型和变量章节。
2.7.2. 运行时接口
2.7.2.1. exit
函数描述：调用exit退出核函数并返回您指定的退出码。
函数类型：同步函数
函数定义：
__device__ void exit(int exit_code)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| exit_code | 输入参数 | int | 您自行定义的退出码。 |
函数返回值：
无
2.7.2.2. abort
函数描述：调用abort退出核函数并返回退出码128。
函数类型：同步函数
函数定义：
__device__ void abort()
函数参数：
无
函数返回值：
无
2.7.2.3. assert
函数描述：调用assert验证指定的条件，在未满足断言时，退出核函数并返回退出码128。
函数类型：同步函数
函数定义：
__device__ void assert(int exp)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| exp | 输入参数 | int | 指定的条件。如果未满足断言，则值为0，退出核函数并返回错误码128。 |
函数返回值：
无
2.7.2.4. memcpy
函数描述：调用memcpy拷贝内存，该函数支持在设备端的本地内存与共享内存之间、共享内存与全局内存之间拷贝内存，支持以指定内存大小或信息格式（通过cpy_config_t描述）的形式拷贝。
函数类型：同步函数
函数定义：
__device__ int memcpy(void* dest, void* src, int len)
__device__ int memcpy(void* dest, void* src, cpy_config_t info)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dest | 输入参数 | void* | 拷贝操作的目的地址，向该指针指向的内存写入数据。 |
| src | 输入参数 | void* | 拷贝操作的源地址，从该指针指向的内存读取数据。 |
| len | 输入参数 | int | 待拷贝内存的大小，单位为字节。 |
| info | 输入参数 | cpy_config_t | 待拷贝内存的信息格式的描述，详细的信息格式说明，请参见cpy_config_t章节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| int | 详细的错误类型含义，请参见npurtError_t章节。 |
2.7.2.5. get_npu_idx
函数描述：获取调用者所在NPU卡的ID。
函数类型：同步函数
函数定义：
int get_npu_idx()
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| int | NPU卡的逻辑ID，范围：0~（NPU卡数量-1）。 |
2.7.2.6. get_npu_nr
函数描述：查询程序可以使用的NPU卡数量。
函数类型：同步函数
函数定义：
int get_npu_nr()
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| int | 返回值为0：非跨卡程序 返回值为其他数值：跨卡程序使用的NPU卡数量 |
2.7.2.7. memcpy_p2p
函数描述：在本地发起一个向远端NPU DDR的数据传输，将本NPU卡的数据传输到远端NPU卡。当远端flag_addr为非0时，可以同时向该远端地址写入一个4字节的flag值。
由于卡卡互联的功能性与主机系统的配置（包括BIOS设置和Linux引导参数）有关系，所以通常情况下，完全关闭CPU的虚拟化功能即可支持P2P传输。若需在开启CPU的虚拟化功能的情况下使用P2P功能，则需设置intel_iommu=on（或amd_iommu=on）和iommu=pt。有时，也可能需要升级主机的BIOS（或CPU Microcode）才可使能P2P功能。
说明：每个NPC均可发起P2P传输任务，传输任务是异步执行的。在该任务完成前，该NPC不能发起另一个任务；此时如果再次调用该接口，会等待上次一任务完成后才能完成新任务提交。
函数类型：同步函数
函数定义：
int memcpy_p2p(uint64_t dst, uint64_t src, uint64_t len, int dst_npu_idx, uint64_t flag_addr, uint32_t flag)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | uint64_t | 远端NPU卡的目的地址，必须4字节对齐。 |
| src | 输入参数 | uint64_t | 本地NPU卡的源地址，必须4字节对齐。 |
| len | 输入参数 | uint64_t | 传输的数据长度，必须为4字节倍数。当len为零时，不传输数据。 |
| dst_npu_idx | 输入参数 | int | 远端NPU卡的ID。 |
| flag_addr | 输入参数 | uint64_t | 远端NPU卡的flag地址。当flag_addr为零时，不传输flag的值。 |
| flag | 输入参数 | uint32_t | flag的值。 |
说明：当len和flag_addr都为零时，为无效传输，系统报错。
函数返回值：
| 类型 | 说明 |
|---|---|
| int | 返回值为0：任务创建成功（始终成功）。 |
2.7.2.8. memcpy_p2p_wait
函数描述：等待本地发起的p2p传输完成。
函数类型：同步函数
函数定义：
void memcpy_p2p_wait()
函数参数：
无
函数返回值：
无
2.7.2.9. wait_flag
函数描述：等待远端NPU卡上的地址值被修改为预期标识值，用于确认远端发起的P2P传输是否完成。
函数类型：同步函数
函数定义：
void wait_flag(uint32_t addr, uint32_t flag)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| addr | 输入参数 | uint32_t | 远端NPU卡的目的地址。 |
| flag | 输入参数 | uint32_t | 预期标识值。 |
函数返回值：
无
2.7.3. 数据类型
2.7.3.1. npurtError_t
数据描述：记录调用设备端运行时接口时返回的错误。支持返回的错误类型如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| npurtSuccess | 0 | 函数调用成功，未返回错误。 |
| npurtErrorSysdmaInvalidDmaID | 1 | 使用SysDMA在共享内存和全局内存间传输数据时，指定了无效的DMA控制器ID。 |
| npurtErrorSysdmaInvalidChannelID | 2 | 指定了无效的DMA通道ID。 |
| npurtErrorSysdmaInvalidDataRow | 3 | 指定了无效的数据行。 |
| npurtErrorSysdmaInvalidDataCol | 4 | 指定了无效的数据列。 |
| npurtErrorSysdmaInvalidDataStride | 5 | 指定了无效的数据Stride。 |
| npurtErrorSysdmaInvalidDataSize | 6 | 指定了无效的数据大小。 |
| npurtErrorSysdmaInvalidState | 7 | 无效的DMA状态。 |
| npurtErrorSysdmaReqFull | 8 | DMA请求队列已满，无法处理更多请求。 |
| npurtErrorSysdmaInvalidAddress | 9 | 指定了无效的地址。 |
| npurtErrorSysdmaInvalidXferType | 10 | 指定了无效的传输类型。 |
| npurtErrorSysdmaBusy | 11 | DMA正在使用中，无法执行其他任务。 |
| npurtErrorSysdmaXferFailure | 12 | 数据传输失败。 |
| npurtMemcpyIncompatibleParamLen | 13 | 拷贝内存时，指定了无法识别的参数。 |
| npurtErrorMemcpyInvalidAddress | 14 | 指定了无效的地址。 |
| npurtErrorMemcpyInvalidConfigRows | 15 | 指定了无效的数据行。 |
| npurtErrorMemcpyInvalidConfigCols | 16 | 指定了无效的数据列。 |
| npurtErrorMemcpyInvalidDataStrides | 17 | 指定了无效的数据Stride。 |
| npurtErrorMemcpyInvalidConfigDtype | 18 | 指定了无效的数据类型。 |
| npurtErrorMemcpyInvalidConfigDestcore | 19 | 指定了无效的目标NPC。 |
| npurtErrorMemcpyInvalidDataSize | 20 | 指定了无效的数据大小。 |
| npurtErrorMemcpyOverflowBoundary | 21 | 超过了源或目的地址空间。 |
| npurtErrorMemcpyGlbmem32Unaligned | 22 | 发现源或目的地址没有32字节对齐。 |
2.7.3.2. cpy_config_t
数据描述：待拷贝内存的信息格式。包含的成员变量如下：
| 成员变量 | 说明 |
|---|---|
| rows | 待拷贝数据的行数。 |
| cols | 待拷贝数据的列数。 |
| stride_cols | 拷贝数据时使用的列数Stride。 |
| dtype | 待拷贝数据的类型。详细的数据类型定义，请参见cpy_dtype_t章节。 |
| dest_core | 拷贝数据时NPC Cluster内目标NPC的标识符，该参数只适用于在不同NPC的本地内存间传输数据。 说明：目前 memcpy仅支持在同一NPC Cluster内拷贝。 |
2.7.3.3. cpy_dtype_t
数据描述：待拷贝数据的类型。支持的数据类型如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| DTYPE_BYTE_1 | 0 | 每个数据元素占一个字节大小。 |
| DTYPE_BYTE_2 | 1 | 每个数据元素占两个字节大小。 |
2.7.4. 变量
2.7.4.1. CoreID
描述：获取NPC在NPC Cluster内的标识符，例如STCP920中CoreID的范围为0 ~ 7。
类型：整型只读
2.7.4.2. CoreNum
描述：获取运行核函数时所使用NPC的个数。
类型：整型只读