华为昇腾CANN算子开发类认证，考试试卷（HCIP/微认证）快速通关方案，附6张试卷，后面三张直接在题目下面有答案

zhangfeng1133

215人浏览 · 2026-06-11 15:41:07

zhangfeng1133 · 2026-06-11 15:41:07 发布

昇腾CANN算子开发类认证（HCIP/微认证）快速通关方案，附6张试卷，后面三张直接在题目下面有答案

核心思路：优先攻克高分模块，先吃透理论保底，再突击实操拿分，最后刷题复盘查漏补缺，零基础也可按步骤落地。
ai生成

视频教程
https://www.bilibili.com/video/BV1C94y1V7ZV?spm_id_from=333.788.videopod.episodes&vd_source=7d3841cb81fdca8b20a31fa9ee5ec8ef&p=7

一、考试模块&分值优先级（按投入产出排序）

结合当前CANN 8.x版本、Ascend C算子开发认证考点，划分优先级，优先拿下占比高、易得分内容：

第一梯队（高占比、易拿分，优先学）
- 昇腾AI处理器架构（14%）：纯理论，背诵即可得分
- 并行计算与Kernel编程（13%）：理论+基础代码，实操核心框架
- 模板库基本概念（13%）：概念+场景题，记忆为主
- 存储与计算单元功能（10%）：理解数据通路，选择题高频考点
第二梯队（中等难度，保底得分）
- 模型迁移适配（10%）、性能分析工具使用（10%）：流程+工具操作，上手快
- 自定义算子开发（10%）：套用固定代码模板即可完成基础开发
第三梯队（偏难，最后攻坚）
- GEMM算子实现与优化（10%）、算子性能调优（10%）：代码+优化思路，零基础先掌握核心逻辑，不用深挖高阶技巧

二、分模块突击学习（直接照着执行）

（一）理论模块（拿下约50%分值，2天完成）

目标：选择、判断题尽量满分，不丢基础分，以背诵+理解为主。

昇腾AI处理器架构（14%）
- 必背核心：CANN五层架构（应用使能层→框架层→算子开发层→运行时层→驱动/硬件层）及每层作用；AI Core三大单元分工（Cube矩阵运算、Vector向量运算、Scalar标量运算）；软硬件分离耦合架构优势、Stream/Queue/Event同步机制、HCCL通信库用途。
- 学习方式：查看官方架构文档，手绘架构简图，早晚各背诵1遍。
存储与计算单元功能（10%）
- 必背核心：标准数据流向 Global Memory(GM) → L2 → L1 → UB；UB读写速度最快，GM最慢；DMA数据搬运逻辑、多级缓存作用。
- 答题技巧：性能类题目，优先选择“数据搬运优化、UB复用”相关选项。
并行计算与Kernel编程（13%）
- 必背核心：SPMD并行模型；GetBlockIdx() 索引用法；CANN 8.x核函数标准声明格式。
- 重点记忆：并行块尺寸设置原则，代码边界判断写法。
模板库基本概念（13%）
- 必背核心：模板库作用（简化开发、代码复用）；CATLASS高性能矩阵模板的应用场景（大模型训练/推理）。
补充理论（第二梯队，半天搞定）
- 模型迁移：牢记迁移全流程（评估→迁移→调试→部署），区分自动迁移、手动迁移的优缺点。
- 性能工具：熟记msprof工具用途（采集性能数据、定位算子瓶颈）。

（二）实操模块（拿下约40%分值，3天完成，CANN 8.x版本）

零基础套用固定模板，先保证功能跑通，再简单优化，不用深究复杂原理。

基础准备
优先使用华为开发者空间在线环境（预装CANN 8.x），规避本地环境搭建报错；牢记：工程目录禁止修改、数据操作必须32字节对齐。
核心模板背诵&练习（重中之重）
熟练掌握向量加法算子完整代码、编译命令，这是实操通用框架，大部分基础算子仅需修改计算API。
- 核函数代码、算子注册语句、ascendc编译命令，全程默写练习。
- 拓展练习：基于加法算子，改成乘法算子，熟悉代码改动逻辑。
自定义算子开发
熟记开发流程：核函数编写→内存搬运→计算逻辑→结果回写→算子注册，严格遵循固定流程。
高阶内容（GEMM&性能调优，浅学即可）
- 不用手写完整复杂代码，只记核心思路：大矩阵分块存入UB、双缓冲技术隐藏搬运延迟、数据Padding对齐提升硬件利用率。
- 性能调优三大通用方案：优化数据搬运、保证内存对齐、优先使用Cube/Vector批量运算。

（三）工具使用（半天掌握）

掌握msprof基础操作：数据采集、查看算子耗时、识别数据搬运/计算瓶颈，应对工具类选择题、简答题。

三、7天完整冲刺计划表（每日任务清晰，零基础适配）

天数	核心任务	达成目标
第1天	架构、存储单元理论学习+背诵，刷对应章节习题	吃透两大高分理论模块，选择题正确率90%以上
第2天	并行编程、模板库、模型迁移理论背诵，整理易混知识点	完成全部理论模块学习，整理专属背诵笔记
第3天	环境熟悉，跑通Add加法算子，理解代码每一行含义	熟练编译、运行，能完整默写基础算子代码
第4天	改写Mul乘法算子，练习索引、边界判断、DataCopy用法	掌握基础算子通用修改逻辑，规避对齐报错
第5天	学习msprof工具使用，熟记GEMM、性能调优核心思路	应对工具题、优化类简答题
第6天	整套模拟题/真题刷题，整理错题本，复盘薄弱点	摸清考试题型、出题规律，补齐知识漏洞
第7天	复盘全部笔记+代码模板，回看错题，考前查漏	巩固全部考点，调整状态，准备考试

四、考场避坑&得分技巧（零基础必看）

理论题：优先作答，利用背诵内容快速选择，不确定的题目标记后回头再看。
实操题
- 原则：先跑通功能，再谈性能，不要一开始就追求极致优化。
- 红线：不改动官方工程目录；所有DataCopy严格保证32字节对齐；必须添加边界判断防止数组越界。
- 格式：核函数、算子注册严格使用CANN 8.x标准写法，拒绝旧版本语法。
时间分配：理论题控制在30分钟内，剩余时间专注代码编写、调试、提交。

五、配套资料&资源建议

官方资料：CANN 8.x 开发文档、昇腾社区配套课程、训练营样例代码（最权威，优先使用）。
练习资料：认证配套章节题、模拟卷，反复刷错题。
环境：全程使用华为在线开发环境，省去本地部署麻烦。

CANN 8.x 标准、可直接复制运行的完整案例：*

向量加法算子 Add*，包含源码、编译命令、运行说明、考试必记要点，零基础直接照着用即可。

一、完整算子源码（AddKernel.cpp）

新建文件 AddKernel.cpp，复制下面全部代码（CANN 8.x 标准写法）：

#include "kernel_operator.h"
using namespace AscendC;

// 核函数声明：CANN 8.x 标准格式
__global__ __aicore__ void AddKernel(
    GM_ADDR const __restrict__ src0,
    GM_ADDR const __restrict__ src1,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    // 1. 获取当前 Block 编号
    uint32_t blockIdx = GetBlockIdx();
    // 单块处理数据量（适配UB，考试常用 1024）
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    // 边界保护：超出总长度直接返回
    if (offset >= totalLen) {
        return;
    }

    // 2. 申请 UB 本地张量（CANN 8.x 推荐 LocalTensor）
    LocalTensor<half> ubSrc0;
    LocalTensor<half> ubSrc1;
    LocalTensor<half> ubDst;

    ubSrc0.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    ubSrc1.SetBuffer((half *)GM2UBAddr(1, 0), eachBlockLen);
    ubDst.SetBuffer((half *)GM2UBAddr(2, 0), eachBlockLen);

    // 3. GM -> UB 数据搬运（必须保证 32B 对齐）
    DataCopy(ubSrc0, src0 + offset, eachBlockLen * sizeof(half));
    DataCopy(ubSrc1, src1 + offset, eachBlockLen * sizeof(half));

    // 4. 向量加法计算（Ascend C 内置API）
    Add(ubDst, ubSrc0, ubSrc1, eachBlockLen);

    // 5. UB -> GM 结果写回
    DataCopy(dst + offset, ubDst, eachBlockLen * sizeof(half));
}

// 算子注册（考试固定写法，不要改）
这个可能有问题，不用考虑 REG_KERNEL(AddKernel, 0, AddFunc);

二、CANN 8.x 编译命令（直接复制执行）

环境：CANN 8.0 / 8.0.RC 系列，芯片以 Ascend910B 为例（考试主流）

ascendc \
--soc-version=ascend910b \
--kernel-name=AddKernel \
--output=AddKernel.o \
AddKernel.cpp

执行成功：生成 AddKernel.o 目标文件，无报错 = 编译通过
常见报错：
1. 命令不存在 → 环境没加载 CANN 8.x
2. 头文件找不到 → 环境版本不对，确认是 8.x

三、极简运行 & 测试流程（考试通用步骤）

编译得到 .o 文件
使用 msOpGen 生成完整算子工程（在线考试平台会自动封装，你只需写核函数）
构造半精度(half)测试数据，调用算子验证结果

线上认证实操题：只要求编写核函数 + 注册语句，不用自己搭完整工程。

四、零基础必背 6 个考点（选择/判断/实操高频）

1. 核函数固定格式（必考）

__global__ __aicore__ void 函数名(参数...)

8.x 必须写全这两个修饰符，少一个直接报错。

2. 索引获取

GetBlockIdx()：获取当前计算块编号（最常用）
不要越界：一定要加 if (offset >= totalLen) return; 边界判断

3. 内存层级（理论题必背）

数据流向：
全局内存 GM → L2缓存 → L1缓存 → 本地缓存 UB

UB：读写速度最快，计算只能在 UB 中完成
GM：存放原始数据，速度最慢

4. DataCopy 铁律（实操80%报错来源）

所有 DataCopy 长度、地址必须 32 字节对齐
考试只要出现拷贝报错，优先查对齐。

5. LocalTensor（8.x 新标准）

替代旧版 __ub__ 数组，认证现在全部考这种写法，不要再用旧数组写法。

6. 算子注册（固定死，直接默写）

这个可能有问题，不用考虑 REG_KERNEL(核函数名, 0, 自定义名称);

第二个参数考试统一填 0。

五、考场实操极简流程（照着走就能拿分）

新建 .cpp 文件
粘贴头文件 + using namespace AscendC;
仿写上面的核函数（改计算逻辑即可，框架不动）
补充边界判断 + DataCopy + 计算API
这个可能有问题，不用考虑最后一行加上 REG_KERNEL 注册
在线平台点击编译 → 提交

六、下一步练习建议（零基础循序渐进）

先把上面 Add 算子跑通、默写下来（保底题）
再改一版 Mul 乘法算子（只把 Add() 改成 Mul()，其余全不变）
熟悉后再看分块、双缓冲、简单矩阵运算

模拟试卷一

我来为你生成一份昇腾CANN算子开发类认证（HCIP/微认证）模拟考试试卷，基于你提供的备考资料中的核心考点。试卷包含理论题和实操题，完全贴合CANN 8.x版本考试要求。

📝 昇腾CANN算子开发类认证模拟考试试卷

考试时间：90分钟　　满分：100分　　及格线：60分

第一部分：选择题（每题2分，共30分）

1. CANN五层架构中，负责算子开发和编译优化的是哪一层？

A. 应用使能层
B. 框架层
C. 算子开发层
D. 运行时层

2. AI Core的三大计算单元中，负责矩阵运算的是？

A. Vector单元
B. Scalar单元
C. Cube单元
D. DMA单元

3. 昇腾AI处理器中，读写速度最快的存储层级是？

A. Global Memory (GM)
B. L2缓存
C. L1缓存
D. Unified Buffer (UB)

4. 在CANN 8.x中，核函数的标准声明格式是？

A. __global__ void kernelName(...)
B. __aicore__ void kernelName(...)
C. __global__ __aicore__ void kernelName(...)
D. __kernel__ void kernelName(...)

5. SPMD并行模型中，获取当前Block索引的API是？

A. GetThreadIdx()
B. GetBlockIdx()
C. GetGridIdx()
D. GetWarpIdx()

6. 下列关于DataCopy的说法，错误的是？

A. DataCopy用于在GM和UB之间搬运数据
B. 所有DataCopy的长度和地址必须32字节对齐
C. DataCopy可以在UB内部直接进行计算
D. 数据搬运是实操题中最常见的报错来源

7. CANN 8.x中，LocalTensor的设置方式是？

A. __ub__ half ubSrc[1024];
B. LocalTensor<half> ubSrc; ubSrc.SetBuffer(...);
C. Tensor<half> ubSrc(1024);
D. Buffer<half> ubSrc = GM2UBAddr(0,0);

8. 算子注册的固定写法是？

A. REGISTER_KERNEL(kernelName, 0, funcName);
B. 这个可能有问题，不用考虑 REG_KERNEL(kernelName, 0, funcName);
C. KERNEL_REG(kernelName, 0, funcName);
D. OP_REGISTER(kernelName, 0, funcName);

9. 模型迁移适配的标准流程顺序是？

A. 评估→迁移→调试→部署
B. 迁移→评估→调试→部署
C. 评估→调试→迁移→部署
D. 迁移→调试→评估→部署

10. msprof工具的主要用途是？

A. 编译算子代码
B. 采集性能数据、定位算子瓶颈
C. 自动生成算子注册代码
D. 管理CANN版本升级

11. CATLASS高性能矩阵模板主要应用于？

A. 图像预处理
B. 大模型训练/推理
C. 数据加密解密
D. 网络通信加速

12. 在自定义算子开发流程中，正确的顺序是？

