设置kernel类型

函数功能

用于用户自定义设置kernel类型，控制算子执行时只启动该类型的核，避免启动不需要工作的核，缩短核启动开销。

函数原型

• 设置全局默认的kernel type, 对所有的tiling key生效。

  当前支持在自定义算子工程和kernel直调工程中使用，其中kernel直调工程仅支持使用该宏使能Vector Core。

  KERNEL_TASK_TYPE_DEFAULT(value)

• 设置某一个具体的tiling key对应的kernel type。

  当前仅支持在自定义算子工程中使用。

  KERNEL_TASK_TYPE(key, value)

参数说明

enum KernelMetaType {
    KERNEL_TYPE_AIV_ONLY,
    KERNEL_TYPE_AIC_ONLY,
    KERNEL_TYPE_MIX_AIV_1_0,
    KERNEL_TYPE_MIX_AIC_1_0,
    KERNEL_TYPE_MIX_AIC_1_1,
    KERNEL_TYPE_MIX_AIC_1_2,
    KERNEL_TYPE_AICORE,
    KERNEL_TYPE_VECTORCORE,
    KERNEL_TYPE_MIX_AICORE,
    KERNEL_TYPE_MIX_VECTOR_CORE,
    KERNEL_TYPE_MAX
};

支持的型号

Atlas A2训练系列产品/Atlas 800I A2推理产品

Atlas推理系列产品AI Core

约束说明

• KERNEL_TASK_TYPE优先级高于KERNEL_TASK_TYPE_DEFAULT，同时设置了全局kernel type和某一个tiling key的kernel type，该tiling key的kernel type以KERNEL_TASK_TYPE设置的为准。
• 没有设置全局默认kernel type的情况下，如果开发者只为其中的某几个tiling key设置kernel type，即部分tiling key没有设置kernel type，会导致算子kernel编译报错。
• 当设置具体的kernel task type时，用户的算子实现需要与kernel type相匹配。比如用户设置kernel type为KERNEL_TYPE_MIX_AIC_1_2, 则算子内部实现应与核配比AIC:AIV为1:2相对应；若用户设置kernel type为KERNEL_TYPE_AIC_ONLY， 则算子内部实现应该为纯cube逻辑，不应该存在vector部分的逻辑。其他的kernel type类似。
• 当纯cube或者纯vec算子强制设定kernel type为MIX类型时，workspace的大小不能设置为0，需要设置一个大于0的值（比如16、32等）。

调用示例

• 示例一：使能VectorCore样例
  1. 完成算子kernel侧开发时，需要通过本接口开启支持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);
     }

• 示例二：设置某一个具体的tiling key对应的kernel type。如下代码为伪代码 ，不可直接运行。

  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_AIV_ONLY); // 设置默认的kernel类型为纯AIV类型
      if (TILING_KEY_IS(1)) {
          KERNEL_TASK_TYPE(1, KERNEL_TYPE_MIX_AIV_1_0); // 设置tiling key=1对应的kernel类型为MIX AIV 1:0
          op.Process1();
      } else if (TILING_KEY_IS(2)) {
          KERNEL_TASK_TYPE(2, KERNEL_TYPE_AIV_ONLY); // 设置tiling key=2对应的kernel类型为纯AIV类型
          op.Process2();
      }
      // ...
  }
  配套的host侧tiling函数示例：
  ge::graphStatus TilingFunc(gert::TilingContext* context)
  {	
      // ...
      if (context->GetInputShape(0) > 10) {
          context->SetTilingKey(1);
      } else if (some condition) {
          context->SetTilingKey(2);
      }
  }