MLU220快速上手指南 – 寒武纪开发者社区
安装硬件驱动和软件工具链，也可以直接使用寒武纪官方开发平台：寒武纪开发平台
本地开发安装完工具需要进行一些配置：
根据包管理器的反馈，CNToolkit安装成功后，还需要进行最后检查和环境配置才能正确工作。
RPM包管理器查询包是否安装：
$ cat /usr/local/neuware/version.txt
$ rpm -qa | grep cncc
$ rpm -qa | grep cnas
$ rpm -qa | grep cncodec
$ rpm -qa | grep cndev
$ rpm -qa | grep cndrv
$ rpm -qa | grep cngdb
$ rpm -qa | grep cnpapi
$ rpm -qa | grep cnperf
$ rpm -qa | grep cnrt
$ rpm -qa | grep cnrtc
$ rpm -qa | grep cnstudio
Deb包管理器查询包是否安装：
$ cat /usr/local/neuware/version.txt
$ dpkg -l | grep cncc
$ dpkg -l | grep cnas
$ dpkg -l | grep cncodec
$ dpkg -l | grep cndev
$ dpkg -l | grep cndrv
$ dpkg -l | grep cngdb
$ dpkg -l | grep cnpapi
$ dpkg -l | grep cnperf
$ dpkg -l | grep cnrt
$ dpkg -l | grep cnrtc
$ dpkg -l | grep cnstudio
Neuware和CNToolkit统一使用环境变量 NEUWARE_HOME 来存储根目录路径。默认如下，也可以自定义为其他已安装目录。
$ export NEUWARE_HOME="/usr/local/neuware"
RPM或Deb包管理器在安装后处理阶段会调用 /sbin/ldconfig 刷新链接缓存，但仍需要为运行时库配置 LD_LIBRARY_PATH 和二进制工具的 PATH 。
$ export PATH="${NEUWARE_HOME}/bin:${PATH}"
$ which cncc
$ which cnas
$ which cngdb
$ which cnperf-cli
$ cncc --version
$ cnas --version
$ cngdb --version
$ cnperf-cli --version
$ export LD_LIBRARY_PATH="${NEUWARE_HOME}/lib64:${LD_LIBRARY_PATH}"
异构编程实现了使用具有不同类型指令集和架构的计算单元。BANG C异构编程模型基于CPU与MLU协同计算，突破CPU开发瓶颈，利用MLU的机器学习能力，有效解决能耗和扩展性问题。
Cambricon BANG C语言兼顾终端和云目标平台。本章结合MLU硬件架构的概念，介绍BANG C语言的并行模型和存储模型及相关概念。
异构计算系统通常由通用处理器和许多特定领域的处理器组成： 通用处理器作为控制设备（简称Host），进行复杂的控制和调度；特定领域的处理器作为子设备（简称 MLU），用于大规模并行计算和特定领域的计算任务。Host 和 MLU 合作完成计算任务。对于异构计算系统，原有的同构并行编程模型不再适用。因此，异构并行编程模型越来越受到学术界和工业界的关注。本章简要介绍了寒武纪 MLU 异构编程。
异构编程包括 Host 和 MLU。对于Host，主要包括设备获取、数据/参数准备、执行流程创建、任务描述、内核启动、输出获取等。Entry函数是MLU上的程序入口，可以调用MLU函数。MLU程序使用异构编程模型的C/C++语言扩展。二进制文件由 MLU 的指定编译器编译。
MLU异构程序的编译和链接过程
MLU异构程序的编译和链接过程如图MLU异构程序的编译和链接过程。采用分离的编程方式，即Host程序和MLU程序位于不同的文件中（即Host和Kernel文件）。Host程序和MLU程序的异构并行程序需要自己的编译器编译。
具体来说，Host程序是一个普通的C/C++程序，用户可以使用GCC、CLANG等任何C/C++编译器。MLU程序是基于C/C++语言的扩展，可以通过以下方式编译寒武纪提供的指定编译器 CNCC。Host Linker 通过将 Host 和 MLU 中的两个目标文件、运行时库和其他文件链接起来，形成一个可执行程序。
下面的例子将l2loss.mlu、l2loss_main.cpp、l2loss_ops.cpp这三个文件编译成可执行程序l2loss。
cncc -c l2loss.mlu --bang-device-only -o l2loss.o
g++ -c l2loss_main.cpp
g++ -c l2loss_ops.cpp -I/usr/local/neuware/include
g++ l2loss.o l2loss_main.o l2loss_ops.o -o l2loss -L/usr/local/neuware/lib64 -lcnrt
使用 cncc -c xx.mlu 时，会生成 .cncode 中间文件。使用cncc xx.mlu时，不带-c选项，会生成.cnbin中间文件。多个 .cnbin 文件将链接到 .cnfatbin 文件中。
Host程序是一个常见的C/C++程序，通过调用CNRT-API来初始化设备，管理设备内存，准备Kernel参数，启动Kernel，释放资源。下面将介绍Host程序调用Kernel程序的主要过程。
头文件
Host程序需要包含运行时头文件cnrt.h，该文件提供了异构编程所需的运行时接口的声明以及Host程序使用的相关数据类型的定义。有关详细信息，请参阅寒武纪 CNRT 开发人员指南。
初始化设备
在启动 Kernel 之前，用户需要调用 CNRT 接口来初始化设备，如下例所示：
cnrtInit(0);
获取设备
初始化设备后，用户可以通过以下方式获取设备：
cnrtDev_t dev;
cnrtGetDeviceHandle(&dev, 0);
cnrtSetCurrentDevice(dev);
准备 MLU 的输入数据
Host端需要准备好MLU程序的输入数据，并将输入数据复制到MLU的指定位置。由于 MLU 支持一些特殊的数据类型，例如 C/C++ 语言目前不支持半类型（半精度浮点，即两个字节的浮点），用户需要在 Host 上转换 MLU 的输入数据. 进行转换，例如将float/double数据转换为half数据并存储在两个字节中，例如float/double数据应该转换为half数据并存储在两个字节中。目前runtime提供了2个接口函数，cnrtConvertDoubleToHalf以及cnrtConvertFloatToHalf，方便用户将double/float数据转换为half数据。有关详细信息，请参阅寒武纪 CNRT 开发人员指南。需要注意的是，这两个数据转换函数都将半数据存储在 uint16_t 数据类型中。具体使用方法请参考以下示例：
typedef uint16_t half;
half* input_half = (half*)(malloc(dims_a * sizeof(half)));
for (int i = 0; i< len; i++) {
cnrtConvertFloatToHalf(input_half+i, input[i]);
}
传输 MLU 的输入数据
由于 Kernel 参数不支持数组和簇类型，因此需要用户将数组类型的输入数据显式复制到 MLU 空间中。用户首先需要调用cnrtMalloc申请一个DDR空间，然后调用cnrtMemcpy将Host中准备好的数据复制到DDR。下面是一个具体的例子。
half* mlu_input;
cnrtMalloc((void**)(&mlu_input), dims_a * sizeof(half));
cnrtMemcpy(mlu_input, input_half, dims_a * sizeof(half),
CNRT_MEM_TRANS_DIR_HOST2DEV);
准备内核参数
Kernel 参数仅支持标量数据。传输Kernel参数时，应该在Host上获取一个cnrtKernelParamsBuffer_t。然后按照Kernel参数定义的顺序，依次调用运行时函数cnrtKernelParamsBufferAddParam将标量参数推入缓冲区。下面是一个具体的例子。
cnrtKernelParamsBuffer_t params;
cnrtGetKernelParamsBuffer(¶ms);
cnrtKernelParamsBufferAddParam(params, &mlu_input, sizeof(half*));
创建队列
与执行 BANG C 程序 Kernel 相关的一个重要概念是队列。同一个队列的任务是串行执行的，不同队列之间的任务是并行执行的。当我们启动 Kernel 时，我们需要指定 Kernel 执行的队列。创建队列的示例如下：
cnrtQueue_t pQueue;
cnrtCreateQueue(&pQueue);
指定内核的任务大小
CNRT 提供的数据类型 cnrtDim3_t 用于指定 Kernel 任务的大小。任务大小有 XYZ 三个维度。任务大小是有一定限制的。详情请参阅任务部。以下示例显示了如何指定任务的执行大小（任务的并行度为 1）。
cnrtDim3_t dim;
dim.x = 1;
dim.y = 1;
dim.z = 1;
指定内核的任务类型
cnrtFunctionType_t用于指定Kernel启动时的计算并行度，其值可以是BLOCK或UNIONn（n = 1,2,4,8）类型。下面是一个简单的例子。
cnrtFunctionType_t ft = CNRT_FUNC_TYPE_BLOCK;
启动内核
Kernel启动时，需要将Kernel的函数名、任务大小、Kernel参数、任务类型、队列作为参数传递给cnrt运行时函数cnrtInvokeKernel_V2。下面是一个具体的例子。
ret = cnrtInvokeKernel_V2((void *)(&L2LossKernel), dim,
params, ft, pQueue);
从 MLU 到主机的数据传输
MLU 的计算结果需要用户显式复制到 Host。用户调用运行时函数cnrtMemcpy，将方向参数设置为CNRT_MEM_TRANS_DIR_DEV2HOST，即可完成从MLU到Host的数据传输。下面是一个例子。
cnrtMemcpy(output_half, mlu_output, sizeof(half),
CNRT_MEM_TRANS_DIR_DEV2HOST);
cnrtConvertHalfToFloat(output, output_half[0]);
需要注意的是，半数据需要在Host上转换为float/double数据才能处理。cnrt运行时提供了两个函数，cnrtConvertHalfToFloat和cnrtConvertHalfToDouble，完成数据转换功能。
发布资源
调用Kernel后，应该释放相关资源。这些资源主要包括Kernel参数buffer、queue、MLU上的DDR数据，以及Host上对应的malloc资源。最后，用户需要调用 cnrtDestroy 来释放 cnrt 运行时资源。有关详细信息，请参阅以下示例：
ret = cnrtDestroyKernelParamsBuffer(params);
ret = cnrtDestroyQueue(pQueue);
cnrtFree(mlu_input);
free(output_half);
cnrtDestroy();
核心
MLU 编程模型基于异构编程模型。在 MLU 上执行任务的程序称为 Kernel。当有足够的资源可用时，MLU 可以同时执行多个并行内核。每个Kernel都有一个Entry函数，Entry函数用来调用Device函数和Func函数。Device 函数和 Func 函数可以相互调用。Kernel 的语句由内置函数语句和 C/C++ 语言语句组成。
入口功能
Entry 函数由 BANG C 语言中的 mlu_entry 指定，如下例所示。
__mlu_entry__ void L2LossKernel(half* input, half* output) {
......
}
调用CNRT-API的cnrtInvokeKernel_V2函数启动Kernel。在 Host 上启动 l2lossKernel 函数的示例代码如下。
ret = cnrtInvokeKernel_V2((void *)(&kernel), dim, params, ft, pQueue);
当前版本的BANG C语言要求.mlu文件只有一个mlu_entry修改函数，并且入口函数所依赖的所有设备函数和数据必须在同一个编译单元中。
设备功能
设备函数是带有函数修饰符 __mlu_device__ 的函数类型。它用于定义递归函数或调用另一个函数。它有一定的函数调用成本（可以由编译选项决定是否使用内联优化）。以下是设备功能的示例。
__mlu_device__ void CreateBox(half* box, half* anchor_,
half* delt_, int A, int W,
int H, half im_w, half im_h)
Entry函数可以调用Device函数，Entry和Device函数中的语句分别使用C/C++语言和C/C++语言扩展的BANG C语言。
函数函数
Func 函数，带有函数修饰符 __mlu_func__，默认为具有 inline 属性的 Device 函数。当不需要递归函数时，可以选择 Func 函数来提高性能。func 函数必须与依赖它的入口函数在同一个编译单元中。下面是一个 Func 函数的例子。
__mlu_func__ void CreateBox(half* box, half* anchor_,
half* delt_, int A, int W,
int H, half im_w, half im_h)
Entry函数可以调用Func函数，Entry和Func函数中的语句分别使用C/C++语言和C/C++语言扩展的BANG C语言。
文件名后缀
MLU上的程序文件后缀为*.mlu，头文件的名称与C语言类似，即*.h。
头文件
MLU 程序必须包含头文件 mlu.h，其中包含 MLU 编程所需数据类型的定义，以及函数接口的声明。
内核程序示例
以下示例定义了一个 MLU 内核函数 L2LossKernel。内核函数与常见的C语言函数很相似，主要区别在于__mlu_entry__用于指定该函数为内核的入口函数，或者内核函数使用了__nram__地址空间的修饰符。有关内核编程规范，请参阅以下部分。
#include "mlu.h"
#define ONELINE 64
__mlu_entry__ void L2LossKernel(half* input, half* output, int32_t len) {
__nram__ int32_t quotient = len / ONELINE;
__nram__ int32_t rem = len % ONELINE;
__nram__ half input_nram[ONELINE];
output[0] = 0;
for (int32_t i = 0; i < quotient; i++) {
__memcpy(input_nram, input + i * ONELINE,
ONELINE * sizeof(half) , GDRAM2NRAM);
__bang_mul(input_nram, input_nram, input_nram, ONELINE);
__bang_mul_const(input_nram, input_nram, 0.5, ONELINE);
for (int32_t j = 0; j < ONELINE; j++) {
output[0] += input_nram[j];
}
}
if (rem != 0) {
__memcpy(input_nram, input + quotient * ONELINE,
ONELINE * sizeof(half), GDRAM2NRAM);
__bang_mul(input_nram, input_nram, input_nram, ONELINE);
__bang_mul_const(input_nram, input_nram, 0.5, ONELINE);
for (int i = 0; i < rem; i++) {
output[0] += input_nram[i];
}
}
}
混合编程模型基于异构编程模型。MLU程序在混合编程模型中同时执行Host端任务和MLU端任务。对于运行混合编程模型，我们建议您将 CNCC 更新到 v2.5.0 或更高版本，将 CNAS 更新到 v2.5.0 或更高版本，将 CNRT 更新到 v4.3.0 或更高版本。
下面的例子介绍了常见的异构编程模型和混合编程模型。
常见的异构编程
MLU 端 .mlu 文件如下。
__mlu_entry__ void foo(half val) {
__bang_printf("%hf\n",val);
}
Host端main.cpp文件如下。
#include "cnrt.h"
typedef unsigned short half;
#ifdef __cplusplus
extern "C" {
#endif
void foo(half);
#ifdef __cplusplus
}
#endif
int main() {
// ...
half val = 0;
cnrtConvertFloatToHalf(&val, 3.14);
cnrtDim3_t dim = {4, 1, 1};
cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;
cnrtQueue_t queue;
CNRT_CHECK(cnrtCreateQueue(&queue));
cnrtKernelParamsBuffer_t params;
CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &val, sizeof(half)));
CNRT_CHECK(cnrtInvokeKernel_V2((void *)&foo, dim, params, ktype, queue));
cnrtSyncQueue(queue);
cnrtDestroy();
return 0;
}
编译选项是：
cncc --bang-device-only --bang-mlu-arch=MLU270 -c foo.mlu -o foo.o
g++ main.cpp foo.o -I $NEUWARE_HOME/include -L $NEUWARE_HOME/lib64 -lcnrt -o a.output
混合编程
在混合编程模型中，.mlu 文件是 Host-MLU 端，如下所示。
#include<bang.h>
#include<stdio.h>
__mlu_entry__ void foo(half val) {
printf("%hf\n", val);
}
int foo() {
// ...
half val = 3.14h;
cnrtDim3_t dim = {4, 1, 1};
cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_UNION1;
cnrtQueue_t queue;
CNRT_CHECK(cnrtCreateQueue(&queue));
foo<<<dim, ktype, queue>>>(val);
cnrtSyncQueue(queue);
cnrtDestroy();
return 0;
}
int main() { return foo(); }
编译选项是：
cncc --bang-mlu-arch=MLU100 --bang-mlu-arch=MLU270 foo.mlu -o a.out
在混合编程中，invoke_kernel 在“<<<”和“>>>”中实现。“<<<”和“>>>”用于将执行参数传递给内核函数。上例中，第一个参数dim为taskdimension（详细请参考并行编程），第二个参数ktype为内核函数类型，第三个参数queue为运行时任务队列。在混合编程模型中，Host 端支持 half 的类型。
使用“<<<”和“>>>”时，必须包含bang.h，而不是mlu.h；
使用“<<<”和“>>>”时，参数类型应为u32、s32、u16、s16、u8、s8、f32，否则会报错；
在混合编程中，允许用户包含以下头文件：stdio.h、stdlib.h、math.h、libio.h、string.h、time.h、sys/time.h、memory.h
在混合编程中， -L $NEUWARE_HOME/lib64 -l cnrt 在编译选项中自动添加
有关更多混合编程示例，请参见示例目录中的示例。
BANG C 语言提供了多种内置变量供用户显式并行编程。本章介绍了 BANG C 并行内置变量和并行编程的基本概念，并举例说明。
对于单芯片中的cluster和core，用户可以使用内置变量clusterDim、clusterId、coreDim、coreId分别表示cluster和core的维度和ID。另外，每个硬件计算核心执行一个任务，用户通过cnrtDim3_t类型指定Kernel任务的大小。任务大小通常具有三个维度：x、y 和 z。用户可以根据场景需求指定任务大小。下面的例子展示了如何在 x 维度上将任务分成四个部分。
cnrtDim3_t dim;
dim.x = 4;
dim.y = 1;
dim.z = 1;
task的内置变量包括：taskDim、taskDimX、taskDimY、taskDimZ、taskIdX、taskIdY、taskIdZ、taskId。taskDim = taskDimX * taskDimY * taskDimZ。在上面的代码中，taskDimX 为 4，taskDimY 为 1，taskDimZ 为 1。taskDim 与 UNION 的数量和集群的数量无关。每个任务都映射到一个计算核，即每个任务由一个核执行，核是MLU的基本处理单元。一个 Kernel 由多个任务组成，这些任务通过 cnrtDim3_t 数据结构在三个维度上进行索引。
在 BANG C 语言中，每个内核都有一个 Block 类型的任务。块是编程模型层的基本调度单元。表示将Kernel中的任务调度到单核上执行。对于每个 Cluster，都有一个 Union1 类型的任务。两个簇形成一个 Union2，以此类推。我们称之为任务类型。任务类型指定了一次Kernel启动所需的硬件核数，即在Kernel执行周期内一直占用的物理核数。块型任务是单核任务，联合型任务是多核并行任务。
以完成16384个长向量相加为例，假设任务的任务类型为Union4，可以设置dim.x = 16。每个任务完成n = 1024个向量相加，总共16个任务后，你得到一个暗淡的向量加法。X * n = 16384 长度。深度学习处理器上的向量加法代码如下例所示。
#define N 1024
__mlu_entry__ add(float* x, float* y, float* z) {
__nram__ float x_tmp[N];
__nram__ float y_tmp[N];
__memcpy(x_tmp, x + taskID * N, N * sizeof(float), GDRAM2NRAM);
__memcpy(y_tmp, y + taskID * N, N * sizeof(float), GDRAM2NRAM);
__bang_add(x_tmp, x_tmp, y_tmp, N);
__memcpy(z + taskID * N, x_tmp, N * sizeof(float), NRAM2GDRAM);
}
笔记
执行模型
MLU270在启动执行任务时按照taskDimX作为1个单位的并行度下发和执行指令，即最小并行粒度为taskDimX = clusterDim * coreDim。对于MLU270，coreDim = 4，用户可以通过指定Union类型来控制并行最小粒度taskDimX。这个限制如下表所示。当用户指定Union类型时，如果指定的taskDimX不是clusterDim * coreDim的正整数倍，CNRT运行时会报错。
| 硬件并行占用 | 软件并行限制 | ||
| 联盟 | 簇 | 核 | 任务暗淡 |
| 1 | 1 | 1*4=4 | taskDimX = N * clusterDim * coreDim
|
| 2 | 2 | 2*4=8 | |
| 4 | 4 | 4*4=16 | |
| 8 | 8 | 8*4=32 | |
MLU270 并行执行模型
如图MLU270 Parallel Executed Model所示，Host 上有 3 个不同的 Kernel 函数被传输到 MLU270 设备执行。具体执行流程如下：
对于 Kernel1，Union 类型和 taskDim 定义如下：
cnrtFunctionType_t Union2 = CNRT_FUNC_TYPE_UNION2;
cnrtDim3_t taskDim = {8, 1, 1}; // for example
cnrtInvokeKernel_V2((void *)&kernel1, taskDim, params, Union2, pQueue);
Host 处的 CNRT 将 Kernel1 的任务大小 taskDim 和任务类型 Union2 传输给任务调度器。调度器会等到 2 个硬件 CLUSTER 空闲（用户在 Kernel1 中得到 clusterId = [0 ~ 1], clusterDim = 2）；
当调度器发现有 2 个 CLUSTER 空闲时，它会启动 2 个 CLUSTER 中的总共 8 个 CORE，这样每个 CORE 开始执行相同的 Kernel1 代码。
Kernel1 软件的总并行大小为 taskDimX * taskDimY * taskDimZ。根据表2中taskDimX * taskDimY * taskDimZ = 8的限制，调度器会占用2个CLUSTER在Time1启动Union2大小的并行任务，直到该MLU上8个CORE的所有指令执行结束。
当Kernel1下发时，如果Kernel2和Kernel1之间没有依赖关系（即Kernel2和Kernel1没有绑定到同一个Queue，两个计算任务可以独立下发，不需要计算输入和结果的相互依赖关系），那么Host 将继续发出 Kernel2 任务。
对于 Kernel2，Union 类型和 taskDim 定义如下：
cnrtFunctionType_t Union1 = CNRT_FUNC_TYPE_UNION1;
cnrtDim3_t taskDim = {4, 1, 1}; // for example
cnrtInvokeKernel_V2((void *)&kernel2, taskDim, params, Union1, pQueue);
调度器会等到一个硬件CLUSTER空闲后，才会下发Kernel2的任务。Kernel2 和 Kernel1 是并行执行的，除了 GDRAM 之外，它们都是相互独立的。
Kernel2在硬件CLUSTER上的执行方式与Kernel1相同，Kernel2软件的总并行大小为taskDimX * taskDimY * taskDimZ = 4。调度器会占用1个CLUSTER在Time1启动Union1并行任务，直到所有的执行完毕此 MLU 上的 4 个 CORE 指令结束。
当Kernel2下发时，如果Kernel3和Kernel1或Kernel1之间没有依赖关系（例如等待同步），那么Host会继续下发Kernel3任务（在图MLU270 Parallel Executed Model中，假设CLUSTER[4 -7] 总是很忙）。
对于 Kernel3，Union 类型和 taskDim 定义如下：
cnrtFunctionType_t Union4 = CNRT_FUNC_TYPE_UNION4;
cnrtDim3_t taskDim = {16, 1, 1}; // for example
cnrtInvokeKernel_V2((void *)&kernel3, taskDim, params, Union4, pQueue);
该程序声明需要执行 4 个硬件 CLUSTER。在Time2，调度器发现逻辑ID相邻的CLUSTER[0-3]是空闲的，于是此时发出Kernel3的Union4任务。
笔记
coreDim
BANG C语言内置变量名称保留关键字，取值等于1个CLUSTER拥有的CORE个数（在MLU220/MLU270上等于4个）。
核心标识
BANG C语言内置变量名称保留关键字，取值等于CLUSTER内各硬件CORE的逻辑ID（MLU220/MLU270取值范围为[0-3]）。
集群昏暗
BANG C语言内置变量名称保留关键字，取值由任务的Union信息决定（MLU220/MLU270最大值为4，MLU100最大值为8）。
集群 ID
BANG C语言内置变量名称保留关键字，取值等于程序运行所在CLUSTER的逻辑ID，取值范围为[0-clusterDim-1]。
任务DimX
BANG C 语言的内置变量名称保留关键字。在调用内核之前，用户需要设置此任务的逻辑大小。逻辑大小有三个维度：{X, Y, Z}，taskDimX的值等于X方向的大小。
任务维度
BANG C 语言的内置变量名称保留关键字。在调用内核之前，用户需要设置此任务的逻辑大小。逻辑大小有三个维度：{X, Y, Z}，taskDimY的值等于Y方向的大小。
任务DimZ
BANG C 语言的内置变量名称保留关键字。在调用内核之前，用户需要设置此任务的逻辑大小。逻辑大小有三个维度：{X, Y, Z}，taskDimZ的值等于Z方向的大小。
任务暗淡
BANG C 语言的内置变量名称保留关键字。在调用内核之前，用户需要设置此任务的逻辑大小。逻辑大小有三个维度：{X,Y,Z}，taskDim的值等于当前任务降维后的逻辑大小的值，即taskDim=taskDimX*taskDimY*taskDimZ。
任务IdX
BANG C语言内置变量名称保留关键字，取值等于程序运行时分配的逻辑大小的X维度的任务ID，取值范围为[0-taskDimX-1]。
任务IDY
BANG C语言内置变量名称保留关键字，取值等于程序运行时分配的逻辑大小的Y维度的任务ID，取值范围为[0-taskDimY-1]。
任务IDZ
BANG C语言内置变量名称保留关键字，取值等于程序运行时分配的逻辑大小的Z维度的任务ID，取值范围为[0-taskDimZ-1]。
任务标识
BANG C语言内置变量名称保留关键字，取值等于程序运行时分配的任务ID，取值范围为[0-taskDim-1]。taskId的值等于降维后逻辑大小的任务ID，即：taskId=taskIdZ*taskDimY *taskDimX + taskIdY *taskDimX+taskIdX
将 BANG C 程序与机器学习框架集成有两种方法：
第一种方式是将具体的BANG C实现文件编译成动态链接库，调用动态链接库中的函数计算框架中的算子，如图BANG C程序集成框架与动态链接图书馆。
BANG C 程序将框架与动态链接库集成
第二种方式是直接将 BANG C 源码集成到机器学习框架中，编译 BANG C 源码，如图BANG C Program Integrates Framework with Source Code所示。
BANG C 程序将框架与源代码集成
我们以 upsample_layer 和 caffe 框架的集成为例来说明具体的集成过程。流程图如图Upsample_layer 与 Caffe Framework 的集成。
upsample_layer 与 Caffe 框架的集成
先写对应的算子。有两个文件：mlu_upsample.cpp、mlu_upsample.mlu。具体代码请参考demo。
将步骤1)中的两个文件编译到动态链接库libmluupsample.so中。编译命令显示在示例中的 Makefile 中。
新建caffe/src/caffe/mlu_layers文件夹，将步骤2)中的libmluupsample.so复制到该文件夹，在caffe/src/caffe/CMakeLists文件中添加动态链接库的链接：
链接目录（./mlu_layers）
链接库（libmluupsample.so）
修改框架，在caffe中调用步骤2)动态链接库libmluupsample.so中的函数；添加新文件 caffe/include/caffe/layers/mlu_upsample_layer.hpp 和 caffe/src/caffe/layers/mlu_upsample_layer.cpp；在文件 caffe/src/caffe/layers/mlu_upsample_layer.cpp 中声明并实现 MLUUpsampleLayer 和 forward_mlu 函数。
修改caffe/src/caffe/layer_factory.cpp，增加GetUpsampleLayer函数。具体代码请参考示例中的示例。
重新编译 caffe。
详情参考：Programming Model — Cambricon BANG C Developer Guide 2.15.0 documentation
矩阵乘Demo执行过程中，用户先输入参数m,k,n，代表要计算的左右矩阵分别为m*k和k*n大小，HOST端对这两个矩阵进行随机赋值，将输入矩阵以及大小相应的参数传入MLU端进行矩阵运算，最后将运算结果传回HOST端，在HOST端打印矩阵乘的硬件处理时间。
HOST端关键代码如下：
1）输入左右矩阵初始化
float *A = (float *)malloc(M * K * sizeof(float));
float *B = (float *)malloc(K * N * sizeof(float));
float *C = (float *)malloc(M * N * sizeof(float));
float *Cmlu = (float *)malloc(M * N * sizeof(float));
// 给A和B矩阵随机赋值
for (int i = 0; i < M; i++)
{
for (int j = 0; j < K; j++)
{
A[i*K +j] = (i + rand()%16)/217.0;
}
}
int p = 0;
for (int i = 0; i < K; i++)
{
for (int j = 0; j < N; j++)
{
B[i*N +j] = ((float)(rand()%20+3))/1003.0;
}
}
2）准备相关参数，启动Kernel，将参数传入MLU端
cnrtDim3_t dim;cnrtFunctionType_t func_type = CNRT_FUNC_TYPE_BLOCK;
dim.x = 1;
dim.y = 1;
dim.z = 1;
CNRT_CHECK(cnrtMalloc((void **)&d_c, sizeof(half) * M * N_align));
CNRT_CHECK(cnrtMalloc((void **)&d_a, sizeof(half) * M * K));
CNRT_CHECK(cnrtMalloc((void **)&d_w, sizeof(half) * K * N_align));
A_half = (half *)malloc(sizeof(half) * M * K);
B_half = (half *)malloc(sizeof(half) * K * N_align);
cnrtKernelInitParam_t init_param;CNRT_CHECK(cnrtCreateKernelInitParam(&init_param));CNRT_CHECK(cnrtInitKernelMemory((const void*)gemm16Kernel, init_param));cnrtKernelParamsBuffer_t params;CNRT_CHECK(cnrtGetKernelParamsBuffer(¶ms)); CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &d_c, sizeof(half *))); CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &d_a, sizeof(half *)));CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &d_w, sizeof(half *)));CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &M, sizeof(uint32_t)));CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &K, sizeof(uint32_t)));CNRT_CHECK(cnrtKernelParamsBufferAddParam(params, &N_align, sizeof(uint32_t)));cnrtNotifier_t notifier_start; // A pointer which points to the struct describing notifier.cnrtNotifier_t notifier_end;CNRT_CHECK(cnrtCreateNotifier(¬ifier_start));CNRT_CHECK(cnrtCreateNotifier(¬ifier_end));float timeTotal = 0.0;//printf("start invoke : \n");gettimeofday(&start, NULL);CNRT_CHECK(cnrtPlaceNotifier(notifier_start, pQueue)); // Places a notifier in specified queueCNRT_CHECK(cnrtInvokeKernel_V3((void *)&gemm16Kernel,init_param, dim, params, func_type, pQueue, NULL)); // Invokes a kernel written in Bang with given params on MLUCNRT_CHECK(cnrtPlaceNotifier(notifier_end, pQueue)); // Places a notifier in specified queue
在后续优化的过程中，HOST端的代码基本不变，我们重点关注MLU端代码的开发和优化过程。
MLU端的BANG C实现，我们分成6个步骤，逐一优化，希望帮助大家理解BANG C的使用和优化方法。
直接在GDRAM上使用循环和标量操作进行计算
无须对输入的矩阵作任何处理，使用矩阵乘公式直接计算，完全没有利用到MLU硬件架构的优势，所以整个计算时间很长。MLU端关键代码如下：
#include "mlu.h"
__mlu_entry__ void gemm16Kernel(half *outputDDR,
half *input1DDR,
half *input2DDR,
uint32_t m,
uint32_t k,
uint32_t n) {
half ret;
__bang_printf("m=%d,k=%d,n=%d\n",m,k,n);
for (uint32_t i = 0; i < m; i++) {
for (uint32_t j = 0; j < n; j++) {
ret = 0;
for (uint32_t t = 0; t < k; t++) {
ret += input1DDR[i*k+t] * input2DDR[t*n+j];
}
outputDDR[i*n+j] = ret;
}
}
}
基本的数据调度
第二步是在第一步的基础上引入NRAM/WRAM的使用，每个core都有自己的NRAM和WRAM，虽然相比于GDRAM空间小，但是可以获得更高的读写带宽和更低的访问时延。（片上存储层次相关介绍可参考BANG C开发者手册）
我们将输入的左右矩阵从GDRAM拷入NRAM中，在NRAM中进行计算，然后拷回GDRAM。需要注意的是，在这个例子中我们假设输入的左右矩阵规模都为256*256，来保证输入的矩阵可以一次性拷入NRAM/WRAM。一旦输入矩阵规模超过NRAM/WRAM的空间大小时，则需要对NRAM/WRAM复用，进行多次拷入和拷出。
MLU端关键代码如下：
#include "mlu.h"
__mlu_entry__ void gemm16Kernel(half *outputDDR, half *input1DDR, half *input2DDR,
uint32_t m, uint32_t k, uint32_t n) {
__nram__ half input1NRAM[256*256];
__nram__ half input2NRAM[256*256];
__nram__ half outputNRAM[256*256];
__memcpy(input1NRAM, input1DDR, m * k * sizeof(half), GDRAM2NRAM); //从 GDRAM拷入NRAM
__memcpy(input2NRAM, input2DDR, k * n * sizeof(half), GDRAM2NRAM);
for (uint32_t i = 0; i < m; i++) {
for (uint32_t j = 0; j < n; j++) {
half ret = 0.0;
half c = 0.0;
for (uint32_t t = 0; t < k; t++) {
half v = input1NRAM[i*k+t] * input2NRAM[t*n+j];
half y = v - c;
half temp = ret + y;
c = ( temp - ret) - y;
ret = temp;
}
outputNRAM[i*n+j] = ret;
}
}
__memcpy(outputDDR, outputNRAM, m * n * sizeof(half), NRAM2GDRAM); //将计算结果拷回GDRAM
}
计算的向量化
第三步在以上的基础上，使用BANG C提供的向量计算指令完成矩阵乘的计算，采用向量计算指令可以更好地发挥MLU硬件性能，减少计算时间。
我们先介绍下后续要解决的矩阵乘中的矩阵规模大小问题，为了方便展示和读者理解，假设的是左矩阵规模大小为256*256，右矩阵规模大小为256*N（N可被256整除）。
在这种情况下，输入的左矩阵可一次性拷入NRAM，在执行卷积指令操作时，将输入的右矩阵拷入WRAM中。在往WRAM拷入前，需要对数据进行量化处理，并且摆放成特定要求的数据摆放格式，使用__bang_conv指令进行计算。由于右矩阵规模较大，将右矩阵分批次拷入WRAM进行计算。
MLU端关键代码如下。其中，all_round表示计算的循环次数，这和右矩阵规模大小相关；dst_stride和src_stride代表调整右矩阵数据摆放格式过程中的步长；total_times表示调整右矩阵数据格式的次数，因为目前MLU270上有64个卷积计算单元，所以需要将原本顺序摆放的数据按照64个为一组间隔摆放。
#include "mlu.h"
#define ROUND 256
__mlu_entry__ void gemm16Kernel(half *outputDDR, int8_t *input1DDR, int8_t *input2DDR,
uint32_t m, uint32_t k, uint32_t n, int16_t pos) {
__nram__ int8_t input1NRAM[256*256];
__nram__ int8_t input2NRAM[256*256];
__nram__ int8_t input2NRAM_tmp[256*256];
__wram__ int8_t input2WRAM[256*256];
__nram__ half outputNRAM[256*256];
__memcpy(input1NRAM, input1DDR, m * k * sizeof(int8_t), GDRAM2NRAM);
//在这里将左矩阵一次性拷入NRAM
int all_round = n / ROUND;
int32_t dst_stride = (ROUND * k / 64) * sizeof(int8_t);
int32_t src_stride = k * sizeof(int8_t);
int32_t size = k * sizeof(int8_t);
int32_t total_times = ROUND / 64;
//__bang_printf("taskDim=%d,clusterId=%d,coreId=%d\n",taskDim,clusterId,coreId);
for(int i = 0; i < all_round; i++) {
__bang_write_zero((half *)input2NRAM_tmp, 256 * 128);
__bang_write_zero((half *)input2NRAM, 256 * 128);
__memcpy(input2NRAM_tmp, input2DDR + i * ROUND * k,
k * ROUND * sizeof(int8_t), GDRAM2NRAM);
for (int j = 0; j < total_times; j++) { //这里将数据摆放成bang_conv可以使用的格式
__memcpy(input2NRAM + j * k, input2NRAM_tmp + j * 64 * k,
size, NRAM2NRAM, dst_stride, src_stride, 64);
}
__memcpy(input2WRAM, input2NRAM, ROUND*k*sizeof(int8_t), NRAM2WRAM);
__bang_conv(outputNRAM, input1NRAM, input2WRAM, k, m, 1, 1, 1, 1, 1, ROUND, pos);
__memcpy(outputDDR + i * ROUND,
outputNRAM,
ROUND * sizeof(half),
NRAM2GDRAM,
n * sizeof(half),
ROUND * sizeof(half),
m-1);
/*for (int j = 0; j < m; j++) { //要对每轮计算的结果进行拼接
__memcpy(outputDDR + i * ROUND + j * n, outputNRAM + j * ROUND,
ROUND * sizeof(half), NRAM2GDRAM);
}*/
}
}
计算任务的多核拆分
在第三步的计算中，我们只使用了1个MLU core进行计算，MLU270上有16个MLU core，这一步可以进一步采用16个core并行运算。根据输入矩阵规模的大小，将输入矩阵拆分成多份并分配给不同的计算core，最后再对计算结果进行合并，提高了计算效率。
MLU端关键代码如下。在实现过程中，我们会用到与并行相关的内置变量：taskDim表示任务规模，taskId表示程序运行时所分配的任务ID,在这步的方法中taskDim=16，taskId范围为[0,15]。更多关于taskDim和taskId的介绍，读者可以参考BANG C用户手册第5章的内容：
#include "mlu.h"
#define ROUND 256
#define NRAM_ARRAY_SIZE 256*256
__mlu_entry__ void gemm16Kernel(half *outputDDR, int8_t *input1DDR, int8_t *input2DDR,
uint32_t m, uint32_t k, uint32_t n, int16_t pos) {
__nram__ int8_t input1NRAM[NRAM_ARRAY_SIZE];
__nram__ int8_t input2NRAM[NRAM_ARRAY_SIZE];
__nram__ int8_t input2NRAM_tmp[NRAM_ARRAY_SIZE];
__wram__ int8_t input2WRAM[NRAM_ARRAY_SIZE];
__nram__ half outputNRAM[NRAM_ARRAY_SIZE];
__memcpy(input1NRAM, input1DDR, m * k * sizeof(int8_t), GDRAM2NRAM);
//在这里将左矩阵一次性拷入NRAM
int all_round = n / ( taskDim * ROUND); //因为现在使用16个核同时运算，所以每个核循环的次数也相应减少
int32_t dst_stride = (ROUND * k / 64) * sizeof(int8_t);
int32_t src_stride = k * sizeof(int8_t);
int32_t size = k * sizeof(int8_t);
int32_t total_times = ROUND / 64;
//__bang_printf("taskDim=%d,taskId=%d\n",taskDim, taskId);
for(int i = 0; i < all_round; i++) {
__memcpy(input2NRAM_tmp, input2DDR + ROUND * (i * taskDim + taskId) * k , //只涉及这个核需要的数据
k * ROUND * sizeof(int8_t), GDRAM2NRAM);
for (int j = 0; j < total_times; j++) {
__memcpy(input2NRAM + j * k, input2NRAM_tmp + j * 64 * k,
size, NRAM2NRAM, dst_stride, src_stride, 64 - 1);
}
__memcpy(input2WRAM, input2NRAM, ROUND*k*sizeof(int8_t), NRAM2WRAM);
__bang_conv(outputNRAM, input1NRAM, input2WRAM, k, m, 1, 1, 1, 1, 1, ROUND, pos);
for (int j = 0; j < m; j++) { //向GDRAM回写的时候也要注意每个核的位置不同
__memcpy(outputDDR + (i * taskDim + taskId) * ROUND + j * n,
outputNRAM + j * ROUND, ROUND * sizeof(half), NRAM2GDRAM);
}
}
}
SRAM的使用
第五步是在第四步的基础上引入Shared-RAM，在MLU270中，一个cluster中的4个core共享一个SRAM。在第四步中，因为使用了4个cluster的16个core进行并行计算，而同1个cluster上的4个core在从GDRAM上拷贝数据到各自的NRAM/WRAM时，会争抢这个cluster到GDRAM的带宽，从而导致数据读取速度降低。所以我们将数据先从GDRAM拷贝到SRAM，再从SRAM分发到NRAM/WRAM中，避免了调度争抢问题，提高了数据读取速度。
特别注意的是，从GDRAM拷入数据到SRAM和从SRAM拷入数据到NRAM这两个操作，是由两种不同功能的core执行（这个会在后文中解释），所以这两个操作是并行的关系。为了避免数据冲突，我们要设置同步功能，保证数据从GDRAM拷入到SRAM之后，才能执行从SRAM拷入到NRAM的过程，在BANG C中我们可以使用内置的__sync_cluster()函数完成同步功能。图示如下：
整个执行过程如下图所示：
MLU端关键代码如下，其中clusterId表示此时执行任务的是哪个cluster，范围为[0,3]：
#include "mlu.h"
#define ROUND 256
__mlu_entry__ void gemm16Kernel(half *outputDDR, int8_t *input1DDR, int8_t *input2DDR,
uint32_t m, uint32_t k, uint32_t n, int16_t pos) {
__nram__ int8_t input1NRAM[256*256];
__nram__ int8_t input2NRAM[256*256];
__nram__ int8_t input2NRAM_tmp[256*256];
__wram__ int8_t input2WRAM[256*256];
__nram__ half outputNRAM[256*256];
__memcpy(input1NRAM, input1DDR, m * k * sizeof(int8_t), GDRAM2NRAM);
//在这里将左矩阵一次性拷入NRAM
int all_round = n / ( taskDim * ROUND); //因为现在使用16个核同时运算，所以每个核循环的次数也相应减少
int32_t dst_stride = (ROUND * k / 64) * sizeof(int8_t);
int32_t src_stride = k * sizeof(int8_t);
int32_t size = k * sizeof(int8_t);
int32_t total_times = ROUND / 64;
__mlu_shared__ int8_t input2SRAM[256*1024];
//_bang_printf("taskDim=%d,clusterId=%d,coreId=%d\n",taskDim,clusterId,coreId);
for(int i = 0; i < all_round; i++)
{
// copy GDRAM2SRAM
__memcpy(input2SRAM, input2DDR + ROUND * (i * taskDim + clusterId * 4) * k ,
k * ROUND * 4 * sizeof(int8_t), GDRAM2SRAM); // 只将右矩阵拷入SRAM中
__sync_cluster(); //设置sync barrier
// copy SRAM2NRAM
__memcpy(input2NRAM_tmp, input2SRAM + ROUND * coreId * k , k * ROUND * sizeof(int8_t), SRAM2NRAM);
// 将数据摆好对应的格式
for (int j = 0; j < total_times; j++) {
__memcpy(input2NRAM + j * k, input2NRAM_tmp + j * 64 * k,
size, NRAM2NRAM, dst_stride, src_stride, 64 - 1);
}
// copy NRAM2WRAM
__memcpy(input2WRAM, input2NRAM, ROUND*k*sizeof(int8_t), NRAM2WRAM);
// compute
__bang_conv(outputNRAM, input1NRAM, input2WRAM, k, m, 1, 1, 1, 1, 1, ROUND, pos);
// copy NRAM2GDRAM
for (int j = 0; j < m; j++) { //向GDRAM回写的时候也要注意每个核的位置不同
__memcpy(outputDDR + (i * taskDim + taskId) * ROUND + j * n,
outputNRAM + j * ROUND, ROUND * sizeof(half), NRAM2GDRAM);
}
__sync_cluster(); //设置sync barrier
}
}
基本的流水优化
MLU270上，每个cluster除了4个普通的计算core之外，还有专门用以管理片上总线和SRAM的memory core。这就是上一步提到的两种不同功能的计算单元，为我们使用流水线优化创造了条件。
第六步在上面的基础上，实现了4个cluster并行计算，且每个cluster中的memory core和其他4个 MLU core构成流水线的计算模式。在每个cluster中，memory core只负责将数据从GDRAM拷入SRAM，其余的每个MLU core则负责从SRAM拷入数据、矩阵乘计算、将数据拷回GDRAM。
我们设置了在SRAM上的两个变量input2SRAM1，input2SRAM2。初始时，memory core从GDRAM上拷入数据到input2SRAM1，当数据拷入完成后，4个core开始工作，它们将自己需要的数据部分从input2SRAM1拷入进行计算。在MLU core工作的同时，memory core不会停止工作，它会将下一次需要计算的数据从GDRAM拷入input2SRAM2，供给4个MLU core在下一次使用，减少了拷入等待时间，input2SRAM1和input2SRAM2交替读写重复上述过程直至所有数据计算完成。
从中可以发现，耗时很长GDRAM到SRAM的这一步拷贝时间被“藏起来”了。和原来相比，在相同的时间内，我们搬运了更多的GDRAM数据到片上并且完成了计算。那么为什么会使用两个SRAM变量对GDRAM上的数据进行拷贝呢？因为在上述过程中，MLU core在从SRAM读取数据的同时，SRAM也会从GDRAM写入数据，如果只使用一个SRAM变量，则很有可能导致MLU core应该读取的数据在读取前被写入覆盖。
有经验的开发者可能已经发现，这里使用的是一种常用的数据流控制的处理技巧，乒乓操作。
整个过程如下图所示:
MLU端关键代码如下：
#include "mlu.h"
#define ROUND 256
__mlu_entry__ void gemm16Kernel(half *outputDDR, int8_t *input1DDR, int8_t *input2DDR,
uint32_t m, uint32_t k, uint32_t n, int16_t pos) {
__nram__ int8_t input1NRAM[256*256];
__nram__ int8_t input2NRAM[256*256];
__nram__ int8_t input2NRAM_tmp[256*256];
__wram__ int8_t input2WRAM[256*256];
__nram__ half outputNRAM[256*256];
__memcpy(input1NRAM, input1DDR, m * k * sizeof(int8_t), GDRAM2NRAM);
//在这里将左矩阵一次性拷入NRAM
int all_round = n / ( taskDim * ROUND); //因为现在使用16个核同时运算，所以每个核循环的次数也相应减少
int32_t dst_stride = (ROUND * k / 64) * sizeof(int8_t);
int32_t src_stride = k * sizeof(int8_t);
int32_t size = k * sizeof(int8_t);
int32_t total_times = ROUND / 64;
__mlu_shared__ int8_t input2SRAM1[256*1024];
__mlu_shared__ int8_t input2SRAM2[256*1024];
__mlu_shared__ int8_t * input2SRAM_read;
__mlu_shared__ int8_t * input2SRAM_write;
input2SRAM_write=input2SRAM1;
// copy GDRAM2SRAM
__memcpy(input2SRAM_write, input2DDR + ROUND * (clusterId * 4) * k,
k * ROUND * 4 * sizeof(int8_t), GDRAM2SRAM); // 只将右矩阵拷入SRAM中
__sync_cluster(); //设置sync barrier
//_bang_printf("taskDim=%d,clusterId=%d,coreId=%d\n",taskDim,clusterId,coreId);
for(int i = 0; i < all_round-1; i++)
{
if (i % 2 == 0)
{
input2SRAM_read=input2SRAM1;
input2SRAM_write=input2SRAM2;
} else {
input2SRAM_read=input2SRAM2;
input2SRAM_write=input2SRAM1;
}
if (coreId == 0x80) {
// copy GDRAM2SRAM
__memcpy(input2SRAM_write, input2DDR + ROUND * ((i+1) * taskDim + clusterId * 4) * k,
k * ROUND * 4 * sizeof(int8_t), GDRAM2SRAM); // 只将右矩阵拷入SRAM中
} else {
// copy SRAM2NRAM
__memcpy(input2NRAM_tmp, input2SRAM_read + ROUND * coreId * k , k * ROUND * sizeof(int8_t), SRAM2NRAM);
// 将数据摆好对应的格式
for (int j = 0; j < total_times; j++) {
__memcpy(input2NRAM + j * k, input2NRAM_tmp + j * 64 * k,
size, NRAM2NRAM, dst_stride, src_stride, 64 - 1);
}
// copy NRAM2WRAM
__memcpy(input2WRAM, input2NRAM, ROUND*k*sizeof(int8_t), NRAM2WRAM);
// compute
__bang_conv(outputNRAM, input1NRAM, input2WRAM, k, m, 1, 1, 1, 1, 1, ROUND, pos);
// copy NRAM2GDRAM
for (int j = 0; j < m; j++) { //向GDRAM回写的时候也要注意每个核的位置不同
__memcpy(outputDDR + (i * taskDim + taskId) * ROUND + j * n,
outputNRAM + j * ROUND, ROUND * sizeof(half), NRAM2GDRAM);
}
}
__sync_cluster(); //设置sync barrier
}
__memcpy(input2NRAM_tmp, input2SRAM_write + ROUND * coreId * k , k * ROUND * sizeof(int8_t), SRAM2NRAM);
// 将数据摆好对应的格式
for (int j = 0; j < total_times; j++) {
__memcpy(input2NRAM + j * k, input2NRAM_tmp + j * 64 * k,
size, NRAM2NRAM, dst_stride, src_stride, 64 - 1);
}
// copy NRAM2WRAM
__memcpy(input2WRAM, input2NRAM, ROUND*k*sizeof(int8_t), NRAM2WRAM);
// compute
__bang_conv(outputNRAM, input1NRAM, input2WRAM, k, m, 1, 1, 1, 1, 1, ROUND, pos);
// copy NRAM2GDRAM
for (int j = 0; j < m; j++) { //向GDRAM回写的时候也要注意每个核的位置不同
__memcpy(outputDDR + ((all_round - 1) * taskDim + taskId) * ROUND + j * n,
outputNRAM + j * ROUND, ROUND * sizeof(half), NRAM2GDRAM);
}
}
具体开发还可参考开发手册：Built-in Function — Cambricon BANG C Developer Guide 2.15.0 documentation
本文对开源yolov5s模型进行寒武纪平台的移植
代码参考链接：GitHub - CambriconECO/Pytorch_Yolov5_Inference at eb36b68e5a4ee3fe23c45ec545bef68dced452bb
整个移植过程分为模型结构转换、添加后处理算子框架代码、模型量化、在线推理和离线推理共五个步骤。
对于原始Yolov5网络的后处理部分的逻辑，Cambricon-PyTorch直接使用一个大的BANGC算子完成后处理的计算，需要对原生的pytorch网络进行修改，将后处理部分的整体计算换成BANGC算子。
具体做法为是将yololayer层替换成了yolov5_detection_output，把三个yololayer的输入传给了yolov5_detection_output。修改部分在yolo.py中，如下：
if x[0].device.type == 'mlu':
for i in range(self.nl):
x[i] = self.m[i](x[i]) # conv
y = x[i].sigmoid()
output.append(y)
detect_out = torch.ops.torch_mlu.yolov5_detection_output(output[0], output[1], output[2],
self.anchors_list,self.nc, self.num_anchors,
self.img_h, self.img_w, self.conf_thres, self.iou_thres, self.maxBoxNum)
return detect_out
if x[0].device.type == 'cpu':
z = []
for i in range(self.nl):
x[i] = self.m[i](x[i]) # conv
bs, _, ny, nx = x[i].shape # x(bs,255,20,20) to x(bs,3,20,20,85)
x[i] = x[i].view(bs, self.na, self.no, ny, nx).permute(0, 1, 3, 4, 2).contiguous()
if not self.training: # inference
if self.grid[i].shape[2:4] != x[i].shape[2:4]:
self.grid[i] = self._make_grid(nx, ny).to(x[i].device)
y = x[i].sigmoid()
y[..., 0:2] = (y[..., 0:2] * 2. - 0.5 + self.grid[i].to(x[i].device)) * self.stride[i] # xy
y[..., 2:4] = (y[..., 2:4] * 2) ** 2 * self.anchor_grid[i] # wh
z.append(y.view(bs, -1, self.no))
return x if self.training else torch.cat(z, 1)
在第一步中我们使用了BANGC编写的Yolov5Detection算子替换掉了原始的后处理逻辑。为了保证能够正确调用到这个算子，需要将该Yolov5Detection算子集成到框架中。
共分成两步：先将算子集成到CNPlugin中，然后将CNPlugin算子集成到Cambricon-Pytorch。该算子的实现与集成到CNPlugin会在另一个教程中详细介绍，在这里只介绍将该算子集成到Cambricon Pytorch这一步骤。
1）声明算子。在catch/torch_mlu/tools/mlu_functions.yaml中声明算子
- name: yolov5_detection_output // 算子名称
use_mlu_dispatcher: custom // 分发类型，unboxed_only为标准化算子，custom为客制化算子
derived_type: cnml // 派生类型
schema_string: torch_mlu::yolov5_detection_output // 用于算子注册
arguments: // 参数
- name: alpha_data // 参数名称
type: const at::Tensor & // 参数类型
- name: beta_data
type: const at::Tensor &
- name: gamma_data
type: const at::Tensor &
- name: anchor_data
type: torch::List<int64_t>
- name: num_classes
type: int64_t
- name: num_anchors
type: int64_t
- name: img_height
type: int64_t
- name: img_width
type: int64_t
- name: conf_thres
type: double
- name: nms_thres
type: double
- name: maxBoxNum
type: int64_t return_type: at::Tensor // 函数返回类型
2）添加OpMethods基类中的CPU实现。
Catch模块中包含AtenMluType标准算子类型和AtenMluCustomType定制化算子类型，AtenMluType和AtenMluCustomType会通过OpMethods下发到推理算子或训练算子。根据模板生成的 op_methods.h算子声明，在catch/torch_mlu/csrc/aten/operators/op_methods.cpp中添加算子的CPU实现。
// op_methods.h
virtual at::Tensor yolov5_detection_output(const at::Tensor & alpha_data, const at::Tensor & beta_data, const at::Tensor & gamma_data, torch::List<int64_t> anchor_data, int64_t num_classes, int64_t num_anchors, int64_t img_height, int64_t img_width, double conf_thres, double nms_thres, int64_t maxBoxNum);
// op_methods.cpp 因为在这里并没有添加CPU实现，直接抛出异常
at::Tensor OpMethods::yolov5_detection_output(
const at::Tensor& alpha_data, const at::Tensor& beta_data,
const at::Tensor& gamma_data, torch::List<int64_t> anchor_data,
int64_t num_classes, int64_t num_anchors, int64_t img_height, int64_t img_width,
double conf_thres, double nms_thres, int64_t maxBoxNum) {
// To do: cpu kernel
throw std::invalid_argument("To do for CPU");
3）添加wrapper
推理算子或训练算子会优先分发到wrapper中。wrapper是对算子kernel的封装，每个算子对应一个wrapper。根据模板生成的wrapper头文件cnml_kernel.h，添加wrapper的实现。
// cnml_kernel.h
at::Tensor cnml_yolov5_detection_output(const at::Tensor & alpha_data, const at::Tensor & beta_data, const at::Tensor & gamma_data, torch::List<int64_t> anchor_data, int64_t num_classes, int64_t num_anchors, int64_t img_height, int64_t img_width, double conf_thres, double nms_thres, int64_t maxBoxNum);
// yolov5_detection_output.cpp
at::Tensor cnml_yolov5_detection_output(const at::Tensor& alpha_data,
const at::Tensor& beta_data,
const at::Tensor& gamma_data,
torch::List<int64_t> anchor_data,
int64_t num_classes,
int64_t num_anchors,
int64_t img_height,
int64_t img_width,
double conf_thres,
double nms_thres,
int64_t maxBoxNum) {
auto alpha_new = alpha_data;
auto beta_new = beta_data;
auto gamma_new = gamma_data;
bool cast_fp32_fp16 = (toCnmlDataType(alpha_data.dtype()) == CNML_DATA_FLOAT32);
if (cast_fp32_fp16) {
alpha_new = cnml_cast_internal(alpha_new, CNML_CAST_FLOAT32_TO_FLOAT16);
beta_new = cnml_cast_internal(beta_new, CNML_CAST_FLOAT32_TO_FLOAT16);
gamma_new = cnml_cast_internal(gamma_new, CNML_CAST_FLOAT32_TO_FLOAT16);
}
return cnml_yolov5_detection_output_internal(alpha_new,
beta_new,
gamma_new,
anchor_data,
num_classes,
num_anchors,
img_height,
img_width,
conf_thres,
nms_thres,
maxBoxNum);
}
4）添加kernel
Wrapper中通过调用kernel实现算子功能。算子的具体实现主要通过调用CNML库来完成。以下是CNML库的简要逻辑。
kernel 实 现 就 是 按 照 上 述 编 程 逻 辑 调 用 CNML 库接 口 完 成 的，在 catch/torch_mlu/csrc/aten/operators/cnml/internal/cnml_internal.h 和 catch/torch_mlu/csrc/aten/operators/cnml/internal/yolov5_detection_output_internal.cpp 中添加 kernel 函数的声明和实现。
// cnml_internal.h
at::Tensor cnml_yolov5_detection_output_internal(const at::Tensor& alpha_data,
const at::Tensor& beta_data,
const at::Tensor& gamma_data,
torch::List<int64_t> anchor_data,
int64_t num_classes,
int64_t num_anchors,
int64_t img_height,
int64_t img_width,
double conf_thres,
double nms_thres,
int64_t maxBoxNum);
// yolov5_detection_output_internal.cpp
at::Tensor cnml_yolov5_detection_output_internal(const at::Tensor& alpha_data,
const at::Tensor& beta_data,
const at::Tensor& gamma_data,
torch::List<int64_t> anchor_data,
int64_t num_classes,
int64_t num_anchors,
int64_t img_height,
int64_t img_width,
double conf_thres,
double nms_thres,
int64_t maxBoxNum) {
int batch_size = alpha_data.size(0);
int inputNum = 3;
int output_num = 2;
int maskGroupNum = 3;
int classNum = num_classes;
int maxbox_num = maxBoxNum;
int net_w = img_width;
int net_h = img_height;
float confidence_thresh = static_cast<float>(conf_thres);
float nms_thresh = static_cast<float>(nms_thres);
cnmlTensor_t cnml_input_ptr[3];
cnmlTensor_t cnml_output_ptr[2];
// prepare input cnml tensor
auto* alpha_impl = getMluTensorImpl(alpha_data);
auto alpha_cnml = alpha_impl->CreateCnmlTensor(CNML_TENSOR,
toCnmlDataType(alpha_data.dtype()));
auto* beta_impl = getMluTensorImpl(beta_data);
auto beta_cnml = beta_impl->CreateCnmlTensor(CNML_TENSOR,
toCnmlDataType(beta_data.dtype()));
auto* gamma_impl = getMluTensorImpl(gamma_data);
auto gamma_cnml = gamma_impl->CreateCnmlTensor(CNML_TENSOR,
toCnmlDataType(gamma_data.dtype()));
// prepare input cnml tensor
auto output = at::empty({batch_size, maxbox_num * 7 + 64, 1, 1},
alpha_data.options());
auto* output_impl = getMluTensorImpl(output);
auto output_cnml = output_impl->CreateCnmlTensor(CNML_TENSOR,
toCnmlDataType(output.dtype()));
// prepare input cnml tensor for multicore
int buf_size = 1024 * (alpha_data.size(2) * alpha_data.size(3) +
beta_data.size(2) * beta_data.size(3) +
gamma_data.size(2) * gamma_data.size(3));
auto temp_buf = at::empty({batch_size, buf_size, 1, 1}, alpha_data.options());
auto* temp_buf_impl = getMluTensorImpl(temp_buf);
auto temp_buf_cnml = temp_buf_impl->CreateCnmlTensor(CNML_TENSOR,
toCnmlDataType(temp_buf.dtype()));
// End the execution flow if not MLU device
CHECK_MLU_DEVICE(output);
cnml_input_ptr[0] = alpha_cnml;
cnml_input_ptr[1] = beta_cnml;
cnml_input_ptr[2] = gamma_cnml;
cnml_output_ptr[0] = output_cnml;
cnml_output_ptr[1] = temp_buf_cnml;
// prepare h_arr
std::vector<int> h_arr_data(64, 1);
h_arr_data[0] = alpha_data.size(2);
h_arr_data[1] = beta_data.size(2);
h_arr_data[2] = gamma_data.size(2);
int h_data[] = {h_arr_data[0], h_arr_data[1], h_arr_data[2]};
// prepare w_arr
std::vector<int> w_arr_data(64, 1);
w_arr_data[0] = alpha_data.size(3);
w_arr_data[1] = beta_data.size(3);
w_arr_data[2] = gamma_data.size(3);
int w_data[] = {w_arr_data[0], w_arr_data[1], w_arr_data[2]};
// prepare bias_arr
std::vector<float> bias_arr_data(64, 1.0);
float bias_data[64];
for (int i = 0; i < num_anchors; i++) {
bias_arr_data[i] = (float)anchor_data[i];
bias_data[i] = bias_arr_data[i];
}
cnmlPluginYolov5DetectionOutputOpParam_t Yolov5params;
TORCH_CNML_CHECK(cnmlCreatePluginYolov5DetectionOutputOpParam(&Yolov5params,
batch_size,
inputNum,
classNum,
maskGroupNum,
maxbox_num,
net_w,
net_h,
confidence_thresh,
nms_thresh,
GET_CORE_VERSION,
w_data,
h_data,
bias_data));
cnmlBaseOp_t yolov5_op;
TORCH_CNML_CHECK(cnmlCreatePluginYolov5DetectionOutputOp(&yolov5_op,
Yolov5params,
cnml_input_ptr,
cnml_output_ptr));
// return to JIT if running mode is fuse
CHECK_RETURN_TO_FUSE(yolov5_op, output);
// get queue and func_param
cnrtInvokeFuncParam_t func_param;
static u32_t affinity = 0x01;
int data_parallelism = 1;
func_param.affinity = &affinity;
func_param.data_parallelism = &data_parallelism;
func_param.end = CNRT_PARAM_END;
auto queue = getCurQueue();
// compile all ops
TORCH_CNML_CHECK(cnmlCompileBaseOp(yolov5_op,
GET_CORE_VERSION,
GET_CORE_NUMBER));
void* input_addrs[3];
void* output_addrs[2];
input_addrs[0] = alpha_impl->raw_mutable_data();
input_addrs[1] = beta_impl->raw_mutable_data();
input_addrs[2] = gamma_impl->raw_mutable_data();
output_addrs[0] = output_impl->raw_mutable_data();
output_addrs[1] = temp_buf_impl->raw_mutable_data();
// compute operator
TORCH_CNML_CHECK(cnmlComputePluginYolov5DetectionOutputOpForward(yolov5_op,
input_addrs,
3,
output_addrs,
2,
&func_param,
queue));
syncQueue(queue);
TORCH_CNML_CHECK(cnmlDestroyPluginYolov5DetectionOutputOpParam(&Yolov5params));
TORCH_CNML_CHECK(cnmlDestroyBaseOp(&yolov5_op));
return output;
}
5)重新编译Cambricon-Pytorch
上述步骤操作完，重新编译Cambricon Pytorch, 进入python环境确认是否集成成功，如下：
Python 3.5.2 (default, Nov 12 2018, 13:43:14)
[GCC 5.4.0 20160609] on linux
Type "help", "copyright", "credits" or "license" for more information.
>>> import torch
>>> import torch_mlu
CNML: 7.7.0 a414883
CNRT: 4.6.0 e158c88
>>> torch.ops.torch_mlu.yolov5_detection_output
<built-in method yolov5_detection_output of PyCapsule object at 0x7f40636e7810>
为什么要量化：量化是将float32的模型转换为int8/int16的模型，可以保证计算精度在目标误差范围内的情况下，显著减少模型占用的存储空间和带宽，加速推理；比如int8模型是指将数值以有符号8位整型数据保存，并提供int8定点数的指数position和缩放因子scale，因此int8模型中每个8位整数i表示的实际值为：value=( i*2^position ) / scale。设备在进行在线推理和生成离线模型时仅支持输入量化后的模型。
操作步骤：对已有的yolov5s_1.3a.pt模型文件进行量化。示例参考yolov5_pytorch_demo/quantize_online目录中python quant.py脚本，量化后即在当前目录生成量化后的模型为yolov5s_int8.pt
参数含义：
其中qconfig_spec包括：{‘iteration’ : 1,‘use_avg’ : False,‘data_scale’ : 1.0,‘mean’ : [0,0,0],‘std’ : [1,1,1], ‘firstconv’ : True, ‘per_channel’ : False}
参考代码：
parser = argparse.ArgumentParser()
parser.add_argument('--cfg', type=str, default='yolov5s.yaml',help='model.yaml')
parser.add_argument('--device', default='cpu',help='cuda device, i.e. 0 or 0,1,2,3 or cpu')
opt = parser.parse_args()
# 获取yolov5网络文件
net = yolo.get_model(opt)
# 在这里设置firstconv参数为False,因为该模型首层为focus算子，非卷积，无法开启first_conv
qconfig={'iteration': 1, 'use_avg':False, 'data_scale':1.0, 'firstconv':False, 'per_channel': False}
# 调用量化接口
quantized_net = mlu_quantize.quantize_dynamic_mlu(net.float(),qconfig_spec=qconfig, dtype='int8', gen_quant=True)
# 设置为推理模式
quantized_net = quantized_net.eval().float()
# 读取图片做预处理
img_mat = Image.open("./images/image.jpg")
if img_mat.mode != 'RGB':
img_mat = img_mat.convert('RGB')
crop = 640
resize = 640
transform = transforms.Compose([
transforms.Resize(resize),
transforms.CenterCrop(crop),
transforms.ToTensor(),
])
img = transform(img_mat)
im_tensor = torch.unsqueeze(img, 0)
im_tensor = im_tensor.float()
# 推理生成量化值
quantized_net(im_tensor)
# 保存量化后的模型
torch.save(quantized_net.state_dict(), './yolov5s_int8.pt')
操作步骤：对步骤2生成的量化后的yolov5s_int8.pt进行在线推理测试。对图片进行推理，画出目标框和标注置信度。在yolov5_pytorch_demo/quantize_online目录中，示例如下：
1) 逐层模式 python detect.py
推理后的图片存储在./results目录下
推理过程：
融合模式：被融合的多个层作为单独的运算（单个 Kernel）在 MLU上运⾏。根据⽹络中的层是否可以被融合，⽹络被拆分为若⼲个⼦⽹络段。 MLU 与 CPU 间的数据拷⻉只在各个⼦⽹络之间发⽣。
逐层模式：逐层模式中，每层的操作都作为单独的运算（单个 Kernel）在 MLU 上运⾏，⽤⼾可以将每层结果导出到 CPU 上，⽅便⽤⼾进⾏调试。
一般来说，在线逐层模式更适用于调试环节，在线融合模式可以查看网络融合情况；
主要步骤：
1）设置：torch.set_grad_enabled(False) # 注意：在运行 MLU 推理融合模式时，这个条件是必须要设置的。
2）获取模型加载权重：
net = yolo.get_empty_model(opt) quantized_net = torch_mlu.core.mlu_quantize.quantize_dynamic_mlu(net) state_dict = torch.load('yolov5s_int8.pt') quantized_net.load_state_dict(state_dict, strict=False)
3）使用逐层模式进行推理： input_imgs = input_imgs.to(ct.mlu_device()) # 该命令将输入tensor转换到mlu上运行 detections = model(input_imgs) # 推理
4）如果要运行在线融合模式，需要在运行前向过程前调用jit.trace()接口生成静态图。首先会对整个网络运行一遍逐层模式，同时构建一个静态图；然后对静态图进行优化（包括去除冗余算子、小算子融、数据块复用等）得到一个优化后的静态图；之后会根据输入数据的设备类型进行基于设备的优化，生成针对当前设备的指令：
randn_input = torch.randn(batch_size, 3,img_height, img_width).float() model = torch.jit.trace(model, randn_input.to(ct.mlu_device()), check_trace = False)
5）最后根据推理结果为图片加框和标记置信度
参考代码：
parser = argparse.ArgumentParser()
parser.add_argument('--cfg', type=str, default='yolov5s.yaml',help='model.yaml')
parser.add_argument('--device', default='cpu',help='cuda device, i.e. 0 or 0,1,2,3 or cpu')
parser.add_argument('--jit',type=bool,help='fusion',default=False)
parser.add_argument('--save',type=bool,default=False,help='selection of save *.cambrcion')
opt = parser.parse_args()
# 获取yolov5网络并加载量化后的权重
net = yolo.get_empty_model(opt)
quantized_net = torch_mlu.core.mlu_quantize.quantize_dynamic_mlu(net)
state_dict = torch.load('yolov5s_int8.pt')
quantized_net.load_state_dict(state_dict, strict=False)
# 设置为推理模式
quantized_net = quantized_net.eval().float()
# 转移到MLU上进行推理
device = ct.mlu_device()
quantized_net.to(ct.mlu_device())
# 读取图片
img_mat = cv2.imread('images/image.jpg')
# 调用预处理函数做预处理
img = letter_box(img_mat)
# 设置融合模式，save选项表示是否生成离线模型，因为在进行在线融合推理时，可以生成离线模型
if opt.jit:
if opt.save:
ct.save_as_cambricon('yolov5s')
torch.set_grad_enabled(False)
ct.set_core_number(4)
trace_input = torch.randn(1, 3, 640, 640, dtype=torch.float)
trace_input=trace_input.to(ct.mlu_device())
quantized_net = torch.jit.trace(quantized_net, trace_input, check_trace = False)
# 推理
detect_out = quantized_net(img.to(ct.mlu_device()))
if opt.jit:
# 关闭生成离线模型
if opt.save:
ct.save_as_cambricon("")
detect_out=detect_out.to(torch.device('cpu'))
# 为原图添加框、检测类别和置信度
box_result = get_boxes(detect_out)
draw_boxes(box_result)
注意：原始的yolov5网络可以对输入的图片做自适应的预处理，使预处理后的图片可以有不同的大小。而在当前demo中，为了能够满足在线融合推理和离线推理必须是固定大小的要求，将预处理都改成了固定大小为640*640。
1）生成离线模型
操作步骤：
生成离线模型与在线推理代码相似，在yolov5_pytorch_demo/quantize_online目录中，示例如下：
python detect.py --jit True --save True
会在当前目录生成离线模型yolov5s.cambricon和离线模型信息文件yolov5s.cambricon_twins
2）离线推理
对一张图片进行离线推理，画出目标框和置信度。
示例图片放置在yolov5_pytorch_demo/offline/yolov5_offline_simple_demo/data目录下，离线模型放置在model目录下。执行make.sh在src目录下生成可执行文件，执行run.sh对一张图片进行推理，在result目录下生成推理后的图片。
参考： PyTorch框架的Yolov5移植教程 – 寒武纪开发者社区
参考：以矩阵乘为例的BANG C编程实验 – 寒武纪开发者社区