设置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参数说明参数输入/输出说明key输入tiling key的key值此参数是正数表示某个核函数的分支。value输入设置的kernel类型可选值范围kernel类型具体说明请参考表2。不同硬件架构支持的参数取值不同具体支持的参数取值请参考kernel type取值约束。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_AIV_ONLY算子执行时仅启动AI Core上的Vector核比如用户在host侧设置numBlocks为10则会启动10个Vector核。KERNEL_TYPE_AIC_ONLY算子执行时仅启动AI Core上的Cube核比如用户在host侧设置numBlocks为10则会启动10个Cube核。KERNEL_TYPE_MIX_AIV_1_0AIC、AIV混合场景下使用了多核控制相关指令时设置核函数的类型为MIX AIV:AIC 1:0带有硬同步算子执行时仅会启动AI Core上的Vector核比如用户在host侧设置numBlocks为10则会启动10个Vector核。硬同步的概念解释如下当不同核之间操作同一块全局内存且可能存在读后写、写后读以及写后写等数据依赖问题时通过调用SyncAll()函数来插入同步语句来避免上述数据依赖时可能出现的数据读写错误问题。目前多核同步分为硬同步和软同步硬同步是利用硬件自带的全核同步指令由硬件保证多核同步。KERNEL_TYPE_MIX_AIC_1_0AIC、AIV混合场景下使用了多核控制相关指令时设置核函数的类型为MIX AIC:AIV 1:0带有硬同步算子执行时仅会启动AI Core上的Cube核比如用户在host侧设置numBlocks为10则会启动10个Cube核。KERNEL_TYPE_MIX_AIC_1_1AIC、AIV混合场景下设置核函数的类型为MIX AIC:AIV 1:1算子执行时会同时启动AI Core上的Cube核和Vector核比如用户在host侧设置numBlocks为10则会启动10个Cube核和10个Vector核。KERNEL_TYPE_MIX_AIC_1_2AIC、AIV混合场景下设置核函数的类型为MIX AIC:AIV 1:2算子执行时会同时启动AI Core上的Cube核和Vector核比如用户在host侧设置numBlocks为10则会启动10个Cube核和20个Vector核。KERNEL_TYPE_AICORE算子执行时仅会启动AI Core比如用户在host侧设置numBlocks为5则会启动5个AI Core。KERNEL_TYPE_VECTORCORE该参数为预留参数当前版本暂不支持。KERNEL_TYPE_MIX_AICORE该参数为预留参数当前版本暂不支持。KERNEL_TYPE_MIX_VECTOR_CORE基于Ascend C开发的矢量计算相关的算子可以运行在Vector Core上调用本接口传入该参数用于启用Vector Core。启用Vector Core后算子执行时会同时启动AI Core和Vector Core用于并行计算。比如用户在host侧设置numBlocks为10则会启动总数为10的AI Core和Vector Core。需要注意的是通过SetBlockDim设置核数时需要大于AI Core的核数否则不会启动VectorCore。约束说明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样例完成算子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(); } // ... }完成算子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 key1对应的kernel类型为MIX AIV 1:0 op.Process1(); } else if (TILING_KEY_IS(2)) { KERNEL_TASK_TYPE(2, KERNEL_TYPE_AIV_ONLY); // 设置tiling key2对应的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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考