concat 算子开发零、前言理解 HPC 组同学在算子开发这块的相关工作，理解NPU内部的基础运算逻辑。在runt

零、前言

理解 HPC 组同学在算子开发这块的相关工作，理解NPU内部的基础运算逻辑。
在runtime 组同学和HPC 组同学沟通中，起到促进互相工作理解的目的。

算子开发的三个层次

基础：

对开发的那个特定的算子有一些基础的认识，知道算子具体的数学运算。
知道host 传递给 device 侧的参数在算子侧怎么参与的运算
能够根据算子(运算规则)和硬件(存储层级)的特点，进行拆分(tiling) 瓦片化
了解基础硬件加速特性，比如 bulk 模式。

中高级：

对硬件了如指掌，清楚的知道算子C语言转换为汇编后的每一条指令的 cycle .（CLA/FMA/SFU/Scalar/AIE）
对总线结构，数据通路的带宽有清晰的认识，明确带宽瓶颈。

目前的级别刚刚算是基础摸了一个门，对算子的理解也远不够深刻。只能说有幸参与。

先看demo

一个算子的开发分成两部分。因为NPU 属于协处理器.所以自然就区分为 host侧和 device 侧。

host 侧

device侧

三件事

准备要执行的 kernel 的指令
给kernel 传入参数，已经参数里面的数据要准备好
给kernel 传入一些配置，让kernel 知道该算多少。该用多少的资源算。
指令

module就加载了 kernel 指令，也就是 mida 语言 (类 C 语言) 通过 XNCC 编译出来的 ELF 二进制文件。

这个二进制文件，就是 bell 这个NPU 架构的 ISA 指令集的一些组合。实现一个具体的运算功能需要的指令。

参数

kernel 是需要参数的，每个kernle 都需要输入输出以及一些针对输入输出需要的一些参数配置等。

这个kernel 传入了 5个参数， size1, output0_dev, size2, input0_dev, input1_dev. 这里就对应了 device 侧代码需要的具体参数。

配置

launch 配置， 决定了这个kernel 需要多少的资源来参与运算。这个资源的多少和这个kernel 要完成的工作密切相关，也和实现kernel 的具体算法相关，总体思想就是并行度尽可能的高，mac 利用率高，带宽 利用率高

这里只暴露了两个参数，gridDim (1,1,1) blockDim(size,1,1)

先说 blockDim ：

是一个包含3个变量的结构体传入了 size, 1, 1 这个3个值，其中用到了x.传到了 device 就变成了

blockDim.x 这个值约束了 threadIdx.x 的范围。也就是说只能取值 0~(size-1）.

再说 gridDim ：

这个 gridDim是一个结构体，有 x ,y, z 三个值，x 就对应 device侧的一个内建变量 gridDim.x 这个值是一个确定的值，它约束了 blockIdx.x 的值。也就是说当 gridDim.x =1 时 blockIdx.x 只能取 0, 当 gridDim.x =10 ,blockIdx.x = 0 ~ 9 每个TB 取各自的编号。