希姆计算异构编程手册
版本历史
| 版本 | 作者 | 日期 | 说明 |
|---|---|---|---|
| V1.2.0 | 希姆计算 | 2022-04-11 | - 文档版本号对齐HPE发版。 - 整篇编辑优化。 |
| V1.0.0 | 希姆计算 | 2021-09-01 | 初始版本。 |
概述
希姆计算推出了异构编程环境HPE（Heterogeneous Programming Engine），安装HPE后您可以使用C/C++语言开发异构程序，方便地使用NPU进行并行计算。
异构程序包括在CPU上运行的主机端程序和在NPU上运行的设备端程序，希姆计算的异构编程环境为开发、编译、运行主机端程序和设备端程序提供了完整的工具链。主要包括以下模块：
| 模块 | 全称 | 说明 |
|---|---|---|
| SHC | Stream Computing Heterogeneous C++ | 异构编程语言，完整兼容C++17标准，并针对异构程序扩展了语法和函数库，例如支持调用核函数的<<< >>>符号。 |
| stcc | Stream Computing Heterogeneous C++ Compiler | 异构程序编译器，统一编译异构程序中的主机端程序和设备端程序，生成可执行文件。 |
| stc-dkms | Stream Computing Dynamic Kernel Module Support | 异构驱动模块，对应有异构驱动stc.ko文件。主机端与设备端使用PCIe总线连接，异构驱动将主机端的操作请求通过PCIe总线传递到设备端。 stc-dkms方便异构驱动适应不同的Linux内核版本，不用逐版本手动编译也能保证异构驱动始终可用，主机端和设备端的交互不会因Linux内核版本变化而产生异常。 |
| stc-smi | Stream Computing System Management Interface | 设备管理工具，管理和监控NPU设备的状态。 |
| stc-prof | Stream Computing Profiler | 性能调优工具，采集异构程序的性能数据，分析异构程序的性能。 说明：stc-prof是命令行工具，希姆计算还提供了接口形式的性能调优工具STCPTI（ Stream Computing Profiling Tool Interface ）。 |
| stc-gdb | Stream Computing Debugger | 异构程序调试工具，监视程序运行状态，获取和修改程序的中间运行结果。stc-gdb完全兼容GDB原生命令，并添加了希姆计算的扩展命令。 |
| hpert | Heterogeneous Programming Engine Runtime | 主机端运行时库，为主机端程序提供控制设备端内存访问、核函数执行等操作的接口。 |
| npurt | NPU Runtime | 设备端运行时库，为设备端程序提供信息打印、内存拷贝等操作的接口。 |
基本概念
| 概念 | 全称 | 说明 |
|---|---|---|
| NPU | Neural-network Processing Unit | 神经网络处理器，专门面向大规模神经元和突触处理的处理器。采用数据驱动并行计算的架构，非常适合并行计算的场景。 |
| NerualScale | NerualScale | 希姆计算自研的NPU神经网络计算加速架构，支持基于RISC-V指令集的扩展指令，具有良好的可编程性。 |
| STCP920 | Stream Computing P920 | 希姆计算推出的云端AI推理计算加速板卡，内含可编程NPU，基于NerualScale架构、软硬件协同设计。 |
| NPC | Neural-network Processing Core | 希姆计算自研NPU中的神经网络处理核心。 每个NPU中包含多个NPC Cluster，例如STCP920中包含4个NPC Cluster，每个NPC Cluster包括8个NPC，因此STCP920中共包含32个NPC。 每个NPC Cluster通过NPC Cluster ID唯一标识，对N个NPC Cluster，异构编程环境中使用0至N-1依次标记，同一个NPU中的所有NPC Cluster的ID是连续的。 |
| 核函数 | Kernel Function | 一个核函数对应一个计算任务，由主机端（CPU）部署到设备端（NPU），并在设备端的多个NPC上并行执行。 |
| 全局内存（DDR） | Double Data Rate Synchronous Dynamic Random-Access Memory | 每个NPC Cluster私有的内存，由NPC Cluster内的NPC共享，访问速度最慢。STCP920中，每个NPC Cluster的全局内存大小为4GiB。 |
| 共享内存（LLB） | Last Level Buffer | 每个NPC Cluster私有的内存，由NPC Cluster内的NPC共享，访问速度较快。STCP920中，每个NPC Cluster的共享内存大小为8MiB。 |
| 本地内存（L1） | L1 | 每个NPC私有的内存，访问速度最快。STCP920中，每个NPC的本地内存大小为1.25MiB。 |
| sysDMA | System Direct Memory Access | 数据传输通道，用于在NPC Cluster内的共享内存和全局内存间传输数据。 |
异构编程模型
在异构编程环境中，设备端与主机端在计算和存储结构上存在差别，因此需要通过不同的方式执行函数和访问内存。异构编程模型中需要考虑核函数、内存管理、异步函数等因素。
核函数
定义核函数
异构程序包括主机端程序和设备端程序，并引入核函数的概念来关联两类程序，核函数是两类代码的桥梁。核函数在主机端调用，由主机端部署到设备端，并在设备端的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;
}
编译核函数
使用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
执行核函数
运行编译生成的二进制文件，在执行到核函数时，会自动转到设备端。
$ ./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上立即执行下一个核函数。
内存管理
内存布局
希姆计算自研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的内存布局完全相同，如下所示：
访问内存
在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刷新缓存会导致无法保障全局内存中数据的正确性。
访问内存用例
下方的用例中，在主机端分配全局内存并写入数据，然后在设备端使用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;
}
异步函数
部分运行时接口的函数类型设计为异步函数，在主机端调用异步函数后会立即返回，不用等待完成异步函数规定的所有操作，有利于主机端和设备端并行处理任务。例如，调用stcLaunchKernel在主机端启动核函数后，主机端无需等待完成核函数的所有操作即可开始处理下一个任务；调用stcMemcpyAsync在主机端和设备端之间拷贝数据后，主机端无需等待拷贝完所有数据即可开始处理下一个任务。
异步函数的返回值不是异步操作的结果。如果您调用了异步函数，但仍然需要等待所有操作完成，可以调用stcDeviceSynchronize或stcStreamSynchronize等待设备或流上的所有操作完成后再处理下一个任务。stcDeviceSynchronize和stcStreamSynchronize的返回值是异步操作的结果，但仅返回最后一个异步操作的结果。
说明：核函数没有返回值，如果核函数在真正开始执行前出错退出，只能通过
stcGetLastError获取错误信息。例如指定的NPC数量超过实际的NPC数量，导致启动核函数后还未实际执行就出错退出。
典型操作
希姆计算提供了丰富的运行时接口，方便您从主机端控制在设备端执行任务。
指定运行设备
主机端访问设备端时需要指定运行设备。在主机端程序中启动核函数后，默认在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。
说明：STCP920的NPU包含4个Cluster，因此在安装多张STCP920时需要依次顺序标记。
$ stc-smi -q NPU: 0 Product Name: STCP920 Chip count: 1 Temperature: 40.51C Fan speed: N/A Power: 33.46W Cluster count: 4 Frequency: 1000M Bus: 0000:65:00.0 Vendor: 23e2 Device: 0100 Current link speed: 16 GT/s Max link speed: 16 GT/s Current link width: 16 Max link width: 16 Write bytes: 0B Read bytes: 9579B HPE version: 1.1.2 Driver: 1.1.1 Chip version: 20200102 MCU firmware: 1.0.3 NPU ctrl firmware: 1.1.5 Cluster 0: Frequency: 1000M Core count: 8 DMA count: 2 Status: WORK Memory Used/Total: 156.00M/4.00G Npc status: IDLE IDLE IDLE IDLE IDLE IDLE IDLE IDLE Cluster 1: Frequency: 1000M Core count: 8 DMA count: 2 Status: WORK Memory Used/Total: 156.00M/4.00G Npc status: IDLE IDLE IDLE IDLE IDLE IDLE IDLE IDLE Cluster 2: Frequency: 1000M Core count: 8 DMA count: 2 Status: WORK Memory Used/Total: 156.00M/4.00G Npc status: IDLE IDLE IDLE IDLE IDLE IDLE IDLE IDLE Cluster 3: Frequency: 1000M Core count: 8 DMA count: 2 Status: WORK Memory Used/Total: 156.00M/4.00G Npc status: IDLE IDLE IDLE IDLE IDLE IDLE IDLE IDLE
在代码中调用
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; }
访问设备端全局内存
访问设备端全局内存的运行时接口如下：
在主机端调用
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; }
并行执行
主机端发送给设备端的请求类型包括执行核函数（Kernel）、从主机端向设备端拷贝数据（H2D）、从设备端向主机端拷贝数据（D2H）。主机端可以并行发起不同类型的请求，但一个NPC Cluster内是顺序处理请求的。希姆计算提供了流和队列的运行时接口，您可以基于流或队列实现一个NPC Cluster内并行处理请求。
流
流定义
流（stream）是由主机端发起、设备端处理的一系列请求。同一个流内的请求顺序处理，不同流间的不同类型请求可以并行处理。 流包括以下类型：
隐式声明流：默认创建的流，只有一个，可以包括同步、异步请求。执行核函数、拷贝数据时默认使用隐式声明流。
显示声明流：您自行创建的流，只能包括异步请求。执行核函数、拷贝数据（异步）时可以使用显示声明流。
说明：在一个NPC Cluster上创建的流只能在该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;
}
流调度
同一流内的请求只能顺序调度，不同流间的不同类型请求可以并行调度。假定有三组请求，均包括H2D、Kernel、D2H请求，且不同组的请求之间没有依赖关系，采取并行调度可以提高性能。顺序调度和并行调度的性能对比如下：
顺序调度和并行调度的代码示例如下：
顺序调度，顺序在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; }
流同步
事件（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;
}
队列
队列（fifo）是一种特殊的流形式。通过队列可以在不退出核函数的情况下，不断读入和写出数据，提高核函数的执行效率。输入数据和输出数据均对应一个队列，队列可以与核函数并行执行。
下方的用例中，核函数一次从队列中读入并写出一个整数，循环执行共8次。在以下情况会出现同步阻塞：
主机端调用
stcFifoPush与设备端调用fifo_push，但队列已满时。主机端调用
stcFifoPop与设备端调用fifo_pop，但队列为空时。
核函数无法预测待处理数据的个数，因此通常使用无限循环体实现。在写出全部数据后，可以调用stcStreamClean退出核函数结束循环。
说明：请确保调用
stcStreamClean前对应的流上只有核函数，否则会将流上的其他请求一并销毁。
#include <hpe.h>
#include <npurt.h>
#include <pthread.h>
#include <stdio.h>
#define DATA_NUM 8
int host_in[DATA_NUM] = {1, 2, 3, 4, 5, 6, 7, 8};
int host_out[DATA_NUM];
__global__ void kernel(void *fifo_in, void *fifo_out) {
while (1) {
int data;
if (CoreID == 0) {
fifo_pop(&data, fifo_in);
fifo_push(&data, fifo_out);
}
}
}
void *thread_in(void *data) {
stcFifo_t fifo_in = *(stcFifo_t *)data;
for (int i = 0; i < DATA_NUM; i++)
stcFifoPush(&host_in[i], fifo_in);
return NULL;
}
void *thread_out(void *data) {
stcFifo_t fifo_out = *(stcFifo_t *)data;
for (int i = 0; i < DATA_NUM; i++)
stcFifoPop(&host_out[i], fifo_out);
// stop kernel function
stcStreamClean(NULL);
return NULL;
}
int main(void) {
stcFifo_t fifo_in, fifo_out;
stcFifoCreate(&fifo_in, sizeof(int), 2);
stcFifoCreate(&fifo_out, sizeof(int), 2);
// Kernel
kernel<<<8>>>(fifo_in, fifo_out);
// H2D
pthread_t tid_in;
pthread_create(&tid_in, NULL, thread_in, &fifo_in);
// D2H
pthread_t tid_out;
pthread_create(&tid_out, NULL, thread_out, &fifo_out);
stcDeviceSynchronize();
pthread_join(tid_in, NULL);
pthread_join(tid_out, NULL);
// print result
for (int i = 0; i < DATA_NUM; i++)
printf("%d, ", host_out[i]);
printf("\n");
stcFifoDestroy(fifo_in);
stcFifoDestroy(fifo_out);
return 0;
}
进程独占NPU
调用NPC Cluster函数接口时，允许多个进程的H2D、Kernel、D2H操作在同一个NPC Cluster上执行，由于同一类型的操作只能顺序执行，无法预测指定进程的指定操作在何时开始执行。如有需求，您可以使用NPU函数接口独占NPU，避免多个进程同时使用同一个NPC Cluster。
NPU函数接口对NPC Cluster函数接口进行了封装，并限制同时只能有一个进程访问一个NPU内所有NPC Cluster。NPU函数接口具有以下特点：
一个进程独占一个或多个NPU。
调用NPU函数接口后会作用到NPU内的所有NPC Cluster上。
核函数在一个NPU内的所有NPC上并行执行。
注意：避免同时使用NPU函数接口和NPC Cluster函数接口，否则会导致运行结果无法预测。希姆计算异构编程环境中支持的NPU函数接口，请参见NPU管理章节。
对使用NPU函数接口的异构程序，必须单独编译设备端程序，然后在主机端程序中读入编译得到的设备端目标程序。编译示例如下：
编译设备端程序得到设备端目标程序。
说明：单独编译设备端程序时，编译器无法识别核函数修饰符
__global__，定义核函数时需要改用__attribute__((force_stack_param))属性。$ cat device.cc #include <npurt.h> __attribute__((force_stack_param)) void kernel(void) { printf("hello world from core %d/%d.\n", CoreID, CoreNum); } $ stc-clang++ --target=riscv32npu -c device.cc $ stc-ld.lld -flavor gnu device.o /usr/local/hpe/riscv32npu/lib/libnpurt_hld.a -r -o fatbin.o
在主机端程序中读入设备端目标程序，然后编译主机端程序并生成可执行文件。
运行异构程序时，主机端读入设备端目标程序，然后实现NPU的每个NPC分别输出一条hello word信息。
$ cat host.hc #include <hpe.h> #include <stdlib.h> #include <sys/stat.h> #include <unistd.h> #include <fcntl.h> __global__ void kernel(void); int main(void) { struct stat st; stcModule_t module; int fd = open("fatbin.o", O_RDONLY); fstat(fd, &st); void *mem = malloc(st.st_size); read(fd, mem, st.st_size); close(fd); stcModuleLoadData(&module, mem, st.st_size); kernel<<<8>>>(); stcDeviceSynchronize(); free(mem); return 0; } $ stcc host.hc -o hello_world
流类型用例
下方的用例中，使用一个NPU中的4个NPC Cluster分别读入并写出一个整数，其中以多流方式处理H2D、Kernel、D2H操作。
$ cat device.cc
#include <npurt.h>
extern "C" {
__attribute__((force_stack_param)) void kernel(int *in, int *out) {
if(CoreID == 0)
*out = *in;
}
}
$ cat host.hc
#include <hpe.h>
#include <stdlib.h>
#include <sys/stat.h>
#include <unistd.h>
#include <fcntl.h>
#include <stdio.h>
int main(void) {
struct stat st;
int fd = open("fatbin.o", O_RDONLY);
fstat(fd, &st);
void *mem = malloc(st.st_size);
read(fd, mem, st.st_size);
close(fd);
int input[] = {1, 2, 3, 4};
int output[] = {0, 0, 0 ,0};
int npu;
stcAcquireNpu(&npu);
stcNpuModule_t module;
stcLoadNpuModule(&module, npu, mem, st.st_size);
stcActiveNpuModule(module);
stcNpuTask_t task;
stcKernelData_t ins[] = {{(char *)input, sizeof(input)}};
stcKernelData_t outs[] = {{(char *)output, sizeof(output)}};
stcEnqueueNpu(&task, npu, "kernel", ins, 1, outs, 1);
stcWaitForNpuTask(task);
printf("output %d:%d:%d:%d\n", output[0], output[1], output[2], output[3]);
stcDestroyNpuTask(task);
free(mem);
return 0;
}
队列类型用例
下方的用例中，使用一个NPU中的4个NPC Cluster分别读入并写出一个整数。
$ cat device.cc
#include <npurt.h>
extern "C" {
__attribute__((force_stack_param)) void kernel(void *fifo_in, void *fifo_out) {
while (1) {
if (CoreID == 0) {
int data;
fifo_pop(&data, fifo_in);
fifo_push(&data, fifo_out);
}
}
}
}
$ cat host.hc
#include <fcntl.h>
#include <hpe.h>
#include <stdio.h>
#include <stdlib.h>
#include <sys/stat.h>
#include <unistd.h>
int main(void) {
struct stat st;
int fd = open("fatbin.o", O_RDONLY);
fstat(fd, &st);
void *mem = malloc(st.st_size);
read(fd, mem, st.st_size);
close(fd);
int input[] = {1, 2, 3, 4};
int output[] = {0, 0, 0, 0};
int npu;
stcAcquireNpu(&npu);
stcNpuModule_t module;
stcLoadNpuModule(&module, npu, mem, st.st_size);
stcActiveNpuModule(module);
stcNpuTask_t task;
stcKernelData_t ins[] = {{(char *)input, sizeof(input)}};
stcKernelData_t outs[] = {{(char *)output, sizeof(output)}};
stcEnqueueNpu(&task, npu, "kernel", ins, 1, outs, 1, stcNpuTaskFifo);
stcWaitForNpuTask(task);
printf("output %d:%d:%d:%d\n", output[0], output[1], output[2], output[3]);
stcDestroyNpuTask(task);
free(mem);
return 0;
}
C++语言扩展接口
SHC完整兼容C++17标准，并针对异构程序扩展了语法和函数库。基于SHC提供的运行时接口编写代码时，您可以使用C++扩展语言接口方便地控制执行代码的逻辑。
调用核函数
在SHC中，调用核函数的方式如下：
kernel_function<<<NCORE, stream, flags>>>(arg0, ...)
命令中配置部分和参数部分的含义如下所示：
| 命令内容 | 说明 |
|---|---|
| << |
指定设备端的配置，配置项含义如下： - NCORE：执行核函数所使用NPC的个数。 - stream：指定执行核函数时所在的流，默认为0代表使用隐式声明流。详细的流使用说明，请参见并行执行章节。 - flags：指定核函数的运行标志，默认为0（stcKernelFlagNone）代表无运行标志。详细的运行标志含义，请参见stcKernelFlag_t章节。 |
| (arg0, ...) | 指定核函数的参数，参数需要满足以下条件： - 参数列表中的变量类型和个数必须和核函数的定义相匹配。 - 每个参数的大小不能超过4字节。 |
函数修饰符
SHC支持函数修饰符__host__、__global__、__device__，用于区分不同用途的函数。函数修饰符以及函数用途说明如下所示：
| 函数用途 | 修饰符 | 说明 |
|---|---|---|
| 主机端函数 | __host__ |
具有以下特点： - 在主机端执行，对设备端程序不可见。 - 主机端函数中可以调用主机端函数、核函数、双边函数，可以使用STL、libc、libc++库，支持递归调用自身。 - 参数来自于用户输入。 与普通C++函数没有任何区别，可以作为函数模板、类方法或匿名函数。 说明：如果函数没有添加任何函数修饰符，默认是主机端函数，编译器会自动添加函数修饰符。 |
| 核函数 | __global__ |
具有以下特点： - 能且只能由主机端函数调用，并在设备端执行。 - 核函数中可以调用设备端函数、双边函数，可以使用libnpurt库。 - 从栈上取参数。 - 能且只能是普通函数。 - 返回类型必须是void。 - 具有以下限制： - 不支持递归调用自身。 - 不能包含 long、longlong或double类型的参数。- 不支持使用其他变量修饰符修饰核函数的参数。 |
| 设备端函数 | __device__ |
具有以下特点： - 在设备端执行，对主机端程序不可见。 - 设备端函数中可以调用设备端函数、双边函数，可以使用libnpurt库，支持递归调用自身。 - 从寄存器取参数。 - 与普通C++函数没有任何区别，可以作为函数模板、类方法或匿名函数。 说明：如果需要函数需要在设备端执行，则不可省略函数修饰符。 |
| 双边函数 | __host__和__device__ |
具有以下特点： - 可以在主机端、设备端执行。 - 能且只能调用双边函数，支持递归调用自身。 - 可以被主机端函数、核函数、双边函数调用，参数来自于调用者。 - 双边函数一般是一些主机端和设备端都会用到的小型辅助函数，例如求数组最大值。 |
说明：不建议使用
__device__和__host__各自修饰类方法，会导致在主机端和设备端看到的类定义不同。
变量修饰符
设备端不支持动态分配内存，但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类型（例如普通变量、数组、结构体）。
主机端编程
接口调用要求
调用主机端运行时接口时，需要包含对应的头文件：
#include <hpe.h>
主机端运行时接口提供以下功能：
设备管理：提供操作设备（NPC Cluster）相关的功能，例如指定待使用的设备、获取设备信息等。
内存管理：提供操作内存相关的功能，例如分配/释放内存、拷贝内存数据等。
执行控制：提供执行目标程序相关的功能，例如注册/释放目标程序、指定运行配置、启动核函数、加载/卸载目标程序等。
流管理：提供操作流相关的功能，例如创建/销毁流、创建/销毁事件、添加事件、创建/销毁队列、数据写入/写出队列等。
NPU管理：提供独占NPU相关的功能，例如进程占用/释放NPU、加载设备端目标程序到所有NPC Cluster、获取NPU占用情况等。
错误处理：提供获取错误信息相关的功能，例如获取错误码、获取错误详情等。
说明：主机端运行时接口涉及的数据类型和环境变量，请参见数据类型和环境变量章节。
运行时接口
设备管理
stcSetDevice
函数描述：调用stcSetDevice设置用于执行设备端程序的NPC Cluster。
函数类型：同步函数
函数定义：
stcError_t stcSetDevice (int device)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| device | 输入参数 | int | NPC Cluster ID，由NPU设备标识符、NPC Cluster设备标识符推导得出。详细的推导说明，请参见指定NPC Cluster ID章节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcGetDevice
函数描述：调用stcGetDevice获取用于执行设备端程序的NPC Cluster。
函数类型：同步函数
函数定义：
stcError_t stcGetDevice (int *device)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| device | 输出参数 | int* | 指向所获取NPC Cluster ID的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
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章节。 |
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章节。 |
stcGetDeviceCount
函数描述：调用stcGetDeviceCount获取所有NPU上NPC Cluster的数量。
函数类型：同步函数
函数定义：
stcError_t stcGetDeviceCount (int *count)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| count | 输出参数 | int* | 指向所获取NPC Cluster数量的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcDeviceSynchronize
函数描述：调用stcDeviceSynchronize等待当前进程的所有设备端操作执行结束。如果核函数执行异常退出，则输出触发异常时核函数的调用栈。
函数类型：同步函数
函数定义：
stcError_t stcDeviceSynchronize (void)
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
内存管理
stcMalloc
函数描述：调用stcMalloc在设备端全局内存的0GiB ~ 3GiB范围动态分配内存。
函数类型：同步函数
函数定义：
stcError_t stcMalloc (void **devPtr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| devPtr | 输出参数 | void** | 指向所分配内存地址的指针。 |
| size | 输入参数 | size_t | 所需分配的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcMallocHigh
函数描述：调用stcMallocHigh在设备端全局内存的3GiB ~ 4GiB范围（也称为高端内存）动态分配内存。
函数类型：同步函数
函数定义：
stcError_t stcMallocHigh(void **devPtr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| devPtr | 输出参数 | void** | 指向所分配内存地址的指针。 |
| size | 输入参数 | size_t | 所需分配的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcFree
函数描述：调用stcFree释放在设备端动态分配的全局内存。
函数类型：同步函数
函数定义：
stcError_t stcFree (void *devPtr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| devPtr | 输入参数 | void* | 指向待释放内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcMallocHost
函数描述：调用stcMallocHost在主机端分配内存，并设置为不会被换出的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcMallocHost (void **ptr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输出参数 | void** | 指向所分配内存地址的指针。 |
| size | 输入参数 | size_t | 所需分配的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcFreeHost
函数描述：调用stcFreeHost释放在主机端分配的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcFreeHost (void *ptr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输入参数 | void* | 指向待释放内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcHostRegister
函数描述：调用stcHostRegister将主机端的内存设置为不会被换出的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcHostRegister (void *ptr, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输入参数 | void* | 指向待锁定内存的指针。 |
| size | 输入参数 | size_t | 所需锁定的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcHostUnregister
函数描述：调用stcHostUnregister解锁主机端的页锁定内存。
函数类型：同步函数
函数定义：
stcError_t stcHostUnregister (void *ptr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| ptr | 输入参数 | void* | 指向待解锁内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
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章节。 |
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章节。 |
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章节。 |
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章节。 |
执行控制
stcRegisterFatBinary
函数描述：调用stcRegisterFatBinary注册在设备端执行的目标程序。
函数类型：同步函数
函数定义：
stcError_t stcRegisterFatBinary (const void *data)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| data | 输入参数 | const void* | 指向目标程序所占内存的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcUnregisterFatBinary
函数描述：调用stcUnregisterFatBinary释放所有已注册的目标程序。
函数类型：同步函数
函数定义：
stcError_t stcUnregisterFatBinary (void)
函数参数：
无
函数返回值：
stcSuccess
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章节。 |
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章节。 |
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章节。 |
stcModuleUnload
函数描述：调用stcModuleUnload从设备端卸载目标程序模块。
函数类型：同步函数
函数定义：
stcError_t stcModuleUnload (stcModule_t module)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输入参数 | stcModule_t | 待释放的目标程序模块。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
流管理
stcStreamCreate
函数描述：调用stcStreamCreate创建一个流。
函数类型：同步函数
函数定义：
stcError_t stcStreamCreate (stcStream_t *pStream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pStream | 输出参数 | stcStream_t* | 指向创建的流标识符的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcStreamDestroy
函数描述：调用stcStreamDestroy销毁指定流。如果流中有未完成的操作，会终止执行并释放相关资源。
函数类型：同步函数
函数定义：
stcError_t stcStreamDestroy (stcStream_t stream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 待销毁流的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcStreamSynchronize
函数描述：调用stcStreamSynchronize等待指定流上的所有操作执行结束。如果核函数执行异常退出，则输出触发异常时核函数的调用栈。
函数类型：同步函数
函数定义：
stcError_t stcStreamSynchronize (stcStream_t stream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 流标识符，等待该流上的所有操作执行结束。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcStreamClean
函数描述：调用stcStreamClean停止处理流上的请求并销毁所有请求，但不会销毁流本身。
函数类型：同步函数
函数定义：
stcError_t stcStreamClean(stcStream_t stream)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 流标识符，停止处理该流的请求，并销毁所有请求。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcEventCreate
函数描述：调用stcEventCreate创建一个事件。
函数类型：同步函数
函数定义：
stcError_t stcEventCreate(stcEvent_t *pEvent)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pEvent | 输出参数 | stcEvent_t* | 指向创建的事件标识符的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcEventDestroy
函数描述：调用stcEventDestroy销毁指定事件。
函数类型：同步函数
函数定义：
stcError_t stcError_t stcEventDestroy(stcEvent_t event)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| event | 输入参数 | stcEvent_t | 待销毁事件的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcEventSynchronize
函数描述：调用stcEventSynchronize等待指定事件进入完成状态。
函数类型：同步函数
函数定义：
stcError_t stcEventSynchronize(stcEvent_t event)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| event | 输入参数 | stcEvent_t | 事件标识符，等待流中该事件前的所有操作执行结束后，才会将事件置为完成状态。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcEventRecord
函数描述：调用stcEventRecord在指定流的当前运行点添加事件。
函数类型：同步函数
函数定义：
stcEventRecord(stcEvent_t event, stcStream_t stream=0)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| event | 输入参数 | stcEvent_t | 待添加事件的标识符。 |
| stream | 输入参数 | stcStream_t | 待添加事件的流的标识符，默认为0，代表在所有流的当前运行点添加事件。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
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章节。 |
stcStreamWaitEvent
函数描述：调用stcStreamWaitEvent在多流场景中建立流之间的同步关系，指定某个流需要等指定事件被置为完成状态后再开始处理请求。
函数类型：同步函数
函数定义：
stcError_t stcStreamWaitEvent(stcStream_t stream, stcEvent_t event)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| stream | 输入参数 | stcStream_t | 流标识符，该流在对应事件被置为完成状态后再开始处理请求。 |
| event | 输入参数 | stcEvent_t | 事件标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcFifoCreate
函数描述：调用stcFifoCreate创建一个队列。
函数类型：同步函数
函数定义：
stcError_t stcFifoCreate(stcFifo_t *pFifo, size_t size, int fifoNr)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| pFifo | 输出参数 | stcFifo_t* | 指向创建的队列标识符的指针。 |
| size | 输入参数 | size_t | 队列中每个数据单元的大小。 |
| fifoNr | 输入参数 | int | 队列中数据单元的最大个数。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcFifoDestroy
函数描述：调用stcFifoDestroy销毁指定队列。
函数类型：同步函数
函数定义：
stcError_t stcFifoDestroy(stcFifo_t fifo)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| fifo | 输入参数 | stcFifo_t | 待销毁队列的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcFifoPush
函数描述：调用stcFifoPush向指定队列存入一个数据单元。
说明：不允许stcFifoPush和stcFifoPop同时操作同一个队列。
函数类型：同步函数
函数定义：
stcError_t stcFifoPush(void *src, stcFifo_t fifo)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| src | 输入参数 | void* | 指向待存入数据的指针。 |
| fifo | 输入参数 | stcFifo_t | 队列的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcFifoPop
函数描述：调用stcFifoPop从队列中取出一个数据单元。
说明：不允许stcFifoPush和stcFifoPop同时操作同一个队列。
函数类型：同步函数
函数定义：
stcError_t stcFifoPop(void *dst, stcFifo_t fifo)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | void* | 指向待取出数据的指针。 |
| fifo | 输入参数 | stcFifo_t | 队列的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
NPU管理
stcAcquireNpu
函数描述：调用stcAcquireNpu占用一个空闲的NPU设备。
说明：一个NPU设备不能同时被多个进程占用。
函数类型：同步函数
函数定义：
stcError_t stcAcquireNpu(int *npu)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| npu | 输出参数 | int* | 指向NPU设备标识符的指针。如果主机中有N个NPU，异构编程环境中使用0至N-1依次标记。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcGetNumOfAcquiredNpus
函数描述：调用stcGetNumOfAcquiredNpus统计已占用的NPU设备个数。
函数类型：同步函数
函数定义：
stcError_t stcGetNumOfAcquiredNpus(int *count)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| count | 输出参数 | int* | 指向统计的NPU设备个数的指针。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcLoadNpuModule
函数描述：调用stcLoadNpuModule将设备端目标程序加载到一个NPU内的所有NPC Cluster。
函数类型：同步函数
函数定义：
stcError_t stcLoadNpuModule(stcNpuModule_t *module, int npu, const void *data, size_t size)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输出参数 | stcNpuModule_t* | 指向加载到NPU设备的目标程序模块的指针。 |
| npu | 输入参数 | int | NPU设备标识符。 |
| data | 输入参数 | const void* | 目标程序的地址。 |
| size | 输入参数 | size_t | 目标程序占用的内存大小，单位为字节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcUnloadNpuModule
函数描述：调用stcUnloadNpuModule卸载一个NPU目标程序模块。如果执行中的NPU任务正在使用该模块，则卸载失败。
函数类型：同步函数
函数定义：
stcError_t stcUnloadNpuModule(stcNpuModule_t module)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输入参数 | stcNpuModule_t | 待卸载的NPU目标程序模块。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 如果执行中的任务正在使用NPU目标程序模块，则卸载失败，返回stcErrorNpuTaskRunning。更多错误类型含义，请参见stcError_t章节。 |
stcNumLoadedNpuModule
函数描述：调用stcNumLoadedNpuModule统计一个NPU上加载的目标程序模块的个数。
函数类型：同步函数
函数定义：
stcError_t stcNumLoadedNpuModule(int *count, int npu)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| count | 输出参数 | int* | 指向统计的NPU目标程序模块个数的指针。 |
| npu | 输入参数 | int | NPU设备标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcActiveNpuModule
函数描述：调用stcActiveNpuModule激活一个已加载的NPU目标程序模块。一个NPU中只能存在一个处于激活状态的NPU目标程序模块，激活新的NPU目标程序模块后，原有处于激活状态的NPU目标程序模块会被置为未激活状态。
函数类型：同步函数
函数定义：
stcError_t stcActiveNpuModule(stcNpuModule_t module)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输入参数 | stcNpuModule_t | 待激活的NPU目标程序模块。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcGetActiveNpuModule
函数描述：调用stcGetActiveNpuModule获取处于激活状态的NPU目标程序模块。
函数类型：同步函数
函数定义：
stcError_t stcGetActiveNpuModule(stcNpuModule_t *module, int npu)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| module | 输出参数 | stcNpuModule_t* | 指向已激活的NPU目标程序模块的指针。 |
| npu | 输入参数 | int | NPU设备标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcEnqueueNpu
函数描述：调用stcEnqueueNpu平分数据，并拷贝到NPU设备中的各个NPC Cluster并行处理。具体来说，将主机端输入和输出的每一个数据区按照一个NPU内NPC Cluster的总数平分，分别拷贝到每个NPC Cluster的全局内存，然后依次做为核函数的参数在每个NPC Cluster上并行处理。
函数类型：异步函数
函数定义：
stcError_t stcEnqueueNpu(stcNpuTask_t *task, int npu, const char *kname, stcKernelData_t *input, int inNum, stcKernelData_t *output, int outNum, stcNpuTaskKind_t kind, unsigned int kflags)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| task | 输入/输出参数 | stcNpuTask_t* | 指向NPU任务的指针，可以作为输入参数或输出参数。 - 如果NPU任务已经存在，并且函数的各个参数与NPU任务匹配，则继续使用NPU任务，作为输入参数。 - 如果新建一个NPU任务，则作为输出参数。 |
| npu | 输入参数 | int | NPU设备标识符。 |
| kname | 输入参数 | const char* | 核函数名。 |
| input | 输入参数 | stcKernelData_t* | 核函数输入参数数组。详细的数据结构定义，请参见stcKernelData_t章节。 |
| inNum | 输入参数 | int | 输入参数数组的大小。 |
| output | 输入参数 | stcKernelData_t* | 核函数输出参数数组。详细的数据结构定义，请参见stcKernelData_t章节。 |
| outNum | 输入参数 | int | 输出参数数组的大小。 |
| kind | 输入参数 | stcNpuTaskKind_t | NPU任务类型，默认为流。详细的任务类型含义，请参见stcNpuTaskKind_t章节。 |
| kflags | 输入参数 | unsigned int | 指定核函数的运行标志，默认无运行标志（stcKernelFlagNone）。详细的运行标志含义，请参见stcKernelFlag_t章节。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcWaitForNpuTask
函数描述：调用stcWaitForNpuTask等待一个NPU任务执行完成。
函数类型：同步函数
函数定义：
stcError_t stcWaitForNpuTask(stcNpuTask_t task)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| task | 输入参数 | stcNpuTask_t | NPU任务的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcDestroyNpuTask
函数描述：调用stcDestroyNpuTask强制销毁一个NPU任务。如果NPU任务正在运行，在销毁NPU任务前会停止执行并释放相关资源。
函数类型：同步函数
函数定义：
stcError_t stcDestroyNpuTask(stcNpuTask_t task)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| task | 输入参数 | stcNpuTask_t | NPU任务的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcReleaseNpu
函数描述：调用stcReleaseNpu释放一个NPU设备。在释放NPU设备前，会在该NPU上销毁所有NPU任务并卸载所有NPU目标程序模块。
函数类型：同步函数
函数定义：
stcError_t stcReleaseNpu(npu)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| npu | 输入参数 | int | NPU设备标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcDestroyNpus
函数描述：调用stcDestroyNpus释放当前进程占用的所有NPU设备。在释放NPU设备前，会在所有NPU上销毁所有NPU任务并卸载所有NPU目标程序模块。
函数类型：同步函数
函数定义：
stcError_t stcDestroyNpus()
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
错误处理
stcGetLastError
函数描述：调用stcGetLastError返回一个线程调用主机端运行时接口时产生的最后一个错误，然后将结果重置为stcSuccess。如果没有错误，则返回stcSuccess。
函数类型：同步函数
函数定义：
stcError_t stcGetLastError(void)
函数参数：
无
函数返回值：
| 类型 | 说明 |
|---|---|
| stcError_t | 详细的错误类型含义，请参见stcError_t章节。 |
stcGetErrorName
函数描述：调用stcGetErrorName从获取的错误码得到错误名称。
函数类型：同步函数
函数定义：
const char* stcGetErrorName (stcError_t error)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| error | 输入参数 | stcError_t | 调用stcGetLastError获取到的错误码。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| const char* | 指向错误名称字符串的指针，字符串以NULL结尾。 |
stcGetErrorString
函数描述：调用stcGetErrorString从获取的错误码得到错误详情。
函数类型：同步函数
函数定义：
const char* stcGetErrorString (stcError_t error)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| error | 输入参数 | stcError_t | 调用stcGetLastError获取到的错误码。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| const char* | 指向NULL结尾的字符串。 |
数据类型
stcKernelParams_t
数据描述：记录一个核函数的所有参数信息。
stcModule_t
数据描述：记录一个NPC Cluster设备目标程序模块的相关信息。
stcStream_t
数据描述：记录一个流的相关信息。
stcEvent_t
数据描述：记录一个事件的相关信息。
stcFifo_t
数据描述：记录一个队列的相关信息。
stcNpuModule_t
数据描述：记录一个NPU设备目标程序模块的相关信息。
stcNpuTask_t
数据描述：记录一个NPU任务的相关信息，包括核函数、参数等。
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 | 事件不可用。 |
| stcErrorInvalidFifo | 14 | 队列不可用。 |
| stcErrorDeviceImageException | 15 | 设备端目标程序运行时出现异常。 |
| stcErrorSyscallFailure | 16 | 主机端的系统调用失败。 |
| stcErrorForkForbidden | 17 | 父进程已调用过运行时接口，禁止子进程再次调用。 |
| stcErrorInvalidCoreNum | 18 | 使用了无效的核数量。 |
| stcErrorGDBFailure | 19 | stc-gdb跟踪失败，无法获取调试信息。 |
| stcErrorDriverMismatch | 20 | NPU驱动版本和NPU设备不匹配。 |
| stcErrorDeviceBreakdown | 21 | NPU设备出现故障，无法继续使用。 |
| stcErrorNoFreeNpu | 22 | 没有可用的NPU设备。 |
| stcErrorNpuNotAcquired | 23 | NPU设备未被占用，但无法访问该NPU设备。 |
| stcErrorInvalidNpuTask | 24 | NPU任务不可用。 |
| stcErrorInvalidImageDataSection | 25 | 设备端目标程序包含了太多可写数据段，导致程序无法正常运行。 |
| stcErrorNpuTaskRunning | 26 | NPU任务正在使用NPU模块，不可以卸载该NPU模块。 |
| stcErrorDriverFailure | 99 | NPU驱动错误。 |
stcMemcpyKind_t
数据描述：设置拷贝内存操作的类型。支持设置的操作类型如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| stcMemcpyHostToHost | 0 | 在同一台主机的不同内存段间拷贝数据。 |
| stcMemcpyHostToDevice | 1 | 从主机端向设备端拷贝数据。 |
| stcMemcpyDeviceToHost | 2 | 从设备端向主机端拷贝数据。 |
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外枚举定义的数量。 |
stcKernelFlag_t
数据描述：设置核函数的运行标志。支持设置的运行标志如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| stcKernelFlagNone | 0 | NPC执行核函数前后均处理DCache，可以视为无运行标志。 |
| stcKernelFlagInputDataBypassDcache | 1 | NPC执行核函数前，不需要处理DCache。 |
| stcKernelFlagOutputDataBypassDache | 2 | NPC执行核函数后，不需要处理DCache。 |
stcKernelData_t
数据描述：核函数输入输出对应的主机端数据区的属性。包含的成员变量如下所示：
| 成员变量 | 说明 |
|---|---|
| data | 数据区的地址。 |
| size | 数据区的大小。 |
stcNpuTaskKind_t
数据描述：执行NPU任务时处理核函数输入输出的方式。包含的成员变量如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| stcNpuTaskStream | 0 | 多流方式处理。 |
| stcNpuTaskFifo | 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。 |
设备端编程
接口调用要求
调用设备端运行时接口时，需要包含对应的头文件：
#include <npurt.h>
说明：设备端运行时接口涉及的数据类型和变量，请参见数据类型和变量章节。
运行时接口
exit
函数描述：调用exit退出核函数并返回您指定的退出码。
函数类型：同步函数
函数定义：
__device__ void exit(int exit_code)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| exit_code | 输入参数 | int | 您自行定义的退出码。 |
函数返回值：
无
abort
函数描述：调用abort退出核函数并返回退出码128。
函数类型：同步函数
函数定义：
__device__ void abort()
函数参数：
无
函数返回值：
无
assert
函数描述：调用assert验证指定的条件，在未满足断言时，退出核函数并返回退出码128。
函数类型：同步函数
函数定义：
__device__ void assert(int exp)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| exp | 输入参数 | int | 指定的条件。如果未满足断言，则值为0，退出核函数并返回错误码128。 |
函数返回值：
无
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章节。 |
fifo_push
函数描述：调用fifo_push向指定队列存入一个数据单元。
函数类型：同步函数
函数定义：
__device__ void fifo_push(void* src, void* fifo_id, bool lock_prot=true, bool clear_dcache=true)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| src | 输入参数 | void* | 指向待存入数据的指针。 |
| fifo_id | 输入参数 | void* | 队列的标识符。 |
| lock_prot | 输入参数 | bool | 队列是否需要锁保护，默认为需要。用于保护多个NPC同时访问队列，保证所有NPC中看到的是同一份数据。取值含义如下： - true：需要锁保护。 - false：不需要锁保护。 |
| clear_dcache | 输入参数 | bool | 队列是否需要消除DCache影响，默认为需要。用于保护多个NPC同时访问队列，保证所有NPC中看到的是同一份数据。取值含义如下： - true：需要消除DCache影响。 - false：不需要消除DCache影响。 |
函数返回值：
无
fifo_pop
函数描述：调用fifo_pop从队列中取出一个数据单元。
函数类型：同步函数
函数定义：
__device__ void fifo_pop(void *dst, void *fifo_id, bool lock_prot=true, bool clear_dcache=true)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| dst | 输入参数 | void* | 指向待取出数据的指针。 |
| fifo_id | 输入参数 | void* | 队列的标识符。 |
| lock_prot | 输入参数 | bool | 队列是否需要锁保护，默认为需要。用于保护多个NPC同时访问队列，保证所有NPC中看到的是同一份数据。取值含义如下： - true：需要锁保护。 - false：不需要锁保护。 |
| clear_dcache | 输入参数 | bool | 队列是否需要消除DCache影响，默认为需要。用于保护多个NPC同时访问队列，保证所有NPC中看到的是同一份数据。取值含义如下： - true：需要消除DCache影响。 - false：不需要消除DCache影响。 |
函数返回值：
无
fifo_size
函数描述：调用fifo_size查看指定队列中已有数据单元的个数。
函数类型：同步函数
函数定义：
__device__ unsigned int fifo_size(void* fifo_id)
函数参数：
| 名称 | 输入/输出 | 类型 | 说明 |
|---|---|---|---|
| fifo_id | 输入参数 | void* | 队列的标识符。 |
函数返回值：
| 类型 | 说明 |
|---|---|
| unsigned int | 已有数据单元的个数。 |
数据类型
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字节对齐。 |
cpy_config_t
数据描述：待拷贝内存的信息格式。包含的成员变量如下：
| 成员变量 | 说明 |
|---|---|
| rows | 待拷贝数据的行数。 |
| cols | 待拷贝数据的列数。 |
| stride_cols | 拷贝数据时使用的列数Stride。 |
| dtype | 待拷贝数据的类型。详细的数据类型定义，请参见cpy_dtype_t章节。 |
| dest_core | 拷贝数据时NPC Cluster内目标NPC的标识符，该参数只适用于在不同NPC的本地内存间传输数据。 说明：目前 memcpy仅支持在同一NPC Cluster内拷贝。 |
cpy_dtype_t
数据描述：待拷贝数据的类型。支持的数据类型如下所示：
| 枚举定义 | 枚举值 | 说明 |
|---|---|---|
| DTYPE_BYTE_1 | 0 | 每个数据元素占一个字节大小。 |
| DTYPE_BYTE_2 | 1 | 每个数据元素占两个字节大小。 |
变量
CoreID
描述：获取NPC在NPC Cluster内的标识符，例如STCP920中CoreID的范围为0 ~ 7。
类型：整型只读
CoreNum
描述：获取运行核函数时所使用NPC的个数。
类型：整型只读