编程艺术寒武纪 BANG C 异构编程方式

Posted 极智视界

tags:

篇首语:本文由小常识网(cha138.com)小编为大家整理,主要介绍了编程艺术寒武纪 BANG C 异构编程方式相关的知识,希望对你有一定的参考价值。

欢迎关注我的公众号 [极智视界],获取我的更多笔记分享

O_o>_<o_OO_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(&params);
cnrtKernelParamsBufferAddParam(params, &param1, 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 这些,下次我再写一篇关于寒武纪架构的文章说明一下。


  好了,收工~


 【公众号传送】

《【编程艺术】寒武纪 BANG C 异构编程方式》



扫描下方二维码即可关注我的微信公众号【极智视界】,获取更多AI经验分享,让我们用极致+极客的心态来迎接AI !

以上是关于编程艺术寒武纪 BANG C 异构编程方式的主要内容,如果未能解决你的问题,请参考以下文章

以编程方式将按钮添加到片段

JavaScript DOM编程艺术 - 读书笔记

JavaScript DOM编程艺术学习笔记

html 将以编程方式附加外部脚本文件的javascript代码片段,并按顺序排列。用于响应式网站,其中ma

使用 Pygments 检测代码片段的编程语言

oracle database 9i/10g/11g 编程艺术 源代码下载