从零开始学昇腾Ascend C算子开发-第三篇：算子开发基础

3.1 算子开发流程

3.1.1 算子需求分析

注：运行前看好自己的版本：

明确算子功能

开发算子之前，先得搞清楚这个算子要干什么。比如要做一个Add算子，那就是两个输入相加得到输出。听起来简单，但实际要考虑的东西还挺多。

输入输出规格：输入有几个？输出有几个？每个输入输出的形状是什么？数据类型是什么？这些都要明确。

边界情况：空输入怎么处理？形状不匹配怎么办？数值溢出怎么处理？这些边界情况都要考虑。

性能要求：这个算子对性能有什么要求？是追求极致性能，还是够用就行？这会影响你的实现方案。

分析计算特点

不同的算子有不同的计算特点，分析清楚才能选对实现方法。

计算密集还是内存密集：如果计算很复杂，那要优化计算部分。如果主要是内存访问，那要优化内存访问模式。

数据依赖：计算之间有没有依赖关系？能不能并行？能不能向量化？

数据重用：中间结果能不能重用？能不能减少重复计算？

3.1.2 算子设计思路

确定实现方案

根据算子特点，确定实现方案。

简单算子：像Add、Mul这种元素级算子，直接用向量化API就行，实现起来比较简单。

复杂算子：像卷积、矩阵乘法这种，可能需要分块（Tiling），需要仔细设计数据流和计算流。

特殊算子：有些算子可能需要特殊的优化技巧，比如算子融合、内存复用这些。

设计数据流

数据流设计很重要，直接影响性能。

数据加载策略：一次加载多少数据？什么时候加载？是提前加载还是按需加载？

计算顺序：先算哪部分？后算哪部分？能不能流水线化？

结果写回策略：什么时候写回结果？是算完一块写一块，还是算完所有再写？

3.1.3 算子实现步骤

第一步：创建工程

在ModelArts Notebook里，先创建工程目录结构：

# 创建工程目录
mkdir my_operator
cd my_operator

# 创建源代码目录
mkdir src include

工程结构大概这样：

• src/：放核函数代码
• include/：放头文件
• build.sh：编译脚本
• CMakeLists.txt：CMake配置文件（如果用CMake的话）

第二步：编写核函数

在src/目录下创建核函数文件，比如add_kernel.cpp：

// 概念性示例
#include "ascendc.h"

extern "C" __global__ __aicore__ void AddKernel(
    GlobalTensor<float> input1,
    GlobalTensor<float> input2,
    GlobalTensor<float> output,
    int32_t total_length
) {
    // 1. 分配LocalTensor
    LocalTensor<float> local_input1;
    LocalTensor<float> local_input2;
    LocalTensor<float> local_output;
    
    local_input1.Alloc(total_length);
    local_input2.Alloc(total_length);
    local_output.Alloc(total_length);
    
    // 2. 数据加载
    DataCopy(local_input1, input1, total_length);
    DataCopy(local_input2, input2, total_length);
    
    // 3. 计算
    Add(local_output, local_input1, local_input2);
    
    // 4. 结果写回
    DataCopy(output, local_output, total_length);
    
    // 5. 释放内存
    local_input1.Free();
    local_input2.Free();
    local_output.Free();
}

这只是个概念性示例，实际的API调用会更复杂一些。

第三步：实现Tiling函数

如果算子需要分块处理，还要实现Tiling函数。Tiling函数在Host端运行，计算每个块的大小和偏移。

// 概念性示例
void TilingFunc(TilingData* tiling) {
    // 计算每个块的大小
    int32_t block_size = 256;
    int32_t total_blocks = (total_length + block_size - 1) / block_size;
    
    tiling->block_size = block_size;
    tiling->total_blocks = total_blocks;
}

第四步：注册算子

在Host端注册算子，让框架知道有这个算子：

// 概念性示例
void RegisterOperator() {
    // 注册算子原型
    RegisterOp("Add", AddKernel, TilingFunc);
}

3.1.4 算子测试验证

CPU侧调试

Ascend C支持CPU侧调试，可以在CPU上模拟NPU的行为，方便调试。

# 概念性命令
# 在CPU模式下运行算子
cpu_run_kernel AddKernel input1 input2 output

CPU侧调试的好处是速度快，不需要NPU硬件，但只能验证逻辑正确性，不能验证性能。

NPU侧验证

逻辑没问题后，要在NPU上实际跑一下，验证性能和正确性。

# 概念性命令
# 在NPU上运行算子
npu_run_kernel AddKernel input1 input2 output

NPU侧验证能看到真实的性能，也能发现一些CPU侧发现不了的问题。

结果验证

验证结果是否正确：