A. 核函数编写→计算逻辑→内存搬运→结果回写→算子注册
B. 核函数编写→内存搬运→计算逻辑→结果回写→算子注册
C. 内存搬运→核函数编写→计算逻辑→结果回写→算子注册
D. 算子注册→核函数编写→内存搬运→计算逻辑→结果回写

13. GEMM算子优化中，"双缓冲技术"的主要作用是？

A. 减少核函数参数数量
B. 隐藏数据搬运延迟，提升计算吞吐
C. 简化算子注册流程
D. 降低UB内存占用

14. 性能调优的三大通用方案不包括？

A. 优化数据搬运
B. 保证内存对齐
C. 优先使用Scalar逐元素运算
D. 优先使用Cube/Vector批量运算

15. 在线考试实操题中，以下哪项是红线要求？

A. 必须手写完整Makefile
B. 可以修改官方工程目录结构
C. 所有DataCopy必须保证32字节对齐
D. 必须使用旧版本__ub__数组写法

第二部分：判断题（每题2分，共20分）

16. CANN五层架构从上到下依次是：应用使能层→框架层→算子开发层→运行时层→驱动/硬件层。（　）

17. Scalar单元主要负责标量运算和控制流处理。（　）

18. 标准数据流向是：UB → L1 → L2 → Global Memory。（　）

19. 在核函数中，if (offset >= totalLen) return; 这样的边界判断是可选的，不影响编译。（　）

20. CANN 8.x版本兼容使用旧版的__ub__数组写法。（　）

21. 自动迁移工具可以处理所有模型的迁移，无需人工干预。（　）

22. HCCL是昇腾的集合通信库，用于多卡分布式训练。（　）

23. 在线考试实操题只要求编写核函数+注册语句，不需要自己搭建完整工程。（　）

24. 大矩阵分块存入UB是GEMM算子优化的核心思路之一。（　）

25. 理论题建议控制在30分钟内完成，剩余时间留给代码编写。（　）

第三部分：简答题（每题5分，共20分）

26. 简述CANN五层架构的各层名称及主要作用。（5分）

27. 请写出CANN 8.x核函数的标准声明格式，并说明__global__和__aicore__两个修饰符的作用。（5分）

28. 简述自定义算子开发的完整流程（5个步骤）。（5分）

29. 请列举msprof工具的三个主要功能或用途。（5分）

第四部分：实操编程题（30分）

30. 向量乘法算子开发（MulKernel）

题目要求：

基于以下Add加法算子的代码框架，完成**向量乘法算子（MulKernel）**的开发。要求：

使用CANN 8.x标准写法
正确处理边界保护
数据搬运保证32字节对齐
完成算子注册

已知条件：

输入：两个半精度(half)数组 src0 和 src1
输出：半精度(half)数组 dst
单块处理数据量：eachBlockLen = 1024
使用 Mul() API进行向量乘法计算

请补全以下代码中的空白处（①~⑥）：

#include "kernel_operator.h"
using namespace AscendC;

// 核函数声明
__global__ __aicore__ void MulKernel(
    GM_ADDR const __restrict__ src0,
    GM_ADDR const __restrict__ src1,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    // ① 获取当前Block编号（1分）
    uint32_t blockIdx = __________;
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    // ② 边界保护（2分）
    __________ {
        return;
    }

    // ③ 申请UB本地张量（3分）
    LocalTensor<half> ubSrc0;
    LocalTensor<half> ubSrc1;
    LocalTensor<half> ubDst;

    ubSrc0.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    __________;
    __________;

    // ④ GM -> UB 数据搬运（3分）
    DataCopy(ubSrc0, src0 + offset, eachBlockLen * sizeof(half));
    __________;

    // ⑤ 向量乘法计算（3分）
    __________;

    // ⑥ UB -> GM 结果写回（3分）
    __________;
}

// 算子注册（5分）
__________;

附加要求（10分）：
请写出编译该算子的ascendc命令（需指定soc-version为ascend910b）。

✅ 参考答案

点击展开查看答案

第一部分：选择题答案

题号	答案	解析
1	C	算子开发层负责算子开发、编译优化
2	C	Cube单元专门负责矩阵运算
3	D	UB是本地缓存，读写速度最快
4	C	CANN 8.x必须写全`__global__ __aicore__`
5	B	`GetBlockIdx()`获取当前Block索引
6	C	DataCopy只负责搬运，计算在UB中用Add/Mul等API完成
7	B	8.x标准写法：LocalTensor + SetBuffer
8	B	这个可能有问题，不用考虑 `REG_KERNEL(核函数名, 0, 自定义名称);`
9	A	标准流程：评估→迁移→调试→部署
10	B	msprof用于性能分析和瓶颈定位
11	B	CATLASS用于大模型训练/推理加速
12	B	正确顺序：核函数→内存搬运→计算逻辑→结果回写→注册
13	B	双缓冲隐藏搬运延迟，提升吞吐
14	C	应优先使用Cube/Vector批量运算，而非Scalar
15	C	32字节对齐是实操红线要求

第二部分：判断题答案

题号	答案	解析
16	√	五层架构顺序正确
17	√	Scalar负责标量运算和控制流
18	×	正确流向：GM → L2 → L1 → UB
19	×	边界判断是必须的，防止数组越界
20	×	8.x必须使用LocalTensor，旧写法废弃
21	×	自动迁移无法处理所有情况，需人工干预
22	√	HCCL是昇腾集合通信库
23	√	实操题只需写核函数+注册
24	√	分块是GEMM优化核心思路
25	√	理论题控制在30分钟内

第三部分：简答题参考答案

26. CANN五层架构（5分）

层级	作用
应用使能层	提供行业应用SDK，简化业务开发
框架层	支持TensorFlow、PyTorch等主流框架
算子开发层	算子开发、编译优化、自定义算子
运行时层	任务调度、内存管理、执行控制
驱动/硬件层	底层驱动、AI Core硬件资源管理

27. 核函数声明格式（5分）

__global__ __aicore__ void 函数名(参数...)

__global__：表示该函数为全局核函数，可从Host端调用
__aicore__：表示该函数在AI Core上执行

28. 自定义算子开发流程（5分）

核函数编写
内存搬运（GM→UB）
计算逻辑实现
结果回写（UB→GM）
算子注册

29. msprof工具功能（5分，任意3点）

采集算子性能数据
查看算子执行耗时
识别数据搬运瓶颈
识别计算瓶颈
定位性能优化方向

第四部分：实操编程题答案

30. 向量乘法算子（MulKernel）

#include "kernel_operator.h"
using namespace AscendC;

// 核函数声明
__global__ __aicore__ void MulKernel(
    GM_ADDR const __restrict__ src0,
    GM_ADDR const __restrict__ src1,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    // ① 获取当前Block编号
    uint32_t blockIdx = GetBlockIdx();
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    // ② 边界保护
    if (offset >= totalLen) {
        return;
    }

    // ③ 申请UB本地张量
    LocalTensor<half> ubSrc0;
    LocalTensor<half> ubSrc1;
    LocalTensor<half> ubDst;

    ubSrc0.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    ubSrc1.SetBuffer((half *)GM2UBAddr(1, 0), eachBlockLen);
    ubDst.SetBuffer((half *)GM2UBAddr(2, 0), eachBlockLen);

    // ④ GM -> UB 数据搬运
    DataCopy(ubSrc0, src0 + offset, eachBlockLen * sizeof(half));
    DataCopy(ubSrc1, src1 + offset, eachBlockLen * sizeof(half));

    // ⑤ 向量乘法计算
    Mul(ubDst, ubSrc0, ubSrc1, eachBlockLen);

    // ⑥ UB -> GM 结果写回
    DataCopy(dst + offset, ubDst, eachBlockLen * sizeof(half));
}

// 算子注册
这个可能有问题，不用考虑 REG_KERNEL(MulKernel, 0, MulFunc);

编译命令：

ascendc \
--soc-version=ascend910b \
--kernel-name=MulKernel \
--output=MulKernel.o \
MulKernel.cpp

📊 评分建议

分数段	评价
90-100分	🟢 优秀，可以直接上考场
75-89分	🟡 良好，再巩固薄弱模块
60-74分	🟠 及格，需要针对性复习错题
<60分	🔴 需重新学习理论+实操基础

第二套模拟试卷（侧重不同考点，比如GEMM优化、性能调优类题目），

好的，以下是第二套模拟试卷，侧重GEMM优化、性能调优、工具使用和边界场景，与第一套形成互补。

📝 昇腾CANN算子开发类认证模拟考试试卷（第二套）

考试时间：90分钟　　满分：100分　　及格线：60分

第一部分：选择题（每题2分，共30分）

1. 在昇腾AI处理器中，负责向量运算的AI Core单元是？

A. Cube
B. Scalar
C. Vector
D. DMA

2. CANN 8.x中，以下哪种写法是错误的核函数声明？

A. __global__ __aicore__ void MyKernel(...)
B. __aicore__ __global__ void MyKernel(...)
C. __global__ void MyKernel(...)
D. __global__ __aicore__ void MyKernel(GM_ADDR src, ...)

3. 关于数据搬运对齐，以下说法正确的是？

A. 只有GM地址需要32字节对齐，长度不需要
B. DataCopy的长度和地址都必须32字节对齐
C. UB内部操作不需要任何对齐
D. 对齐要求只在CANN 7.x中存在，8.x已取消

4. GEMM算子实现中，"大矩阵分块存入UB"的主要目的是？

A. 减少核函数参数数量
B. 适配UB有限容量，实现分块计算
C. 简化算子注册语句
D. 避免使用LocalTensor

5. 以下关于双缓冲（Double Buffer）技术的描述，正确的是？

A. 双缓冲会占用双倍GM空间
B. 双缓冲通过并行执行数据搬运和计算，隐藏搬运延迟
C. 双缓冲技术只能在L2缓存中使用
D. 双缓冲会降低计算吞吐量

6. GM2UBAddr(0, 0) 中两个参数的含义是？

A. (Block编号, 偏移量)
B. (Buffer编号, 偏移量)
C. (GM地址, UB地址)
D. (线程编号, 数据长度)

7. 性能调优时，以下哪种做法不能有效提升算子性能？

A. 将频繁访问的数据驻留在UB中复用
B. 使用Scalar单元逐元素处理大量数据
C. 采用双缓冲隐藏数据搬运延迟
D. 对数据进行Padding对齐以提升硬件利用率

8. 关于Stream/Queue/Event同步机制，以下说法错误的是？

A. Stream是任务执行的队列
B. 同一Stream内的任务按顺序执行
C. Event用于不同Stream之间的同步
D. Queue和Stream是完全相同的概念，可以互换使用

9. 在算子性能分析中，若msprof显示"数据搬运耗时占比80%"，优先优化方向是？

A. 增加计算指令数量
B. 优化数据搬运策略，提升数据复用率
C. 减少核函数Block数量
D. 更换更高精度的数据类型

10. 以下关于自动迁移和手动迁移的说法，正确的是？

A. 自动迁移适用于所有场景，效果总是优于手动迁移
B. 手动迁移灵活性高，但需要开发者深入理解模型和框架差异
C. 自动迁移和手动迁移不能结合使用
D. 手动迁移不需要评估阶段，直接修改代码即可

11. CANN 8.x中，LocalTensor相比旧版__ub__数组的优势不包括？

A. 更安全的内存管理
B. 支持动态大小调整
C. 更符合现代C++编程范式
D. 编译器可以更好地进行优化分析

12. 在SPMD并行模型中，若总数据量为4096，每个Block处理1024，则需要的Block数量是？

A. 2
B. 4
C. 8
D. 16

13. 数据Padding对齐的主要作用是？

A. 减少数据总量
B. 使数据长度满足硬件对齐要求，提升访存效率
C. 加密保护数据安全
D. 压缩数据存储空间

14. 以下关于GetBlockIdx()的说法，错误的是？

A. 返回当前Block的全局索引
B. 不同Block返回的索引值一定不同
C. 索引从0开始递增
D. 可以用于计算当前Block处理的数据偏移

15. 在线考试实操题中，编译成功后会生成的文件是？

A. .exe 可执行文件
B. .o 目标文件
C. .so 动态库文件
D. .py Python脚本

第二部分：判断题（每题2分，共20分）

16. AI Core的Cube单元既可以进行矩阵运算，也可以进行向量运算。（　）

17. 在CANN 8.x中，这个可能有问题，不用考虑 REG_KERNEL的第二个参数考试时必须填写0。（　）

18. 数据从GM搬运到UB的速度比从L1搬运到UB的速度更快。（　）

19. 边界判断if (offset >= totalLen)中的条件可以写成if (offset > totalLen)而不影响正确性。（　）

20. 性能调优时，应优先保证功能正确性，再考虑性能优化。（　）

21. DataCopy的第三个参数表示要搬运的字节数。（　）

22. 双缓冲技术需要至少两个独立的UB Buffer来回切换使用。（　）

23. 模型迁移适配的"评估"阶段主要是评估模型是否可以在昇腾上运行，以及预期的性能表现。（　）

24. 在核函数中，可以直接对GM地址进行读写操作，不需要通过DataCopy搬运到UB。（　）

25. AscendC::Add() API支持在UB中执行向量加法运算。（　）

第三部分：简答题（每题5分，共20分）

26. 请简述AI Core三大计算单元（Cube、Vector、Scalar）的分工，并各举一个适用场景。（5分）

27. 解释什么是"数据搬运优化"和"UB复用"，并说明它们在算子性能调优中的作用。（5分）

28. 简述GEMM算子优化的三个核心思路（大矩阵分块、双缓冲、数据Padding），并说明每个思路解决的主要问题。（5分）

