基于VectorCore编程
基于Ascend C开发的矢量计算相关的算子可以运行在Vector Core上,本节主要介绍如何基于Vector Core架构进行算子编程,该特性仅支持Atlas 推理系列产品。
学习本节内容之前,请先参考算子实现、Kernel直调算子开发、工程化算子开发学习基于AI Core的算子端到端的开发流程,下文内容会重点介绍基于VectorCore进行算子编程的差异点。具体差异点如下:
- 完成算子kernel侧开发时,需要通过宏KERNEL_TASK_TYPE_DEFAULT开启支持Vector Core,算子执行时会同时启动AI Core和Vector Core, 此时AI Core会当成Vector Core使用。如下的代码样例展示了使能Vector Core的方法:
extern "C" __global__ __aicore__ void add_custom(__gm__ uint8_t *x, __gm__ uint8_t *y, __gm__ uint8_t *z, __gm__ uint8_t *workspace, __gm__ uint8_t *tiling) { GET_TILING_DATA(tilingData, tiling); if (workspace == nullptr) { return; } GM_ADDR usr = AscendC::GetUserWorkspace(workspace); KernelAdd op; op.Init(x, y, z, tilingData.blockDim, tilingData.totalLength, tilingData.tileNum); KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_VECTOR_CORE); // 使能VectorCore if (TILING_KEY_IS(1)) { op.Process1(); } else if (TILING_KEY_IS(2)) { op.Process2(); } // ... }
- 完成host侧tiling开发时,设置的blockDim代表的是AI Core和Vector Core的总数,比如用户在host侧设置blockDim为10,则会启动总数为10的AI Core和Vector Core;为保证启动Vector Core,设置数值应大于AI Core的核数。您可以通过GetCoreNumAic接口获取AI Core的核数,GetCoreNumVector接口获取Vector Core的核数。 如下代码片段,分别为使用kernel直调工程和自定义算子工程时的设置样例,此处设置为AI Core和Vector Core的总和,表示所有AI Core和Vector Core都启动。
- kernel直调工程
auto ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance(); auto totalCoreNum = ascendcPlatform.GetCoreNumAic(); // ASCENDXXX请替换为实际的版本型号 if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCENDXXX) { totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector(); } ... kernel_name<<<totalCoreNum , l2ctrl, stream>>>(argument list);
- 自定义算子工程
// 配套的host侧tiling函数示例: ge::graphStatus TilingFunc(gert::TilingContext* context) { // 使能VectorCore,将blockDim置为AI Core中vector核数 + Vector Core中的vector核数 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo); auto totalCoreNum = ascendcPlatform.GetCoreNumAic(); // ASCENDXXX请替换为实际的版本型号 if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCENDXXX) { totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector(); } context->SetBlockDim(totalCoreNum); }
- kernel直调工程
父主题: 附录