智能编程模型和智能编程语言

两个实验

包含ai生成文本

智能编程语言算子开发与集成实验

实验目的：

掌握使用智能编程语言 BCL 进行算子开发、编译扩展高性能库算子,并集成到 PyTorch 框架中的方法和流程。能够用 BCL 实现 Sigmoid 算子,并集成进 PyTorch 的推理网络高效 地运行在 MLU 硬件上。

实验环境：

硬件平台:MLU 云平台环境。

软件环境:AI 框架 PyTorch、计算加速库 CNNL、云端开发工具集 CNToolkit、Python3 及相关的扩展库。

运行环境：

镜像收藏：mlu370_ubuntu22.04-for-student ：v7.1

目录：/ opt / code_chap7 / exp_7_2_matmul_opt /

代码文件组织：

Sigmoid 自定义算子实验根目录
├── README.md               # 描述算子功能的说明文档
├── mlu_custom_ext          # 生成的 module 模块，用于在 python 层导入
│   ├── __init__.py         # python 包固有文件
│   ├── mlu                 # mlu 代码文件，根据实际情况自己创建，在 setup.py 中修改即可
│   │   ├── include         # 头文件目录（头文件和实现分离，属于代码习惯，建议采用此布局）
│   │   │   ├── bang_sigmoid_sample.h   # 实现对 mlu 函数的封装
│   │   │   ├── kernel.h                # BCL 代码中的宏，良好组织代码的需要
│   │   │   └── custom_ops.h            # 算子对外头文件
│   │   └── src
│   │       ├── bang_sigmoid.cpp        # 对 PyTorch 层面 Tensor 的封装，和自定义算子中 xxx_internal 的实现类似
│   │       └── bang_sigmoid_sample.mlu # 核心 BangC 实现
│   └── mlu_functions       # 合理的组织自己的代码，方便后续调用
│       ├── __init__.py     # 包必备文件
│       └── mlu_functions.py # 对 C++ 代码的封装
├── setup.py                # 构建包的脚本
└── tests
    └── test_sigmoid.py     # 对绑定代码的 python 侧测试

代码补全

补全mlu_custom_ext/mlu/src/bang_sigmoid_sample.mlu、mlu_custom_ext/mlu/src/bang_sigmoid.cpp/bang_sigmoid.cpp、/mlu_custom_ext/mlu_functions/mlu_functions.py、tests/test_sigmoid.py、tests/test_sigmoid_benchmark.py一系列文件

实现 Sigmoid 主程序

mlu_custom_ext/mlu/src/bang_sigmoid.cpp/bang_sigmoid.cpp

实现 Sigmoid 核函数

在 mlu_custom_ext/mlu/src/bang_sigmoid_sample.mlu 文件中实现 bang_sigmoid_kernel_entry 函数并通过头文件对 外暴露。

头文件bang_sigmoid_sample.h声明：

#pragma once
#include <cnrt.h>
template <typename T>
void bang_sigmoid_kernel_entry(
    cnrtQueue* queue,
    T* d_dst,
    T* d_src,
    int elem_count);

通过 pybind11 暴露 Op 接口

customed_ops . h

这俩是写好了的

使用 setuptools 编译和安装

setup . py

精度对比测试

tests/test_sigmoid.py

性能测试

tests/test_sigmoid_benchmark.py

mlu_functions.py

实验运行

先编译

python setup.py clean --all
python setup.py install

进入tests目录

python test_sigmoid.py
python test_sigmoid_benchmark.py

智能编程语言性能优化实验

实验目的：

掌握使用智能编程语言优化算法性能的原理,掌握智能编程语言的调试和调优方法, 能够使用智能编程语言在 MLU 上加速矩阵乘的计算。

实验环境：

硬件平台:MLU 云平台环境。

软件环境:云端开发工具集 CNToolkit

运行环境：

镜像收藏：mlu370_ubuntu22.04-for-student ：v7.1

目录：/opt/code_chap7/exp_7_1_custom_pytorch_mlu_op/

实验内容：

(1) 实现一个 CPU 版本的标量矩阵乘,然后将标量矩阵乘实现迁移到 MLU 上,做性 能对比;

