Cambricon BANG C编程入门

rebekk

于 2024-07-21 22:44:46 发布

阅读量1k

点赞数 20

CC 4.0 BY-SA版权

分类专栏：寒武纪文章标签：开发语言寒武纪 BANG C

本文链接：https://blog.csdn.net/qq_43527718/article/details/140589185

寒武纪专栏收录该内容

3 篇文章

订阅专栏

Cambricon BANG异构并行编程模型利用CPU和MLU协同计算。CPU作为主机侧的控制设备，用于完成复杂的控制和任务调度；设备侧的MLU则用于大规模并行计算和领域相关的计算任务。

使用Cambricon BANG C编写程序时，需要同时编写主机侧和设备侧的代码。其中，主机侧程序通过调用CNRT（Cambricon Runtime Library，寒武纪运行时库）或者CNDrv（Cambricon Driver API，寒武纪软件栈驱动接口）接口来完成设备初始化、设备内存管理、主机端与设备端的数据拷贝、启动Kernel、释放设备资源等工作；设备端程序由多个Kernel构成。不同的Kernel之间可以并行执行，也可以串行执行，用户可以调用CNRT或者CNDrv接口进行控制。

在具体编程过程中，用户需要先将一个完整的计算任务拆解为一系列可以并行的Task，所有的Task构成一个三维网格。这个三维网格的维度信息由用户做任务拆分时确定。在由Task构成的三维网格中，每个Task都有唯一的坐标。每个任务除了一个三维坐标外，还有一个全局唯一的线性ID。在实际执行时，每个Task会映射到一个物理MLU Core上执行。MLU Core在执行一个Task的过程中不会发生切换，只有一个Task执行完毕，另一个Task才能开始执行。

为了描述上述三维网格式的任务规模，Cambricon BANG C编程语言中引入了cnrtDim3_t数据类型，如下面的配置所构建的三维任务网格如下图所示：

cnrtDim3_t dim;
dim.x = 8;
dim.y = 8;
dim.z = 4;

请添加图片描述

在Cambricon BANG异构并行编程模型中，一个Kernel描述了一个Task的行为。在MLU上执行的程序称为Kernel，每个task都执行一次对应的Kernel函数，在MLU上可以同时执行多个并行的Kernel。

结合代码来看