29. 在实操题中，如果编译报错"头文件kernel_operator.h找不到"，请列出至少两种可能的原因及解决方法。（5分）

第四部分：实操编程题（30分）

30. 带边界保护的通用向量加法算子

题目要求：

以下代码是一个不完整的向量加法算子，存在多处错误和遗漏。请找出问题并修正，使其成为符合CANN 8.x标准的正确代码。

已知条件：

输入：两个半精度(half)数组 src0 和 src1
输出：半精度(half)数组 dst
总长度 totalLen 不一定是1024的整数倍
单块处理数据量：eachBlockLen = 1024
关键要求：必须正确处理最后一块的边界情况（数据量不足1024时）

待修正代码：

#include "kernel_operator.h"
using namespace AscendC;

__global__ void AddKernel(
    GM_ADDR const __restrict__ src0,
    GM_ADDR const __restrict__ src1,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    uint32_t blockIdx = GetBlockIdx();
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    LocalTensor<half> ubSrc0;
    LocalTensor<half> ubSrc1;
    LocalTensor<half> ubDst;

    ubSrc0.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    ubSrc1.SetBuffer((half *)GM2UBAddr(1, 0), eachBlockLen);
    ubDst.SetBuffer((half *)GM2UBAddr(2, 0), eachBlockLen);

    DataCopy(ubSrc0, src0 + offset, eachBlockLen * sizeof(half));
    DataCopy(ubSrc1, src1 + offset, eachBlockLen * sizeof(half));

    Add(ubDst, ubSrc0, ubSrc1, eachBlockLen);

    DataCopy(dst + offset, ubDst, eachBlockLen * sizeof(half));
}

 这个可能有问题，不用考虑 REG_KERNEL(AddKernel, 1, AddFunc);

请回答以下问题：

(1) 找出并修正核函数声明中的错误。（3分）

(2) 当前代码缺少边界保护，请补充正确的边界判断逻辑，并说明为什么需要处理totalLen不是1024整数倍的情况。（5分）

(3) 当最后一块数据量不足1024时，DataCopy和Add的长度参数应该如何处理？请写出修正后的核心代码段。（10分）

(4) 算子注册语句中存在什么错误？请写出正确的注册语句。（2分）

(5) 请写出完整的、修正后的核函数代码。（10分）

✅ 参考答案

点击展开查看答案

第一部分：选择题答案

题号	答案	解析
1	C	Vector单元负责向量运算
2	C	缺少`__aicore__`修饰符，CANN 8.x必须两个都写
3	B	DataCopy的长度和地址都必须32字节对齐
4	B	分块是为了适配UB有限容量
5	B	双缓冲并行执行搬运和计算，隐藏延迟
6	B	(Buffer编号, 偏移量)
7	B	Scalar逐元素处理效率低，应使用Vector批量处理
8	D	Queue和Stream概念相关但不完全相同
9	B	搬运耗时高应优化搬运策略
10	B	手动迁移灵活但需要深入理解
11	B	LocalTensor大小在SetBuffer时确定，不支持动态调整
12	B	4096/1024 = 4个Block
13	B	Padding使数据满足对齐要求，提升访存效率
14	B	若Block数量超过数据量，部分Block可能不工作，但索引仍不同（此题有歧义，核心考点是索引从0开始）
15	B	编译生成`.o`目标文件

第14题补充说明：实际上不同Block的索引值确实不同（从0递增），但如果Block数量设置过多，超出数据范围的Block会提前return。此题主要考察对边界情况的理解。

第二部分：判断题答案

题号	答案	解析
16	×	Cube专用于矩阵运算，向量运算由Vector处理
17	√	考试统一填0
18	×	L1比GM更接近计算单元，搬运更快
19	×	`>=`包含等于情况，若用`>`会漏掉offset恰好等于totalLen的边界
20	√	先功能正确，再性能优化
21	√	第三个参数是字节数
22	√	双缓冲需要两个Buffer交替使用
23	√	评估阶段判断可行性和预期性能
24	×	GM不能直接读写，必须通过DataCopy到UB
25	√	Add()是UB中的向量加法API

第三部分：简答题参考答案

26. AI Core三大计算单元分工（5分）

单元	分工	适用场景
Cube	矩阵运算（GEMM等）	大模型训练中的矩阵乘法
Vector	向量运算（逐元素操作）	向量加法、激活函数计算
Scalar	标量运算、控制流	循环控制、条件判断、索引计算

27. 数据搬运优化与UB复用（5分）

数据搬运优化：减少GM和UB之间的数据搬运次数，通过合理的数据排布和访问模式降低搬运开销。
UB复用：将频繁访问的数据驻留在UB中，避免重复从GM搬运，提升数据局部性。
作用：数据搬运是算子性能的主要瓶颈之一，优化搬运和复用UB数据可以显著提升算子执行效率。

28. GEMM优化三个核心思路（5分）

思路	解决的问题
大矩阵分块存入UB	UB容量有限，无法一次性加载大矩阵，分块后逐块计算
双缓冲技术	数据搬运和计算串行执行导致计算单元空闲，双缓冲并行化两者
数据Padding对齐	原始数据长度不满足硬件对齐要求，Padding后提升访存和计算效率

29. 头文件找不到的原因及解决（5分，任意两种）

原因	解决方法
CANN环境未正确加载	执行`source /usr/local/Ascend/ascend-toolkit/set_env.sh`
环境版本不是CANN 8.x	确认环境版本，升级到CANN 8.x
头文件路径未包含	检查编译命令是否包含正确的头文件搜索路径
使用在线环境时未进入正确目录	确认当前工作目录在CANN工程目录下

第四部分：实操编程题答案

30. 带边界保护的通用向量加法算子

(1) 核函数声明修正（3分）

原代码：__global__ void AddKernel(...)

错误：缺少__aicore__修饰符。

修正：

__global__ __aicore__ void AddKernel(...)

(2) 边界保护逻辑（5分）

需要补充的边界判断：

if (offset >= totalLen) {
    return;
}

为什么需要处理：当totalLen不是1024的整数倍时，最后一个Block处理的数据量会少于1024。如果不处理，会导致：

访问越界（读取/写入超出数组范围）
DataCopy长度不对齐或超出实际数据量
计算结果错误

(3) 最后一块数据处理（10分）

需要计算实际处理长度：

uint32_t actualLen = eachBlockLen;
if (offset + eachBlockLen > totalLen) {
    actualLen = totalLen - offset;
}

然后所有使用eachBlockLen的地方替换为actualLen，并确保actualLen * sizeof(half)满足32字节对齐要求（考试中若数据量保证对齐可简化处理）。

(4) 算子注册修正（2分）

原代码：这个可能有问题，不用考虑 REG_KERNEL(AddKernel, 1, AddFunc);

错误：第二个参数应为0。

修正：

这个可能有问题，不用考虑 REG_KERNEL(AddKernel, 0, AddFunc);

(5) 完整修正代码（10分）

#include "kernel_operator.h"
using namespace AscendC;

__global__ __aicore__ void AddKernel(
    GM_ADDR const __restrict__ src0,
    GM_ADDR const __restrict__ src1,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    uint32_t blockIdx = GetBlockIdx();
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    // 边界保护：超出范围直接返回
    if (offset >= totalLen) {
        return;
    }

    // 计算当前Block实际处理的数据长度
    uint32_t actualLen = eachBlockLen;
    if (offset + eachBlockLen > totalLen) {
        actualLen = totalLen - offset;
    }

    // 申请UB本地张量
    LocalTensor<half> ubSrc0;
    LocalTensor<half> ubSrc1;
    LocalTensor<half> ubDst;

    ubSrc0.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    ubSrc1.SetBuffer((half *)GM2UBAddr(1, 0), eachBlockLen);
    ubDst.SetBuffer((half *)GM2UBAddr(2, 0), eachBlockLen);

    // GM -> UB 数据搬运（使用actualLen）
    DataCopy(ubSrc0, src0 + offset, actualLen * sizeof(half));
    DataCopy(ubSrc1, src1 + offset, actualLen * sizeof(half));

    // 向量加法计算（使用actualLen）
    Add(ubDst, ubSrc0, ubSrc1, actualLen);

    // UB -> GM 结果写回（使用actualLen）
    DataCopy(dst + offset, ubDst, actualLen * sizeof(half));
}

// 算子注册（第二个参数固定为0）
这个可能有问题，不用考虑 REG_KERNEL(AddKernel, 0, AddFunc);

注意：实际考试中若数据量保证32字节对齐，可直接使用actualLen。若不确定，可能需要对actualLen进行向上取整到32字节对齐。

📊 两套试卷考点对比

考点模块	第一套侧重	第二套侧重
架构理论	✅ 五层架构、AI Core单元	✅ 单元分工、同步机制
存储层级	✅ 数据流向、速度对比	✅ 各级缓存特性、对齐要求
核函数格式	✅ 标准写法	✅ 常见错误、边界场景
并行编程	✅ SPMD、GetBlockIdx	✅ Block数量计算、索引理解
模板库	✅ CATLASS应用场景	-
数据搬运	✅ DataCopy铁律	✅ 边界处理、实际长度计算
GEMM优化	✅ 核心思路简述	✅ 分块、双缓冲、Padding深入
性能调优	✅ 三大通用方案	✅ 瓶颈分析、优化方向判断
工具使用	✅ msprof用途	✅ 报错排查、结果分析
模型迁移	✅ 流程、自动/手动对比	✅ 评估阶段、结合使用
实操难度	⭐⭐⭐ 基础算子（Mul）	⭐⭐⭐⭐ 边界处理+代码纠错

*第三套**（侧重真题风格，全是综合应用题），

好的，以下是第三套模拟试卷，采用真题综合应用风格，题目更贴近实际考试场景，融合多个知识点进行考察。

📝 昇腾CANN算子开发类认证模拟考试试卷（第三套 · 真题风格）

考试时间：90分钟　　满分：100分　　及格线：60分

第一部分：单选题（每题3分，共30分）

1. 某算子在msprof分析中显示"Cube利用率15%，Vector利用率85%，DMA带宽利用率90%"，优先优化方向是？

A. 减少Vector计算量，增加Cube矩阵运算
B. 优化数据搬运策略，降低DMA压力
C. 增加Scalar控制逻辑复杂度
D. 扩大UB缓冲区容量

2. 以下核函数代码片段存在什么问题？

__global__ __aicore__ void MyKernel(GM_ADDR src, GM_ADDR dst, uint32_t len) {
    uint32_t idx = GetBlockIdx();
    uint32_t offset = idx * 1024;
    LocalTensor<half> ub;
    ub.SetBuffer((half *)GM2UBAddr(0, 0), 2048);
    DataCopy(ub, src + offset, 1024 * sizeof(half));
    Add(ub, ub, ub, 1024);
    DataCopy(dst + offset, ub, 1024 * sizeof(half));
}

A. 缺少边界判断，且Add的源和目的相同
B. 核函数声明格式错误
C. UB缓冲区申请太小
D. 没有使用双缓冲技术

3. 在CANN 8.x中，将PyTorch模型迁移到昇腾平台，以下步骤正确的是？

A. 直接修改模型源码中的算子实现 → 编译运行 → 调试
B. 评估模型算子支持度 → 使用ATC工具自动迁移 → 手动修复不支持的算子 → 精度调试 → 性能优化
C. 先用msprof分析原模型性能 → 直接部署到昇腾 → 对比性能差异
D. 将模型转换为ONNX格式 → 直接调用昇腾推理接口

4. 某GEMM算子M=2048, N=2048, K=2048，UB可用空间为64KB（half精度），最佳分块策略是？

A. 不分块，一次性加载全部数据
B. 将K维度分块，每次加载部分K到UB计算
C. 将M和N维度分块，每次计算一个小块矩阵
D. 只将M维度分块，N和K全量加载

5. 以下关于DataCopy对齐的说法，正确的是？

A. 地址对齐即可，长度可以任意
B. 长度对齐即可，地址可以任意
C. 地址和长度都必须满足32字节对齐
D. 只有在搬运half类型时才需要对齐

6. 自定义算子开发中，以下哪项不属于算子注册的必要信息？

A. 核函数名称
B. 算子输入输出Tensor描述
C. 第二个固定参数0
D. 自定义算子名称

7. 在SPMD并行模型中，若设置BlockDim=8，但数据总量只需4个Block即可处理完，会发生什么？

A. 程序报错，Block数量必须等于数据块数
B. 后4个Block的offset >= totalLen，通过边界判断直接return
C. 后4个Block会重复计算前4个Block的数据
D. 系统自动调整BlockDim为4

8. 以下代码中，actualLen的计算方式是否正确？（totalLen=2500, eachBlockLen=1024）

uint32_t actualLen = eachBlockLen;
if (offset + eachBlockLen > totalLen) {
    actualLen = totalLen - offset;
}

A. 正确，能处理所有边界情况
B. 错误，actualLen可能不满足32字节对齐
C. 错误，应该使用>=而不是>
D. 错误，offset计算方式不对

9. 性能调优中，"计算隐藏搬运"的核心思想是？

A. 让计算单元在数据搬运时保持空闲以节省功耗
B. 通过流水线并行，使计算和数据搬运同时进行
C. 减少计算量以降低对数据的需求
D. 增加搬运次数以分散带宽压力

10. 以下关于CATLASS模板库的说法，错误的是？

A. CATLASS提供了高性能矩阵运算的抽象接口
B. CATLASS主要用于大模型训练/推理场景
C. CATLASS可以自动生成最优的GEMM分块策略
D. CATLASS是昇腾自研的，与CUDA生态无关

第二部分：多选题（每题4分，共20分，漏选得2分，错选不得分）

11. 以下哪些情况会导致算子编译失败？

