一、动态 Shape 算子核心定义与应用场景

（一）核心定义

动态 Shape 算子指输入张量的维度（如 N、C、H、W）可在一定范围内灵活变化（如 N 支持 1~1024、H 支持 256~2048），算子通过 Ascend C 提供的 API 动态推导输入输出 Shape、分块参数及资源配置，无需硬编码固定维度。相较于固定 Shape 算子，动态 Shape 算子的核心优势在于通用性强、框架兼容性好（支持 PyTorch/TensorFlow 动态组网）、开发效率高（一次开发覆盖多场景）。

（二）典型应用场景

• 动态 Batch 场景：训练或推理过程中，Batch Size 根据硬件资源或业务需求动态调整（如 16、32、64）。
• 多分辨率适配场景：计算机视觉任务中，输入图像分辨率不固定（如 256×256、512×512、1024×1024）。
• 弹性组网场景：分布式训练中，不同节点处理的数据分片维度可能不同，需算子自适应适配。

二、动态 Shape 算子开发关键技术

（一）Shape 动态推导

通过 Ascend C 的GetShape、GetDim、GetDimNum等 API，从输入张量描述（TensorDesc）中提取维度信息，动态确定输入输出 Shape 及分块规则。核心 API 如下：

API 接口	功能说明	示例用法
TensorDesc::GetShape	获取张量的 Shape 对象	auto shape = input_desc.GetShape();
Shape::GetDim	获取指定索引的维度大小（索引从 0 开始）	int32_t h = shape.GetDim(2);（获取 H 维度大小）
Shape::GetDimNum	获取张量的维度数量	int32_t dim_num = shape.GetDimNum();（如 4 维张量返回 4）
TensorDesc::GetDataType	获取张量的数据类型	auto dtype = input_desc.GetDataType();

（二）资源动态配置

根据动态推导的 Shape 信息，灵活分配 Local Memory 缓冲区、线程配置、分块参数等资源，避免固定资源配置导致的内存浪费或溢出。核心要点包括：

1. Local Memory 缓冲区动态分配：根据分块大小计算所需内存，通过new或AllocTensor动态申请，避免硬编码固定大小。
2. 分块参数动态计算：结合 Local Memory 容量、输入维度大小，动态推导分块维度（tile_h/tile_w）与分块数量（tile_num_h/tile_num_w）。
3. 边界分块自适应处理：当输入维度不能被分块大小整除时，动态计算边界分块的实际数据量，避免越界访问或数据丢失。

（三）Tiling 技术动态适配

Tiling 技术是动态 Shape 算子开发的核心，通过将大规模张量分块处理，适配 Local Memory 有限容量。动态 Shape 场景下的 Tiling 设计需满足：

1. 分块大小自适应：根据输入维度与 Local Memory 容量动态调整分块大小（如小 Shape 用 32×32 分块，大 Shape 用 64×64 分块）。
2. 分块数量动态计算：通过(dim + tile_dim - 1) / tile_dim（向上取整）计算分块数量，确保覆盖全部数据。
3. 分块参数传递：Host 侧通过SetTilingData将动态计算的 Tiling 参数传递给 Kernel 侧，Kernel 侧通过__gm__修饰符接收。

三、ReLU 动态 Shape 算子开发实战

以 ReLU 算子（数学公式：f(x) = max(0, x)）为例，详细讲解动态 Shape 算子的完整开发流程，输入张量支持 Shape[N, C, H, W]动态变化，数据类型为float16，存储格式为NHWC。

（一）Step 1：算子原型定义（支持动态维度）

在relu_dynamic.json原型文件中，将维度定义为动态参数（N、C、H、W），不指定固定值：

{
    "opName": "ReluDynamic",
    "inputDesc": [
        {
            "name": "x",
            "dtype": "float16",
            "format": "NHWC",
            "shape": ["N", "C", "H", "W"]
        }
    ],
    "outputDesc": [
        {
            "name": "y",
            "dtype": "float16",
            "format": "NHWC",
            "shape": ["N", "C", "H", "W"]
        }
    ]
}

（二）Step 2：Tiling 结构体定义与动态分块实现

1. Tiling 结构体定义（relu_dynamic_tiling.h）

包含分块维度、分块数量、输入实际维度等核心参数，支持动态推导：

struct ReluDynamicTiling {
    int32_t tile_h;      // H维度分块大小（动态推导）
    int32_t tile_w;      // W维度分块大小（动态推导）
    int32_t tile_c;      // C维度分块大小（动态推导）
    int32_t tile_num_h;  // H维度分块数量
    int32_t tile_num_w;  // W维度分块数量
    int32_t tile_num_c;  // C维度分块数量
    int32_t total_tile;  // 总分块数量
    int32_t input_n;     // 输入N维度实际大小
    int32_t input_c;     // 输入C维度实际大小
    int32_t input_h;     // 输入H维度实际大小
    int32_t input_w;     // 输入W维度实际大小
};

