在 Ascend C 算子开发中，Tiling 技术是解决 Local Memory 空间有限的核心方案，通过数据分块实现大规模张量计算的并行处理。本文从原理、实现流程、场景适配三个维度，结合代码案例详解 Tiling 技术的开发要点。

一、Tiling 技术核心原理

（一）技术背景

昇腾 AI 芯片的 Local Memory（片上内存）空间较小（如单 Core 约 512KB），而实际算子输入张量（如 [N, 1024, 1024] 的 float32 张量）占用内存远超 Local Memory 容量。Tiling 技术通过将输入数据按固定规则分块，逐块加载至 Local Memory 进行计算，最终合并结果，实现大规模数据的高效处理。

（二）核心概念

• 分块规则：由 Tiling 结构体定义，包含分块维度（如 H、W 维度分块）、分块数量、单块大小等参数。
• 数据调度：Host 侧负责分块逻辑处理，Kernel 侧接收分块信息，从 Global Memory 加载单块数据至 Local Memory 计算。
• 并行优化：分块后的数据可通过多 Core 并行处理，提升算子计算效率。

二、Tiling 技术实现流程（Host 侧 + Kernel 侧）

（一）Host 侧实现：分块逻辑定义

1. 定义 Tiling 结构体：描述分块规则，示例代码如下：

c

运行

// add_custom_tiling.h
struct AddCustomTiling {
    int32_t tile_h;  // H维度分块大小
    int32_t tile_w;  // W维度分块大小
    int32_t tile_num; // 总分块数量
};

1. 实现 Tiling 函数：根据输入张量 shape 推导分块参数，适配固定 / 动态 Shape 场景：

c

运行

// add_custom.cpp
Status AddCustomTilingFunc(const ge::Operator &op, AddCustomTiling &tiling) {
    // 获取输入张量shape
    auto input_shape = op.GetInputDesc(0).GetShape();
    int32_t h = input_shape.GetDim(2);
    int32_t w = input_shape.GetDim(3);
    
    // 固定Shape场景：直接指定分块大小
    tiling.tile_h = 32;
    tiling.tile_w = 32;
    
    // 动态Shape场景：根据输入shape动态计算分块数量
    tiling.tile_num = (h / tiling.tile_h) * (w / tiling.tile_w);
    return SUCCESS;
}

1. 算子注册：将 Tiling 函数与算子绑定，确保执行时自动调用分块逻辑。

（二）Kernel 侧实现：分块数据计算

Kernel 函数接收 Host 侧传递的 Tiling 信息，加载单块数据至 Local Memory 并执行计算：

c

运行

// add_custom_kernel.cpp
__global__ void add_custom(__gm__ const float *A, __gm__ const float *B, __gm__ float *C,
                           __gm__ AddCustomTiling tiling, int32_t idx) {
    // 根据分块索引获取当前块的偏移量
    int32_t tile_h_offset = (idx / (tiling.tile_w)) * tiling.tile_h;
    int32_t tile_w_offset = (idx % (tiling.tile_w)) * tiling.tile_w;
    
    // 加载单块数据至Local Memory
    __local__ float local_A[32][32];
    __local__ float local_B[32][32];
    ld_matrix(local_A, A + tile_h_offset * W + tile_w_offset, tiling.tile_h, tiling.tile_w);
    ld_matrix(local_B, B + tile_h_offset * W + tile_w_offset, tiling.tile_h, tiling.tile_w);
    
    // 并行计算
    int32_t local_h = threadIdx.x / 32;
    int32_t local_w = threadIdx.x % 32;
    local_A[local_h][local_w] += local_B[local_h][local_w];
    
    // 结果写回Global Memory
    st_matrix(C + tile_h_offset * W + tile_w_offset, local_A, tiling.tile_h, tiling.tile_w);
}

三、固定与动态 Shape 场景适配

（一）固定 Shape 场景

• 特点：输入张量 shape 固定（如 [N=1, C=3, H=512, W=512]），分块参数可硬编码。
• 实现要点：Tiling 结构体参数直接赋值，无需动态推导，代码简洁，计算效率高。
• 局限：仅适配固定输入规模，灵活性差。

（二）动态 Shape 场景

• 特点：输入张量 shape 可动态变化（如 H、W 维度支持 [256, 512, 1024]），通过 API 推导分块参数。
• 实现要点：调用GetInputDesc、GetShape等 API 获取输入 shape，动态计算分块大小与数量，示例代码如下：

c

运行

// 动态Shape的Tiling函数调整
Status AddCustomTilingFunc(const ge::Operator &op, AddCustomTiling &tiling) {
    auto input_desc = op.GetInputDesc(0);
    int32_t h = input_desc.GetShape().GetDim(2);
    int32_t w = input_desc.GetShape().GetDim(3);
    
    // 动态适配分块大小（确保单块不超过Local Memory容量）
    tiling.tile_h = (h > 64) ? 64 : h;
    tiling.tile_w = (w > 64) ? 64 : w;
    tiling.tile_num = (h + tiling.tile_h - 1) / tiling.tile_h * (w + tiling.tile_w - 1) / tiling.tile_w;
    return SUCCESS;
}

• 优势：适配多场景输入，通用性强，符合生产级算子需求。

四、实战避坑指南

1. 分块大小校验：单块数据占用内存需小于 Local Memory 容量（如 float32 类型的 32×32 矩阵占用 4KB，需预留部分空间给临时变量）。
2. 边界处理：当输入 shape 不能被分块大小整除时，需处理剩余数据块（如补充零值或单独计算）。
3. 性能优化：分块大小需结合芯片 Core 数量调整（如 64×64 分块适配多 Core 并行，提升计算吞吐量）。

   训练营简介：

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

   报名链接：

   https://www.hiascend.com/developer/activities/cann20252?tab=overview