A. 核函数缺少__aicore__修饰符
B. 使用了CANN 7.x的__ub__数组写法
C. 这个可能有问题，不用考虑 REG_KERNEL第二个参数写成了1
D. 头文件kernel_operator.h未找到

12. 在昇腾AI处理器中，以下关于存储层级的说法正确的是？

A. GM容量最大，速度最慢
B. L1缓存可以直接被AI Core计算单元访问
C. UB是AI Core内部的本地缓存，速度最快
D. L2缓存是所有AI Core共享的

13. 以下哪些属于算子性能调优的有效手段？

A. 使用双缓冲隐藏数据搬运延迟
B. 将数据Padding到32字节对齐
C. 优先使用Vector批量运算替代Scalar逐元素运算
D. 减少不必要的GM到UB的数据搬运

14. 模型迁移适配过程中，"评估阶段"需要完成的工作包括？

A. 分析模型中各算子在昇腾平台的支持度
B. 预估模型在昇腾上的性能表现
C. 直接开始修改模型源码
D. 识别需要自定义开发的算子

15. 以下关于LocalTensor的说法正确的是？

A. 必须通过SetBuffer方法绑定UB内存地址
B. 可以动态调整绑定的缓冲区大小
C. 是CANN 8.x推荐的标准写法
D. 支持多种数据类型（half、float、int32等）

第三部分：综合应用题（共50分）

16. 算子开发综合分析题（20分）

某开发者需要实现一个向量乘加算子（FusedMulAdd），功能为：dst[i] = src0[i] * src1[i] + src2[i]，数据类型为half。

已知条件：

输入：三个GM数组 src0、src1、src2
输出：一个GM数组 dst
总长度 totalLen 任意（不一定是1024整数倍）
单Block处理1024个元素
UB可用空间充足

请完成以下任务：

(1) 画出该算子的数据流向图（从GM到UB再到GM），标注关键步骤。（5分）

(2) 写出完整的核函数代码，要求：

使用CANN 8.x标准写法
正确处理边界情况
数据搬运和计算长度使用实际长度
包含算子注册语句
（10分）

(3) 如果要对该算子进行性能优化，请给出至少2个优化方向（无需写代码，说明思路即可）。（5分）

17. 性能分析与调优实战题（15分）

某算子经msprof分析后得到以下性能数据：

指标	数值
总执行时间	120μs
Cube计算时间	10μs
Vector计算时间	15μs
DMA搬运时间（GM→UB）	45μs
DMA搬运时间（UB→GM）	40μs
Scalar控制时间	10μs

请回答：

(1) 计算DMA搬运时间占总执行时间的百分比，并判断性能瓶颈所在。（3分）

(2) 针对该瓶颈，提出至少3个具体的优化方案。（9分）

(3) 优化后预期DMA占比降低到30%以下，估算优化后的总执行时间（假设其他时间不变）。（3分）

18. 代码纠错与完善题（15分）

以下是一个开发者提交的ReduceSum算子代码（对向量求和），存在多处错误。请找出所有错误并给出修正后的完整代码。

原始代码：

#include "kernel_operator.h"
using namespace AscendC;

__global__ void ReduceSum(GM_ADDR src, GM_ADDR dst, uint32_t len) {
    uint32_t blockIdx = GetBlockIdx();
    uint32_t offset = blockIdx * 1024;
    
    if (offset > len) {
        return;
    }
    
    LocalTensor<half> ubSrc;
    LocalTensor<half> ubDst;
    
    ubSrc.SetBuffer((half *)GM2UBAddr(0, 0), 1024);
    ubDst.SetBuffer((half *)GM2UBAddr(1, 0), 1);
    
    DataCopy(ubSrc, src + offset, 1024);
    
    half sum = 0;
    for (int i = 0; i < 1024; i++) {
        sum += ubSrc.GetValue(i);
    }
    
    ubDst.SetValue(0, sum);
    DataCopy(dst + blockIdx, ubDst, 1);
}

这个可能有问题，不用考虑 
REG_KERNEL(ReduceSum, 1, ReduceSumFunc);

错误提示：

该算子意图实现：每个Block对其负责的1024个元素求和，结果写入dst对应位置
但代码存在多处编译或逻辑错误

请完成：

(1) 找出至少5处错误，并说明错误原因。（8分）

(2) 写出修正后的完整代码。（7分）

✅ 参考答案

点击展开查看答案

第一部分：单选题答案

题号	答案	解析
1	B	DMA带宽利用率90%是主要瓶颈，应优化数据搬运
2	A	缺少边界判断，且Add(src, src, src)语义错误（应为三个不同Tensor）
3	B	标准迁移流程：评估→自动迁移→手动修复→调试→优化
4	C	M×N×K=8GB远超UB容量，必须对M和N分块
5	C	地址和长度都必须32字节对齐
6	B	算子注册只需核函数名、固定参数0、自定义名称，输入输出描述在算子原型中
7	B	多余Block通过边界判断直接return
8	B	2500-2048=452，452×2=904字节，不满足32字节对齐（实际考试数据通常保证对齐）
9	B	流水线并行，计算和搬运同时进行
10	D	CATLASS借鉴了CUTLASS设计理念，并非完全无关

第二部分：多选题答案

题号	答案	解析
11	ABD	A缺少修饰符编译失败；B旧写法8.x不支持；D头文件找不到编译失败；C第二个参数1不会导致编译失败（但不符合考试规范）
12	ACD	A正确；B错误，L1需通过UB中转；C正确；D正确
13	ABCD	全部正确
14	ABD	C是迁移阶段的工作，不是评估阶段
15	ACD	B错误，SetBuffer后大小固定，不能动态调整

第三部分：综合应用题答案

16. 算子开发综合分析题（20分）

(1) 数据流向图（5分）

┌─────────┐     DataCopy     ┌─────────┐
│  src0   │ ───────────────→ │  ubSrc0 │
│  (GM)   │                  │  (UB)   │
└─────────┘                  └─────────┘
                                    ↘
┌─────────┐     DataCopy     ┌─────────┐    Mul    ┌─────────┐
│  src1   │ ───────────────→ │  ubSrc1 │ ────────→ │  ubTmp  │
│  (GM)   │                  │  (UB)   │           │  (UB)   │
└─────────┘                  └─────────┘           └────┬────┘
                                                        │ Add
┌─────────┐     DataCopy     ┌─────────┐              ↓
│  src2   │ ───────────────→ │  ubSrc2 │ ─────────→ │  ubDst  │
│  (GM)   │                  │  (UB)   │              │  (UB)  │
└─────────┘                  └─────────┘              └──┬────┘
                                                         │ DataCopy
                                                         ↓
                                                    ┌─────────┐
                                                    │   dst   │
                                                    │  (GM)   │
                                                    └─────────┘

关键步骤：

GM→UB：搬运src0、src1、src2到UB
UB计算：Mul(ubSrc0, ubSrc1) → ubTmp；Add(ubTmp, ubSrc2) → ubDst
UB→GM：将ubDst写回dst

(2) 完整核函数代码（10分）

#include "kernel_operator.h"
using namespace AscendC;

__global__ __aicore__ void FusedMulAddKernel(
    GM_ADDR const __restrict__ src0,
    GM_ADDR const __restrict__ src1,
    GM_ADDR const __restrict__ src2,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    uint32_t blockIdx = GetBlockIdx();
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    // 边界保护
    if (offset >= totalLen) {
        return;
    }

    // 计算实际处理长度
    uint32_t actualLen = eachBlockLen;
    if (offset + eachBlockLen > totalLen) {
        actualLen = totalLen - offset;
    }

    // 申请UB张量
    LocalTensor<half> ubSrc0;
    LocalTensor<half> ubSrc1;
    LocalTensor<half> ubSrc2;
    LocalTensor<half> ubDst;

    ubSrc0.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    ubSrc1.SetBuffer((half *)GM2UBAddr(1, 0), eachBlockLen);
    ubSrc2.SetBuffer((half *)GM2UBAddr(2, 0), eachBlockLen);
    ubDst.SetBuffer((half *)GM2UBAddr(3, 0), eachBlockLen);

    // GM -> UB 数据搬运
    DataCopy(ubSrc0, src0 + offset, actualLen * sizeof(half));
    DataCopy(ubSrc1, src1 + offset, actualLen * sizeof(half));
    DataCopy(ubSrc2, src2 + offset, actualLen * sizeof(half));

    // 计算：dst = src0 * src1 + src2
    // 先乘后加，使用ubDst作为中间结果和最终结果
    Mul(ubDst, ubSrc0, ubSrc1, actualLen);
    Add(ubDst, ubDst, ubSrc2, actualLen);

    // UB -> GM 结果写回
    DataCopy(dst + offset, ubDst, actualLen * sizeof(half));
}

// 算子注册
这个可能有问题，不用考虑 REG_KERNEL(FusedMulAddKernel, 0, FusedMulAddFunc);

(3) 优化方向（5分，任意2个）

优化方向	说明
算子融合	将Mul和Add融合为单个FMA指令（如果硬件支持），减少一次UB读写
双缓冲	使用两组UB Buffer，当前Block计算时预取下一批数据，隐藏搬运延迟
数据复用	若src2为广播数据，避免重复搬运
向量化加载	使用更大的向量宽度（如float16x8）提升访存效率

17. 性能分析与调优实战题（15分）

(1) DMA占比计算（3分）

DMA总时间 = 45 + 40 = 85μs
DMA占比 = 85 / 120 × 100% = 70.8%

性能瓶颈：DMA数据搬运（占比70.8%，严重偏高）

(2) 优化方案（9分，任意3个）

方案	具体措施	预期效果
① 双缓冲优化	使用两组UB Buffer交替进行数据搬运和计算，实现流水线并行	隐藏搬运延迟，降低DMA等待时间
② 数据复用	分析数据访问模式，将重复使用的数据驻留UB，减少GM访问次数	减少DMA搬运总量
③ 增大分块尺寸	在UB容量允许范围内，增大eachBlockLen，减少Block数量和启动开销	减少分块带来的额外搬运
④ 数据对齐优化	确保所有DataCopy的地址和长度满足32字节对齐，避免硬件额外处理	提升单次搬运效率
⑤ 融合相邻算子	若该算子前后有相邻算子，考虑算子融合，减少中间结果的GM读写	减少整体DMA搬运

(3) 优化后总执行时间估算（3分）

设优化后DMA占比为30%，其他时间不变 = 10+15+10 = 35μs

设优化后总时间为T：
DMA时间 = 0.3T
其他时间 = 0.7T = 35μs
T = 35 / 0.7 = 50μs

或直接用比例计算：
DMA时间从85μs降到x，x/(35+x) = 0.3 → x = 15μs
总时间 = 35 + 15 = 50μs

18. 代码纠错与完善题（15分）

(1) 错误找出（8分，至少5处）

序号	错误位置	错误原因	修正
1	`__global__ void ReduceSum`	缺少`__aicore__`修饰符	`__global__ __aicore__ void ReduceSum`
2	`if (offset > len)`	应为`>=`，否则offset==len时越界	`if (offset >= len)`
3	`DataCopy(ubSrc, src + offset, 1024)`	长度单位错误，应为字节数	`DataCopy(ubSrc, src + offset, 1024 * sizeof(half))`
4	`for (int i = 0; ...)`	未处理最后一块不足1024的情况	使用actualLen控制循环
5	`sum += ubSrc.GetValue(i)`	LocalTensor没有GetValue方法，应使用`ubSrc[i]`或提取到临时变量	使用`half val = ubSrc.GetValue(i)` 或改用Vector API
6	`ubDst.SetValue(0, sum)`	LocalTensor没有SetValue方法	使用`ubDst[0] = sum` 或直接用Scalar写入
7	`DataCopy(dst + blockIdx, ubDst, 1)`	地址偏移应为`dst + blockIdx * sizeof(half)`，且长度应为字节	`DataCopy(dst + blockIdx * sizeof(half), ubDst, sizeof(half))`
8	`这个可能有问题，不用考虑 REG_KERNEL(ReduceSum, 1, ...)`	第二个参数应为0	`REG_KERNEL(ReduceSum, 0, ReduceSumFunc)`

注：LocalTensor的元素访问在CANN 8.x中通常通过operator[]或特定API，具体以官方文档为准。此处按常见写法修正。

(2) 修正后完整代码（7分）

#include "kernel_operator.h"
using namespace AscendC;

__global__ __aicore__ void ReduceSum(
    GM_ADDR const __restrict__ src,
    GM_ADDR __restrict__ dst,
    uint32_t totalLen)
{
    uint32_t blockIdx = GetBlockIdx();
    const uint32_t eachBlockLen = 1024;
    uint32_t offset = blockIdx * eachBlockLen;

    // 边界保护
    if (offset >= totalLen) {
        return;
    }

    // 计算实际处理长度
    uint32_t actualLen = eachBlockLen;
    if (offset + eachBlockLen > totalLen) {
        actualLen = totalLen - offset;
    }

    // 申请UB张量
    LocalTensor<half> ubSrc;
    LocalTensor<half> ubDst;

    ubSrc.SetBuffer((half *)GM2UBAddr(0, 0), eachBlockLen);
    ubDst.SetBuffer((half *)GM2UBAddr(1, 0), 1);

    // GM -> UB 数据搬运
    DataCopy(ubSrc, src + offset, actualLen * sizeof(half));

    // 求和计算（使用Scalar累加）
    half sum = static_cast<half>(0);
    for (uint32_t i = 0; i < actualLen; i++) {
        sum = sum + ubSrc[i];  // 或使用对应API
    }

    // 结果写入UB
    ubDst[0] = sum;

    // UB -> GM 结果写回（每个Block输出一个sum）
    DataCopy(dst + blockIdx * sizeof(half), ubDst, sizeof(half));
}

