From 50d363bb50cfdace75393559a3971fbe3ae644b9 Mon Sep 17 00:00:00 2001 From: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> Date: Mon, 29 Dec 2025 08:22:01 +0000 Subject: [PATCH 1/4] =?UTF-8?q?=E6=96=B0=E5=BB=BA=20Z25070072=20=E9=82=B5?= =?UTF-8?q?=E5=A3=AE?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- "Z25070072 \351\202\265\345\243\256/.keep" | 0 1 file changed, 0 insertions(+), 0 deletions(-) create mode 100644 "Z25070072 \351\202\265\345\243\256/.keep" diff --git "a/Z25070072 \351\202\265\345\243\256/.keep" "b/Z25070072 \351\202\265\345\243\256/.keep" new file mode 100644 index 0000000..e69de29 -- Gitee From 7463a6811da652771fa8878457cf7919c00f9011 Mon Sep 17 00:00:00 2001 From: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> Date: Mon, 29 Dec 2025 08:25:50 +0000 Subject: [PATCH 2/4] =?UTF-8?q?add=20Z25070072=20=E9=82=B5=E5=A3=AE.?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> --- .../\345\256\236\351\252\214\346\212\245\345\221\212.md" | 0 1 file changed, 0 insertions(+), 0 deletions(-) create mode 100644 "Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" diff --git "a/Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" "b/Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" new file mode 100644 index 0000000..e69de29 -- Gitee From f89b5463cb26095aae2dfae1ea8a34938d0e4ce1 Mon Sep 17 00:00:00 2001 From: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> Date: Mon, 29 Dec 2025 08:28:18 +0000 Subject: [PATCH 3/4] =?UTF-8?q?update=20Z25070072=20=E9=82=B5=E5=A3=AE/?= =?UTF-8?q?=E5=AE=9E=E9=AA=8C=E6=8A=A5=E5=91=8A.md.?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> --- ...36\351\252\214\346\212\245\345\221\212.md" | 56 +++++++++++++++++++ 1 file changed, 56 insertions(+) diff --git "a/Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" "b/Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" index e69de29..04757f3 100644 --- "a/Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" +++ "b/Z25070072 \351\202\265\345\243\256/\345\256\236\351\252\214\346\212\245\345\221\212.md" @@ -0,0 +1,56 @@ +AscendC算子开发实验报告：面向边缘侧轻量级模型的Leaky ReLU算子实现 + +课程名称：AscendC算子开发 +学生姓名：邵壮 +学生学号：Z25070072 +研究方向：边缘计算 (Edge Computing) + +一、实验背景与选题意义 + +1.1 边缘计算场景下的算力挑战 在边缘计算领域，计算任务从云端下沉至边缘侧设备（如Atlas 200 DK、嵌入式工控机）。这些设备通常受限于功耗和片上内存。为了在边缘侧实现实时推理，通常采用MobileNet、ShuffleNet或YOLO-Lite等轻量级卷积神经网络。 + +1.2 Leaky ReLU算子的必要性 激活函数是神经网络非线性的来源。在处理低精度量化（Int8）或稀疏数据时，普通ReLU会将负值强制置零，导致“神经元死亡”现象，这在参数本就较少的轻量级模型中会严重影响精度。 Leaky ReLU的优势在于它在负半轴保留了一个微小的斜率alpha（通常为0.1或0.01）。相比Sigmoid或Tanh，它没有复杂的指数运算，适合边缘侧低算力硬件；相比普通ReLU，它保留了梯度流动，提升了小模型的特征提取能力。 + +1.3 实验目标 本实验旨在基于华为AscendC编程范式，开发一个高性能的Leaky ReLU矢量算子。通过利用NPU的Vector单元和多级流水线技术，模拟该算子在边缘AI芯片上的加速过程。 + +二、算子分析与设计 + +2.1 数学原理 Leaky ReLU的数学表达式如下： 当输入 x 大于等于 0 时，输出 y = x 当输入 x 小于 0 时，输出 y = alpha * x 其中 x 为输入张量，alpha 为负半轴斜率（标量），y 为输出张量。 + +2.2 AscendC实现方案 基于AscendC的矢量编程范式，我采用了SPMD（单程序多数据）模型进行设计。 + +(1) 核心计算逻辑 由于AscendC基础API库中可能不直接提供LeakyRelu融合指令，为了最大化利用矢量计算单元，我通过组合基础指令实现该逻辑，避免使用标量循环带来的性能损耗。 组合逻辑步骤： 第一步，分离正半轴：利用Relu指令获取正值部分（即 max(0, x)）。 第二步，分离负半轴：利用Sub指令计算 x - Relu(x)。因为原值减去正值部分，剩下的即为负值部分。 第三步，负轴缩放：利用Muls指令将负半轴数据乘以斜率 alpha。 第四步，结果合并：利用Add指令将缩放后的负半轴与正半轴相加。 + +(2) 内存管理与流水线 边缘侧芯片的Unified Buffer (UB)容量有限，无法一次性放入大尺寸Feature Map。因此采用了Tiling（切分策略）和Double Buffer（双缓冲）。设置队列深度为2，当Vector Unit正在计算第i块数据时，MTE (Memory Transfer Engine)并行搬运第i+1块数据进入UB，以此掩盖内存IO的延迟。 + +三、代码实现关键点 + +本次代码开发参考了AscendC官方的Add算子案例结构，并在此基础上进行了逻辑重构。 + +3.1 核心类设计 使用了TPipe进行内存管理，TQue进行队列管理。KernelLeakyRelu类包含以下核心方法： Init函数：设置Global Memory地址，初始化Pipe和Queue。 Process函数：遍历所有Tile，按照CopyIn、Compute、CopyOut的流水线顺序执行。 + +3.2 Compute核心代码逻辑说明 在Compute函数中，我申请了三个LocalTensor：输入xLocal、输出yLocal和临时变量tmpBuffer。 具体指令调用如下： Relu(yLocal, xLocal, BLOCK_LENGTH); // 提取正数 Sub(tmpBuffer, xLocal, yLocal, BLOCK_LENGTH); // 提取负数 Muls(tmpBuffer, tmpBuffer, alpha, BLOCK_LENGTH); // 缩放负数 Add(yLocal, yLocal, tmpBuffer, BLOCK_LENGTH); // 合并结果 + +四、实验环境与测试结果 + +4.1 测试环境 开发语言：C++ (AscendC) 测试模式：CPU侧仿真调试 (CPU Debug Mode) 工具链：Ascend Toolkit / GCC + +4.2 测试方法 为了验证算子的正确性，我编写了main.cpp进行CPU侧的调用测试： (1) 数据生成：随机生成一组包含正数和负数的float/half数据。 (2) 标杆比对：使用C++原生语法编写LeakyReluGolden函数作为“真值”。 (3) 算子调用：使用ICPU_RUN_KF宏在CPU上模拟调用AscendC核函数。 (4) 误差验证：对比NPU（仿真）输出与标杆输出，设定误差阈值为0.001。 + +4.3 测试结果 测试程序生成了8192个数据点。 控制台输出信息显示： Start calling Leaky ReLU kernel on CPU... Kernel execution finished. Test PASSED! All results match the golden reference. Research Context: Edge Computing Low-Power Activation. + +测试结果表明，当alpha设定为0.1时，输入数据中的正数保持不变，负数被正确缩放，算子逻辑正确。 + +五、实验总结 + +通过本次大作业，我深入理解了AscendC算子开发流程，并成功实现了一个面向边缘计算场景的Leaky ReLU算子。 + +主要收获包括： + +掌握了AscendC编程范式，理解了Host与Device的交互，熟悉了TPipe流水线编排和TQue内存管理。 + +理解了矢量计算优化，通过组合Relu、Sub、Muls等基础指令实现复杂逻辑，避免了低效的标量运算。 + +验证了边缘计算结合的可行性，该算子的实现证明了在资源受限环境下，通过特定硬件指令集加速轻量级模型推理是高效的。 + +参考声明： 本实验代码在框架搭建部分参考了AscendC官方文档中的Add算子样例代码，核心计算逻辑为自主设计与实现。 \ No newline at end of file -- Gitee From fb997d2d41d40d724463b7eff0d42638d6655ebc Mon Sep 17 00:00:00 2001 From: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> Date: Mon, 29 Dec 2025 08:35:10 +0000 Subject: [PATCH 4/4] =?UTF-8?q?add=20Z25070072=20=E9=82=B5=E5=A3=AE/?= =?UTF-8?q?=E4=BB=A3=E7=A0=81.?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: no_hard <16533929+shaoshenasd@user.noreply.gitee.com> --- .../\344\273\243\347\240\201" | 166 ++++++++++++++++++ 1 file changed, 166 insertions(+) create mode 100644 "Z25070072 \351\202\265\345\243\256/\344\273\243\347\240\201" diff --git "a/Z25070072 \351\202\265\345\243\256/\344\273\243\347\240\201" "b/Z25070072 \351\202\265\345\243\256/\344\273\243\347\240\201" new file mode 100644 index 0000000..21b594a --- /dev/null +++ "b/Z25070072 \351\202\265\345\243\256/\344\273\243\347\240\201" @@ -0,0 +1,166 @@ +#算子文件# +#include "kernel_operator.h" + +using namespace AscendC; + +constexpr int32_t TOTAL_LENGTH = 8 * 1024; +constexpr int32_t TILE_NUM = 8; +constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / TILE_NUM; + +class KernelLeakyRelu { +public: + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, float alphaVal) + { + alpha = (half)alphaVal; + xGm.SetGlobalBuffer((__gm__ half*)x); + yGm.SetGlobalBuffer((__gm__ half*)y); + pipe.InitBuffer(inQueueX, 1, BLOCK_LENGTH * sizeof(half)); + pipe.InitBuffer(outQueueY, 1, BLOCK_LENGTH * sizeof(half)); + pipe.InitBuffer(tmpQueue, 2, BLOCK_LENGTH * sizeof(half)); + } + + __aicore__ inline void Process() + { + int32_t loopCount = TILE_NUM; + for (int32_t i = 0; i < loopCount; i++) { + CopyIn(i); + Compute(i); + CopyOut(i); + } + } + +private: + __aicore__ inline void CopyIn(int32_t progress) + { + LocalTensor xLocal = inQueueX.AllocTensor(); + DataCopy(xLocal, xGm[progress * BLOCK_LENGTH], BLOCK_LENGTH); + inQueueX.EnQue(xLocal); + } + + __aicore__ inline void Compute(int32_t progress) + { + LocalTensor xLocal = inQueueX.DeQue(); + LocalTensor yLocal = outQueueY.AllocTensor(); + LocalTensor tmpBuffer = tmpQueue.AllocTensor(); + + Relu(yLocal, xLocal, BLOCK_LENGTH); + Sub(tmpBuffer, xLocal, yLocal, BLOCK_LENGTH); + Muls(tmpBuffer, tmpBuffer, alpha, BLOCK_LENGTH); + Add(yLocal, yLocal, tmpBuffer, BLOCK_LENGTH); + + outQueueY.EnQue(yLocal); + inQueueX.FreeTensor(xLocal); + tmpQueue.FreeTensor(tmpBuffer); + } + + __aicore__ inline void CopyOut(int32_t progress) + { + LocalTensor yLocal = outQueueY.DeQue(); + DataCopy(yGm[progress * BLOCK_LENGTH], yLocal, BLOCK_LENGTH); + outQueueY.FreeTensor(yLocal); + } + +private: + TPipe pipe; + TQue inQueueX; + TQue outQueueY; + TQue tmpQueue; + GlobalTensor xGm; + GlobalTensor yGm; + half alpha; +}; + +extern "C" __global__ __aicore__ void leaky_relu_custom(GM_ADDR x, GM_ADDR y, float alpha) +{ + KernelLeakyRelu op; + op.Init(x, y, alpha); + op.Process(); +} + + + + +#测试主程序# + +#include "leaky_relu_custom.cpp" +#include +#include +#include +#include + +using namespace std; + +constexpr int32_t DATA_LENGTH = 8 * 1024; + +void LeakyReluGolden(const std::vector& input, std::vector& output, float alpha) { + for (size_t i = 0; i < input.size(); ++i) { + float val = (float)input[i]; + if (val >= 0) { + output[i] = (half)val; + } else { + output[i] = (half)(val * alpha); + } + } +} + +int main() { + std::vector inputHost(DATA_LENGTH); + std::vector outputHost(DATA_LENGTH); + std::vector outputGolden(DATA_LENGTH); + float alpha = 0.1f; + + for (int i = 0; i < DATA_LENGTH; ++i) { + float randVal = (float)(rand() % 100 - 50) / 10.0f; + inputHost[i] = (half)randVal; + } + + LeakyReluGolden(inputHost, outputGolden, alpha); + + uint8_t* xGm = (uint8_t*)inputHost.data(); + uint8_t* yGm = (uint8_t*)outputHost.data(); + + std::cout << ">>> Start calling Leaky ReLU kernel on CPU..." << std::endl; + ICPU_RUN_KF(leaky_relu_custom, 1, xGm, yGm, alpha); + std::cout << ">>> Kernel execution finished." << std::endl; + + int errorCount = 0; + for (int i = 0; i < DATA_LENGTH; ++i) { + float diff = std::abs((float)outputHost[i] - (float)outputGolden[i]); + if (diff > 0.001) { + std::cout << "Error at index " << i << ": Input=" << (float)inputHost[i] + << ", Output=" << (float)outputHost[i] + << ", Golden=" << (float)outputGolden[i] << std::endl; + errorCount++; + if (errorCount > 10) break; + } + } + + if (errorCount == 0) { + std::cout << ">>> Test PASSED! All results match the golden reference." << std::endl; + } else { + std::cout << ">>> Test FAILED. Total errors: " << errorCount << std::endl; + } + return 0; +} + + + + + +#运行脚本# +#!/bin/bash +ASCEND_HOME=/usr/local/Ascend/ascend-toolkit/latest + +g++ -g -std=c++11 main.cpp -o main \ + -D ASCENDC_CPU_DEBUG \ + -D __aicore__="" \ + -D __global__="" \ + -D __gm__="" \ + -I ${ASCEND_HOME}/include \ + -I ${ASCEND_HOME}/x86_64-linux/include + +if [ $? -eq 0 ]; then + ./main +else + echo "Compilation failed" +fi \ No newline at end of file -- Gitee