(2) 利用 BCL 的 NRAM 地址空间加速标量版本的矩阵乘,对比性能提升;

(3) 利用 BCL 提供的向量化接口 bang_matmul 加速矩阵乘,对比性能提升;

(4) 利用 BCL 提供的 Block 任务,使用 MLU 的多核做并行加速,对比性能提升;

(5) 利用 BCL 提供的异步拷贝接口 memcpy_async 做三级流水优化,对比性能提升;

(6) 利用 BCL 提供的 Union 任务和 SRAM 地址空间,做五级流水优化,对比性能提升;

代码补全

矩阵乘标量实现

01 _ s c a l a r . mlu

这段代码实现了一个矩阵乘法的标量（scalar）版本，分别在 CPU 和 MLU（寒武纪人工智能加速卡）上执行，并对比两者的计算结果与运行时间。下面按执行顺序解释主要步骤。

1. 定义矩阵尺寸与辅助函数

• M = 128, N = 256, K = 128：定义左矩阵 A 尺寸为 M×N，右矩阵 B 尺寸为 N×K，结果矩阵 C 尺寸为 M×K。
• relativeError(a, b)：计算两个浮点数的相对误差，用于验证 MLU 与 CPU 结果是否一致。
• generateRandomFloat(min, max)：在 [min, max] 区间生成随机浮点数（调试模式下固定返回 1.1）。

2. 主机端（CPU）矩阵乘法函数

void multiplyMatricesCPU(float *left, float *right, float *result, int m, int n, int k)

• 三重循环实现经典标量乘法：

  for i in 0..m-1:
    for j in 0..k-1:
      sum = 0
      for x in 0..n-1:
        sum += left[i*n + x] * right[x*k + j]
      result[i*k + j] = sum

• 时间复杂度 O(M*N*K)。

3. 设备端（MLU）核函数

__mlu_entry__ void multiplyMatricesMLU(...)

• 使用与 CPU 完全相同的三重循环逻辑，但运行在 MLU 加速卡上。
• 注释中提示“标量矩阵乘计算部分”，已补充了与 CPU 相同的累加代码。

4. 主程序执行流程

4.1 分配内存

• 主机内存：left, right, result（CPU 结果），result_actual（接收 MLU 结果）。
• MLU 设备内存：left_mlu, right_mlu, result_mlu（通过 cnrtMalloc 分配）。

4.2 初始化矩阵数据

• 左矩阵 left[M][N] 每个元素随机生成 [1.0, 1.1]。
• 右矩阵 right[N][K] 同理。
• 通过 cnrtMemcpy 将数据从主机拷贝到 MLU 设备。

4.3 运行 CPU 矩阵乘法并计时

• gettimeofday 记录开始和结束时间，计算耗时（毫秒）。
• 结果保存在 result 中。

4.4 运行 MLU 矩阵乘法并计时

• 创建 CNRT 通知器（notifier）和队列（queue）。
• 开始计时：cnrtPlaceNotifier(st_mlu, queue)。
• 启动核函数：

  multiplyMatricesMLU<<<dim, func_type, queue>>>(left_mlu, right_mlu, result_mlu, m, n, k);

• 结束计时：cnrtPlaceNotifier(et_mlu, queue)。
• 同步队列 cnrtQueueSync(queue) 确保核函数执行完毕。
• cnrtNotifierElapsedTime 计算两个通知器之间的时间差（毫秒），得到 MLU 耗时。

4.5 将 MLU 结果拷贝回主机

• cnrtMemcpy(result_actual, result_mlu, ..., cnrtMemcpyDevToHost)

4.6 正确性验证

• 遍历结果矩阵每个元素，计算 CPU 结果与 MLU 结果的相对误差。
• 若任一元素相对误差 > 0.1，则打印失败位置和数值，最终输出 01FAILED；否则输出 01PASSED。

4.7 资源释放

• 释放主机内存、MLU 设备内存、销毁通知器和队列。

5. 关键点说明

