基于VectorCore编程

基于Ascend C开发的矢量计算相关的算子可以运行在Vector Core上，本节主要介绍如何基于Vector Core架构进行算子编程，该特性仅支持Atlas 推理系列产品。

学习本节内容之前，请先参考算子调用了解基于AI Core的算子端到端的开发流程，下文内容会重点介绍基于VectorCore进行算子编程的差异点。

具体差异点如下：

1. 完成算子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;
       }
       SetSysWorkspace(workspace);
       GM_ADDR usr = 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();
       }
       // ...
   }

2. 完成算子host侧Tiling开发时，设置的block_dim代表的是AI Core和Vector Core的总数，比如用户在host侧设置blockdim为10，则会启动总数为10的AI Core和Vector Core；为保证启动Vector Core，设置数值应大于AI Core的核数。您可以通过GetCoreNumAic接口获取AI Core的核数，GetCoreNumVector接口获取Vector Core的核数。 如下代码片段，展示了block_dim的设置方法，此处设置为AI Core和Vector Core的总和，表示所有AI Core和Vector Core都启动。

   // 配套的host侧tiling函数示例：
   ge::graphStatus TilingFunc(gert::TilingContext* context)
   {	
       // 使能VectorCore，将block_dim置为AI Core中vector核数 + Vector Core中的vector核数
       auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
       auto totalCoreNum = ascendcPlatform.GetCoreNumAiv();
       // ASCENDXXX请替换为实际的版本型号
       if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCENDXXX) {
          totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector();
       }
       context->SetBlockDim(totalCoreNum);
   }