NPU侧调用方式

NPU侧的运行验证当前分为三种方式，使用Kernel Launch接口调用，使用<<<>>>内核调用符调用或者使用Pybind调用。开发者完成算子核函数的开发和Tiling实现后，即可通过任意一种方式以及AscendCL运行时接口，完成算子的调用并实现自己的推理应用；同时提供简易的kernel开发工程，开发者仅需提供kernel侧实现，基于工程框架可以快速实现任意一种方式的调用。

Kernel Launch方式

当前版本，Kernel Launch开放式编程为试用特性，后续版本会存在变更，不支持应用于商用产品中。
当前版本暂不支持获取用户workspace特性。

ACLRT_LAUNCH_KERNEL调用接口的使用方法如下：

ACLRT_LAUNCH_KERNEL(kernel_name)(blockDim, stream, argument list);

kernel_name：算子核函数的名称。
blockDim：规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx。
stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在Device上执行。stream创建等管理接口请参考Stream管理。
argument list：参数列表，与核函数的参数列表保持一致。

考虑名为add_custom的核函数调用的例子，该函数实现两个矢量的相加，调用示例如下：

// blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数，该核函数的参数列表为x，y，z。
ACLRT_LAUNCH_KERNEL(add_custom)(8, stream, x, y, z)

内核调用符方式

核函数可以使用内核调用符<<<...>>>这种语法形式，来规定核函数的执行配置：

kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);

blockDim，规定了核函数将会在几个核上执行。每个执行该核函数的核会被分配一个逻辑ID，即block_idx，可以在核函数的实现中调用GetBlockIdx来获取block_idx；
blockDim是逻辑核的概念，取值范围为[1,65535]。为了充分利用硬件资源，一般设置为物理核的核数或其倍数。对于耦合架构和分离架构，blockDim在运行时的意义和设置规则有一些区别，具体说明如下：
- 耦合架构：由于其Vector、Cube单元是集成在一起的，blockDim用于设置启动多个AICore核实例执行，不区分Vector、Cube。AI Core的核数可以通过GetCoreNumAiv或者GetCoreNumAic获取。
- 分离架构
  - 针对仅包含Vector计算的算子，blockDim用于设置启动多少个Vector（AIV）实例执行，比如某款AI处理器上有40个Vector核，建议设置为40。
  - 针对仅包含Cube计算的算子，blockDim用于设置启动多少个Cube（AIC）实例执行，比如某款AI处理器上有20个Cube核，建议设置为20。
  - 针对Vector/Cube融合计算的算子，启动时，按照AIV和AIC组合启动，blockDim用于设置启动多少个组合执行，比如某款AI处理器上有40个Vector核和20个Cube核，一个组合是2个Vector核和1个Cube核，建议设置为20，此时会启动20个组合，即40个Vector核和20个Cube核。注意：该场景下，设置的blockDim逻辑核的核数不能超过物理核（2个Vector核和1个Cube核组合为1个物理核）的核数。
  - AIC/AIV的核数分别通过GetCoreNumAic和GetCoreNumAiv接口获取。
l2ctrl，保留参数，暂时设置为固定值nullptr，开发者无需关注；
stream，类型为aclrtStream，stream用于维护一些异步操作的执行顺序，确保按照应用程序中的代码调用顺序在device上执行。stream创建等管理接口请参考Stream管理。

考虑名为add_custom的核函数调用的例子，该函数实现两个矢量的相加，调用示例如下：

// blockDim设置为8表示在8个核上调用了add_custom核函数，每个核都会独立且并行地执行该核函数，该核函数的参数列表为x，y，z。
add_custom<<<8, nullptr, stream>>>(x, y, z);

内核调用符进行核函数调用的详细示例可参见核函数调用。

核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机端，可以调用以下aclrtSynchronizeStream函数来强制主机端程序等待所有核函数执行完毕。

aclError aclrtSynchronizeStream(aclrtStream stream);

aclrtSynchronizeStream的具体用法参考《AscendCL应用软件开发指南 (C&C++)》“AscendCL API参考” - 同步等待 - aclrtSynchronizeStream章节。

父主题： Kernel直调