企业官网建设流程全解析

设置Kernel类型

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit

产品支持情况

• Ascend 950PR/Ascend 950DT：支持

• Atlas A3 训练系列产品/Atlas A3 推理系列产品：支持

• Atlas A2 训练系列产品/Atlas A2 推理系列产品：支持

• Atlas 200I/500 A2 推理产品：不支持

• Atlas 推理系列产品AI Core：支持

• Atlas 推理系列产品Vector Core：不支持

• Atlas 训练系列产品：不支持

功能说明

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

函数原型

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

  当前支持在自定义算子工程和Kernel直调工程中使用。

  KERNEL_TASK_TYPE_DEFAULT(value)

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

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

  KERNEL_TASK_TYPE(key, value)

参数说明

表1参数说明

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 };

表2kernel type取值说明

约束说明

• kernel type取值约束

  • Ascend 950PR/Ascend 950DT，支持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。
  • Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持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。
  • Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持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。
  • Atlas 推理系列产品，支持KERNEL_TYPE_AICORE、KERNEL_TYPE_MIX_VECTOR_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等）。

• 使用Tiling模板编程时，需要通过ASCENDC_TPL_KERNEL_TYPE_SEL设置Kernel类型即可，无需再通过该接口进行设置，本接口不生效。

调用示例

• 示例一：启用VectorCore样例

  1. 完成算子kernel侧开发时，需要通过本接口启用Vector Core，算子执行时会同时启动AI Core和Vector Core，此时AI 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; } KernelAdd op; op.Init(x, y, z, tilingData.numBlocks, 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开发时，设置的numBlocks代表的是AI Core和Vector Core的总数，比如用户在host侧设置numBlocks为10，则会启动总数为10的AI Core和Vector Core；为保证启动Vector Core，设置数值应大于AI Core的核数。您可以通过GetCoreNumAic接口获取AI Core的核数，GetCoreNumVector接口获取Vector Core的核数。如下代码片段，展示了numBlocks的设置方法，此处设置为AI Core和Vector Core的总和，表示所有AI Core和Vector Core都启动。

     // 配套的host侧tiling函数示例： ge::graphStatus TilingFunc(gert::TilingContext* context) { // 启用VectorCore，将numBlocks置为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; } KernelAdd op; op.Init(x, y, z, tilingData.numBlocks, 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); } }

【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit