训练营简介
2025年昇腾CANN训练营第二季，基于CANN开源开放全场景，推出0基础入门系列、码力全开特辑、开发者案例等专题课程，助力不同阶段开发者快速提升算子开发技能。获得Ascend C算子中级认证，即可领取精美证书，完成社区任务更有机会赢取华为手机、平板、开发板等大奖。

在参加训练营的过程中，我深入探索了CANN CV算子架构的奥秘。CANN作为昇腾AI的核心基础软件平台，搭起了AI框架与昇腾硬件的桥梁，始终以 “使能每一位创新者” 为目标。本文将分享我在学习CANN CV算子架构过程中的理解、实战经验以及遇到的问题解决方案。
 

Ascend C算子开发能力认证是昇腾社区为开发者提供的重要技能认证，其中Tiling计算是核心考核内容之一。本文将深入解析Tiling的基本概念，并详细介绍动态shape场景下Tiling的实现方法，帮助开发者系统掌握这一关键技术。

一、Tiling基本概念：为何需要数据分块

在Ascend C算子开发中，Tiling是一个基础而重要的概念。大多数情况下，NPU中AI Core的Local Memory存储空间有限，无法完全容纳算子的输入与输出的所有数据。这就需要采用分块处理的策略：每次只搬运一部分输入数据进行计算，然后将结果搬出，再搬运下一部分输入数据，重复这个过程直到得到完整结果。

这种数据切分和分块计算的过程就称为Tiling（数据分块）。其中涉及几个关键术语：

• Tiling块：每次搬运的那一部分数据块

• Tiling算法/策略：根据算子中不同输入形状确定搬入基本块大小的相关算法

• Tiling函数：算子中实现Tiling算法的函数，一般定义在host侧的tiling头文件中

从硬件层面看，Ascend AI处理器内部有Global Memory和Local Memory两级存储。Local Memory靠近计算单元，带宽高但容量有限；Global Memory容量大但带宽较低。因此，合理的Tiling策略能最大程度利用Local Memory的高带宽特性，提升计算效率。

二、Tiling的实现方式：固定shape与动态shape对比

在实际开发中，Tiling的实现根据shape是否可变分为两种场景：固定shape和动态shape，它们有显著的差异。

1. 固定shape场景

固定shape场景下，输入大小在编译时就已经确定。这种情况下，每次搬运的数据量以及总共搬运次数都可以在编译时直接计算出来。

优势：

• 实现难度低，只需考虑固定shape的逻辑处理

• 优化难度低

劣势：

• 灵活性差，不同shape需要重新编译算子，会产生大量的算子二进制文件

• 无法适应实际应用中输入尺寸多变的场景

固定shape的核函数实现通常直接使用常量定义数据尺寸：

#include "add_custom_unalign_tiling.h"
#include "register/op_def_registry.h"

namespace optiling {
constexpr uint32_t BLOCK_DIM = 8;
constexpr uint32_t SIZE_OF_HALF = 2;
constexpr uint32_t BLOCK_SIZE = 32;
// shape需要对齐到的最小单位
constexpr uint32_t ALIGN_NUM = BLOCK_SIZE / SIZE_OF_HALF;

2. 动态shape场景

动态shape场景下，算子的形状可以通过核函数的入参传入核函数内部，参与内部逻辑计算。这种场景实现难度较高，需要考虑不同逻辑分支处理，优化难度也相应增加。

优势：

• 灵活性高，能适应不同shape的使用场景

• 一个算子能处理多种尺寸的输入，减少算子二进制文件数量

劣势：

• 实现复杂度高，需要考虑各种边界情况

• 优化难度大，需要保证不同shape下的性能表现

动态shape的核函数实现较为复杂：

#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
{
    GET_TILING_DATA(tilingData, tiling);
    KernelAdd op;
    op.Init(x, y, z, tilingData.totalLength, tilingData.tileNum);
    if (TILING_KEY_IS(1)) {
        op.Process();
    }
}

三、动态shape场景Tiling实现详解

1. Tiling结构体设计

在Ascend C中，Tiling的直接表示形式是结构体（struct），简称Tiling结构体。这个结构体定义了如何对输入数据进行切分，以及决定了计算过程的细节。

动态shape场景的Tiling结构体通常包含以下关键信息：

• TOTAL_LENGTH：总共需要计算的数据个数

• TILE_NUM：每个核上计算数据分块的个数

• 其他与数据切分相关的参数

Tiling结构体在host侧实例化，并通过指针传入kernel函数中。这种方式使得host侧可以根据实际的shape信息动态计算Tiling参数，然后将这些参数传递给device侧使用。

2. Tiling数据的传递过程

Tiling数据从host侧到device侧的传递需要以下几个步骤：

1. 在host侧为Tiling结构体申请空间：aclrtMallocHost((void**)(&tilingHost), tilingSize)

2. 在device侧为Tiling结构体申请空间：aclrtMalloc((void**)&tilingDevice, tilingSize, ACL_MEM_MALLOC_HUGE_FIRST)

3. 执行内存拷贝，将Tiling结构体从host侧搬运到device侧：aclrtMemcpy(tilingDevice, tilingSize, tilingHost, tilingSize, ACL_MEMCPY_HOST_TO_DEVICE)

3. 核函数中的Tiling解析

在核函数中，需要通过特定的宏函数来解析Tiling数据：

• 使用GET_TILING_DATA宏函数获取tiling结构体

• 对于CPU模式和NPU模式之间的差异，使用CONVERT_TILING_DATA宏函数将__ubuf_uint8_t*转化为__ubuf__tilingstruct*

• 使用INIT_TILING_DATA宏函数区分tiling_data在不同的初始化过程

四、多核并行计算与Tiling策略

在实际的算子开发中，通常需要实现多核并行计算以提高算子的执行效率。这意味着需要对输入数据进行切分，将不同的数据块分配到不同的核上处理。

数据切分原则：

• 将长度为TOTAL_LENGTH的算子输入分配到多个核上进行计算

• 每个核上计算的数据长度为BLOCK_LENGTH

• 对于每个核的计算数据，基于Local Memory的大小进一步切分

• 切分数据块的个数为TILE_NUM，每个数据块的长度为TILE_LENGTH

根据数据切分的均衡性，Tiling策略可以分为四种场景：

1. 核间均分，核内均分：每个核处理的数据量相同，核内每个数据块的数据量相同

2. 核间均分，核内不均分：每个核处理的数据量相同，核内各数据块的数据量不完全相同

3. 核间不均分，核内均分：每个核处理的数据量不同，核内每个数据块的数据量相同

4. 核间不均分，核内不均分：每个核处理的数据量不同，核内各数据块的数据量不完全相同

五、动态shape Tiling的实现要点

1. 固定shape算子改装为动态shape

将现有的固定shape算子改装为动态shape算子，主要改动在于将控制形状的变量（如BLOCK_DIM、TOTAL_LENGTH、TILE_NUM）的来源从常量改为依靠外界输入。在核函数中需要额外传入一个tiling参数，它包含控制核函数逻辑处理的关键变量。

2. 硬件限制与对齐原则

由于硬件限制，在对输入数据进行切分时应遵循以下原则：

• Unified Buffer上的数据存储空间必须保持32字节对齐

• 输入数据不满足32字节对齐时，需要取输入数据长度向上对齐到32字节的长度作为输入数据总长度

• 进行Tiling有关计算时，以32字节为最小单位进行计算

• 尽可能最大利用Unified Buffer空间，减少从Global Memory上搬运数据的次数

• 均衡利用多核计算能力，将计算均衡分配到多个AI Core上

3. 动态shape算子调用流程

动态shape场景下的算子调用流程与固定shape有所不同：

1. 进行算子实现，交付件仅包括算子代码实现文件（.py）

2. 为了提高性能，可以针对不同的shape范围，制定不同的Tiling策略，形成多个算子实现文件

3. 将算子实现文件直接编译成二进制文件（.o）

4. 自定义算子选择器，在算子执行时针对传入的shape，选择使用合适的Tiling策略

5. 用户App通过AscendCL接口，根据算子选择器选择的Tiling策略执行对应的算子二进制文件

六、学习建议与认证准备

对于准备参加Ascend C算子开发能力认证的开发者，建议如下：

1. 理论学习重点

• 深入理解Tiling的基本概念和实现原理

• 掌握固定shape与动态shape的差异和适用场景

• 学习多核并行计算原理和Tiling策略

2. 实践技能培养

• 亲自动手实现简单的固定shape算子

• 尝试将固定shape算子改装为动态shape算子

• 练习使用Ascend C的调试工具进行孪生调试

3. 认证考试重点
根据官方认证大纲，Tiling计算部分主要考核：

• Tiling的基本概念

• 动态shape场景Tiling的实现方法

• 固定与动态shape场景的差异

• 如何构建及调用动态shape的自定义算子核函数调用工程

结语

Tiling计算是Ascend C算子开发的核心技术，掌握动态shape场景下的Tiling实现方法，对于开发高性能、高灵活性的AI算子至关重要。通过系统学习Tiling的基本概念、实现原理和多核并行策略，开发者能够为后续的算子开发工作打下坚实基础，顺利通过Ascend C算子开发能力认证。

值得一提的是，2025年昇腾CANN训练营全新升级，定制化四大专题直播课程，帮助不同阶段的开发者快速入门和提升Ascend C算子开发技术。完成Ascend C算子中级认证和实操挑战，不仅可以领取结业证书，还有机会赢得开发板等大奖，是学习和认证的优质途径。

报名链接：https://www.hiascend.com/developer/activities/cann20252#cann-camp-2502-intro