下载
中文
注册
我要评分
文档获取效率
文档正确性
内容完整性
文档易理解
在线提单
论坛求助
昇腾小AI

基于VectorCore编程

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

学习本节内容之前,请先参考算子实现Kernel直调算子开发工程化算子开发学习基于AI Core的算子端到端的开发流程,下文内容会重点介绍基于VectorCore进行算子编程的差异点。具体差异点如下:

  1. 完成算子kernel侧开发时,需要通过宏KERNEL_TASK_TYPE_DEFAULT开启支持Vector Core,算子执行时会同时启动AI CoreVector 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 CoreVector Core的总数,比如用户在host侧设置blockDim为10,则会启动总数为10的AI CoreVector Core;为保证启动Vector Core,设置数值应大于AI Core的核数。您可以通过GetCoreNumAic接口获取AI Core的核数,GetCoreNumVector接口获取Vector Core的核数。 如下代码片段,分别为使用kernel直调工程和自定义算子工程时的设置样例,此处设置为AI CoreVector Core的总和,表示所有AI CoreVector 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);
      }
  • 请参考Ascend C API中具体API支持的型号,来判断API接口是否支持Atlas推理系列产品Vector Core
  • 支持VectorCore后,因为AICore和VectorCore会分别执行,通过不同的任务进行调度,所以不支持核间同步指令,如IBSet、IBWait、SyncAll等。
  • 算子计算溢出(输入inf/nan或计算结果超出范围)时,需注意AICore和VectorCore结果表现不一致,AICore仅支持饱和模式,VectorCore仅支持inf/nan模式。
搜索结果
找到“0”个结果

当前产品无相关内容

未找到相关内容,请尝试其他搜索词