• 标量实现：指没有使用任何向量化指令或并行优化，纯粹用最内层循环逐元素相乘累加。
• MLU 核函数：虽然代码写的是三重循环，但每个核函数只由一个线程块（dim = {1,1,1}）执行，本质上仍然是串行执行，没有发挥 MLU 的并行优势。这可能是教学示例，用于对比“标量版本”与后续优化的“向量化/并行版本”。
• 随机数范围：[1.0, 1.1] 且调试模式下固定为 1.1，避免浮点溢出或极端值导致误差过大。
• 时间测量：CPU 使用 gettimeofday，MLU 使用硬件通知器，得到精确的设备端执行时间。

6. 输出示例

程序会打印矩阵尺寸、CPU 耗时、MLU 耗时，以及最终 PASSED 或 FAILED。若开启 DEBUG 宏，还会打印完整的 CPU 和 MLU 结果矩阵。

矩阵乘标量 NRAM 实现

02 _scalar_nram . mlu

与第一个示例的主要区别：在第一个示例（01_scalar.mlu）中，核函数 multiplyMatricesMLU 直接在全局内存（GDRAM）上通过三重循环计算，每次访问 left[i*n+x] 和 right[x*k+j] 都要从片外存储器读取，效率很低。而本示例在核函数内部声明了三个 __nram__ 数组：

__nram__ float left_nram[M * N];
__nram__ float right_nram[N * K];
__nram__ float result_nram[M * K];

__nram__ 表示变量存放在 MLU 的片上高速存储器（NRAM）中，访问延迟远低于全局内存。核函数首先通过 __memcpy 将左矩阵和右矩阵从 GDRAM 一次性拷贝到 NRAM：

__memcpy(left_nram, left, M * N * sizeof(float), GDRAM2NRAM);
__memcpy(right_nram, right, N * K * sizeof(float), GDRAM2NRAM);

然后使用与 CPU 完全相同的三重循环在 NRAM 上完成计算，最后将结果矩阵从 NRAM 拷贝回 GDRAM：

__memcpy(result, result_nram, M * K * sizeof(float), NRAM2GDRAM);

这样做的好处是：整个计算过程中，左、右矩阵的数据已经位于片上高速缓存，无需反复访问慢速的 GDRAM，理论上能大幅提升执行效率（虽然本示例中核函数仍然是串行的单线程块执行）。

主程序的流程与前一个示例几乎一致：分配主机内存和 MLU 设备内存，用随机数填充矩阵，拷贝数据到设备，分别运行 CPU 函数和 MLU 核函数并计时，最后比较结果的相对误差。计时部分使用了 cnrtPlaceNotifier 和 cnrtNotifierDuration（示例中写的是 cnrtNotifierElapsedTime 的别名或不同版本）。核函数启动时仍然使用单个线程块 dim = {1,1,1}，因此并没有利用 MLU 的并行计算能力，只是测试了将数据搬到 NRAM 后再计算的效果。

性能预期：由于 M=128, N=256, K=128，整个矩阵数据量较小（左矩阵 128×256×4=128KB，右矩阵 256×128×4=128KB），NRAM 容量足够容纳全部矩阵（寒武纪 MLU 的 NRAM 通常有几百 KB 到几 MB）。因此相比全局内存版本，NRAM 版本会减少大量片外访存，耗时应该显著降低。不过因为核函数仍然是纯标量串行计算，计算量本身并未改变，所以绝对时间依然远大于优化的并行版本。

矩阵乘向量 NRAM 实现

03 _vector_nram . mlu

这段代码是寒武纪 MLU 编程系列中的第三个示例，文件名 03_vector_nram.mlu 表明它在之前的基础上引入了向量化指令（__bang_matmul）来计算矩阵乘法，而不是手动编写三重循环。与前两个示例相比，核函数内部不再使用标量循环，而是直接调用硬件加速的矩阵乘内置函数，充分利用 MLU 的并行计算能力。

核函数的具体步骤：

1. 在片上存储器中声明三个数组：

   • __nram__ float left_nram[M * N]; 左矩阵，存放在 NRAM（高速缓存）。
   • __wram__ float right_wram[N * K]; 右矩阵，存放在 WRAM（另一种片上存储，可能专门为矩阵乘法优化）。
   • __nram__ float result_nram[M * K]; 结果矩阵，存放在 NRAM。