Kernel：
- 设备侧的Kernel是一个带有__mlu_entry__属性的函数，该函数描述一个Task需要执行的所有操作。在Kernel内部还可以通过taskId等内建变量获得每个Task唯一的ID，从而实现不同Task的差异化处理。此外，类似的内建变量还包括clusterId、taskIdX、taskIdY等
- 在主机侧使用<<<...>>>这种语法来启动一个Kernel，如下所示，<<<dim, ktype, pQueue>>>中的dim表示任务规模、pQueue表示该Kernel将会放到哪个任务队列中执行、ktype表示任务类型，即Kernel执行需要的硬件资源数量。在主机侧使用该语法糖启动的Kernel会异步执行（即主机侧不需要等待Kernel执行完毕即可继续执行后续的代码），Cambricon BANG异构并行计算平台会将对应的Kernel插入对应的执行队列中，并在设备侧有资源空闲时调度Kernel到硬件上执行
```
#include "bang.h"

int main(){
	...
	Kernel<<<dim, ktype, pQueue>>>(mlu_result, mlu_source1, mlu_source2);
	...
}
```
任务规模
Cambricon BANG C语言为用户提供了一系列内置变量来显式并行编程。其中，与任务规模相关的内置变量包括：taskDim、taskDimX、taskDimY、taskDimZ
- taskDimX、taskDimY、taskDimZ分别对应任务规模的三个维度：dim.x、dim.y、dim.z
- taskIdX、taskIdY、taskIdZ的取值范围为[0, taskDimX-1、[0, taskDimY-1、[0, taskDimZ-1
- taskDim等于taskDimX、taskDimY、taskDimZ三者的乘积
- taskId = taskIdZ * taskDimY * taskDimX + taskIdY * taskDimX + taskIdX。taskId的取值范围是[0, taskDim-1]
Notifier
Notifier可以帮助用户对Queue中的任务执行进行管理与检查。Notifier可以记录当前Queue还未被执行的任务，或是作为一个等待任务被放置到队列中执行。相比计算任务Notifier不执行实际的硬件操作。当Notifier记录了某个Queue未被执行的任务状态后，可以在另一个Queue中等待该Queue的未被执行的任务完成。Notifier会在记录的Queue任务都完成后记录一个时间，可通过两个Notifier来实现对计算任务耗时的统计
BANG C程序执行流程：
- 通过CNRT接口选择硬件设备
- 在主机侧准备输入数据，并为输出数据分配空间
- 在主机侧调用CNRT接口分配设备内存，并将输入数据拷贝到设备内存
- 设置Kernel的任务规模（使用CNRT定义的cnrtDim3_t数据类型）
- 设置Kernel的任务类型：CNRT定义了cnrtFunctionType_t数据类型来设置任务类型，其值可以是Block或UnionN
- 通过CNRT接口创建任务队列
- 向任务队列添加Kernel
- 调用CNRT接口等待任务队列执行完成
- 调用CNRT接口将计算结果拷贝至主机侧
- 释放主机侧和设备侧的各类资源。这些资源主要包括任务队列、设备侧内存、主机侧内存等

一个BANG C代码示例：
BANG C程序的文件后缀是*.mlu，BANG C异构程序必须包含头文件bang.h，该头文件包含了混合编程必需的数据类型的定义以及函数接口声明。下面是一个简单的向量加法示例

#include <bang.h>

#define EPS 1e-7
#define LEN 1024

__mlu_entry__ void Kernel(float* dst, float* source1, float* source2){
	__nram__ float dest[LEN];
	__nram__ float src1[LEN];
	__nram__ float src2[LEN];
	__memcpy(src1, source1, LEN*sizeof(float), GDRAM2NRAM);
	__memcpy(src2, source2, LEN*sizeof(float), GDRAM2NRAM);
	__bang_add(dest, src1, src2, LEN);
	__memcpy(dst, dest, LEN*sizeof(float), NRAM2GDRAM);
}

int main(void){
	cnrtQueue_t queue;
	// 通过CNRT接口选择硬件设备
	CNRT_CHECK(cnrtSetDevice(0));
	
	// 通过CNRT接口创建任务队列
	CNRT_CHECK(cnrtQueueCreate(&queue));

	// 设置Kernel的任务规模（使用CNRT定义的cnrtDim3_t数据类型）
	cnrtDim3_t dim = {1, 1, 1};
	// 设置Kernel的任务类型（Block或UnionN）
	cnrtFunctionType_t ktype = CNRT_FUNC_TYPE_BLOCK;

	// 创建Notifier
	cnrtNotifier_t start, end;
	CNRT_CHECK(cnrtNotifierCreate(&start));
	CNRT_CHECK(cnrtNotifierCreate(&end));
	
	// 在主机侧准备输入数据，并为输出数据分配空间
	float* host_dst = (float*)malloc(LEN*sizeof(float));
	float* host_src1 = (float*)malloc(LEN*sizeof(float));
	float* host_src2 = (float*)malloc(LEN*sizeof(float));
	for (int i=0; i<LEN; i++){
		host_src1[i] = i;
		host_src2[i] = i;
	}

	// 在主机侧调用CNRT接口分配设备内存，并将输入数据拷贝到设备内存
	float* mlu_dst;
	float* mlu_src1;
	float* mlu_src2;
	CNRT_CHECK(cnrtMalloc((void**)&mlu_dst, LEN*sizeof(float)));
	CNRT_CHECK(cnrtMalloc((void**)&mlu_src1, LEN*sizeof(float)));
	CNRT_CHECK(cnrtMalloc((void**)&mlu_src2, LEN*sizeof(float)));
	CNRT_CHECK(cnrtMemcpy(mlu_src1, host_src1, LEN*sizeof(float), cnrtMemcpyHostTodev));
	CNRT_CHECK(cnrtMemcpy(mlu_src2, host_src2, LEN*sizeof(float), cnrtMemcpyHostTodev));

	CNRT_CHECK(cnrtPlaceNotifier(start, queue));
	// 向任务队列添加Kernel
	Kernel<<<dim, ktype, queue>>>(mlu_dst, mlu_src1, mlu_src2);
	CNRT_CHECK(cnrtPlaceNotifier(end, queue));

	// 调用CNRT接口等待任务队列执行完成。cnrtQueueSync会等待参数queue的所有操作均执行完成。当Queue中的操作发生异常后，该Queue会立即返回并返回失败；并且此时Queue不再能够下发任务。
	// 然后调用CNRT接口将计算结果从设备侧拷贝至主机侧
	cnrtQueueSync(queue);
	CNRT_CHECK(cnrtMemcpy(host_dst, mlu_dst, LEN*sizeof(float), cnrtMemcpyDevToHost));

	// 因为执行的是加法，所以应该的结果是2i，这里的作用就是判断一下host_dst的结果是否是正确的，具体做法就是查看和正确值相减后的结果是否小于1e-7（因为是浮点数必不可能相等的？）
	for (int i=0; i<LEN; i++){
		if (fabsf(host_dst[i]-2*i)>EPS){
			printf("%f expected, but %f got!\n", (float)(2*i), host_dst[i]);  
		}
	}

	float timeTotal;
	CNRT_CHECK(cnrtNotifierDuration(start, end, &timeTotal));
	printf("Total time: %.3f ms\n", timeTotal/1000.0);

	// 释放主机侧和设备侧的各类资源（任务队列、设备侧内存、主机侧内存）
	CNRT_CHECK(cnrtQueueDestroy(queue));
	cnrtFree(mlu_dst);
	cnrtFree(mlu_src1);
	cnrtFree(mlu_src2);
	free(host_dst);
	free(host_src1);
	free(host_src2);

	return 0;