希姆计算异构编程手册

版本历史

版本 作者 日期 说明
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的内存布局完全相同,如下所示:

../_images/mem-layout.png

访问内存

在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读写设备端本地内存、共享内存、全局内存。

  • 在主机端调用stcMemcpystcMemcpyAsync读写设备端全局内存。

说明:在设备端访问全局内存时(访问局部变量除外),请确保访问区域不会被其他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在主机端和设备端之间拷贝数据后,主机端无需等待拷贝完所有数据即可开始处理下一个任务。

异步函数的返回值不是异步操作的结果。如果您调用了异步函数,但仍然需要等待所有操作完成,可以调用stcDeviceSynchronizestcStreamSynchronize等待设备或流上的所有操作完成后再处理下一个任务。stcDeviceSynchronizestcStreamSynchronize的返回值是异步操作的结果,但仅返回最后一个异步操作的结果。

说明:核函数没有返回值,如果核函数在真正开始执行前出错退出,只能通过stcGetLastError获取错误信息。例如指定的NPC数量超过实际的NPC数量,导致启动核函数后还未实际执行就出错退出。

典型操作

希姆计算提供了丰富的运行时接口,方便您从主机端控制在设备端执行任务。

指定运行设备

主机端访问设备端时需要指定运行设备。在主机端程序中启动核函数后,默认在NPC Cluster 0上执行核函数,您也可以提前通过调用stcSetDevice或修改STC_SET_DEVICES指定其他NPC Cluster。示例如下:

  1. 推导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
    
  2. 在代码中调用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)。

  • 在主机端调用stcMemcpystcMemcpyAsync读写设备端全局内存。

在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上使用。

下方的用例中,创建了一个显示声明流来执行核函数。如果需要等待流上的所有操作完成后再处理下一个请求,调用stcStreamSynchronizestcDeviceSynchronize即可。

#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请求,且不同组的请求之间没有依赖关系,采取并行调度可以提高性能。顺序调度和并行调度的性能对比如下:

../_images/stream-scheduling.png

顺序调度和并行调度的代码示例如下:

  • 顺序调度,顺序在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函数接口的异构程序,必须单独编译设备端程序,然后在主机端程序中读入编译得到的设备端目标程序。编译示例如下:

  1. 编译设备端程序得到设备端目标程序。

    说明:单独编译设备端程序时,编译器无法识别核函数修饰符__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
    
  2. 在主机端程序中读入设备端目标程序,然后编译主机端程序并生成可执行文件。

    运行异构程序时,主机端读入设备端目标程序,然后实现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。
- 具有以下限制:
- 不支持递归调用自身。
- 不能包含longlonglongdouble类型的参数。
- 不支持使用其他变量修饰符修饰核函数的参数。
设备端函数 __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的个数。

类型:整型只读