2. 通过 __memcpy 将数据从全局内存（GDRAM）搬运到片上：

   __memcpy(left_nram, left, M * N * sizeof(float), GDRAM2NRAM);
   __memcpy(right_wram, right, N * K * sizeof(float), GDRAM2WRAM);

3. 调用向量化矩阵乘法函数一次性完成计算：

   __bang_matmul(result_nram, left_nram, right_wram, m, n, k);

   这个函数会利用 MLU 的硬件矩阵运算单元（可能包含脉动阵列或多核并行）高效地执行 result = left × right，避免了显式的循环。

4. 将结果从 NRAM 拷贝回全局内存：

   __memcpy(result, result_nram, M * K * sizeof(float), NRAM2GDRAM);

主程序流程与前两个示例完全一致：分配主机和设备内存，随机生成左、右矩阵（元素范围 1.0~1.1），CPU 执行标量三重循环计算并计时，MLU 执行上述核函数并计时（使用 cnrtPlaceNotifier 和 cnrtNotifierDuration），最后比较 CPU 和 MLU 结果的相对误差，误差阈值 0.1，输出 03PASSED 或 03FAILED。若开启 DEBUG 宏，还会打印详细的矩阵内容。

与前两个示例的关键区别：

• 01_scalar.mlu：直接在全局内存上用标量三重循环计算，没有利用片上缓存，也没有并行。
• 02_scalar_nram.mlu：将数据预取到 NRAM，但计算仍然是标量三重循环，只是减少了访存延迟。
• 03_vector_nram.mlu：不仅将数据预取到片上，还使用了硬件加速的 __bang_matmul 指令，计算过程是向量化且高度并行的，执行效率远超前两个版本。

性能预期：由于 M=128, N=256, K=128，矩阵规模较小，CPU 标量版本耗时可能在几十毫秒量级，而 MLU 向量化版本由于硬件加速，耗时可能降到亚毫秒甚至微秒级。不过实际加速比取决于 MLU 的架构和 __bang_matmul 的实现。

矩阵乘多核向量 NRAM 实现

04 _vector_nram_blocks . mlu

这段代码是第四个示例，在向量化（__bang_matmul）的基础上，加入了多块并行和循环分块，以处理远超片上存储容量的超大矩阵。

矩阵规模：通过宏定义，最终左矩阵的行数 M 被扩展到 524288，列数 N=256；右矩阵 256×128；结果矩阵 524288×128。由于左矩阵太大，无法一次性放入 NRAM，因此需要拆分计算。

核函数内部逻辑：

• 右矩阵较小（256×128），一次性从全局内存拷贝到 WRAM，之后不再重复搬运。
• 左矩阵按行分块：每个计算块（taskId）负责连续的若干行（m_per_block = 总行数 / 块数），每个块内部再通过循环（LOOPS 次）每次处理一个更小的行块（m_per_loop 行），将这一小块左矩阵拷贝到 NRAM。
• 对每个小块左矩阵，调用 __bang_matmul 与右矩阵相乘，得到对应小块的结果矩阵（m_per_loop × K）。
• 将结果小块拷贝回全局结果矩阵的对应位置。

并行方式：主程序启动 BLOCKS 个并行块（dim.x = BLOCKS），每个块独立处理不同行范围，互不干扰。这样充分利用 MLU 的多个计算核心。

主程序：与之前类似，分配主机/设备内存、随机初始化、CPU 三重循环计算、MLU 执行核函数并计时、比较误差输出 PASS/FAIL。由于矩阵极大，CPU 计算会非常缓慢，而 MLU 通过并行和向量化获得巨大加速。

优化思路演进：从全局内存标量三重循环（01）→ 数据搬入 NRAM 但仍是标量（02）→ 使用向量化指令 __bang_matmul（03）→ 多块并行 + 循环分块以处理大矩阵（04）。

矩阵乘多核向量 NRAM 三级流水实现

05 _vector_nram_blocks_pipe3 . mlu

第五个示例 05_vector_nram_blocks_pipe3.mlu 在第四个示例（多块并行+循环分块）的基础上引入了三级流水线（双缓冲 + 异步拷贝），通过重叠数据传输与计算来隐藏内存访问延迟，进一步提高执行效率。核心变化如下：