功能正确性：输出结果对不对？可以用参考实现（比如NumPy）对比一下。

边界情况：空输入、异常输入这些边界情况都测试一下。

性能指标：算子的执行时间、内存使用、NPU利用率这些指标都要看看。

---

3.2 向量化编程

3.2.1 向量化指令集

什么是向量化指令

向量化指令就是一条指令同时处理多个数据的指令。比如Add指令，可以同时把两个向量的所有元素加起来，而不是循环一个一个加。

Ascend C提供了丰富的向量化指令，覆盖了大部分常见的运算。

向量加载和存储指令

Load指令：从LocalTensor加载数据到向量。

// 概念性示例
Vector<float, 256> vec;
vec.Load(local_tensor, offset);  // 从local_tensor的offset位置加载256个元素

Store指令：把向量数据存回LocalTensor。

// 概念性示例
Vector<float, 256> vec;
// ... 对vec进行计算 ...
vec.Store(local_tensor, offset);  // 存回local_tensor的offset位置

向量运算指令

向量运算指令有很多，常用的有：

算术运算：Add、Sub、Mul、Div，这些是基础的加减乘除。

数学函数：Exp、Ln、Sqrt、Abs，这些是数学函数。

比较运算：Max、Min，找最大值最小值。

融合指令：FusedMulAdd，乘加融合，一条指令完成乘法和加法，性能更好。

3.2.2 向量加载和存储

加载策略

数据加载有几种策略：

顺序加载：按顺序一块一块地加载，简单直接。

预取加载：提前加载下一块数据，让加载和计算重叠。

批量加载：一次加载多块数据，减少加载次数。

选择哪种策略要看具体情况，通常预取加载效果比较好。

存储策略

存储也有策略：

立即存储：算完一块就存一块，简单但可能影响流水线。

延迟存储：算完所有再统一存储，可能更高效。

批量存储：一次存储多块数据。

对齐要求

加载和存储都要注意对齐。数据不对齐的话，可能要多访问一次内存，或者访问跨缓存行，都很慢。

// 概念性示例
// 确保offset是向量长度的倍数
int32_t aligned_offset = (offset + vector_length - 1) / vector_length * vector_length;
vec.Load(local_tensor, aligned_offset);

3.2.3 向量运算操作

元素级运算

元素级运算就是对向量的每个元素分别做运算：

// 概念性示例
Vector<float, 256> a, b, c;

// 向量加法：c[i] = a[i] + b[i]
Add(c, a, b);

// 向量乘法：c[i] = a[i] * b[i]
Mul(c, a, b);

// 向量最大值：c[i] = max(a[i], b[i])
Max(c, a, b);

标量运算

标量和向量的运算：

// 概念性示例
Vector<float, 256> vec;
Scalar<float> s = 3.14;

// 向量加标量：vec[i] = vec[i] + s
Adds(vec, vec, s);

// 向量乘标量：vec[i] = vec[i] * s
Muls(vec, vec, s);

归约运算

归约运算就是把向量的所有元素归约成一个值：

// 概念性示例
Vector<float, 256> vec;
Scalar<float> sum;

// 求和
ReduceSum(sum, vec);

// 求最大值
ReduceMax(max_val, vec);

3.2.4 向量化优化技巧

循环向量化

把标量循环改成向量化循环：

// 标量版本（慢）
for (int i = 0; i < n; i++) {
    c[i] = a[i] + b[i];
}

// 向量化版本（快）
for (int i = 0; i < n; i += vector_length) {
    Vector<float, 256> va, vb, vc;
    va.Load(a, i);
    vb.Load(b, i);
    Add(vc, va, vb);
    vc.Store(c, i);
}

数据重用

如果数据要用多次，尽量在向量里多待一会儿：

// 概念性示例
Vector<float, 256> vec;
vec.Load(local_tensor, offset);

// 用vec做多个运算
Add(result1, vec, vec);
Mul(result2, vec, vec);
// vec不用重复加载

融合操作

尽量用融合指令，减少指令数：

// 分开写（两条指令）
Mul(temp, a, b);
Add(c, temp, c);

// 融合指令（一条指令，更快）
FusedMulAdd(c, a, b, c);  // c = a * b + c

---

3.3 并行计算

3.3.1 多核并行计算

并行执行模型

昇腾处理器有多个AI Core，可以并行执行。每个Core执行同一个核函数，但处理不同的数据块。

// 概念性示例
// 核函数会被多个Core并行执行
// 每个Core处理不同的数据块
extern "C" __global__ __aicore__ void MyKernel(...) {
    // 获取当前Core的ID
    int32_t core_id = GetCoreId();
    
    // 根据Core ID计算要处理的数据范围
    int32_t start = core_id * block_size;
    int32_t end = start + block_size;
    
    // 处理这个范围的数据
    ProcessBlock(start, end);
}