// 算子注册
这个可能有问题，不用考虑 REG_KERNEL(ReduceSum, 0, ReduceSumFunc);

说明：ReduceSum的精确实现可能因CANN版本和API差异有所不同。核心要点：边界处理、对齐、正确的Tensor访问方式。

📊 三套试卷难度与考点覆盖总览

维度	第一套	第二套	第三套
难度	⭐⭐⭐ 基础	⭐⭐⭐⭐ 进阶	⭐⭐⭐⭐⭐ 综合
题型	选择+判断+简答+编程	选择+判断+简答+纠错编程	单选+多选+综合应用
编程题	补全代码（Mul）	纠错+边界处理	完整开发+性能分析+纠错
侧重	基础语法记忆	边界场景处理	综合应用与性能优化
最适阶段	第3-4天学习	第5-6天巩固	第6-7天冲刺

💡 使用建议

第一套：学完理论+跑通Add算子后做，检验基础掌握度
第二套：熟悉边界处理后做，重点练习代码纠错能力
第三套：考前1-2天做，模拟真实考试节奏和难度

试卷3

算子开发者认证（入门级）试卷三

考试时间：90分钟 | 满分：100分 | 题量：50题

答题说明

单选题：每题2分，共30题，计60分
多选题：每题3分，共10题，计30分（少选得1分，多选/错选不得分）
判断题：每题1分，共10题，计10分

昇腾AI处理器架构

1. (单选题，2分) 以下哪项不属于AI Core的计算单元？

A. Cube
B. Vector
C. Scalar
D. DMA

【答案】D

【解析】DMA是直接内存访问单元，属于数据搬运部件，不属于计算单元。计算单元为Cube、Vector和Scalar。

2. (单选题，2分) CANN运行时层(Runtime)的主要功能是？

A. 编译计算图
B. 管理设备资源、任务调度和内存分配
C. 提供算子库
D. 绘制计算图

【答案】B

【解析】运行时层负责设备资源管理、任务调度、内存分配等，是程序运行时的支撑层。

3. (单选题，2分) 以下关于昇腾AI处理器分离架构的优势，正确的是？

A. 降低了芯片面积
B. AI Core可独立扩展和升级
C. 减少了功耗
D. 提高了集成度

【答案】B

【解析】分离架构中AI Core与CPU独立部署，AI Core可独立扩展和升级，灵活性更高。

4. (单选题，2分) 软硬件同步机制中，【信号量】的主要用途是？

A. 数据加密
B. 流水线中不同单元间的同步通知
C. 内存分配
D. 指令译码

【答案】B

【解析】信号量用于流水线中不同执行单元（如计算与搬运）之间的同步通知，确保数据依赖正确。

5. (多选题，3分) 以下关于达芬奇架构的描述，正确的有哪些？

A. 采用Cube+Vector+Scalar三种计算单元
B. Cube单元专用于矩阵乘法
C. 仅支持推理场景
D. 具有多级存储层次

【答案】ABD

【解析】达芬奇架构包含Cube、Vector、Scalar三种计算单元和多级存储层次，同时支持训练和推理。

6. (多选题，3分) 以下属于CANN生态组件的有哪些？

A. ATC(Ascend Tensor Compiler)
B. AscendCL(Ascend Computing Language)
C. TensorFlow框架
D. Graph Engine

【答案】ABD

【解析】CANN生态包括ATC编译器、AscendCL运行时接口、Graph Engine计算图引擎等，TensorFlow是第三方框架。

7. (判断题，1分) Ascend 310主要面向推理场景，Ascend 910主要面向训练场景。

正确
错误

【答案】正确

【解析】Ascend 310定位边缘/推理场景，Ascend 910定位云端训练场景，两者面向不同应用需求。

8. (判断题，1分) CANN驱动层直接与昇腾AI处理器硬件交互，管理硬件资源。

正确
错误

【答案】正确

【解析】驱动层是CANN最底层，直接与硬件交互，管理NPU设备资源，为上层提供硬件抽象。

存储与计算单元功能

9. (单选题，2分) HBM（High Bandwidth Memory）在AI Core存储层次中的位置是？

A. 最快的缓存
B. 片上L1缓存
C. 芯片外部的高带宽内存
D. 寄存器

【答案】C

【解析】HBM是AI Core外部的高带宽内存，容量最大但访问速度最慢，是存储层次的最低层。

10. (单选题，2分) 计算、存储与搬运单元之间的协作关系是？

A. 串行执行，互不重叠
B. 通过流水线和同步机制实现重叠协作
C. 完全并行，无依赖
D. 仅计算单元工作

【答案】B

【解析】三者通过流水线机制实现重叠执行，搬运单元准备数据供计算单元使用，通过同步信号协调。

11. (单选题，2分) Queue同步机制中，【入队】操作的含义是？

A. 数据写入完成，通知后续阶段可读取
B. 删除数据
C. 分配内存
D. 执行计算

【答案】A

【解析】入队操作表示当前阶段数据处理完成，将数据放入队列通知后续阶段可以读取处理。

12. (多选题，3分) 以下属于AI Core中数据搬运路径的有哪些？

A. HBM与L1之间的搬运
B. L1与L0之间的搬运
C. UB与HBM之间的搬运
D. Cube单元内部的计算

【答案】ABC

【解析】数据搬运路径包括HBM与L1、L1与L0、UB与HBM等存储层次间的传输，Cube内部计算不是搬运。

13. (判断题，1分) 多级缓存设计的目的是通过数据复用减少对外部存储的访问次数。

正确
错误

【答案】正确

【解析】多级缓存通过将常用数据保留在片上高速缓存中复用，减少对慢速HBM的访问，提升整体性能。

并行计算与Kernel编程

14. (单选题，2分) 在Ascend C Kernel编程中，Tiling的含义是？

A. 数据分块/切分策略
B. 数据类型转换
C. 编译优化
D. 内存分配

【答案】A

【解析】Tiling指将输入数据切分为多个数据块(Tile)，分配给不同核并行处理，是昇腾算子开发的核心概念。

15. (单选题，2分) 以下关于blockDim取值的最佳实践，正确的是？

A. 越大越好
B. 等于或接近物理核数，充分利用硬件并行度
C. 固定为1
D. 固定为32

【答案】B

【解析】blockDim设置为等于或接近AI Core物理核数，可充分利用硬件并行能力，避免多轮调度开销。

16. (单选题，2分) 核函数调用时，传入参数通常存储在哪里？

A. HBM中
B. 通过GM(Global Memory)传入，Kernel内通过指针访问
C. 寄存器中
D. 硬编码在Kernel中

【答案】B

【解析】核函数参数（输入输出数据地址等）通过Global Memory传入，Kernel内通过指针访问这些数据。

17. (单选题，2分) 经典并行策略中，【流水线并行】的核心思想是？

A. 不同核处理不同数据
B. 将任务分为多个阶段，不同核处理不同阶段，形成流水线
C. 所有核做相同工作
D. 串行处理

【答案】B

【解析】流水线并行将任务分为多个阶段，不同核负责不同阶段，数据在阶段间流转，形成流水线。

18. (多选题，3分) 以下关于核函数定义的描述，正确的有哪些？

A. 使用__global__关键字声明
B. 在Device侧(AI Core)执行
C. 由Host侧调用启动
D. 可以直接调用标准库函数

【答案】ABC

【解析】核函数用__global__声明，在Device侧执行，由Host侧调用启动。Device侧不能直接调用Host侧的标准库函数。

19. (判断题，1分) SPMD模型中，所有Block执行的Kernel代码完全相同，仅通过block_idx区分处理不同数据。

正确
错误

【答案】正确

【解析】SPMD模型的核心是同一程序处理不同数据，所有Block运行相同Kernel，通过block_idx索引区分数据。

模板库基本概念

20. (单选题，2分) 以下关于模板库价值的描述，最准确的是？

A. 完全替代手写算子
B. 提供高性能可复用模板，降低开发门槛并保证性能
C. 仅用于教学
D. 增加开发难度

【答案】B

【解析】模板库提供经优化的可复用模板，开发者无需从零编写，降低开发门槛同时保证性能。

21. (单选题，2分) CATLASS与CUTLASS的关系是？

A. 完全相同
B. CATLASS是昇腾版的高性能线性代数模板库，类似GPU上的CUTLASS
C. 互不相关
D. CATLASS是CUTLASS的子集

【答案】B

【解析】CATLASS是昇腾平台的高性能线性代数模板库，定位类似NVIDIA GPU上的CUTLASS，但针对达芬奇架构优化。

22. (单选题，2分) 模板库中实现层的主要职责是？

A. 对外提供API
B. 实现具体的计算逻辑和优化策略
C. 管理硬件资源
D. 编译代码

【答案】B

【解析】实现层负责具体的计算逻辑实现和性能优化策略，为接口层提供底层支持。

23. (单选题，2分) 模板库的调度层主要负责什么？

A. 用户接口管理
B. 根据硬件特性自动选择最优执行路径
C. 数据存储
D. 编译优化

【答案】B

【解析】调度层根据硬件特性和输入参数自动选择最优的执行路径和策略，实现性能自适应。

24. (多选题，3分) 以下关于模板库分层Block复用的描述，正确的有哪些？

A. 不同层次的算子可共享底层Block
B. 减少代码冗余
C. 提高代码可维护性
D. 降低算子性能

【答案】ABC

【解析】分层Block复用让不同层次算子共享底层实现，减少冗余、提高可维护性，不影响性能。

25. (判断题，1分) CATLASS模板库仅支持float16数据类型。

正确
错误

【答案】错误

【解析】CATLASS支持多种数据类型，包括float16、float32、int8等，覆盖常见的AI计算精度需求。

自定义算子开发

26. (单选题，2分) 自定义算子开发中，Tiling实现的文件是？

A. ${op_name}_def.cpp
B. ${op_name}_tiling.cpp
C. ${op_name}_kernel.cpp
D. ${op_name}_infershape.cpp

【答案】B

【解析】Tiling逻辑实现在${op_name}_tiling.cpp中，负责数据切分策略和tiling data的生成。

27. (单选题，2分) ACLNN接口的作用是？

A. 编译算子
B. 提供算子对外的调用接口，使框架能调用自定义算子
C. 管理硬件
D. 进行精度测试

【答案】B

【解析】ACLNN(Ascend C Language for Neural Network)接口是算子对外的调用接口，使上层框架能调用算子。

28. (单选题，2分) 通信原语AllReduce的功能是？

A. 将数据广播到所有设备
B. 对所有设备上的数据进行规约（如求和），结果分发到所有设备
C. 收集所有设备的数据
D. 分发数据到不同设备

【答案】B

【解析】AllReduce对所有设备上的数据进行规约操作（如求和、求最大值），并将结果分发到所有设备。

29. (多选题，3分) 以下属于自定义算子开发交付物的有哪些？

A. 算子工程代码
B. 测试用例
C. 设计文档
D. 模型训练脚本

【答案】ABC

【解析】自定义算子交付物包括算子工程代码、测试用例和设计文档，模型训练脚本不属于算子交付物。

30. (判断题，1分) Vector编程范式的基本流程包括：CopyIn到Compute到CopyOut。

正确
错误

【答案】正确

【解析】Vector编程范式的标准流程为数据搬入(CopyIn)到计算(Compute)到数据搬出(CopyOut)，形成完整的处理循环。

性能分析工具使用与性能评估

31. (单选题，2分) msprof采集的性能数据通常以什么形式存储？

A. 仅文本文件
B. 结构化的性能数据文件，可用可视化工具打开
C. 图片
D. 可执行文件

【答案】B

【解析】msprof采集性能数据生成结构化文件，可通过MindStudio或msprof可视化工具打开分析。

32. (单选题，2分) 性能分析中，【内存带宽利用率】反映的是？

A. 内存容量使用情况
B. 实际内存带宽与峰值带宽的比值
C. 内存错误率
D. 内存条数量

【答案】B

【解析】内存带宽利用率是实际使用的内存带宽与硬件峰值带宽的比值，反映访存效率。

33. (单选题，2分) 以下哪种情况最可能导致算力利用率偏低？

A. 计算量过大
B. 数据搬运成为瓶颈，计算单元等待数据
C. 代码行数过多
D. 注释太多

【答案】B

【解析】当数据搬运速度跟不上计算速度时，计算单元空闲等待，导致算力利用率偏低。

34. (多选题，3分) 以下属于msprof工具能力的有哪些？

A. 采集AI Core执行时间
B. 采集内存搬运数据
C. 分析流水线利用情况
D. 修改算子代码

【答案】ABC

【解析】msprof能采集执行时间、搬运数据、流水利用情况等，但不能修改算子代码，代码修改需开发者完成。

35. (判断题，1分) 性能分析中，量化指标比定性分析更能精确定位性能瓶颈。

正确
错误

【答案】正确

【解析】量化指标（如利用率百分比、延迟数值）能精确衡量性能问题严重程度，比定性分析更准确。

模型迁移、适配与精度调试

36. (单选题，2分) 模型迁移前进行可行性评估的首要目标是？

A. 确定迁移所需时间
B. 判断模型能否在目标平台运行及潜在风险
C. 计算迁移成本
D. 选择开发团队

【答案】B

【解析】可行性评估首要判断模型能否在目标平台正常运行，识别不支持算子和潜在风险。

37. (单选题，2分) 以下哪种迁移方式需要开发者对模型结构和算子实现有深入理解？

A. 自动迁移
B. 工具辅助迁移
C. 手工迁移
D. 一键迁移

【答案】C

【解析】手工迁移需要开发者深入理解模型结构和算子实现，逐个适配不支持的算子。

38. (单选题，2分) 模型迁移后精度对齐测试的基准是？

