编程艺术寒武纪 BANG C 异构编程方式
Posted 极智视界
tags:
篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了编程艺术寒武纪 BANG C 异构编程方式相关的知识,希望对你有一定的参考价值。
欢迎关注我的公众号 [极智视界],获取我的更多笔记分享
O_o
>_<
o_O
O_o
~_~
o_O
本文介绍一下寒武纪 BANG C 的编程方式。
相信接触过 CUDA C 的同学比较多, 但接触过 BANG C 的同学会少很多,这主要源于 GPU 的使用人群量是 MLU 不可比拟的,不过多学习一些异构编程方面的东西还是好的。这里分享一下寒武纪 BANG C 的编程方式,看起来和 CUDA C 会比较类似。我们知道一般的异构编程会分为 Host 端和 Device 端,在 Host 端做逻辑判断,在 Device 端做并行计算。在 BANG C 编程中也是一样,分为 Host 端和 MLU 端,下面从这两端编程方式分别进行介绍。
1、Host 端
Host 端调用 CNRT 接口,CNRT 是 MLU 的推理运行时库,需要包 cnrt.h 头:
#include "cnrt.h"
进行设备初始化:
cnrtInit(0);
cnrtDev_t dev;
cnrtGetDeviceHandle(&dev, 0);
cnrtSetCurrentDevice(dev);
创建 Queue:
cnrtQueue_t queue;
cnrtCreateQueue(&queue);
指定任务规模和调度类型:
cnrtDim3_t dim;
dim.x = 1;
dim.y = 1;
dim.z = 1;
cnrtFunctionType_t func_type = CNRT_FUNC_TYPE_BLOCK;
准备 Kernel 的参数:
cnrtKernelParamsBuffer_t params;
cnrtKernelInitParam_t init_param;
cnrtCreateKernelInitParam(&init_param);
cnrtInitKernelMemory(kernel, init_param);
...
uint32_t param1 = SIZE_INPUT;
cnrtGetKernelParamsBuffer(¶ms);
cnrtKernelParamsBufferAddParam(params, ¶m1, sizeof(uint32_t));
...
在 GDRAM 上为输入数据分配内存:
half *ptr_mlu_input;
cnrtMalloc((void **)&ptr_mlu_input, N * sizeof(half));
输入数据拷入 GDRAM:
cnrtMemcpy(&ptr_mlu_input prt_host_input, CNRT_MEM_TRANS_DIR_HOST2DEV);
启动 Kernel,并绑定到 Queue:
cnrtInvokeKernel_V3((void *)&kernel, init_param, dim, params, func_type, queue, NULL);
同步 Queue:
cnrtSyncQueue(queue);
数据拷出:
cnrtMemcpy(ptr_host_output, ptr_mlu_output, SIZE_OUTPUT, CNRT_MEM_TRANS_DIR_DEV2HOST);
资源释放:
cnrtDestroyQueue(queue);
cnrtDestoryKernelParamsBuffer(params);
cnrtDestroy();
2、MLU 端
MLU 端的代码就是基于 BANG C 来写的了,类似 CUDA C,是基于 C 语言的扩展,其文件名后缀为 .mlu。
MLU 端程序需要包含 mlu.h 头:
#inlcude "mlu.h"
再包含一些其他需要的头:
#include "macro.h"
#define s1 N1
#define s2 N2
BANG C Kernel 的入口函数需要用 __mlu_entry__
标记,类似于 CUDA C 中的 __global__
:
__mlu__entry__ void kernel(uint32_t size1, uint32_t size2, half* c, half* a, half* b)
{
/// __nram__ 是标记 NRAM 空间的变量
__nram__ half a_tmp[s1];
__nram__ half b_tmp[s2];
__nram__ half c_tmp[s3];
/// 将数据从 GRAM 拷贝到 NRAM 上
__memcpy(a_tmp, a, s1 * sizeof(half), GDRAM2NRAM);
__memcpy(b_tmp, b, s2 * sizeof(half), GDRAM2NRAM);
/// 将 NRAM 上数据进行循环 add 计算
__bang_cycle_add(c_tmp, b_tmp, a_tmp, size2, size1);
/// 将计算结果从 NRAM 拷贝到 GDRAM,HOST端可以从 RDRAM 拿数据
__memcpy(c, c_tmp, s2 * sizeof(half), NRAM2GDRAM);
}
以上分享了一下 BANG C 异构编程的方式,关于什么是 NRAM、GDRAM 这些,下次我再写一篇关于寒武纪架构的文章说明一下。
好了,收工~
【公众号传送】
扫描下方二维码即可关注我的微信公众号【极智视界】,获取更多AI经验分享,让我们用极致+极客的心态来迎接AI !
以上是关于编程艺术寒武纪 BANG C 异构编程方式的主要内容,如果未能解决你的问题,请参考以下文章