新增参数：#define STAGES 2，表示使用两个缓冲区（双缓冲）。

片上存储加倍：左矩阵 NRAM 数组大小变为 M_PER_LOOP * N * STAGES，结果矩阵 NRAM 数组大小变为 M_PER_LOOP * K * STAGES，即每个缓冲区独立存放一块行数据，允许在执行计算的同时异步加载下一块。

核函数循环次数扩展：循环次数从 LOOPS 变为 LOOPS + STAGES（即 34 次），用于处理流水线启动（前 STAGES 次）和排空（后 STAGES 次）阶段。

循环体内三级操作：每次迭代同时触发最多三个阶段的异步操作，通过取模 % STAGES 切换缓冲区：

1. 左矩阵异步加载（当 loop < LOOPS 时）：
   使用 __memcpy_async 将左矩阵的第 loop 块拷贝到 left_nram 的 (loop % STAGES) 号缓冲区。

2. 矩阵乘计算（当 loop >= 1 && loop <= LOOPS 时）：
   对 (loop-1) % STAGES 号缓冲区中的左矩阵块，调用 __bang_matmul 计算乘积，结果存入结果 NRAM 的同号缓冲区。

3. 结果异步写回（当 loop >= STAGES 时）：
   将 (loop-2) % STAGES 号缓冲区中的计算结果，通过 __memcpy_async 写回全局结果矩阵的对应位置。

同步点：每次循环末尾调用 __sync()，确保当前循环中发起的异步内存操作（加载、计算、写回）在进入下一次循环前已完成，保证数据依赖正确。

右矩阵拷贝：也改为 __memcpy_async 异步拷贝，进一步流水线化。

主程序：与第四个示例完全相同（仍启动 128 个块，矩阵规模一致），唯一区别是核函数内部实现了三级流水线。

优化效果：在第四个示例中，每个循环是“拷贝左块 → 等待完成 → 矩阵乘 → 等待完成 → 写回结果 → 等待完成”的串行步骤；而本示例通过双缓冲和异步操作，使得第 t 次计算可以和第 t+1 次加载、第 t-1 次写回同时进行，有效隐藏了访存延迟，提升了整体吞吐量。

矩阵乘多核向量 SRAM 五级流水实现

06 _vector_sram_unions_pipe5 . mlu

第六个示例 06_vector_sram_blocks_pipe5.mlu 引入了更复杂的存储层次和并行结构，在前五个示例的基础上增加了 SRAM（共享内存） 和 多级流水线（文件名中的 pipe5 暗示五级流水线）。代码不再仅仅使用 NRAM 和 WRAM，而是加入了 __mlu_shared__ 声明的 SRAM，它可以被同一个计算簇（cluster）内的多个计算核（core）共享。

首先看宏定义的变化：M_PER_LOOP 从 128 减小到 64，新增了 M_PER_LOOP_BLOCK = M_PER_LOOP * STAGES = 128 和 M_PER_LOOP_UNION = M_PER_LOOP_BLOCK * 4 = 512（4 是隐含的 coreDim）。最终 M 仍然是 524288，但分块粒度更细。并行层次上，主程序启动核函数时使用 cnrtDim3_t dim = {CNRT_FUNC_TYPE_UNION1, 256, 1}，func_type = CNRT_FUNC_TYPE_UNION1，表示启动 256 个 union（每个 union 包含多个 core，这里 coreDim 为 4）。这样总共有 256×4=1024 个计算核同时工作。

核函数 multiplyMatricesMLU 中，声明了三组 SRAM 数组：left_sram、right_sram、result_sram，大小分别为 M_PER_LOOP_UNION * N * STAGES、N * K、M_PER_LOOP_UNION * K * STAGES。右矩阵先异步从 GDRAM 拷贝到 SRAM（代码中写了两次，可能是笔误），然后进入主循环。主循环将左矩阵按行分块：每次循环从 GDRAM 异步拷贝 m_per_loop_union 行（即 512 行）到 left_sram 的一个缓冲区中，然后调用 multiplyMatrices 函数（这是一个 __mlu_func__，会在每个 core 上执行）处理该块，最后将结果从 SRAM 异步写回 GDRAM。