2. 动态 Tiling 函数实现（relu_dynamic.cpp）

Host 侧通过 Tiling 函数，根据输入 Shape 与 Local Memory 容量，动态计算分块参数：

#include "relu_dynamic_tiling.h"
#include "ge_api.h"
using namespace ge;

Status ReluDynamicTilingFunc(const Operator &op, ReluDynamicTiling &tiling) {
    // 1. 获取输入张量描述与Shape
    auto input_desc = op.GetInputDesc(0);
    auto input_shape = input_desc.GetShape();
    
    // 2. 动态提取输入维度实际大小
    tiling.input_n = input_shape.GetDim(0);
    tiling.input_c = input_shape.GetDim(1);
    tiling.input_h = input_shape.GetDim(2);
    tiling.input_w = input_shape.GetDim(3);
    
    // 3. 动态推导分块大小（适配512KB Local Memory，float16占2字节）
    const int32_t MAX_LOCAL_BYTES = 512 * 1024;
    const int32_t ELEM_BYTES = 2;
    int32_t max_tile_elem = MAX_LOCAL_BYTES / ELEM_BYTES;  // 最大分块元素数：262144
    
    // 按H/W/C均衡分块，避免单维度过大
    tiling.tile_h = std::min(64, tiling.input_h);  // H维度分块最大64
    tiling.tile_w = std::min(64, tiling.input_w);  // W维度分块最大64
    int32_t current_elem = tiling.tile_h * tiling.tile_w;
    tiling.tile_c = std::min(max_tile_elem / current_elem, tiling.input_c);  // C维度分块自适应
    
    // 4. 动态计算分块数量（向上取整，覆盖全部数据）
    tiling.tile_num_h = (tiling.input_h + tiling.tile_h - 1) / tiling.tile_h;
    tiling.tile_num_w = (tiling.input_w + tiling.tile_w - 1) / tiling.tile_w;
    tiling.tile_num_c = (tiling.input_c + tiling.tile_c - 1) / tiling.tile_c;
    tiling.total_tile = tiling.tile_num_h * tiling.tile_num_w * tiling.tile_num_c;
    
    return SUCCESS;
}

Status ReluDynamic::InferShape(const Operator &op, vector<TensorDesc> &output_desc) {
    // 1. 获取输入张量描述
    auto input_desc = op.GetInputDesc(0);
    auto input_shape = input_desc.GetShape();
    
    // 2. 动态校验输入合法性
    // 校验维度数量（需为4维[N,C,H,W]）
    if (input_shape.GetDimNum() != 4) {
        GE_LOGE("Input shape dim num error! Expected 4, got %d", input_shape.GetDimNum());
        return PARAM_INVALID;
    }
    // 校验数据类型（需为float16）
    if (input_desc.GetDataType() != DT_FLOAT16) {
        GE_LOGE("Input dtype error! Expected float16, got %d", input_desc.GetDataType());
        return PARAM_INVALID;
    }
    // 校验维度大小（避免无效维度）
    for (int32_t i = 0; i < 4; i++) {
        if (input_shape.GetDim(i) <= 0) {
            GE_LOGE("Input dim %d is invalid! Value: %d", i, input_shape.GetDim(i));
            return PARAM_INVALID;
        }
    }
    
    // 3. 动态推导输出Shape（与输入一致）
    output_desc.push_back(input_desc);
    return SUCCESS;
}

（四）Step 4：Kernel 类实现（动态资源适配）

Kernel 类需支持动态初始化资源、适配动态分块、处理边界分块，核心包含Init（资源初始化）、Process（流程控制）、Compute（核心计算）方法：

（三）Step 3：Shape 推导函数实现

动态校验输入参数合法性，推导输出 Shape，为后续计算与资源分配提供依据：

#include "relu_dynamic_tiling.h"
#include "ascendc/pipe.h"
#include "ascendc/queue.h"
#include "ascendc/tensor.h"
using namespace AscendC;

class KernelReluDynamic {
public:
    // 动态初始化资源（根据Tiling参数分配Local Memory）
    __aicore__ inline Status Init(__gm__ const ReluDynamicTiling &tiling) {
        // 保存Tiling参数
        tiling_ = tiling;
        
        // 动态计算Local Memory缓冲区大小（按最大分块分配）
        int32_t max_tile_elem = tiling_.tile_h * tiling_.tile_w * tiling_.tile_c;
        local_input_ = new (std::nothrow) float16_t[max_tile_elem];
        if (local_input_ == nullptr) {
            GE_LOGE("Local memory allocation failed! Tile elem: %d", max_tile_elem);
            return MEMALLOC_FAILED;
        }
        
        // 初始化Pipe与Queue（支持Double Buffer）
        constexpr int32_t BUFFER_NUM = 2;
        pipe_.InitBuffer(inQueue_, 4, max_tile_elem * sizeof(float16_t));
        pipe_.InitBuffer(outQueue_, 4, max_tile_elem * sizeof(float16_t));
        
        return SUCCESS;
    }
    