A. 原始框架在原始硬件上的输出
B. 理论值
C. 人工设定阈值
D. 随机数据

【答案】A

【解析】精度对齐测试以原始框架在原始硬件上的输出为基准，比对迁移后输出的差异。

39. (多选题，3分) 以下属于模型迁移适配工作的有哪些？

A. 不支持算子的适配开发
B. 输入输出格式转换
C. 环境部署和配置
D. 模型结构重新设计

【答案】ABC

【解析】迁移适配包括算子适配、格式转换、环境部署等，模型结构重新设计不属于迁移范畴。

40. (判断题，1分) 自动迁移工具能自动处理所有类型的模型，无需人工干预。

正确
错误

【答案】错误

【解析】自动迁移工具主要适用于标准模型，对于包含自定义算子或特殊结构的模型仍需人工干预。

GEMM类算子实现与优化

41. (单选题，2分) GEMM通用分块计算中，分块大小选择的依据是？

A. 随机选择
B. 匹配片上缓存容量，最大化数据复用
C. 越大越好
D. 越小越好

【答案】B

【解析】分块大小应匹配L1/L0缓存容量，使分块数据能在片上缓存中复用，最大化缓存命中率。

42. (单选题，2分) 双缓冲技术中，两个缓冲区的工作状态是？

A. 同时计算
B. 一个用于计算时，另一个用于搬运数据
C. 交替存储结果
D. 互为备份

【答案】B

【解析】双缓冲中一个缓冲区供计算单元使用时，另一个缓冲区同时进行数据搬运，实现重叠。

43. (单选题，2分) 数据格式优化中，NZ格式相比ND格式的优势是？

A. 节省内存
B. 数据在L1缓存中按分块排布，提升Cube单元访存效率
C. 降低精度
D. 简化代码

【答案】B

【解析】NZ格式将数据按分块排布，使Cube单元读取数据时更加连续，减少Bank冲突，提升访存效率。

44. (多选题，3分) 以下属于GEMM分块计算中需要确定的参数有哪些？

A. M方向分块大小
B. N方向分块大小
C. K方向分块大小
D. 代码缩进风格

【答案】ABC

【解析】GEMM分块需要确定M、N、K三个维度的分块大小，代码缩进风格不影响计算性能。

45. (判断题，1分) Padding对齐可能会增加少量额外的计算量，但能提升整体访存效率。

正确
错误

【答案】正确

【解析】Padding补齐数据到对齐边界会引入少量额外计算，但访存效率提升的收益通常大于额外计算开销。

算子性能调优

46. (单选题，2分) 内存使用优化中，【UB空间复用】的含义是？

A. 增加UB容量
B. 不同阶段的数据复用同一UB空间，减少总UB占用
C. 取消UB使用
D. 使用外部存储替代UB

【答案】B

【解析】UB空间复用指不同计算阶段的数据共用同一块UB区域，减少UB总占用，提高空间利用率。

47. (单选题，2分) 资料可能有误，不用考虑，以下哪种映射策略将矩阵按维度拆分到不同核？

A. Block映射
B. Interleave映射
C. Split映射
D. Mode映射

【答案】C

【解析】Split映射将矩阵按特定维度拆分(Split)到不同核，每个核处理矩阵的一部分。

48. (单选题，2分) 算子性能调优的第一步通常是？

A. 重写代码
B. 通过性能分析工具定位瓶颈
C. 增加并行度
D. 修改数据类型

【答案】B

【解析】性能调优应先通过性能分析工具(msprof等)定位瓶颈，再有针对性地优化，避免盲目修改。

49. (多选题，3分) 以下属于数据搬运优化策略的有哪些？

A. 双缓冲(Double Buffering)
B. 合并小搬运
C. 数据预取
D. 增加代码注释

【答案】ABC

【解析】数据搬运优化包括双缓冲、合并小搬运、数据预取等策略，增加注释不影响搬运性能。

50. (判断题，1分) 在算子性能调优中，内存API的选择会影响数据搬运效率。

正确
错误

【答案】正确

【解析】不同内存API有不同的性能特性，选择合适的API（如DataCopy的参数配置）直接影响搬运效率。

算子开发者认证（入门级）试卷五

考试时间：90分钟 | 满分：100分 | 题量：50题

答题说明

单选题：每题2分，共30题，计60分
多选题：每题3分，共10题，计30分（少选得1分，多选/错选不得分）
判断题：每题1分，共10题，计10分

昇腾AI处理器架构

1. (单选题，2分) 昇腾AI处理器AI Core中，负责指令译码和分发的是哪个单元？

A. 存储单元
B. 控制单元
C. 计算单元
D. 搬运单元

【答案】B

【解析】控制单元负责指令的译码、分发和同步控制，是AI Core的调度中枢。

2. (单选题，2分) 以下关于CANN的描述，正确的是？

A. CANN是一个深度学习框架
B. CANN是连接昇腾硬件与上层AI框架的软件栈
C. CANN仅支持推理
D. CANN仅支持训练

【答案】B

【解析】CANN（Compute Architecture for Neural Networks）是昇腾AI计算架构，连接底层硬件与上层框架，同时支持训练和推理。

3. (单选题，2分) 达芬奇架构中Cube单元的主要功能是？

A. 执行标量运算
B. 执行矩阵乘法运算
C. 管理内存分配
D. 进行指令译码

【答案】B

【解析】Cube单元是达芬奇架构中的矩阵计算单元，专门执行矩阵乘法（GEMM）运算，提供高算力。

4. (单选题，2分) 以下哪项是耦合架构的典型特征？

A. AI Core与CPU完全独立封装
B. AI Core与控制CPU集成在同一芯片
C. 仅支持推理场景
D. 不可扩展

【答案】B

【解析】耦合架构将AI Core与控制CPU集成在同一SoC芯片上，具有低延迟、高集成度的特点。

5. (多选题，3分) 以下关于HCCL的描述，正确的有哪些？

A. 支持AllReduce通信原语
B. 支持Broadcast通信原语
C. 仅用于单卡场景
D. 用于多卡/多设备间集合通信

【答案】ABD

【解析】HCCL支持AllReduce、Broadcast等集合通信原语，用于多卡/多设备间的高效通信，不用于单卡场景。

6. (多选题，3分) 昇腾AI处理器中，以下属于控制单元功能的有哪些？

A. 指令译码
B. 指令分发
C. 同步信号管理
D. 矩阵乘法计算

【答案】ABC

【解析】控制单元负责指令译码、分发和同步信号管理，矩阵乘法由Cube计算单元执行。

7. (判断题，1分) CANN编译器层负责将计算图编译为昇腾硬件可执行的指令。

正确
错误

【答案】正确

【解析】CANN编译器层（ATC等工具）将计算图优化并编译为硬件可执行的离线模型。

8. (判断题，1分) 分离架构的AI Core和CPU之间的通信延迟低于耦合架构。

正确
错误

【答案】错误

【解析】分离架构中AI Core和CPU独立部署，通信延迟通常高于耦合架构（集成在同一芯片上）。

存储与计算单元功能

9. (单选题，2分) L0 Buffer包括L0A、L0B、L0C，其中L0C的作用是？

A. 存储矩阵A
B. 存储矩阵B
C. 存储矩阵乘法结果
D. 存储指令

【答案】C

【解析】L0A存储矩阵A输入，L0B存储矩阵B输入，L0C存储矩阵乘法（Cube计算）的输出结果。

10. (单选题，2分) Tile层数据搬运与计算流水编排的核心目标是？

A. 减少代码量
B. 实现搬运与计算的重叠，隐藏访存延迟
C. 降低精度
D. 简化编程

【答案】B

【解析】流水编排通过将数据搬运与计算重叠执行，最大化利用硬件资源，隐藏访存延迟。

11. (单选题，2分) Scalar计算单元的主要功能是？

A. 执行矩阵运算
B. 执行地址计算、循环控制等标量运算
C. 搬运数据
D. 管理通信

【答案】B

【解析】Scalar单元负责标量运算，如地址计算、循环控制、条件判断等，辅助Cube和Vector计算。

12. (多选题，3分) 以下属于MTE数据搬运路径的有哪些？

A. HBM到L1（MTE1）
B. L1到L0（MTE2）
C. UB到HBM/L1（MTE3）
D. L0C到UB

【答案】ABC

【解析】MTE1负责HBM到L1搬运，MTE2负责L1到L0搬运，MTE3负责UB到HBM/L1搬运。L0C到UB使用的是固定搬运通路。

13. (判断题，1分) L1 Buffer是AI Core片上存储，其访问速度高于HBM但低于L0 Buffer。

正确
错误

【答案】正确

【解析】存储层次从快到慢为：寄存器/L0 > L1 > UB > HBM，L1比HBM快但比L0慢。

并行计算与Kernel编程

14. (单选题，2分) 在SPMD模型中，不同Block之间的关系是？

A. 共享同一数据
B. 执行相同代码但处理不同数据
C. 执行不同代码
D. 串行执行

【答案】B

【解析】SPMD模型中所有Block执行相同的Kernel代码，但通过block_idx区分处理不同的数据块。

15. (单选题，2分) blockDim设置过大可能导致什么问题？

A. 编译错误
B. 多轮调度增加，影响性能
C. 精度下降
D. 内存泄漏

【答案】B

【解析】blockDim超过物理核数时，硬件需要多轮调度执行，可能增加调度开销，影响性能。

16. (单选题，2分) Ascend C中核函数的入口函数通过什么关键字标识？

A. global
B. kernel
C. extern C
D. extern_kernel

【答案】A

【解析】Ascend C使用__global__关键字声明核函数入口，类似CUDA的编程模型。

17. (单选题，2分) 通用调优策略中，【数据对齐】的主要目的是？

A. 提高代码可读性
B. 优化访存效率，避免Bank冲突
C. 减少代码量
D. 降低功耗

【答案】B

【解析】数据对齐使访存地址按特定边界对齐，减少Bank冲突，提高内存访问效率。

18. (多选题，3分) 以下关于block_idx的描述，正确的有哪些？

A. block_idx用于区分不同Block
B. block_idx从0开始编号
C. block_idx在Kernel运行时由硬件自动赋值
D. block_idx需要手动传入

【答案】ABC

【解析】block_idx由硬件自动赋值，从0开始编号，用于区分不同Block，开发者无需手动传入。

19. (判断题，1分) 核函数（Kernel）是在AI Core上执行的，由Host侧调用启动。

正确
错误

【答案】正确

【解析】核函数在Device侧的AI Core上执行，由Host侧通过运行时API调用启动。

模板库基本概念

20. (单选题，2分) 以下关于模板库5层架构的描述，正确的是？

A. 每层只提供一种功能
B. 从上到下依次封装，上层调用下层
C. 各层之间完全独立
D. 仅底层可被用户调用

【答案】B

【解析】模板库5层架构从上到下逐层封装，上层调用下层接口，提供不同抽象层次的开发能力。

21. (单选题，2分) CATLASS模板库不适用于以下哪种场景？

A. GEMM算子开发
B. 矩阵分解运算
C. 图像分类模型部署
D. 线性代数计算

【答案】C

【解析】CATLASS是线性代数模板库，适用于GEMM、矩阵分解等计算场景，不用于模型部署。

22. (单选题，2分) 模板库中接口层的主要职责是？

A. 执行硬件指令
B. 对外提供统一的算子开发API
C. 管理内存
D. 进行编译优化

【答案】B

【解析】接口层是模板库的最上层，对外提供统一的、易用的算子开发API，屏蔽底层实现细节。

23. (单选题，2分) 使用模板库开发算子的主要优势是？

A. 代码无法修改
B. 减少重复编码，提升开发效率和性能
C. 增加编译时间
D. 降低性能

【答案】B

【解析】模板库提供经过优化的通用模板，开发者可直接复用，减少重复编码并保证性能。

24. (多选题，3分) 以下属于模板库支持的数据类型的有哪些？

A. float16
B. float32
C. int8
D. float128

【答案】ABC

【解析】模板库支持float16、float32、int8等常见AI计算数据类型，float128不是标准AI精度类型。

25. (判断题，1分) 模板库的分层Block复用机制允许不同层次的算子共享底层实现代码。

正确
错误

【答案】正确

【解析】分层Block复用让不同层次的计算逻辑共享底层Block实现，提高代码复用率和可维护性。

自定义算子开发

26. (单选题，2分) 自定义算子接入注册时，算子原型信息注册在哪个文件中？

A. tiling.cpp
B. def.cpp
C. kernel.cpp
D. CMakeLists.txt

【答案】B

【解析】算子原型信息（输入输出描述、属性等）注册在def.cpp文件中，tiling.cpp负责Tiling逻辑。

27. (单选题，2分) 通信原语按功能分类，不包含以下哪项？

A. AllReduce
B. Broadcast
C. AllGather
D. Compile

【答案】D

【解析】通信原语包括AllReduce、Broadcast、AllGather等，Compile不是通信原语。

28. (单选题，2分) Ascend C中Vector编程范式使用的数据搬入接口是？

A. CopyIn
B. DataCopy
C. MTECopy
D. Memcpy

【答案】B

【解析】Ascend C Vector编程范式中使用DataCopy接口进行UB与GM之间的数据搬运。

29. (多选题，3分) 以下属于GEMV在AIV上实现时需要考虑的要素有哪些？

A. Vector单元的数据加载策略
B. 分块策略
C. Cube单元的矩阵乘法
D. 数据类型选择

【答案】ABD

【解析】GEMV在AIV(Vector Core)上实现需要考虑数据加载、分块和数据类型，不涉及Cube单元矩阵乘法。

30. (判断题，1分) MIX算子的代码框架需要同时包含Cube和Vector的计算逻辑。

正确
错误

【答案】正确

