Star375
2459
代码
代码
Issues167
Pull Requests713
Actions
项目讨论
Wiki
分析
Star375
2459
已合并
【代码侦探Challenge05】补全gcw_93rlw6ed的MulCustom逐元素乘法算子 #2433
gcw_93rlw6ed创建于 9月1日
【代码侦探Challenge05】补全gcw_93rlw6ed的MulCustom逐元素乘法算子 #2433
已合并
gcw_93rlw6ed创建于 9月1日
从gcw_93rlw6ed/cann-outreachchallenge05-mulcustom合入到cann/cann-outreachmaster
共 3 个文件变更+395-0
A2026/CANN-Code-Detective/Challenge05-MulCustom/gcw_93rlw6ed/MulCustom/CMakeLists.txt+11-0
@@ -0,0 +1,11 @@
1+# mul_custom.asc 逐元素乘法算子构建脚本（与模板一致，无需修改）
2+cmake_minimum_required(VERSION 3.16)
3+find_package(ASC REQUIRED)
4+project(kernel_samples LANGUAGES ASC CXX)
5+ 
6+add_executable(mul_test
7+ mul_custom.asc
8+)
9+target_compile_options(mul_test PRIVATE
10+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
11+)
A2026/CANN-Code-Detective/Challenge05-MulCustom/gcw_93rlw6ed/MulCustom/mul_custom.asc+375-0
@@ -0,0 +1,375 @@
1+// mul_custom.asc
2+// 逐元素乘法算子：z[i] = x[i] * y[i]
3+//
4+// 学习自第2章 Add 算子示例，实现包含以下部分：
5+// 1. KernelMul 类成员变量声明
6+// 2. Init —— 初始化 TPipe、TQue、GlobalTensor
7+// 3. Process —— 主循环调用 CopyIn → Compute → CopyOut
8+// 4. CopyIn —— 从 GM 搬入数据到 UB
9+// 5. Compute —— 执行 Mul 计算
10+// 6. CopyOut —— 将结果从 UB 搬回 GM
11+// 7. mul_custom kernel 入口函数
12+// 8. kernel_mul host 侧函数
13+// 9. main 函数中调用 kernel_mul
14+ 
15+#include <cstdint>
16+#include <iostream>
17+#include <vector>
18+#include <algorithm>
19+#include <iterator>
20+#include "acl/acl.h"
21+#include "kernel_operator.h"
22+ 
23+constexpr uint32_t BUFFER_NUM = 2; // Double Buffer 队列深度
24+ 
25+struct MulCustomTilingData
26+{
27+ uint32_t totalLength;
28+ uint32_t tileNum;
29+};
30+ 
31+class KernelMul {
32+public:
33+ __aicore__ inline KernelMul() {}
34+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
35+ {
36+ // 1. 单核处理长度：totalLength 均分到每个 AI Core
37+ this->blockLength = totalLength / AscendC::GetBlockNum();
38+ this->tileNum = tileNum;
39+ // 2. 分片长度：单核数据按 tileNum 切分，每份再拆成 BUFFER_NUM 块构成双缓冲
40+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
41+ 
42+ // 3. 每个核通过 GetBlockIdx() 偏移，只处理自己的那段 Global Memory
43+ uint32_t blockOffset = this->blockLength * AscendC::GetBlockIdx();
44+ xGm.SetGlobalBuffer((__gm__ float *)x + blockOffset, this->blockLength);
45+ yGm.SetGlobalBuffer((__gm__ float *)y + blockOffset, this->blockLength);
46+ zGm.SetGlobalBuffer((__gm__ float *)z + blockOffset, this->blockLength);
47+ 
48+ // 4. 为三个队列各分配 BUFFER_NUM 块 UB 缓冲，每块 tileLength 个 float
49+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
50+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
51+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
52+ }
53+ __aicore__ inline void Process()
54+ {
55+ // 主循环：tileNum * BUFFER_NUM 次迭代，三段流水通过队列解耦
56+ int32_t loopCount = this->tileNum * BUFFER_NUM;
57+ for (int32_t i = 0; i < loopCount; ++i) {
58+ CopyIn(i);
59+ Compute(i);
60+ CopyOut(i);
61+ }
62+ }
63+ 
64+private:
65+ __aicore__ inline void CopyIn(int32_t progress)
66+ {
67+ // 从 GM 搬入第 progress 个分片到 UB，并入队
68+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
69+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
70+ 
71+ AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
72+ AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
73+ 
74+ inQueueX.EnQue(xLocal);
75+ inQueueY.EnQue(yLocal);
76+ }
77+ __aicore__ inline void Compute(int32_t progress)
78+ {
79+ // 出队输入，执行逐元素乘法 zLocal = xLocal * yLocal，结果入队
80+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
81+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
82+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
83+ 
84+ AscendC::Mul(zLocal, xLocal, yLocal, this->tileLength);
85+ 
86+ outQueueZ.EnQue<float>(zLocal);
87+ 
88+ inQueueX.FreeTensor(xLocal);
89+ inQueueY.FreeTensor(yLocal);
90+ }
91+ __aicore__ inline void CopyOut(int32_t progress)
92+ {
93+ // 出队结果，从 UB 搬回 GM
94+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
95+ 
96+ AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
97+ 
98+ outQueueZ.FreeTensor(zLocal);
99+ }
100+ 
101+private:
102+ AscendC::TPipe pipe;
103+ 
104+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
105+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;
106+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
107+ 
108+ AscendC::GlobalTensor<float> xGm;
109+ AscendC::GlobalTensor<float> yGm;
110+ AscendC::GlobalTensor<float> zGm;
111+ 
112+ uint32_t blockLength;
113+ uint32_t tileNum;
114+ uint32_t tileLength;
115+};
116+ 
117+// Kernel 入口函数
118+__global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, MulCustomTilingData tiling)
119+{
120+ // 显式指定 Kernel 任务类型为 AIV-only（Vector 核）
121+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
122+ KernelMul op;
123+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
124+ op.Process();
125+}
126+ 
127+// ----- Host 侧：kernel 直调封装 -----
128+ 
129+std::vector<float> kernel_mul(std::vector<float> &x, std::vector<float> &y)
130+{
131+ if (x.size() != y.size()) {
132+ std::cerr << "Error: input vectors must have the same size." << std::endl;
133+ return {};
134+ }
135+ if (x.empty()) {
136+ std::cerr << "Error: input vectors must not be empty." << std::endl;
137+ return {};
138+ }
139+ 
140+ // 1~2. 多核与分片参数：8 核并行，每核数据再分 8 个 tile（×BUFFER_NUM 双缓冲）
141+ const uint32_t totalLength = static_cast<uint32_t>(x.size());
142+ constexpr uint32_t blockDim = 8;
143+ constexpr uint32_t tileNum = 8;
144+ const size_t byteSize = totalLength * sizeof(float);
145+ 
146+ std::vector<float> output(totalLength, 0.0f);
147+ 
148+ // 3. 初始化 ACL 运行环境
149+ auto ret = aclInit(nullptr);
150+ if (ret != ACL_SUCCESS) {
151+ std::cerr << "aclInit failed. ERROR: " << ret << std::endl;
152+ return output;
153+ }
154+ 
155+ int32_t deviceId = 0;
156+ ret = aclrtSetDevice(deviceId);
157+ if (ret != ACL_SUCCESS) {
158+ std::cerr << "aclrtSetDevice failed. ERROR: " << ret << std::endl;
159+ aclFinalize();
160+ return output;
161+ }
162+ 
163+ aclrtStream stream = nullptr;
164+ ret = aclrtCreateStream(&stream);
165+ if (ret != ACL_SUCCESS) {
166+ std::cerr << "aclrtCreateStream failed. ERROR: " << ret << std::endl;
167+ aclrtResetDevice(deviceId);
168+ aclFinalize();
169+ return output;
170+ }
171+ 
172+ float *xHost = nullptr;
173+ float *yHost = nullptr;
174+ float *zHost = nullptr;
175+ 
176+ uint8_t *xDevice = nullptr;
177+ uint8_t *yDevice = nullptr;
178+ uint8_t *zDevice = nullptr;
179+ 
180+ // 4. 申请 Host/Device 内存
181+ ret = aclrtMallocHost(reinterpret_cast<void **>(&xHost), byteSize);
182+ if (ret != ACL_SUCCESS) {
183+ std::cerr << "aclrtMallocHost for xHost failed. ERROR: " << ret << std::endl;
184+ aclrtDestroyStream(stream);
185+ aclrtResetDevice(deviceId);
186+ aclFinalize();
187+ return output;
188+ }
189+ ret = aclrtMallocHost(reinterpret_cast<void **>(&yHost), byteSize);
190+ if (ret != ACL_SUCCESS) {
191+ std::cerr << "aclrtMallocHost for yHost failed. ERROR: " << ret << std::endl;
192+ aclrtFreeHost(xHost);
193+ aclrtDestroyStream(stream);
194+ aclrtResetDevice(deviceId);
195+ aclFinalize();
196+ return output;
197+ }
198+ ret = aclrtMallocHost(reinterpret_cast<void **>(&zHost), byteSize);
199+ if (ret != ACL_SUCCESS) {
200+ std::cerr << "aclrtMallocHost for zHost failed. ERROR: " << ret << std::endl;
201+ aclrtFreeHost(xHost);
202+ aclrtFreeHost(yHost);
203+ aclrtDestroyStream(stream);
204+ aclrtResetDevice(deviceId);
205+ aclFinalize();
206+ return output;
207+ }
208+ 
209+ std::copy(x.begin(), x.end(), xHost);
210+ std::copy(y.begin(), y.end(), yHost);
211+ 
212+ ret = aclrtMalloc(reinterpret_cast<void **>(&xDevice), byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
213+ if (ret != ACL_SUCCESS) {
214+ std::cerr << "aclrtMalloc for xDevice failed. ERROR: " << ret << std::endl;
215+ aclrtFreeHost(xHost);
216+ aclrtFreeHost(yHost);
217+ aclrtFreeHost(zHost);
218+ aclrtDestroyStream(stream);
219+ aclrtResetDevice(deviceId);
220+ aclFinalize();
221+ return output;
222+ }
223+ ret = aclrtMalloc(reinterpret_cast<void **>(&yDevice), byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
224+ if (ret != ACL_SUCCESS) {
225+ std::cerr << "aclrtMalloc for yDevice failed. ERROR: " << ret << std::endl;
226+ aclrtFree(xDevice);
227+ aclrtFreeHost(xHost);
228+ aclrtFreeHost(yHost);
229+ aclrtFreeHost(zHost);
230+ aclrtDestroyStream(stream);
231+ aclrtResetDevice(deviceId);
232+ aclFinalize();
233+ return output;
234+ }
235+ ret = aclrtMalloc(reinterpret_cast<void **>(&zDevice), byteSize, ACL_MEM_MALLOC_HUGE_FIRST);
236+ if (ret != ACL_SUCCESS) {
237+ std::cerr << "aclrtMalloc for zDevice failed. ERROR: " << ret << std::endl;
238+ aclrtFree(xDevice);
239+ aclrtFree(yDevice);
240+ aclrtFreeHost(xHost);
241+ aclrtFreeHost(yHost);
242+ aclrtFreeHost(zHost);
243+ aclrtDestroyStream(stream);
244+ aclrtResetDevice(deviceId);
245+ aclFinalize();
246+ return output;
247+ }
248+ 
249+ // 5. 输入数据 Host -> Device
250+ ret = aclrtMemcpy(xDevice, byteSize, xHost, byteSize, ACL_MEMCPY_HOST_TO_DEVICE);
251+ if (ret != ACL_SUCCESS) {
252+ std::cerr << "aclrtMemcpy to xDevice failed. ERROR: " << ret << std::endl;
253+ aclrtFree(xDevice);
254+ aclrtFree(yDevice);
255+ aclrtFree(zDevice);
256+ aclrtFreeHost(xHost);
257+ aclrtFreeHost(yHost);
258+ aclrtFreeHost(zHost);
259+ aclrtDestroyStream(stream);
260+ aclrtResetDevice(deviceId);
261+ aclFinalize();
262+ return output;
263+ }
264+ ret = aclrtMemcpy(yDevice, byteSize, yHost, byteSize, ACL_MEMCPY_HOST_TO_DEVICE);
265+ if (ret != ACL_SUCCESS) {
266+ std::cerr << "aclrtMemcpy to yDevice failed. ERROR: " << ret << std::endl;
267+ aclrtFree(xDevice);
268+ aclrtFree(yDevice);
269+ aclrtFree(zDevice);
270+ aclrtFreeHost(xHost);
271+ aclrtFreeHost(yHost);
272+ aclrtFreeHost(zHost);
273+ aclrtDestroyStream(stream);
274+ aclrtResetDevice(deviceId);
275+ aclFinalize();
276+ return output;
277+ }
278+ 
279+ // 6. 组装 tiling 并启动 kernel：8 核 × tileNum 8 × BUFFER_NUM 2
280+ MulCustomTilingData tiling;
281+ tiling.totalLength = totalLength;
282+ tiling.tileNum = tileNum;
283+ 
284+ mul_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
285+ 
286+ // 7. 同步流，等待 kernel 执行完成
287+ ret = aclrtSynchronizeStream(stream);
288+ if (ret != ACL_SUCCESS) {
289+ std::cerr << "aclrtSynchronizeStream failed. ERROR: " << ret << std::endl;
290+ aclrtFree(xDevice);
291+ aclrtFree(yDevice);
292+ aclrtFree(zDevice);
293+ aclrtFreeHost(xHost);
294+ aclrtFreeHost(yHost);
295+ aclrtFreeHost(zHost);
296+ aclrtDestroyStream(stream);
297+ aclrtResetDevice(deviceId);
298+ aclFinalize();
299+ return output;
300+ }
301+ 
302+ // 8. 结果 Device -> Host
303+ ret = aclrtMemcpy(zHost, byteSize, zDevice, byteSize, ACL_MEMCPY_DEVICE_TO_HOST);
304+ if (ret != ACL_SUCCESS) {
305+ std::cerr << "aclrtMemcpy to zHost failed. ERROR: " << ret << std::endl;
306+ aclrtFree(xDevice);
307+ aclrtFree(yDevice);
308+ aclrtFree(zDevice);
309+ aclrtFreeHost(xHost);
310+ aclrtFreeHost(yHost);
311+ aclrtFreeHost(zHost);
312+ aclrtDestroyStream(stream);
313+ aclrtResetDevice(deviceId);
314+ aclFinalize();
315+ return output;
316+ }
317+ 
318+ // 9. 构造返回值
319+ std::copy(zHost, zHost + totalLength, output.begin());
320+ 
321+ // 10. 释放资源，去初始化
322+ aclrtFree(xDevice);
323+ aclrtFree(yDevice);
324+ aclrtFree(zDevice);
325+ 
326+ aclrtFreeHost(xHost);
327+ aclrtFreeHost(yHost);
328+ aclrtFreeHost(zHost);
329+ 
330+ aclrtDestroyStream(stream);
331+ aclrtResetDevice(deviceId);
332+ aclFinalize();
333+ 
334+ return output;
335+}
336+ 
337+// ----- 验证与主函数 -----
338+ 
339+uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
340+{
341+ auto printTensor = [](std::vector<float> &tensor, const char *name) {
342+ constexpr size_t maxPrintSize = 20;
343+ std::cout << name << ": ";
344+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
345+ std::ostream_iterator<float>(std::cout, " "));
346+ if (tensor.size() > maxPrintSize) {
347+ std::cout << "...";
348+ }
349+ std::cout << std::endl;
350+ };
351+ printTensor(output, "Output");
352+ printTensor(golden, "Golden");
353+ if (std::equal(golden.begin(), golden.end(), output.begin())) {
354+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
355+ return 0;
356+ } else {
357+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
358+ return 1;
359+ }
360+}
361+ 
362+int32_t main(int32_t argc, char *argv[])
363+{
364+ constexpr uint32_t totalLength = 8 * 2048;
365+ constexpr float valueX = 1.2f;
366+ constexpr float valueY = 2.3f;
367+ std::vector<float> x(totalLength, valueX);
368+ std::vector<float> y(totalLength, valueY);
369+ 
370+ // 调用 kernel_mul 执行逐元素乘法
371+ std::vector<float> output = kernel_mul(x, y);
372+ 
373+ std::vector<float> golden(totalLength, valueX * valueY);
374+ return VerifyResult(output, golden);
375+}
A2026/CANN-Code-Detective/Challenge05-MulCustom/gcw_93rlw6ed/MulCustom/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+./mul_test