multiplyMatrices 函数在 core 级别运行，它进一步利用 coreDim=4 将传入的 512 行数据拆分为每个 core 处理 128 行（因为 m / coreDim = 512/4=128），然后通过内层双缓冲流水线将左矩阵从 SRAM 异步拷贝到 NRAM、执行 __bang_matmul、再将结果从 NRAM 异步写回 SRAM。内层循环的 loops 固定为 STAGES（2），每个 core 每次处理 64 行（128/2=64），与 M_PER_LOOP 一致。内层使用了 __sync_io() 和 __sync_move() 来保证异步操作完成。

整个数据流动形成五级流水线：GDRAM → SRAM（cluster 级） → NRAM（core 级） → 计算 → NRAM → SRAM → GDRAM。右矩阵虽然从 GDRAM 拷贝到了 SRAM，但 multiplyMatrices 函数中使用的 right_wram 并未在入口函数中显式填充（可能假设已经存在或漏写），不过这并不影响理解其多层次存储的优化意图。

主程序的其他部分（初始化、CPU 计算、计时、结果比较）与之前示例基本相同，只是输出 06PASSED 或 06FAILED。这个示例展示了如何通过组合 SRAM、WRAM、NRAM 以及多级流水线，在 MLU 上实现大规模矩阵乘法的高效计算。

优化思路总结

从 01_scalar.mlu 到 06_vector_sram_blocks_pipe5.mlu，代码逐步展示了在寒武纪 MLU 上优化矩阵乘法的完整演进路径。整体优化思路可以概括为：从串行标量到向量化指令，从全局内存直接访问到多级片上存储（NRAM → WRAM → SRAM），从单线程块到多块并行，再到多核并行（union + core），并通过异步拷贝和双缓冲构建深度流水线，以重叠数据传输与计算，隐藏访存延迟，最终达到高吞吐、大规模矩阵的高效计算。

具体阶段如下：

1. 基线（01_scalar）：直接在全局内存（GDRAM）上用三重循环完成标量乘加，未使用片上存储，计算与访存串行，性能极低。

2. 数据预取（02_scalar_nram）：将左、右矩阵显式拷贝到片上 NRAM，减少访问 GDRAM 的次数。但计算仍为标量串行，访存延迟虽降，算力未释放。

3. 向量化指令（03_vector_nram）：用 Bang 语言内置的 __bang_matmul 硬件加速指令替换手工三重循环，将矩阵乘法交给专用计算单元，大幅提升计算密度。

4. 多块并行 + 分块（04_vector_nram_blocks）：面对超大左矩阵（M=524288），将 M 维度切分为多个独立块（128 个 block），每个块处理一行段，充分利用多个计算核心。同时每个块内部再将行段分成多轮循环（分块 tiling），以适应有限的 NRAM 容量，实现“右矩阵常驻 WRAM，左矩阵分批次搬运”。

5. 异步流水线（05_vector_nram_blocks_pipe3）：引入双缓冲（STAGES=2）和异步内存拷贝（__memcpy_async），将数据加载、计算、结果写回三级操作重叠。循环次数扩展为 LOOPS+STAGES，同一时刻加载第 t+1 块、计算第 t 块、写回第 t-1 块，隐藏了访存延迟。

6. 多级存储与多核协同（06_vector_sram_blocks_pipe5）：增加 __mlu_shared__ 声明的 SRAM（簇内共享），引入更细粒度的并行层次：任务维度为 union（每个 union 含 4 个 core）。数据流变为 GDRAM → SRAM（cluster 级） → NRAM（core 级） → 计算 → NRAM → SRAM → GDRAM，形成五级流水线。外部循环负责 GDRAM↔SRAM 的大块搬运，内部 __mlu_func__ 在 core 上运行，进一步将 SRAM 中的数据分给每个 core，并再次使用双缓冲流水线进行 NRAM 级别的计算。这样通过融合多级存储、多核并行和深度流水，最大限度压榨 MLU 的带宽与算力。

整体优化思路始终围绕 减少访存延迟、提高计算并行度、重叠通信与计算 这三大原则，从最简单的标量实现逐步演进到适合生产环境的高性能矩阵乘法实现。

实验运行

bash test.sh