基于VectorCore编程

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

学习本节内容之前，请先参考算子实现、Kernel直调算子开发、工程化算子开发学习基于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;
       }
       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();
       }
       // ...
   }

2. 完成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();
     if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCEND310P) {
        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();
         if (ascendcPlatform.GetSocVersion() == platform_ascendc::SocVersion::ASCEND310P) {
            totalCoreNum = totalCoreNum + ascendcPlatform.GetCoreNumVector();
         }
         context->SetBlockDim(totalCoreNum);
     }