【解析】MIX算子同时使用Cube和Vector单元，代码框架需要包含两者的计算逻辑及数据交互。

性能分析工具使用与性能评估

31. (单选题，2分) 使用msprof采集性能数据时，以下哪个信息无法直接获取？

A. 核函数执行时间
B. 内存搬运量
C. 模型精度
D. 流水线利用率

【答案】C

【解析】msprof采集性能数据（执行时间、搬运量、流水利用率等），模型精度需要通过精度工具分析。

32. (单选题，2分) 仿真分析波形中，【计算-搬运重叠】的理想状态是？

A. 搬运完全在计算之后
B. 搬运与计算尽量重叠，减少空闲
C. 计算完全在搬运之后
D. 两者交替串行

【答案】B

【解析】理想状态下搬运与计算尽量重叠，隐藏访存延迟，最大化硬件利用率。

33. (单选题，2分) 以下哪个指标直接反映AI Core的计算资源利用程度？

A. 算力利用率
B. 内存占用率
C. 功耗
D. 代码行数

【答案】A

【解析】算力利用率反映AI Core计算单元的实际使用比例，直接体现计算资源利用程度。

34. (多选题，3分) 以下属于性能分析中【缺陷识别】常见目标的有哪些？

A. 流水气泡（空闲等待）
B. Bank冲突
C. 不合理的数据搬运
D. 代码注释不足

【答案】ABC

【解析】缺陷识别关注流水气泡、Bank冲突、不合理搬运等影响性能的问题，注释不足不是性能缺陷。

35. (判断题，1分) msprof工具只能用于推理场景的性能分析，不支持训练场景。

正确
错误

【答案】错误

【解析】msprof同时支持推理和训练场景的性能数据采集与分析。

模型迁移、适配与精度调试

36. (单选题，2分) 手工迁移适用于以下哪种场景？

A. 标准的ResNet模型
B. 包含大量自定义算子的复杂模型
C. 简单的线性模型
D. 标准的BERT模型

【答案】B

【解析】包含大量自定义算子的复杂模型通常无法自动迁移，需要手工逐个适配算子。

37. (单选题，2分) 模型迁移后环境部署不需要以下哪个步骤？

A. 安装CANN工具套件
B. 配置NPU驱动
C. 安装GPU驱动
D. 配置环境变量

【答案】C

【解析】昇腾平台使用NPU，不需要安装GPU驱动，需要安装CANN工具套件和NPU驱动。

38. (单选题，2分) 工具辅助迁移相比手工迁移的优势是？

A. 能处理所有自定义算子
B. 自动化程度更高，减少人工工作量
C. 性能一定更好
D. 不需要任何人工干预

【答案】B

【解析】工具辅助迁移通过自动化工具完成大部分转换工作，减少人工工作量，但不一定能处理所有自定义算子。

39. (多选题，3分) 以下属于模型迁移可行性评估内容的有哪些？

A. 模型中算子在目标平台的支持情况
B. 目标硬件性能是否满足需求
C. 模型输入输出格式
D. 开发团队的编程语言偏好

【答案】ABC

【解析】可行性评估包括算子支持度、硬件性能、模型格式等，编程语言偏好不影响技术可行性。

40. (判断题，1分) 模型迁移后精度下降时，应首先检查算子实现是否与原始框架一致。

正确
错误

【答案】正确

【解析】精度下降通常由算子实现差异引起，应首先通过分层比对和算子级Dump定位问题算子。

GEMM类算子实现与优化

41. (单选题，2分) GEMM分块计算中，L1 Cache中缓存的通常是？

A. 完整的矩阵A和B
B. 矩阵A和B的分块(Tile)
C. 最终结果矩阵
D. 指令

【答案】B

【解析】L1 Cache容量有限，缓存的是矩阵A和B的分块(Tile)，通过循环加载不同分块完成完整计算。

42. (单选题，2分) Padding对齐优化的目的是？

A. 增加数据量
B. 使数据尺寸满足分块对齐要求，提升访存效率
C. 降低精度
D. 减少代码量

【答案】B

【解析】Padding补齐数据尺寸到分块整数倍，避免不规则访存，提升硬件利用效率。

43. (单选题，2分) 在GEMM优化中，Swizzle排布主要解决什么问题？

A. 内存不足
B. L0/L1缓存中数据的Bank冲突和复用
C. 编译错误
D. 精度损失

【答案】B

【解析】Swizzle通过调整数据访问顺序，减少Bank冲突，提高L0/L1缓存中数据的复用率。

44. (多选题，3分) 以下关于双缓冲(Double Buffering)的描述，正确的有哪些？

A. 需要两份缓冲区
B. 实现搬运与计算重叠
C. 增加内存使用量
D. 降低计算精度

【答案】ABC

【解析】双缓冲使用两份缓冲区，实现搬运与计算重叠，会增加少量内存使用，但不影响精度。

45. (判断题，1分) GEMM分块大小越大，性能一定越好。

正确
错误

【答案】错误

【解析】分块大小需要匹配片上缓存容量，过大导致缓存溢出，过小降低复用率，存在最优值。

算子性能调优

46. (单选题，2分) 数据搬运优化中，【合并小搬运】的策略是指？

A. 将多个小数据搬运合并为一次大搬运
B. 减少搬运次数
C. 增加搬运次数
D. 跳过搬运

【答案】A

【解析】合并小搬运将多次小数据传输合并为一次大数据传输，减少搬运启动开销，提高带宽利用率。

47. (单选题，2分) 资料可能有误，不用考虑，以下哪种映射策略通过交替分配数据给不同核，实现负载均衡？

A. Block映射
B. Interleave映射
C. Split映射
D. Round-robin映射

【答案】B

【解析】Interleave映射通过交替（交织）方式分配数据，使各核负载更均衡，适合不规则数据。

48. (单选题，2分) 算子性能调优中，选择高性能API的核心原则是？

A. 选择最复杂的API
B. 根据计算场景选择最匹配的高性能接口
C. 选择最新API
D. 选择最少参数的API

【答案】B

【解析】不同API针对不同场景优化，应根据实际计算需求选择最匹配的高性能接口。

**49. (多选题，3分)**资料可能有误，不用考虑，以下属于4种映射策略的有哪些？

A. Block映射
B. Interleave映射
C. Split映射
D. Random映射

【答案】ABC

【解析】4种映射策略通常指Block映射、Interleave映射、Split映射和Mode映射，Random不是映射策略。

50. (判断题，1分) Cube核的矩阵计算能力通常高于Vector核的矩阵计算能力。

正确
错误

【答案】正确

【解析】Cube核专用于矩阵乘法，具有专门的矩阵计算单元，矩阵计算能力远高于Vector核。

算子开发者认证（入门级）试卷六

考试时间：90分钟 | 满分：100分 | 题量：50题

答题说明

单选题：每题2分，共30题，计60分
多选题：每题3分，共10题，计30分（少选得1分，多选/错选不得分）
判断题：每题1分，共10题，计10分

昇腾AI处理器架构

1. (单选题，2分) 昇腾AI处理器的核心计算架构名称是什么？

A. Volta架构
B. 达芬奇(DaVinci)架构
C. Ampere架构
D. Turing架构

【答案】B

【解析】昇腾AI处理器采用华为自研的达芬奇(DaVinci)架构，专为AI计算场景优化。

2. (单选题，2分) CANN软件体系不包含以下哪一层？

A. 算子库层
B. 计算图层
C. 深度学习框架层
D. 驱动层

【答案】C

【解析】CANN五层体系为：算子库层、计算图层、编译器层、运行时层、驱动层，不包含深度学习框架层。

3. (单选题，2分) AI Core中的控制单元主要负责什么功能？

A. 执行矩阵运算
B. 指令译码、分发和同步控制
C. 数据搬运
D. 存储管理

【答案】B

【解析】控制单元负责指令的译码、分发和同步控制，协调整个AI Core的执行流程。

4. (单选题，2分) HCCL的全称是什么？

A. Huawei Cloud Communication Library
B. Huawei Collective Communication Library
C. High-performance Computing Communication Layer
D. Hardware Control Communication Link

【答案】B

【解析】HCCL全称Huawei Collective Communication Library，是华为提供的集合通信库，用于多卡间高效通信。

5. (多选题，3分) 以下关于分离架构与耦合架构的描述，正确的有哪些？

A. 分离架构中AI Core与控制CPU相互独立
B. 耦合架构将AI Core与CPU集成在同一芯片
C. 分离架构灵活性更高，可独立扩展
D. 耦合架构仅适用于推理场景

【答案】ABC

【解析】分离架构中AI Core与CPU独立部署，耦合架构将二者集成。分离架构灵活性更高，耦合架构并非仅用于推理，Ascend 910耦合架构也用于训练。

6. (多选题，3分) CANN五层体系中，以下属于其组成部分的有哪些？

A. 算子库(AscendC/ATC)
B. 计算图引擎(Graph Engine)
C. 深度学习框架(MindSpore)
D. 运行时(Runtime)

【答案】ABD

【解析】CANN五层包括算子库层、计算图层、编译器层、运行时层和驱动层。MindSpore是独立的AI框架，不属于CANN体系。

7. (判断题，1分) HCCL主要用于单核内部的指令同步。

正确
错误

【答案】错误

【解析】HCCL用于多设备/多卡之间的集合通信（如AllReduce、Broadcast等），不是单核内部同步机制。

8. (判断题，1分) 昇腾AI处理器的软硬件同步机制通过信号量和同步指令实现计算单元与搬运单元之间的协同。

正确
错误

【答案】正确

【解析】昇腾AI Core通过硬件信号量和同步指令实现计算单元和搬运单元之间的流水线协同。

存储与计算单元功能

9. (单选题，2分) AI Core的三大计算单元不包括以下哪项？

A. Cube单元（矩阵计算）
B. Vector单元（向量计算）
C. Scalar单元（标量计算）
D. Tensor单元（张量计算）

【答案】D

【解析】AI Core计算单元包括Cube（矩阵）、Vector（向量）和Scalar（标量），不存在独立的Tensor单元。

10. (单选题，2分) Unified Buffer (UB) 的主要作用是什么？

A. 存储权值数据
B. 存储Vector/Scalar计算的中间数据
C. 存储HBM中的全局数据
D. 存储指令

【答案】B

【解析】UB是Unified Buffer，用于存储Vector和Scalar计算单元的输入输出数据，是片上高速缓冲。

11. (单选题，2分) MTE（Memory Transfer Engine）的主要功能是？

A. 执行矩阵乘法
B. 实现不同存储层级间的数据搬运
C. 进行指令译码
D. 管理电源

【答案】B

【解析】MTE负责AI Core内部不同存储层次之间的数据搬运，如HBM与L1、L1与L0之间的数据传输。

12. (多选题，3分) 以下属于AI Core存储层次结构的有哪些？

A. HBM（高带宽内存）
B. L1 Buffer
C. L0 Buffer (L0A/L0B/L0C)
D. Unified Buffer (UB)

【答案】ABCD

【解析】AI Core存储层次包括HBM、L1 Buffer、L0 Buffer(L0A/L0B/L0C)和Unified Buffer，构成多级缓存体系。

13. (判断题，1分) Queue同步机制用于管理不同流水阶段之间的数据依赖关系。

正确
错误

【答案】正确

【解析】Queue同步通过硬件队列管理计算、搬运等不同流水阶段之间的数据依赖，确保执行顺序正确。

并行计算与Kernel编程

14. (单选题，2分) SPMD编程模型的含义是？

A. Single Program Multiple Data
B. Single Process Multiple Device
C. System Program Multi-Data
D. Stream Program Multi-Dimension

【答案】A

【解析】SPMD（Single Program Multiple Data）指同一程序处理不同数据，各Block执行相同代码但处理不同数据块。

15. (单选题，2分) blockDim参数的作用是什么？

A. 设置数据块大小
B. 设置并行执行的Block数量
C. 设置缓冲区大小
D. 设置线程数量

【答案】B

【解析】blockDim用于设置Kernel并行执行的Block数量，每个Block在不同核上并行执行相同代码。

16. (单选题，2分) 在Ascend C中，获取当前Block索引应使用哪个变量？

A. thread_idx
B. block_idx
C. core_id
D. rank_id

【答案】B

【解析】在Ascend C的SPMD模型中，使用block_idx获取当前Block的索引，用于区分不同数据块。

17. (单选题，2分) Ascend C核函数注册使用的宏是？

A. REGISTER_KERNEL
B. KERNEL_REGISTER
C. ACLNN_REGISTER
D. OP_KERNEL_REG

【答案】A

【解析】Ascend C中使用REGISTER_KERNEL宏注册核函数，将Kernel与算子关联。

18. (多选题，3分) 以下属于经典并行优化策略的有哪些？

A. 数据并行
B. 流水线并行
C. 模型并行
D. 串行执行

【答案】ABC

【解析】经典并行策略包括数据并行、流水线并行和模型并行，串行执行不属于并行策略。

19. (判断题，1分) blockDim的取值不能超过AI Core的物理核数。

正确
错误

【答案】错误

【解析】blockDim可以超过物理核数，硬件会自动进行多轮调度，但通常建议不超过物理核数以获得最佳性能。

模板库基本概念

20. (单选题，2分) Ascend C模板库的主要目标是什么？

A. 替代手写Kernel
B. 提供可复用的高性能算子模板，降低开发难度
C. 管理模型训练
D. 进行模型推理部署

【答案】B

【解析】模板库提供可复用的高性能算子模板，开发者基于模板快速开发自定义算子，降低开发门槛。

21. (单选题，2分) CATLASS的定位是什么？

A. 深度学习框架
B. 高性能线性代数模板库
C. 模型转换工具
D. 性能分析工具

【答案】B