    // 核心流程控制：CopyIn→Compute→CopyOut
    __aicore__ inline Status Process(__gm__ const float16_t *input, __gm__ float16_t *output,
                                     int32_t tile_idx) {
        // 动态计算当前分块的偏移量
        int32_t tile_num_wc = tiling_.tile_num_w * tiling_.tile_num_c;
        int32_t h_idx = tile_idx / tile_num_wc;
        int32_t wc_idx = tile_idx % tile_num_wc;
        int32_t w_idx = wc_idx / tiling_.tile_num_c;
        int32_t c_idx = wc_idx % tiling_.tile_num_c;
        
        int32_t h_offset = h_idx * tiling_.tile_h;
        int32_t w_offset = w_idx * tiling_.tile_w;
        int32_t c_offset = c_idx * tiling_.tile_c;
        
        // 动态处理边界分块（计算实际分块大小）
        int32_t actual_h = std::min(tiling_.tile_h, tiling_.input_h - h_offset);
        int32_t actual_w = std::min(tiling_.tile_w, tiling_.input_w - w_offset);
        int32_t actual_c = std::min(tiling_.tile_c, tiling_.input_c - c_offset);
        int32_t actual_elem = actual_h * actual_w * actual_c;
        
        // 数据搬入（Global→Local）
        CopyIn(input, h_offset, w_offset, c_offset, actual_h, actual_w, actual_c);
        
        // 核心计算（ReLU激活）
        Compute(actual_elem);
        
        // 数据搬出（Local→Global）
        CopyOut(output, h_offset, w_offset, c_offset, actual_h, actual_w, actual_c);
        
        return SUCCESS;
    }
    
private:
    // 数据搬入：从Global Memory搬运当前分块数据至Local Memory
    __aicore__ inline void CopyIn(__gm__ const float16_t *input, int32_t h_offset, int32_t w_offset,
                                  int32_t c_offset, int32_t actual_h, int32_t actual_w, int32_t actual_c) {
        // 计算Global Memory偏移量（NHWC格式）
        int32_t input_n_stride = tiling_.input_c * tiling_.input_h * tiling_.input_w;
        int32_t input_c_stride = tiling_.input_h * tiling_.input_w;
        int32_t input_h_stride = tiling_.input_w;
        int32_t global_offset = 0 * input_n_stride + c_offset * input_c_stride + 
                               h_offset * input_h_stride + w_offset;
        
        // 分配LocalTensor并搬运数据
        LocalTensor<float16_t> local_tensor = inQueue_.AllocTensor<float16_t>();
        DataCopy(local_tensor, input + global_offset, actual_h * actual_w * actual_c);
        inQueue_.EnQue(local_tensor);
    }
    
    // 核心计算：ReLU激活（f(x) = max(0, x)）
    __aicore__ inline void Compute(int32_t actual_elem) {
        LocalTensor<float16_t> local_in = inQueue_.DeQue<float16_t>();
        LocalTensor<float16_t> local_out = outQueue_.AllocTensor<float16_t>();
        
        // 并行执行ReLU计算（Vector指令加速）
        for (int32_t i = 0; i < actual_elem; i++) {
            local_out[i] = (local_in[i] > 0) ? local_in[i] : 0;
        }
        
        outQueue_.EnQue(local_out);
        inQueue_.FreeTensor(local_in);
    }
    
    // 数据搬出：将计算结果从Local Memory搬运至Global Memory
    __aicore__ inline void CopyOut(__gm__ float16_t *output, int32_t h_offset, int32_t w_offset,
                                   int32_t c_offset, int32_t actual_h, int32_t actual_w, int32_t actual_c) {
        // 计算Global Memory偏移量
        int32_t output_n_stride = tiling_.input_c * tiling_.input_h * tiling_.input_w;
        int32_t output_c_stride = tiling_.input_h * tiling_.input_w;
        int32_t output_h_stride = tiling_.input_w;
        int32_t global_offset = 0 * output_n_stride + c_offset * output_c_stride + 
                               h_offset * output_h_stride + w_offset;
        
        // 搬出数据
        LocalTensor<float16_t> local_out = outQueue_.DeQue<float16_t>();
        DataCopy(output + global_offset, local_out, actual_h * actual_w * actual_c);
        outQueue_.FreeTensor(local_out);
    }
    
