【代码侦探Challenge01-DivCustom】补全逐元素除法算子实现+PR 经验证可正确运行。-cann-outreach-AtomGit

已开启

【代码侦探Challenge01-DivCustom】补全逐元素除法算子实现+PR 经验证可正确运行。 #1839

zhangyatao9创建于 8月2日从zhangyatao9/cann-outreach： master 合入到 cann/cann-outreach： master

【代码侦探Challenge01-DivCustom】补全逐元素除法算子实现+PR 经验证可正确运行。 #1839

已开启

zhangyatao9创建于 8月2日

从zhangyatao9/cann-outreach master合入到cann/cann-outreach master

共 3 个文件变更+398-0

A2026/CANN-Code-Detective/Challenge01-DivCustom/Template/DivCustom-zhangyatao9/DivCustom/CMakeLists.txt+10-0

@@ -0,0 +1,10 @@
	1	+cmake_minimum_required(VERSION 3.16)
	2	+find_package(ASC REQUIRED)
	3	+project(kernel_samples LANGUAGES ASC CXX)
	4	+
	5	+add_executable(div_test
	6	+ div_custom.asc
	7	+)
	8	+target_compile_options(div_test PRIVATE
	9	+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
	10	+)

A2026/CANN-Code-Detective/Challenge01-DivCustom/Template/DivCustom-zhangyatao9/DivCustom/div_custom.asc+379-0