数据分块策略

数据怎么分块很重要：

均匀分块：每个Core处理相同大小的块，简单但可能不够灵活。

负载均衡：根据数据特点，让每个Core的负载尽量均衡。

边界处理：分块的时候要注意边界，确保所有数据都被处理，不重复不遗漏。

3.3.2 数据并行和任务并行

数据并行（Data Parallelism）

数据并行就是不同的Core处理不同的数据：

// 概念性示例
// 每个Core处理不同的数据块
void ProcessData() {
    int32_t core_id = GetCoreId();
    int32_t total_cores = GetTotalCores();
    
    // 计算这个Core要处理的数据范围
    int32_t block_size = total_data / total_cores;
    int32_t start = core_id * block_size;
    int32_t end = (core_id == total_cores - 1) ? total_data : start + block_size;
    
    // 处理这个范围
    for (int i = start; i < end; i++) {
        ProcessElement(i);
    }
}

数据并行适合数据之间没有依赖的情况，比如元素级算子。

任务并行（Task Parallelism）

任务并行就是不同的Core执行不同的任务：

// 概念性示例
// 不同的Core执行不同的任务阶段
void ProcessPipeline() {
    int32_t core_id = GetCoreId();
    
    if (core_id == 0) {
        // Core 0负责数据加载
        LoadData();
    } else if (core_id == 1) {
        // Core 1负责计算
        Compute();
    } else if (core_id == 2) {
        // Core 2负责结果写回
        WriteBack();
    }
}

任务并行适合流水线场景，但实现起来复杂一些。

3.3.3 线程同步机制

为什么需要同步

多个Core并行执行的时候，有时候需要同步。比如一个Core要等另一个Core的结果，或者多个Core要协调访问共享资源。

同步原语

Ascend C提供了同步原语：

Barrier（屏障）：等待所有Core都到达这个点，再继续执行。

// 概念性示例
void SynchronizedProcess() {
    // 每个Core执行自己的计算
    DoLocalCompute();
    
    // 等待所有Core完成
    Barrier();
    
    // 所有Core都完成后，继续执行
    DoGlobalCompute();
}

原子操作：对共享变量进行原子操作，避免竞争。

// 概念性示例
// 原子加法
AtomicAdd(shared_counter, value);

3.3.4 负载均衡

负载不均衡的问题

如果负载不均衡，有些Core很忙，有些Core很闲，整体性能就上不去。比如数据分块不均匀，或者不同块的计算复杂度不同。

负载均衡策略

动态分配：根据Core的负载情况，动态分配任务。负载轻的Core多分点任务。

工作窃取：空闲的Core可以从忙碌的Core那里"偷"一些任务来做。

自适应分块：根据数据特点，动态调整分块大小，让每个Core的负载尽量均衡。

性能监控

监控每个Core的负载情况，找出瓶颈：

// 概念性示例
void MonitorLoad() {
    int32_t core_id = GetCoreId();
    int32_t start_time = GetTime();
    
    // 执行计算
    DoCompute();
    
    int32_t end_time = GetTime();
    int32_t duration = end_time - start_time;
    
    // 记录每个Core的执行时间
    LogCoreTime(core_id, duration);
}

根据监控结果，调整分块策略，优化负载均衡。

---

学习检查点

学完这一篇，你应该能做到这些：

理解算子开发的完整流程，从需求分析到测试验证。掌握向量化编程，知道怎么用向量指令，怎么优化向量化代码。理解并行计算，知道数据并行和任务并行的区别，知道怎么同步和负载均衡。能够实现一个简单的算子，比如Add或Mul，完成从编写到测试的全流程。

实践练习

实现Add算子：在ModelArts Notebook中创建一个工程，实现一个Add算子。两个输入相加得到输出，用向量化API实现。完成编译、部署、测试的全流程。

向量化优化：实现一个向量加法的算子，先用标量循环实现，再改成向量化实现，对比性能差异。理解向量化的优势。

并行计算实验：实现一个需要多Core并行的算子，比如大矩阵的加法。理解数据分块、Core同步这些概念。

性能调优：对一个算子进行性能调优，比如调整分块大小、优化内存访问、使用融合指令等，看看性能能提升多少。

---

下一步：掌握了算子开发的基础后，就可以学习常用算子的实现了。下一章会讲各种常见算子的实现方法，比如元素级算子、规约算子、矩阵运算算子等，到时候你就能实现更复杂的算子了。

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

报名链接:https://www.hiascend.com/developer/activities/cann20252

社区地址：https://www.hiascend.com/developer