【解析】CATLASS是Ascend C的高性能线性代数模板库，类似GPU上的CUTLASS，提供GEMM等计算模板。

22. (单选题，2分) Ascend C模板库的分层架构不包含以下哪层？

A. 接口层
B. 实现层
C. 硬件层
D. 调度层

【答案】C

【解析】模板库分层架构一般包含接口层、实现层、调度层等，硬件层属于芯片架构，不属于模板库分层。

23. (单选题，2分) 模板库中分层Block复用机制的主要优势是？

A. 减少代码重复，提高可维护性
B. 增加代码行数
C. 降低执行效率
D. 增加编译时间

【答案】A

【解析】分层Block复用让不同层次的计算逻辑可复用，减少代码冗余，提高可维护性和开发效率。

24. (多选题，3分) 以下关于CATLASS模板库的描述，正确的有哪些？

A. 提供GEMM等矩阵运算模板
B. 支持多种数据类型
C. 是一个深度学习推理框架
D. 可用于高性能算子开发

【答案】ABD

【解析】CATLASS是线性代数模板库，提供GEMM等计算模板，支持多种数据类型，用于高性能算子开发，不是推理框架。

25. (判断题，1分) 模板库支持的数据类型包括float16、float32等常见精度。

正确
错误

【答案】正确

【解析】Ascend C模板库支持float16、float32、int8等多种数据类型，覆盖常见的AI计算精度需求。

自定义算子开发

26. (单选题，2分) Vector编程范式主要适用于什么类型的计算？

A. 矩阵乘法
B. 向量级运算（如ReLU、Add等）
C. 标量运算
D. 通信操作

【答案】B

【解析】Vector编程范式适用于逐元素的向量运算，如ReLU、Add、激活函数等Element-wise操作。

27. (单选题，2分) Cube编程范式主要适用于什么类型的计算？

A. 向量运算
B. 矩阵乘法（GEMM/MatMul）
C. 数据搬运
D. 标量运算

【答案】B

【解析】Cube编程范式利用Cube单元进行矩阵乘法运算，适用于GEMM、MatMul等矩阵计算场景。

28. (单选题，2分) MIX算子的特点是什么？

A. 仅使用Vector单元
B. 同时使用Cube和Vector单元
C. 仅使用Scalar单元
D. 不使用任何计算单元

【答案】B

【解析】MIX算子同时利用Cube和Vector计算单元，实现矩阵运算与向量运算的混合计算。

29. (多选题，3分) 以下属于自定义算子开发流程关键步骤的有哪些？

A. 算子原型注册
B. 核函数实现
C. Tiling实现
D. 模型训练

【答案】ABC

【解析】自定义算子开发包括算子原型注册、核函数实现、Tiling实现等步骤，模型训练不属于算子开发流程。

30. (判断题，1分) GEMV算子在AIV和AIC上的实现方案存在差异，需要根据硬件特性分别设计。

正确
错误

【答案】正确

【解析】AIV（AI Vector Core）和AIC（AI Cube Core）的硬件特性不同，GEMV在这两种核上的实现方案需要分别设计。

性能分析工具使用与性能评估

31. (单选题，2分) msprof工具的主要用途是什么？

A. 模型训练
B. 性能数据采集与分析
C. 代码编译
D. 模型转换

【答案】B

【解析】msprof是昇腾性能分析工具，用于采集AI Core运行时的性能数据，辅助性能瓶颈定位。

32. (单选题，2分) 仿真分析中的波形图主要用于展示什么？

A. 模型精度曲线
B. 各流水阶段的执行时间线
C. 内存占用
D. 功耗变化

【答案】B

【解析】波形图展示计算、搬运等各流水阶段在时间轴上的执行情况，帮助识别流水气泡和瓶颈。

33. (单选题，2分) 性能分析中发现的【流水气泡】指的是什么？

A. 内存泄漏
B. 流水线中的空闲等待时段
C. 数据溢出
D. 编译错误

【答案】B

【解析】流水气泡指流水线中某些阶段空闲等待，未充分利用硬件资源，是性能优化的重要关注点。

34. (多选题，3分) 以下属于性能评估关键指标的有哪些？

A. 算力利用率
B. 内存带宽利用率
C. 流水线利用率
D. 代码行数

【答案】ABC

【解析】性能评估关键指标包括算力利用率、内存带宽利用率、流水线利用率等，代码行数不是性能指标。

35. (判断题，1分) 可视化分析工具可以定性识别性能瓶颈，但无法量化性能指标。

正确
错误

【答案】错误

【解析】现代可视化分析工具既能定性识别瓶颈，也能通过性能数据量化展示各项指标。

模型迁移、适配与精度调试

36. (单选题，2分) 以下哪种迁移方式自动化程度最高？

A. 手工迁移
B. 工具辅助迁移
C. 自动迁移
D. 重新开发

【答案】C

【解析】自动迁移由工具自动完成模型转换和适配，自动化程度最高，适用于标准模型。

37. (单选题，2分) 模型迁移可行性评估不需要考虑以下哪个因素？

A. 算子支持度
B. 性能需求
C. 开发者编程偏好
D. 硬件兼容性

【答案】C

【解析】可行性评估关注算子支持度、性能需求、硬件兼容性等，开发者编程偏好不影响迁移可行性。

38. (单选题，2分) 模型迁移到昇腾平台后，精度调试的常用手段是？

A. 重新设计模型结构
B. 分层精度对比和算子级Dump比对
C. 增加训练数据
D. 更换硬件

【答案】B

【解析】精度调试通常通过分层精度对比和算子级Dump数据比对，定位精度异常的算子。

39. (多选题，3分) 以下属于模型迁移主要挑战的有哪些？

A. 自定义算子适配
B. 精度对齐
C. 性能优化
D. 模型结构创新

【答案】ABC

【解析】模型迁移面临自定义算子适配、精度对齐和性能优化等挑战，模型结构创新不属于迁移挑战。

40. (判断题，1分) 自动迁移适用于所有类型的模型，包括包含大量自定义算子的模型。

正确
错误

【答案】错误

【解析】自动迁移适用于标准模型，包含大量自定义算子的模型通常需要手工迁移或工具辅助迁移。

GEMM类算子实现与优化

41. (单选题，2分) GEMM分块计算的主要目的是什么？

A. 减少代码量
B. 提高数据局部性和缓存利用率
C. 简化编程模型
D. 降低精度

【答案】B

【解析】GEMM分块将大矩阵切分为小块计算，提高数据在片上缓存中的复用率，减少对外部存储的访问。

42. (单选题，2分) 双缓冲(Double Buffering)技术的核心思想是？

A. 使用两份代码
B. 数据搬运与计算重叠执行
C. 双倍存储空间
D. 两个核同时计算

【答案】B

【解析】双缓冲在计算当前数据块的同时搬运下一数据块，实现搬运与计算的重叠，隐藏访存延迟。

43. (单选题，2分) Swizzle排布优化的主要目标是？

A. 减少代码行数
B. 优化L0/L1缓存的数据访问模式
C. 增加并行度
D. 降低功耗

【答案】B

【解析】Swizzle通过重排数据访问顺序，优化Bank冲突和缓存命中，提升数据访问效率。

44. (多选题，3分) 以下属于GEMM优化手段的有哪些？

A. 分块计算(Tiling)
B. 双缓冲(Double Buffering)
C. Swizzle排布
D. 随机访问

【答案】ABC

【解析】GEMM常用优化包括分块计算、双缓冲和Swizzle排布，随机访问会降低性能，不是优化手段。

45. (判断题，1分) 数据格式（如ND格式与NZ格式）会影响GEMM算子的访存效率。

正确
错误

【答案】正确

【解析】不同数据格式影响数据在内存中的排布方式，直接影响访存效率和Bank冲突情况。

算子性能调优

46. (单选题，2分) 自定义算子性能优化的首要关注点是？

A. 代码风格
B. 数据搬运效率
C. 注释完整性
D. 变量命名

【答案】B

【解析】AI Core的计算受限于访存带宽，数据搬运效率通常是性能瓶颈的首要关注点。

47. (单选题，2分) Cube核和Vector核的主要区别是？

A. Cube核做矩阵运算，Vector核做向量运算
B. Cube核做向量运算，Vector核做矩阵运算
C. 两者完全相同
D. Cube核仅用于推理

【答案】A

【解析】Cube核专用于矩阵乘法运算，Vector核用于向量级运算，两者计算特性不同。

48. (单选题，2分) 资料可能有误，不用考虑，以下哪种映射策略将数据按连续地址分给同一核？

A. Block映射
B. Interleave映射
C. Split映射
D. Broadcast映射

【答案】A

【解析】Block映射将连续地址的数据块分配给同一核，有利于连续访存，是最基本的映射策略。

49. (多选题，3分) 以下属于算子性能优化策略的有哪些？

A. 数据搬运优化（如Double Buffering）
B. 内存使用优化
C. API选择优化
D. 增加代码注释

【答案】ABC

【解析】算子性能优化策略包括数据搬运优化、内存使用优化和API选择优化，增加注释不影响性能。

50. (判断题，1分) 资料可能有误，不用考虑，不同的映射策略会影响访存模式，进而影响算子性能。

正确
错误

【答案】正确

【解析】映射策略决定数据到核的分配方式，直接影响访存连续性和Bank冲突，从而影响性能。

ai生成

人工智能6S服务平台

作为“人工智能6S店”的官方数字引擎，为AI开发者与企业提供一个覆盖软硬件全栈、一站式门户。

更多推荐

【共创季稿事节】鸿蒙原生 ArkTS 布局深度解析：Grid 中 GridItem 的自适应尺寸

人工智能6S服务平台

【共创季稿事节】鸿蒙原生 ArkTS 布局深度解析：Grid 中 GridItem 的自适应尺寸

人工智能6S服务平台

GSV5600@ACP# 多协议一体化便携 AI 终端互联芯片

人工智能6S服务平台

所有评论(0)

查看更多评论

zhangfeng1133

@zhangfeng1133

已为社区贡献14条内容

华为昇腾CANN算子开发类认证，考试试卷（HCIP/微认证）快速通关方案，附6张试卷，后面三张直接在题目下面有答案

zhangfeng1133

昇腾CANN算子开发类认证（HCIP/微认证）快速通关方案，附6张试卷，后面三张直接在题目下面有答案

一、考试模块&分值优先级（按投入产出排序）

二、分模块突击学习（直接照着执行）

（一）理论模块（拿下约50%分值，2天完成）

（二）实操模块（拿下约40%分值，3天完成，CANN 8.x版本）

（三）工具使用（半天掌握）

三、7天完整冲刺计划表（每日任务清晰，零基础适配）

四、考场避坑&得分技巧（零基础必看）

五、配套资料&资源建议

CANN 8.x 标准、可直接复制运行 的完整案例：*

一、完整算子源码（AddKernel.cpp）

二、CANN 8.x 编译命令（直接复制执行）

三、极简运行 & 测试流程（考试通用步骤）

四、零基础必背 6 个考点（选择/判断/实操高频）

1. 核函数固定格式（必考）

2. 索引获取

3. 内存层级（理论题必背）

4. DataCopy 铁律（实操80%报错来源）

5. LocalTensor（8.x 新标准）

6. 算子注册（固定死，直接默写）

五、考场实操极简流程（照着走就能拿分）

六、下一步练习建议（零基础循序渐进）

模拟试卷一

📝 昇腾CANN算子开发类认证 模拟考试试卷

第一部分：选择题（每题2分，共30分）

第二部分：判断题（每题2分，共20分）

第三部分：简答题（每题5分，共20分）

第四部分：实操编程题（30分）

✅ 参考答案

第一部分：选择题答案

第二部分：判断题答案

第三部分：简答题参考答案

第四部分：实操编程题答案

📊 评分建议

第二套模拟试卷（侧重不同考点，比如GEMM优化、性能调优类题目），

📝 昇腾CANN算子开发类认证 模拟考试试卷（第二套）

第一部分：选择题（每题2分，共30分）

第二部分：判断题（每题2分，共20分）

第三部分：简答题（每题5分，共20分）

第四部分：实操编程题（30分）

✅ 参考答案

第一部分：选择题答案

第二部分：判断题答案

第三部分：简答题参考答案

第四部分：实操编程题答案

📊 两套试卷考点对比

第三套*（侧重真题风格，全是综合应用题），

📝 昇腾CANN算子开发类认证 模拟考试试卷（第三套 · 真题风格）

第一部分：单选题（每题3分，共30分）

第二部分：多选题（每题4分，共20分，漏选得2分，错选不得分）

第三部分：综合应用题（共50分）

✅ 参考答案

第一部分：单选题答案

第二部分：多选题答案

第三部分：综合应用题答案

16. 算子开发综合分析题（20分）

17. 性能分析与调优实战题（15分）

18. 代码纠错与完善题（15分）

📊 三套试卷难度与考点覆盖总览

💡 使用建议

试卷3

算子开发者认证（入门级）试卷三

答题说明

昇腾AI处理器架构

存储与计算单元功能

并行计算与Kernel编程

模板库基本概念

自定义算子开发

性能分析工具使用与性能评估

模型迁移、适配与精度调试

GEMM类算子实现与优化

算子性能调优

算子开发者认证（入门级）试卷五

答题说明

昇腾AI处理器架构

存储与计算单元功能

并行计算与Kernel编程

模板库基本概念

CANN 8.x 标准、可直接复制运行的完整案例：*

📝 昇腾CANN算子开发类认证模拟考试试卷

📝 昇腾CANN算子开发类认证模拟考试试卷（第二套）

*第三套**（侧重真题风格，全是综合应用题），

📝 昇腾CANN算子开发类认证模拟考试试卷（第三套 · 真题风格）