@@ -0,0 +1,379 @@
	1	+// div_custom.asc
	2	+//
	3	+// Challenge01 · DivCustom
	4	+// 逐元素除法算子: z[i] = x[i] / y[i]
	5	+//
	6	+// 整体流程（一句话）：
	7	+// 把 16384 个 float 分成 8 份给 8 个 AI Core，
	8	+// 每个核再把数据切成小片，用双缓冲流水线完成：
	9	+// 从全局内存搬到片上内存 → 做除法 → 存回全局内存
	10	+//
	11	+// 数据流图：
	12	+// GM(x,y) --CopyIn--> UB(xLocal,yLocal) --Compute(Div)--> UB(zLocal) --CopyOut--> GM(z)
	13	+
	14	+#include <cstdint>
	15	+#include <iostream>
	16	+#include <vector>
	17	+#include <algorithm>
	18	+#include <iterator>
	19	+#include "acl/acl.h" // AscendCL 运行时库（Host 侧用）
	20	+#include "kernel_operator.h" // AscendC 算子开发库（Device 侧用）
	21	+
	22	+// ------------------------------------------------------------
	23	+// 常量定义
	24	+// ------------------------------------------------------------
	25	+// BUFFER_NUM = 2 表示每个队列有 2 个槽位，实现双缓冲（Double Buffer）
	26	+// 这样第 n 次循环在计算时，第 n+1 次循环可以同时搬运数据，隐藏搬运延迟
	27	+constexpr uint32_t BUFFER_NUM = 2;
	28	+
	29	+// ------------------------------------------------------------
	30	+// Tiling 参数结构体
	31	+// 作用：告诉 kernel 总共有多少数据、要切成几份来处理
	32	+// 注意：这个结构体在 Host 侧赋值，然后按值传给 Device 侧 kernel
	33	+// ------------------------------------------------------------
	34	+struct DivCustomTilingData
	35	+{
	36	+ uint32_t totalLength; // 总数据量（元素个数），这里是 16384
	37	+ uint32_t tileNum; // 每个核内部再切成几份，这里是 8
	38	+};
	39	+
	40	+// ------------------------------------------------------------
	41	+// KernelDiv 类：算子的核心实现
	42	+// 职责：管理全局内存、片上内存、队列、搬运和计算
	43	+// ------------------------------------------------------------
	44	+class KernelDiv {
	45	+public:
	46	+ // 构造函数：什么都不做
	47	+ __aicore__ inline KernelDiv() {}
	48	+
	49	+ // --------------------------------------------------------
	50	+ // Init：初始化函数
	51	+ // 输入参数：
	52	+ // x, y, z — 全局内存（GM）上的三个数组首地址
	53	+ // totalLength — 总数据量
	54	+ // tileNum — 每个核内部的切分数
	55	+ //
	56	+ // 做的事情：
	57	+ // 1. 算出当前核负责多少数据（总长 ÷ 核数）
	58	+ // 2. 算出每次循环搬运/计算多少数据
	59	+ // 3. 让 GlobalTensor 指向当前核负责的那段 GM
	60	+ // 4. 在片上内存（UB）上分配队列缓冲区
	61	+ // --------------------------------------------------------
	62	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z,
	63	+ uint32_t totalLength, uint32_t tileNum)
	64	+ {
	65	+ // step 1: 每个核处理的数据量 = 总数据量 ÷ AI Core 数量
	66	+ // AscendC::GetBlockNum() 返回当前启用的核数（这里是 8）
	67	+ this->blockLength = totalLength / AscendC::GetBlockNum();
	68	+
	69	+ // step 2: 记录核内切分数
	70	+ this->tileNum = tileNum;
	71	+
	72	+ // step 3: 每次循环处理的数据量
	73	+ // 公式：核内数据量 ÷ 切分数 ÷ 双缓冲槽位数
	74	+ // 例如：2048 ÷ 8 ÷ 2 = 128，每次处理 128 个 float
	75	+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
	76	+
	77	+ // step 4: 让 GlobalTensor 指向当前核负责的那段全局内存
	78	+ // AscendC::GetBlockIdx() 返回当前核的编号（0 ~ 7）
	79	+ // 每个核的起始位置 = 数组首地址 + 当前核编号 × 每个核的数据量
	80	+ xGm.SetGlobalBuffer(
	81	+ (__gm__ float)x + this->blockLength AscendC::GetBlockIdx(),
	82	+ this->blockLength);
	83	+ yGm.SetGlobalBuffer(
	84	+ (__gm__ float)y + this->blockLength AscendC::GetBlockIdx(),
	85	+ this->blockLength);
	86	+ zGm.SetGlobalBuffer(
	87	+ (__gm__ float)z + this->blockLength AscendC::GetBlockIdx(),
	88	+ this->blockLength);
	89	+
	90	+ // step 5: 在片上内存（Unified Buffer）上分配队列缓冲区
	91	+ // 每个队列有 BUFFER_NUM 个槽位，每个槽位能放 tileLength 个 float
	92	+ // pipe.InitBuffer(队列名, 槽位数, 每个槽位的字节数)
	93	+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
	94	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
	95	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
	96	+ }
	97	+
	98	+ // --------------------------------------------------------
	99	+ // Process：主处理循环
	100	+ // 循环次数 = tileNum × BUFFER_NUM
	101	+ // 为什么是这个数？
	102	+ // - tileNum = 8，表示把核内数据切成 8 个大块
	103	+ // - BUFFER_NUM = 2，每个大块再分成 2 份来实现流水
	104	+ // - 所以总共 8×2 = 16 次循环
	105	+ // 每次循环执行：CopyIn → Compute → CopyOut
	106	+ // --------------------------------------------------------
	107	+ __aicore__ inline void Process()
	108	+ {
	109	+ const int32_t loopCount = this->tileNum * BUFFER_NUM;
	110	+ for (int32_t i = 0; i < loopCount; ++i) {
	111	+ CopyIn(i); // 从 GM 搬到 UB
	112	+ Compute(i); // 在 UB 上做除法
	113	+ CopyOut(i); // 从 UB 搬回 GM
	114	+ }
	115	+ }
	116	+
	117	+private:
	118	+ // --------------------------------------------------------
	119	+ // CopyIn：数据搬入
	120	+ // 作用：从全局内存（GM）把当前进度对应的数据块搬到片上内存（UB）
	121	+ // progress：当前是第几次循环（0 ~ 15）
	122	+ //
	123	+ // 步骤：
	124	+ // 1. 从输入队列分配一块 UB 空间（AllocTensor）
	125	+ // 2. 把 GM 上对应位置的数据拷贝过来（DataCopy）
	126	+ // 3. 把这块 UB 空间入队，通知下游“数据准备好了”（EnQue）
	127	+ // --------------------------------------------------------
	128	+ __aicore__ inline void CopyIn(int32_t progress)
	129	+ {
	130	+ // 从 x 输入队列分配一块 UB 空间
	131	+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
	132	+ // 从 y 输入队列分配一块 UB 空间
	133	+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
	134	+
	135	+ // 把 GM 上第 progress 块数据拷贝到 UB
	136	+ // xGm[progress * this->tileLength] 是 GM 上的起始位置
	137	+ // this->tileLength 是要拷贝的元素个数
	138	+ AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
	139	+ AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
	140	+
	141	+ // 入队：告诉 Compute 阶段数据已就绪
	142	+ inQueueX.EnQue(xLocal);
	143	+ inQueueY.EnQue(yLocal);
	144	+ }
	145	+
	146	+ // --------------------------------------------------------
	147	+ // Compute：计算
	148	+ // 作用：从队列取出输入数据，执行逐元素除法，结果放入输出队列
	149	+ //
	150	+ // 步骤：
	151	+ // 1. 从输入队列取出 x 和 y 的 UB 数据（DeQue）
	152	+ // 2. 从输出队列分配一块 UB 空间存放结果
	153	+ // 3. 执行 Div 指令：zLocal = xLocal / yLocal
	154	+ // 4. 把结果入队（EnQue）
	155	+ // 5. 释放输入数据的 UB 空间（FreeTensor），供后续循环复用
	156	+ // --------------------------------------------------------
	157	+ __aicore__ inline void Compute(int32_t progress)
	158	+ {
	159	+ // 从输入队列取出已经搬运好的数据
	160	+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
	161	+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
	162	+
	163	+ // 从输出队列分配一块空间存放计算结果
	164	+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
	165	+
	166	+ // 执行逐元素除法：zLocal[i] = xLocal[i] / yLocal[i]
	167	+ // 第三个参数是元素个数（this->tileLength）
	168	+ AscendC::Div(zLocal, xLocal, yLocal, this->tileLength);
	169	+
	170	+ // 把计算结果入队，通知 CopyOut 阶段可以搬出了
	171	+ outQueueZ.EnQue(zLocal);
	172	+
	173	+ // 释放输入数据的 UB 空间
	174	+ // 如果不释放，下次循环就没空间分配了
	175	+ inQueueX.FreeTensor(xLocal);
	176	+ inQueueY.FreeTensor(yLocal);
	177	+ }
	178	+
	179	+ // --------------------------------------------------------
	180	+ // CopyOut：结果搬出
	181	+ // 作用：从输出队列取出计算结果，搬回全局内存
	182	+ //
	183	+ // 步骤：
	184	+ // 1. 从输出队列取出结果数据（DeQue）
	185	+ // 2. 把 UB 上的结果拷贝回 GM 上对应的位置（DataCopy）
	186	+ // 3. 释放 UB 空间（FreeTensor）
	187	+ // --------------------------------------------------------
	188	+ __aicore__ inline void CopyOut(int32_t progress)
	189	+ {
	190	+ // 从输出队列取出计算结果
	191	+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
	192	+
	193	+ // 把结果从 UB 搬回 GM
	194	+ // 注意：DataCopy 的第一个参数是目标地址，第二个是源地址
	195	+ AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
	196	+
	197	+ // 释放 UB 空间
	198	+ outQueueZ.FreeTensor(zLocal);
	199	+ }
	200	+
	201	+private:
	202	+ // --------------------------------------------------------
	203	+ // 硬件资源（这些是 AscendC 的核心概念）
	204	+ // --------------------------------------------------------
	205	+
	206	+ // TPipe：管道，统一管理片上内存（UB）的分配和回收
	207	+ AscendC::TPipe pipe;
	208	+
	209	+ // TQue：数据队列
	210	+ // 模板参数1：队列绑定的硬件位置
	211	+ // VECIN — 向量计算单元的输入侧
	212	+ // VECOUT — 向量计算单元的输出侧
	213	+ // 模板参数2：队列深度（槽位数），这里是 2（双缓冲）
	214	+ // inQueueX / inQueueY：输入数据的队列（GM → UB）
	215	+ // outQueueZ：输出结果的队列（UB → GM）
	216	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
	217	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;
	218	+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	219	+
	220	+ // GlobalTensor：全局内存张量，用来读写 GM 上的数据
	221	+ // 相当于一个“带偏移的指针”，知道从哪里读、读多长
	222	+ AscendC::GlobalTensor<float> xGm; // 输入 x
	223	+ AscendC::GlobalTensor<float> yGm; // 输入 y
	224	+ AscendC::GlobalTensor<float> zGm; // 输出 z
	225	+
	226	+ // 以下是 Init 阶段算出来的参数
	227	+ uint32_t blockLength; // 当前核负责的总元素数（2048）
	228	+ uint32_t tileNum; // 核内切分数（8）
	229	+ uint32_t tileLength; // 每次循环处理的元素数（128）
	230	+};
	231	+
	232	+// ------------------------------------------------------------
	233	+// div_custom：kernel 入口函数
	234	+// 这是 Device 侧运行的函数，8 个核各跑一份
	235	+// 参数：
	236	+// x, y, z — GM 上的数组首地址（类型是 GM_ADDR = __gm__ uint8_t*）
	237	+// tiling — 按值传递的 Tiling 参数
	238	+//
	239	+// 执行流程：
	240	+// 1. 实例化 KernelDiv 对象
	241	+// 2. 调用 Init 完成初始化
	242	+// 3. 调用 Process 进入主循环
	243	+// ------------------------------------------------------------
	244	+__global__ __aicore__ void div_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z,
	245	+ DivCustomTilingData tiling)
	246	+{
	247	+ KernelDiv op;
	248	+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
	249	+ op.Process();
	250	+}
	251	+
	252	+// ------------------------------------------------------------
	253	+// kernel_div：Host 侧的封装函数
	254	+// 作用：完成 Host → Device 的数据搬运、启动 kernel、Device → Host 的结果回拷
	255	+//
	256	+// 完整流程：
	257	+// 1. 初始化 AscendCL 运行环境
	258	+// 2. 在 Device 上申请三段全局内存（x, y, z）
	259	+// 3. 把输入数据从 Host 拷贝到 Device
	260	+// 4. 组装 Tiling 参数，以 8 个核启动 kernel
	261	+// 5. 等待 kernel 执行完成
	262	+// 6. 把结果从 Device 拷回 Host
	263	+// 7. 释放资源，清理环境
	264	+// ------------------------------------------------------------
	265	+std::vector<float> kernel_div(std::vector<float>& x, std::vector<float>& y)
	266	+{
	267	+ // 并行参数：8 个核，每个核内部切 8 份
	268	+ constexpr uint32_t blockDim = 8;
	269	+ constexpr uint32_t tileNum = 8;
	270	+ const uint32_t totalLength = static_cast<uint32_t>(x.size());
	271	+
	272	+ // -------- step 1: 初始化 AscendCL 运行环境 --------
	273	+ aclInit(nullptr); // 初始化 ACL
	274	+ aclrtSetDevice(0); // 使用 0 号设备
	275	+ aclrtStream stream = nullptr;
	276	+ aclrtCreateStream(&stream); // 创建 Stream（任务队列）
	277	+
	278	+ // -------- step 2: 在 Device 上申请全局内存 --------
	279	+ const uint32_t dataSize = totalLength * sizeof(float); // 总字节数
	280	+
	281	+ // 保持 uint8_t* 类型以匹配 GM_ADDR（Device 侧是 __gm__ uint8_t*）
	282	+ uint8_t* xDevice = nullptr;
	283	+ uint8_t* yDevice = nullptr;
	284	+ uint8_t* zDevice = nullptr;
	285	+
	286	+ // aclrtMalloc 的参数是 void**，所以需要显式转型
	287	+ aclrtMalloc(reinterpret_cast<void**>(&xDevice), dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
	288	+ aclrtMalloc(reinterpret_cast<void**>(&yDevice), dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
	289	+ aclrtMalloc(reinterpret_cast<void**>(&zDevice), dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
	290	+
	291	+ // -------- step 3: Host → Device 拷贝输入数据 --------
	292	+ aclrtMemcpy(xDevice, dataSize, x.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
	293	+ aclrtMemcpy(yDevice, dataSize, y.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
	294	+
	295	+ // -------- step 4: 组装 Tiling 参数，启动 kernel --------
	296	+ DivCustomTilingData tiling{totalLength, tileNum};
	297	+
	298	+ // <<<blockDim, nullptr, stream>>> 的含义：
	299	+ // blockDim = 8 — 启用 8 个 AI Core
	300	+ // nullptr — 不使用共享内存
	301	+ // stream — 在这个 Stream 上执行
	302	+ // 直接把 uint8_t* 传给 GM_ADDR 参数，类型天然匹配
	303	+ div_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
	304	+
	305	+ // 等待 kernel 执行完成
	306	+ aclrtSynchronizeStream(stream);
	307	+
	308	+ // -------- step 5: Device → Host 拷贝结果 --------
	309	+ std::vector<float> z(totalLength);
	310	+ aclrtMemcpy(z.data(), dataSize, zDevice, dataSize, ACL_MEMCPY_DEVICE_TO_HOST);
	311	+
	312	+ // -------- step 6: 释放资源 --------
	313	+ aclrtFree(xDevice);
	314	+ aclrtFree(yDevice);
	315	+ aclrtFree(zDevice);
	316	+ aclrtDestroyStream(stream);
	317	+ aclrtResetDevice(0);
	318	+ aclFinalize();
	319	+
	320	+ return z;
	321	+}
	322	+
	323	+// ------------------------------------------------------------
	324	+// VerifyResult：验证函数
	325	+// 比较 output 和 golden 是否一致，打印前 20 个值
	326	+// ------------------------------------------------------------
	327	+uint32_t VerifyResult(std::vector<float>& output, std::vector<float>& golden)
	328	+{
	329	+ // 打印前 20 个值的 lambda 函数
	330	+ auto printTensor = [](std::vector<float>& tensor, const char* name) {
	331	+ constexpr size_t maxPrintSize = 20;
	332	+ std::cout << name << ": ";
	333	+ std::copy(tensor.begin(),
	334	+ tensor.begin() + std::min(tensor.size(), maxPrintSize),
	335	+ std::ostream_iterator<float>(std::cout, " "));
	336	+ if (tensor.size() > maxPrintSize) {
	337	+ std::cout << "...";
	338	+ }
	339	+ std::cout << std::endl;
	340	+ };
	341	+
	342	+ printTensor(output, "Output");
	343	+ printTensor(golden, "Golden");
	344	+
	345	+ // 逐元素比较
	346	+ if (std::equal(golden.begin(), golden.end(), output.begin())) {
	347	+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
	348	+ return 0;
	349	+ } else {
	350	+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
	351	+ return 1;
	352	+ }
	353	+}
	354	+
	355	+// ------------------------------------------------------------
	356	+// main 函数
	357	+// 1. 构造 16384 个 1.2 和 2.3
	358	+// 2. 调用 kernel_div 执行逐元素除法
	359	+// 3. 用 golden（1.2/2.3 = 0.521739）验证结果
	360	+// ------------------------------------------------------------
	361	+int32_t main(int32_t argc, char* argv[])
	362	+{
	363	+ constexpr uint32_t totalLength = 8 * 2048; // 16384
	364	+ constexpr float valueX = 1.2f;
	365	+ constexpr float valueY = 2.3f;
	366	+
	367	+ // 构造输入数据：16384 个 1.2 和 16384 个 2.3
	368	+ std::vector<float> x(totalLength, valueX);
	369	+ std::vector<float> y(totalLength, valueY);
	370	+
	371	+ // 调用 kernel 得到输出
	372	+ std::vector<float> output = kernel_div(x, y);
	373	+
	374	+ // 构造期望结果：16384 个 1.2/2.3 = 0.521739
	375	+ std::vector<float> golden(totalLength, valueX / valueY);
	376	+
	377	+ // 验证并返回
	378	+ return VerifyResult(output, golden);
	379	+}

A2026/CANN-Code-Detective/Challenge01-DivCustom/Template/DivCustom-zhangyatao9/DivCustom/run.sh+9-0

@@ -0,0 +1,9 @@
	1	+#!/bin/bash
	2	+# 激活cann环境，可根据实际情况修改，在线环境一般不用修改，
	3	+source "$ASCEND_TOOLKIT_HOME/set_env.sh"
	4	+mkdir -p build
	5	+export ASC_DIR="$ASCEND_HOME_PATH/aarch64-linux/tikcpp/ascendc_kernel_cmake/"
	6	+cd build/ && \
	7	+cmake .. && \
	8	+make && \
	9	+./div_test