    // 成员变量
    ReluDynamicTiling tiling_;          // Tiling参数
    TPipe pipe_;                        // 资源管理对象
    TQue<TQuePosition::VECIN, 2> inQueue_;   // 输入队列
    TQue<TQuePosition::VECOUT, 2> outQueue_;  // 输出队列
    float16_t *local_input_;            // Local Memory缓冲区
};

// 核函数定义（入口）
extern "C" __global__ __aicore__ void relu_dynamic(__gm__ const float16_t *input,
                                                   __gm__ float16_t *output,
                                                   __gm__ const ReluDynamicTiling &tiling,
                                                   int32_t tile_idx) {
    KernelReluDynamic op;
    if (op.Init(tiling) != SUCCESS) {
        GE_LOGE("Kernel init failed! Tile idx: %d", tile_idx);
        return;
    }
    op.Process(input, output, tile_idx);
}

// 核函数调用封装
void relu_dynamic_do(uint32_t blockDim, void* stream, __gm__ const float16_t *input,
                     __gm__ float16_t *output, __gm__ const ReluDynamicTiling &tiling) {
    for (int32_t tile_idx = 0; tile_idx < tiling.total_tile; tile_idx++) {
        relu_dynamic<<<blockDim, nullptr, stream>>>(input, output, tiling, tile_idx);
    }
}

（五）Step 5：算子注册与编译部署

1. 算子注册

REG_OP(ReluDynamic)
    .INPUT(x, TensorType({DT_FLOAT16}))
    .OUTPUT(y, TensorType({DT_FLOAT16}))
    .OPERATOR_CLASS(ReluDynamic)
    .TILING_FUNC(ReluDynamicTilingFunc);

2. 编译配置（CMakeLists.txt）

set(ASCEND_CANN_PACKAGE_PATH "/usr/local/Ascend/ascend-toolkit/latest")
include_directories(${ASCEND_CANN_PACKAGE_PATH}/include)
link_directories(${ASCEND_CANN_PACKAGE_PATH}/lib64)
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++17 -O2 -fPIC -DDYNAMIC_SHAPE")
target_link_libraries(ReluDynamic_opp ascendcl ge_compiler ge_common pthread)

3. 部署与验证

编译生成算子包后，通过部署脚本安装至昇腾 OP 库目录，使用不同 Shape 的输入数据验证算子兼容性：

# 部署算子
sudo ./ReluDynamic_opp_linux_x86_64.run

# 验证不同Shape输入（示例）
# Shape=[1, 3, 256, 256]
./test_relu_dynamic --shape 1,3,256,256
# Shape=[2, 3, 512, 512]
./test_relu_dynamic --shape 2,3,512,512
# Shape=[4, 3, 1024, 1024]
./test_relu_dynamic --shape 4,3,1024,1024

四、动态 Shape 算子优化技巧

（一）分块大小自适应优化

根据输入维度动态调整分块大小，平衡计算效率与内存占用。例如：

• 小 Shape 输入（如 H=128、W=128）：采用 32×32 分块，减少分块数量与调度开销。
• 大 Shape 输入（如 H=1024、W=1024）：采用 64×64 分块，充分利用 Vector 单元并行能力。

（二）Local Memory 复用优化

在 Kernel 类中复用 Local Memory 缓冲区，避免频繁分配与释放内存。例如，将缓冲区初始化放在Init方法中，整个 Kernel 生命周期内仅分配一次内存。

（三）边界分块单独优化

对不能被分块大小整除的边界分块，单独处理实际数据量，避免无效计算。例如，通过std::min(tile_dim, input_dim - offset)计算实际分块大小，仅对有效数据执行计算。

（四）线程配置动态适配

根据分块数量与硬件并行能力，动态调整线程块数量（dim3 block），充分利用 AI Core 算力。例如，按block(tile_h * tile_w)配置线程块，确保每个线程处理一个数据元素。

五、常见问题与解决方案

（一）分块大小超出 Local Memory 容量

问题：动态 Shape 场景下，输入维度过大导致分块大小超出 Local Memory 容量，引发内存溢出。解决方案：在 Tiling 函数中增加分块大小校验，根据 Local Memory 容量动态限制最大分块元素数，公式为max_tile_elem = MAX_LOCAL_BYTES / ELEM_BYTES。

（二）边界分块数据丢失

问题：输入维度不能被分块大小整除时，未处理边界分块，导致部分数据未参与计算。解决方案：通过向上取整计算分块数量（(dim + tile_dim - 1) / tile_dim），并在 Kernel 中动态计算边界分块的实际数据量。

（三）Shape 推导失败

问题：输入维度数量或数据类型不满足要求，导致 Shape 推导函数返回错误。解决方案：在 Shape 推导函数中增加严格的合法性校验，包括维度数量、维度大小、数据类型等，输出明确的错误日志便于定位。

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

报名链接
https://www.hiascend.com/developer/activities/cann20252?tab=overview