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);
考虑名为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的具体用法参考《CANN AscendCL应用软件开发指南(C&C++)》“AscendCL API参考” - 同步等待 - aclrtSynchronizeStream章节。
父主题: Kernel直调