Ascend C 中级认证核心考点：Kernel 类与 Tiling 结构体深度解析

本文围绕认证考点，详细讲解 Kernel 类的结构设计、Tiling 结构体的定义与使用，结合真题考点分析助力备考。掌握 Kernel 类的结构（如Init、Process、Compute方法），理解固定 / 动态 Shape 场景下的实现差异，能够编写符合规范的 Kernel 类代码。掌握 Tiling 结构体的定义规范，理解分块维度、分块数量等参数的含义，能够根据算子需求设计合理的分块规则。C

2401_89199769

832人浏览 · 2025-10-25 21:54:33

2401_89199769 · 2025-10-25 21:54:33 发布

Ascend C 中级认证中，Kernel 类实现与 Tiling 技术是核心考点，占比约 40%。本文围绕认证考点，详细讲解 Kernel 类的结构设计、Tiling 结构体的定义与使用，结合真题考点分析助力备考。

一、Tiling 结构体：分块规则的核心定义

（一）考点要求

掌握 Tiling 结构体的定义规范，理解分块维度、分块数量等参数的含义，能够根据算子需求设计合理的分块规则。

（二）结构体定义规范

核心参数：

分块维度参数：如tile_h（H 维度分块大小）、tile_w（W 维度分块大小）。
分块数量参数：如tile_num（总分块数量）、tile_idx（当前分块索引）。
辅助参数：如data_size（单块数据大小）、pad_size（边界填充大小）。

定义示例（Sigmoid 算子 Tiling 结构体）：

// sigmoid_custom_tiling.h
struct SigmoidCustomTiling {
    // 分块维度：H、W维度分块大小
    int32_t tile_h;
    int32_t tile_w;
    // 分块数量：H、W维度分块数
    int32_t tile_num_h;
    int32_t tile_num_w;
    // 总分块数量
    int32_t total_tile_num;
    // 输入张量总大小
    int32_t input_size;
};

（三）Tiling 数据传递

Host 侧通过getTilingData方法将 Tiling 结构体数据传递给 Kernel 侧，Kernel 侧通过__gm__修饰符接收全局内存中的 Tiling 信息，示例：

// Host侧传递Tiling数据
ge::Tensor tiling_tensor;
tiling_tensor.SetData((uint8_t *)&tiling, sizeof(SigmoidCustomTiling));
op.SetTilingData(tiling_tensor);

// Kernel侧接收Tiling数据
__global__ void sigmoid_custom(__gm__ const float *input, __gm__ float *output,
                               __gm__ SigmoidCustomTiling tiling, int32_t tile_idx) {
    // 使用tiling参数获取分块信息
    int32_t tile_h = tiling.tile_h;
    int32_t tile_w = tiling.tile_w;
    // ...
}

（四）认证考点分析

考点 1：Tiling 结构体参数设计，需根据输入 shape 和 Local Memory 容量计算分块大小。
考点 2：Tiling 数据在 Host 与 Kernel 之间的传递方式，掌握SetTilingData和__gm__修饰符的使用。
考点 3：边界分块处理，当输入 shape 不能被分块大小整除时，如何设计分块规则避免数据丢失。

二、Kernel 类实现：算子计算逻辑的封装

（一）考点要求

掌握 Kernel 类的结构（如Init、Process、Compute方法），理解固定 / 动态 Shape 场景下的实现差异，能够编写符合规范的 Kernel 类代码。

（二）Kernel 类核心结构

以 Sigmoid 算子为例，Kernel 类需包含以下核心方法：

Init方法：初始化资源（如 Local Memory 缓冲区、线程配置），根据 Tiling 信息或输入 shape 确定计算规模。
Process方法：数据加载与结果写回，将 Global Memory 中的数据加载至 Local Memory，计算完成后写回 Global Memory。
Compute方法：核心计算逻辑，实现 Sigmoid 函数（1/(1+exp (-x))）的并行计算。

（三）实现示例（动态 Shape 场景）

// sigmoid_custom_kernel.cpp
class KernelSigmoid {
public:
    // 初始化资源
    Status Init(__gm__ const SigmoidCustomTiling &tiling, int32_t tile_idx) {
        // 获取当前分块的偏移量
        tile_h_offset_ = (tile_idx / tiling.tile_num_w) * tiling.tile_h;
        tile_w_offset_ = (tile_idx % tiling.tile_num_w) * tiling.tile_w;
        // 初始化Local Memory缓冲区
        local_input_ = new (std::nothrow) float[tiling.tile_h * tiling.tile_w];
        if (local_input_ == nullptr) {
            GE_LOGE("Local memory allocation failed!");
            return MEMALLOC_FAILED;
        }
        return SUCCESS;
    }
    
    // 数据加载与写回
    Status Process(__gm__ const float *input, __gm__ float *output, 
                   __gm__ const SigmoidCustomTiling &tiling) {
        // 加载当前分块数据至Local Memory
        ld_matrix(local_input_, input + tile_h_offset_ * W + tile_w_offset_, 
                  tiling.tile_h, tiling.tile_w);
        // 执行计算
        Compute(tiling.tile_h, tiling.tile_w);
        // 写回计算结果
        st_matrix(output + tile_h_offset_ * W + tile_w_offset_, local_input_, 
                  tiling.tile_h, tiling.tile_w);
        return SUCCESS;
    }
    
    // 核心计算逻辑
    void Compute(int32_t tile_h, int32_t tile_w) {
        int32_t idx = threadIdx.x;
        for (int32_t h = 0; h < tile_h; h++) {
            for (int32_t w = 0; w < tile_w; w++) {
                int32_t local_idx = h * tile_w + w;
                if (local_idx == idx) {
                    // Sigmoid计算：1/(1+exp(-x))
                    local_input_[local_idx] = 1.0f / (1.0f + expf(-local_input_[local_idx]));
                }
            }
        }
    }

private:
    float *local_input_;       // Local Memory缓冲区
    int32_t tile_h_offset_;    // H维度分块偏移量
    int32_t tile_w_offset_;    // W维度分块偏移量
};