希姆计算异构编程手册
版本历史
版本 | 作者 | 日期 | 说明 |
---|---|---|---|
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的个数。
类型:整型只读