Star372
2457
代码
代码
Issues214
Pull Requests758
Actions
项目讨论
Wiki
分析
Star372
2457
已开启
【代码侦探Challenge05】实现 MulCustom 逐元素乘法算子 #3563
FineL1创建于 9月9日
【代码侦探Challenge05】实现 MulCustom 逐元素乘法算子 #3563
已开启
FineL1创建于 9月9日
从FineL1/cann-outreach_6062challenge05-finel1合入到cann/cann-outreachmaster
共 5 个文件变更+862-0
A2026/CANN-Code-Detective/Challenge05-MulCustom/FineL1/MulCustom/.ipynb_checkpoints/mul_custom-checkpoint.asc+417-0
@@ -0,0 +1,417 @@
1+// mul_custom.asc
2+// 逐元素乘法算子：z[i] = x[i] * y[i]
3+ 
4+#include <cstdint>
5+#include <iostream>
6+#include <vector>
7+#include <algorithm>
8+#include <iterator>
9+#include "acl/acl.h"
10+#include "kernel_operator.h"
11+ 
12+constexpr uint32_t BUFFER_NUM = 2; // Double Buffer 队列深度
13+ 
14+struct MulCustomTilingData
15+{
16+ uint32_t totalLength;
17+ uint32_t tileNum;
18+};
19+ 
20+class KernelMul {
21+public:
22+ __aicore__ inline KernelMul() {}
23+ 
24+ __aicore__ inline void Init(
25+ GM_ADDR x,
26+ GM_ADDR y,
27+ GM_ADDR z,
28+ uint32_t totalLength,
29+ uint32_t tileNum)
30+ {
31+ // 每个 Core 负责的数据量
32+ this->blockLength =
33+ totalLength / AscendC::GetBlockNum();
34+ 
35+ this->tileNum = tileNum;
36+ 
37+ // 双缓冲下，每个 Tile 实际处理的数据量
38+ this->tileLength =
39+ this->blockLength / tileNum / BUFFER_NUM;
40+ 
41+ // 当前 Core 在 GM 中的数据偏移
42+ uint32_t blockOffset =
43+ this->blockLength * AscendC::GetBlockIdx();
44+ 
45+ // 设置当前 Core 对应的 GM 地址
46+ xGm.SetGlobalBuffer(
47+ (__gm__ float *)x + blockOffset,
48+ this->blockLength);
49+ 
50+ yGm.SetGlobalBuffer(
51+ (__gm__ float *)y + blockOffset,
52+ this->blockLength);
53+ 
54+ zGm.SetGlobalBuffer(
55+ (__gm__ float *)z + blockOffset,
56+ this->blockLength);
57+ 
58+ // 为输入输出队列分配 UB
59+ pipe.InitBuffer(
60+ inQueueX,
61+ BUFFER_NUM,
62+ this->tileLength * sizeof(float));
63+ 
64+ pipe.InitBuffer(
65+ inQueueY,
66+ BUFFER_NUM,
67+ this->tileLength * sizeof(float));
68+ 
69+ pipe.InitBuffer(
70+ outQueueZ,
71+ BUFFER_NUM,
72+ this->tileLength * sizeof(float));
73+ }
74+ 
75+ __aicore__ inline void Process()
76+ {
77+ // tileNum=8，BUFFER_NUM=2
78+ // 所以每个 Core 一共循环 16 次
79+ int32_t loopCount =
80+ this->tileNum * BUFFER_NUM;
81+ 
82+ for (int32_t i = 0; i < loopCount; ++i) {
83+ CopyIn(i);
84+ Compute(i);
85+ CopyOut(i);
86+ }
87+ }
88+ 
89+private:
90+ __aicore__ inline void CopyIn(int32_t progress)
91+ {
92+ AscendC::LocalTensor<float> xLocal =
93+ inQueueX.AllocTensor<float>();
94+ 
95+ AscendC::LocalTensor<float> yLocal =
96+ inQueueY.AllocTensor<float>();
97+ 
98+ // GM -> UB
99+ AscendC::DataCopy(
100+ xLocal,
101+ xGm[progress * this->tileLength],
102+ this->tileLength);
103+ 
104+ AscendC::DataCopy(
105+ yLocal,
106+ yGm[progress * this->tileLength],
107+ this->tileLength);
108+ 
109+ inQueueX.EnQue(xLocal);
110+ inQueueY.EnQue(yLocal);
111+ }
112+ 
113+ __aicore__ inline void Compute(int32_t progress)
114+ {
115+ AscendC::LocalTensor<float> xLocal =
116+ inQueueX.DeQue<float>();
117+ 
118+ AscendC::LocalTensor<float> yLocal =
119+ inQueueY.DeQue<float>();
120+ 
121+ AscendC::LocalTensor<float> zLocal =
122+ outQueueZ.AllocTensor<float>();
123+ 
124+ // 核心计算：z[i] = x[i] * y[i]
125+ AscendC::Mul(
126+ zLocal,
127+ xLocal,
128+ yLocal,
129+ this->tileLength);
130+ 
131+ outQueueZ.EnQue(zLocal);
132+ 
133+ inQueueX.FreeTensor(xLocal);
134+ inQueueY.FreeTensor(yLocal);
135+ }
136+ 
137+ __aicore__ inline void CopyOut(int32_t progress)
138+ {
139+ AscendC::LocalTensor<float> zLocal =
140+ outQueueZ.DeQue<float>();
141+ 
142+ // UB -> GM
143+ AscendC::DataCopy(
144+ zGm[progress * this->tileLength],
145+ zLocal,
146+ this->tileLength);
147+ 
148+ outQueueZ.FreeTensor(zLocal);
149+ }
150+ 
151+private:
152+ AscendC::TPipe pipe;
153+ 
154+ AscendC::TQue<
155+ AscendC::TPosition::VECIN,
156+ BUFFER_NUM> inQueueX;
157+ 
158+ AscendC::TQue<
159+ AscendC::TPosition::VECIN,
160+ BUFFER_NUM> inQueueY;
161+ 
162+ AscendC::TQue<
163+ AscendC::TPosition::VECOUT,
164+ BUFFER_NUM> outQueueZ;
165+ 
166+ AscendC::GlobalTensor<float> xGm;
167+ AscendC::GlobalTensor<float> yGm;
168+ AscendC::GlobalTensor<float> zGm;
169+ 
170+ uint32_t blockLength;
171+ uint32_t tileNum;
172+ uint32_t tileLength;
173+};
174+ 
175+ 
176+// Kernel 入口函数
177+__global__ __aicore__ void mul_custom(
178+ GM_ADDR x,
179+ GM_ADDR y,
180+ GM_ADDR z,
181+ MulCustomTilingData tiling)
182+{
183+ // 本算子只使用 AI Vector Core
184+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
185+ 
186+ KernelMul op;
187+ 
188+ op.Init(
189+ x,
190+ y,
191+ z,
192+ tiling.totalLength,
193+ tiling.tileNum);
194+ 
195+ op.Process();
196+}
197+ 
198+ 
199+// ----- Host 侧：kernel 直调封装 -----
200+ 
201+std::vector<float> kernel_mul(
202+ std::vector<float> &x,
203+ std::vector<float> &y)
204+{
205+ constexpr uint32_t blockDim = 8;
206+ constexpr uint32_t tileNum = 8;
207+ constexpr int32_t deviceId = 0;
208+ 
209+ uint32_t totalLength =
210+ static_cast<uint32_t>(x.size());
211+ 
212+ size_t totalByteSize =
213+ static_cast<size_t>(totalLength) * sizeof(float);
214+ 
215+ MulCustomTilingData tiling {
216+ totalLength,
217+ tileNum
218+ };
219+ 
220+ aclrtStream stream = nullptr;
221+ 
222+ float *xHost = nullptr;
223+ float *yHost = nullptr;
224+ float *zHost = nullptr;
225+ 
226+ uint8_t *xDevice = nullptr;
227+ uint8_t *yDevice = nullptr;
228+ uint8_t *zDevice = nullptr;
229+ 
230+ // 1. ACL 初始化
231+ aclInit(nullptr);
232+ 
233+ // 2. 设置 Device
234+ aclrtSetDevice(deviceId);
235+ 
236+ // 3. 创建 Stream
237+ aclrtCreateStream(&stream);
238+ 
239+ // 4. Host 内存申请
240+ aclrtMallocHost(
241+ reinterpret_cast<void **>(&xHost),
242+ totalByteSize);
243+ 
244+ aclrtMallocHost(
245+ reinterpret_cast<void **>(&yHost),
246+ totalByteSize);
247+ 
248+ aclrtMallocHost(
249+ reinterpret_cast<void **>(&zHost),
250+ totalByteSize);
251+ 
252+ // 把 vector 数据写入 Host 内存
253+ std::copy(
254+ x.begin(),
255+ x.end(),
256+ xHost);
257+ 
258+ std::copy(
259+ y.begin(),
260+ y.end(),
261+ yHost);
262+ 
263+ // 5. Device 内存申请
264+ aclrtMalloc(
265+ reinterpret_cast<void **>(&xDevice),
266+ totalByteSize,
267+ ACL_MEM_MALLOC_HUGE_FIRST);
268+ 
269+ aclrtMalloc(
270+ reinterpret_cast<void **>(&yDevice),
271+ totalByteSize,
272+ ACL_MEM_MALLOC_HUGE_FIRST);
273+ 
274+ aclrtMalloc(
275+ reinterpret_cast<void **>(&zDevice),
276+ totalByteSize,
277+ ACL_MEM_MALLOC_HUGE_FIRST);
278+ 
279+ // 6. Host -> Device
280+ aclrtMemcpy(
281+ xDevice,
282+ totalByteSize,
283+ xHost,
284+ totalByteSize,
285+ ACL_MEMCPY_HOST_TO_DEVICE);
286+ 
287+ aclrtMemcpy(
288+ yDevice,
289+ totalByteSize,
290+ yHost,
291+ totalByteSize,
292+ ACL_MEMCPY_HOST_TO_DEVICE);
293+ 
294+ // 7. Kernel 直调
295+ mul_custom<<<blockDim, nullptr, stream>>>(
296+ xDevice,
297+ yDevice,
298+ zDevice,
299+ tiling);
300+ 
301+ // 8. 等待 Kernel 执行完成
302+ aclrtSynchronizeStream(stream);
303+ 
304+ // 9. Device -> Host
305+ aclrtMemcpy(
306+ zHost,
307+ totalByteSize,
308+ zDevice,
309+ totalByteSize,
310+ ACL_MEMCPY_DEVICE_TO_HOST);
311+ 
312+ // 10. 构造返回结果
313+ std::vector<float> output(
314+ zHost,
315+ zHost + totalLength);
316+ 
317+ // 11. 释放 Device 内存
318+ aclrtFree(xDevice);
319+ aclrtFree(yDevice);
320+ aclrtFree(zDevice);
321+ 
322+ // 12. 释放 Host 内存
323+ aclrtFreeHost(xHost);
324+ aclrtFreeHost(yHost);
325+ aclrtFreeHost(zHost);
326+ 
327+ // 13. 清理运行时资源
328+ aclrtDestroyStream(stream);
329+ aclrtResetDevice(deviceId);
330+ aclFinalize();
331+ 
332+ return output;
333+}
334+ 
335+ 
336+// ----- 验证与主函数 -----
337+ 
338+uint32_t VerifyResult(
339+ std::vector<float> &output,
340+ std::vector<float> &golden)
341+{
342+ auto printTensor = [](
343+ std::vector<float> &tensor,
344+ const char *name)
345+ {
346+ constexpr size_t maxPrintSize = 20;
347+ 
348+ std::cout << name << ": ";
349+ 
350+ std::copy(
351+ tensor.begin(),
352+ tensor.begin() +
353+ std::min(tensor.size(), maxPrintSize),
354+ std::ostream_iterator<float>(
355+ std::cout,
356+ " "));
357+ 
358+ if (tensor.size() > maxPrintSize) {
359+ std::cout << "...";
360+ }
361+ 
362+ std::cout << std::endl;
363+ };
364+ 
365+ printTensor(output, "Output");
366+ printTensor(golden, "Golden");
367+ 
368+ if (std::equal(
369+ golden.begin(),
370+ golden.end(),
371+ output.begin()))
372+ {
373+ std::cout
374+ << "[Success] Case accuracy is verification passed."
375+ << std::endl;
376+ 
377+ return 0;
378+ }
379+ else
380+ {
381+ std::cout
382+ << "[Failed] Case accuracy is verification failed!"
383+ << std::endl;
384+ 
385+ return 1;
386+ }
387+}
388+ 
389+ 
390+int32_t main(int32_t argc, char *argv[])
391+{
392+ constexpr uint32_t totalLength =
393+ 8 * 2048;
394+ 
395+ constexpr float valueX = 1.2f;
396+ constexpr float valueY = 2.3f;
397+ 
398+ std::vector<float> x(
399+ totalLength,
400+ valueX);
401+ 
402+ std::vector<float> y(
403+ totalLength,
404+ valueY);
405+ 
406+ // 调用 Ascend C Kernel
407+ std::vector<float> output =
408+ kernel_mul(x, y);
409+ 
410+ std::vector<float> golden(
411+ totalLength,
412+ valueX * valueY);
413+ 
414+ return VerifyResult(
415+ output,
416+ golden);
417+}
A2026/CANN-Code-Detective/Challenge05-MulCustom/FineL1/MulCustom/.ipynb_checkpoints/run-checkpoint.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
A2026/CANN-Code-Detective/Challenge05-MulCustom/FineL1/MulCustom/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(mul_test
6+ mul_custom.asc
7+)
8+target_compile_options(mul_test PRIVATE
9+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
10+)
A2026/CANN-Code-Detective/Challenge05-MulCustom/FineL1/MulCustom/mul_custom.asc+417-0
@@ -0,0 +1,417 @@
1+// mul_custom.asc
2+// 逐元素乘法算子：z[i] = x[i] * y[i]
3+ 
4+#include <cstdint>
5+#include <iostream>
6+#include <vector>
7+#include <algorithm>
8+#include <iterator>
9+#include "acl/acl.h"
10+#include "kernel_operator.h"
11+ 
12+constexpr uint32_t BUFFER_NUM = 2; // Double Buffer 队列深度
13+ 
14+struct MulCustomTilingData
15+{
16+ uint32_t totalLength;
17+ uint32_t tileNum;
18+};
19+ 
20+class KernelMul {
21+public:
22+ __aicore__ inline KernelMul() {}
23+ 
24+ __aicore__ inline void Init(
25+ GM_ADDR x,
26+ GM_ADDR y,
27+ GM_ADDR z,
28+ uint32_t totalLength,
29+ uint32_t tileNum)
30+ {
31+ // 每个 Core 负责的数据量
32+ this->blockLength =
33+ totalLength / AscendC::GetBlockNum();
34+ 
35+ this->tileNum = tileNum;
36+ 
37+ // 双缓冲下，每个 Tile 实际处理的数据量
38+ this->tileLength =
39+ this->blockLength / tileNum / BUFFER_NUM;
40+ 
41+ // 当前 Core 在 GM 中的数据偏移
42+ uint32_t blockOffset =
43+ this->blockLength * AscendC::GetBlockIdx();
44+ 
45+ // 设置当前 Core 对应的 GM 地址
46+ xGm.SetGlobalBuffer(
47+ (__gm__ float *)x + blockOffset,
48+ this->blockLength);
49+ 
50+ yGm.SetGlobalBuffer(
51+ (__gm__ float *)y + blockOffset,
52+ this->blockLength);
53+ 
54+ zGm.SetGlobalBuffer(
55+ (__gm__ float *)z + blockOffset,
56+ this->blockLength);
57+ 
58+ // 为输入输出队列分配 UB
59+ pipe.InitBuffer(
60+ inQueueX,
61+ BUFFER_NUM,
62+ this->tileLength * sizeof(float));
63+ 
64+ pipe.InitBuffer(
65+ inQueueY,
66+ BUFFER_NUM,
67+ this->tileLength * sizeof(float));
68+ 
69+ pipe.InitBuffer(
70+ outQueueZ,
71+ BUFFER_NUM,
72+ this->tileLength * sizeof(float));
73+ }
74+ 
75+ __aicore__ inline void Process()
76+ {
77+ // tileNum=8，BUFFER_NUM=2
78+ // 所以每个 Core 一共循环 16 次
79+ int32_t loopCount =
80+ this->tileNum * BUFFER_NUM;
81+ 
82+ for (int32_t i = 0; i < loopCount; ++i) {
83+ CopyIn(i);
84+ Compute(i);
85+ CopyOut(i);
86+ }
87+ }
88+ 
89+private:
90+ __aicore__ inline void CopyIn(int32_t progress)
91+ {
92+ AscendC::LocalTensor<float> xLocal =
93+ inQueueX.AllocTensor<float>();
94+ 
95+ AscendC::LocalTensor<float> yLocal =
96+ inQueueY.AllocTensor<float>();
97+ 
98+ // GM -> UB
99+ AscendC::DataCopy(
100+ xLocal,
101+ xGm[progress * this->tileLength],
102+ this->tileLength);
103+ 
104+ AscendC::DataCopy(
105+ yLocal,
106+ yGm[progress * this->tileLength],
107+ this->tileLength);
108+ 
109+ inQueueX.EnQue(xLocal);
110+ inQueueY.EnQue(yLocal);
111+ }
112+ 
113+ __aicore__ inline void Compute(int32_t progress)
114+ {
115+ AscendC::LocalTensor<float> xLocal =
116+ inQueueX.DeQue<float>();
117+ 
118+ AscendC::LocalTensor<float> yLocal =
119+ inQueueY.DeQue<float>();
120+ 
121+ AscendC::LocalTensor<float> zLocal =
122+ outQueueZ.AllocTensor<float>();
123+ 
124+ // 核心计算：z[i] = x[i] * y[i]
125+ AscendC::Mul(
126+ zLocal,
127+ xLocal,
128+ yLocal,
129+ this->tileLength);
130+ 
131+ outQueueZ.EnQue(zLocal);
132+ 
133+ inQueueX.FreeTensor(xLocal);
134+ inQueueY.FreeTensor(yLocal);
135+ }
136+ 
137+ __aicore__ inline void CopyOut(int32_t progress)
138+ {
139+ AscendC::LocalTensor<float> zLocal =
140+ outQueueZ.DeQue<float>();
141+ 
142+ // UB -> GM
143+ AscendC::DataCopy(
144+ zGm[progress * this->tileLength],
145+ zLocal,
146+ this->tileLength);
147+ 
148+ outQueueZ.FreeTensor(zLocal);
149+ }
150+ 
151+private:
152+ AscendC::TPipe pipe;
153+ 
154+ AscendC::TQue<
155+ AscendC::TPosition::VECIN,
156+ BUFFER_NUM> inQueueX;
157+ 
158+ AscendC::TQue<
159+ AscendC::TPosition::VECIN,
160+ BUFFER_NUM> inQueueY;
161+ 
162+ AscendC::TQue<
163+ AscendC::TPosition::VECOUT,
164+ BUFFER_NUM> outQueueZ;
165+ 
166+ AscendC::GlobalTensor<float> xGm;
167+ AscendC::GlobalTensor<float> yGm;
168+ AscendC::GlobalTensor<float> zGm;
169+ 
170+ uint32_t blockLength;
171+ uint32_t tileNum;
172+ uint32_t tileLength;
173+};
174+ 
175+ 
176+// Kernel 入口函数
177+__global__ __aicore__ void mul_custom(
178+ GM_ADDR x,
179+ GM_ADDR y,
180+ GM_ADDR z,
181+ MulCustomTilingData tiling)
182+{
183+ // 本算子只使用 AI Vector Core
184+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
185+ 
186+ KernelMul op;
187+ 
188+ op.Init(
189+ x,
190+ y,
191+ z,
192+ tiling.totalLength,
193+ tiling.tileNum);
194+ 
195+ op.Process();
196+}
197+ 
198+ 
199+// ----- Host 侧：kernel 直调封装 -----
200+ 
201+std::vector<float> kernel_mul(
202+ std::vector<float> &x,
203+ std::vector<float> &y)
204+{
205+ constexpr uint32_t blockDim = 8;
206+ constexpr uint32_t tileNum = 8;
207+ constexpr int32_t deviceId = 0;
208+ 
209+ uint32_t totalLength =
210+ static_cast<uint32_t>(x.size());
211+ 
212+ size_t totalByteSize =
213+ static_cast<size_t>(totalLength) * sizeof(float);
214+ 
215+ MulCustomTilingData tiling {
216+ totalLength,
217+ tileNum
218+ };
219+ 
220+ aclrtStream stream = nullptr;
221+ 
222+ float *xHost = nullptr;
223+ float *yHost = nullptr;
224+ float *zHost = nullptr;
225+ 
226+ uint8_t *xDevice = nullptr;
227+ uint8_t *yDevice = nullptr;
228+ uint8_t *zDevice = nullptr;
229+ 
230+ // 1. ACL 初始化
231+ aclInit(nullptr);
232+ 
233+ // 2. 设置 Device
234+ aclrtSetDevice(deviceId);
235+ 
236+ // 3. 创建 Stream
237+ aclrtCreateStream(&stream);
238+ 
239+ // 4. Host 内存申请
240+ aclrtMallocHost(
241+ reinterpret_cast<void **>(&xHost),
242+ totalByteSize);
243+ 
244+ aclrtMallocHost(
245+ reinterpret_cast<void **>(&yHost),
246+ totalByteSize);
247+ 
248+ aclrtMallocHost(
249+ reinterpret_cast<void **>(&zHost),
250+ totalByteSize);
251+ 
252+ // 把 vector 数据写入 Host 内存
253+ std::copy(
254+ x.begin(),
255+ x.end(),
256+ xHost);
257+ 
258+ std::copy(
259+ y.begin(),
260+ y.end(),
261+ yHost);
262+ 
263+ // 5. Device 内存申请
264+ aclrtMalloc(
265+ reinterpret_cast<void **>(&xDevice),
266+ totalByteSize,
267+ ACL_MEM_MALLOC_HUGE_FIRST);
268+ 
269+ aclrtMalloc(
270+ reinterpret_cast<void **>(&yDevice),
271+ totalByteSize,
272+ ACL_MEM_MALLOC_HUGE_FIRST);
273+ 
274+ aclrtMalloc(
275+ reinterpret_cast<void **>(&zDevice),
276+ totalByteSize,
277+ ACL_MEM_MALLOC_HUGE_FIRST);
278+ 
279+ // 6. Host -> Device
280+ aclrtMemcpy(
281+ xDevice,
282+ totalByteSize,
283+ xHost,
284+ totalByteSize,
285+ ACL_MEMCPY_HOST_TO_DEVICE);
286+ 
287+ aclrtMemcpy(
288+ yDevice,
289+ totalByteSize,
290+ yHost,
291+ totalByteSize,
292+ ACL_MEMCPY_HOST_TO_DEVICE);
293+ 
294+ // 7. Kernel 直调
295+ mul_custom<<<blockDim, nullptr, stream>>>(
296+ xDevice,
297+ yDevice,
298+ zDevice,
299+ tiling);
300+ 
301+ // 8. 等待 Kernel 执行完成
302+ aclrtSynchronizeStream(stream);
303+ 
304+ // 9. Device -> Host
305+ aclrtMemcpy(
306+ zHost,
307+ totalByteSize,
308+ zDevice,
309+ totalByteSize,
310+ ACL_MEMCPY_DEVICE_TO_HOST);
311+ 
312+ // 10. 构造返回结果
313+ std::vector<float> output(
314+ zHost,
315+ zHost + totalLength);
316+ 
317+ // 11. 释放 Device 内存
318+ aclrtFree(xDevice);
319+ aclrtFree(yDevice);
320+ aclrtFree(zDevice);
321+ 
322+ // 12. 释放 Host 内存
323+ aclrtFreeHost(xHost);
324+ aclrtFreeHost(yHost);
325+ aclrtFreeHost(zHost);
326+ 
327+ // 13. 清理运行时资源
328+ aclrtDestroyStream(stream);
329+ aclrtResetDevice(deviceId);
330+ aclFinalize();
331+ 
332+ return output;
333+}
334+ 
335+ 
336+// ----- 验证与主函数 -----
337+ 
338+uint32_t VerifyResult(
339+ std::vector<float> &output,
340+ std::vector<float> &golden)
341+{
342+ auto printTensor = [](
343+ std::vector<float> &tensor,
344+ const char *name)
345+ {
346+ constexpr size_t maxPrintSize = 20;
347+ 
348+ std::cout << name << ": ";
349+ 
350+ std::copy(
351+ tensor.begin(),
352+ tensor.begin() +
353+ std::min(tensor.size(), maxPrintSize),
354+ std::ostream_iterator<float>(
355+ std::cout,
356+ " "));
357+ 
358+ if (tensor.size() > maxPrintSize) {
359+ std::cout << "...";
360+ }
361+ 
362+ std::cout << std::endl;
363+ };
364+ 
365+ printTensor(output, "Output");
366+ printTensor(golden, "Golden");
367+ 
368+ if (std::equal(
369+ golden.begin(),
370+ golden.end(),
371+ output.begin()))
372+ {
373+ std::cout
374+ << "[Success] Case accuracy is verification passed."
375+ << std::endl;
376+ 
377+ return 0;
378+ }
379+ else
380+ {
381+ std::cout
382+ << "[Failed] Case accuracy is verification failed!"
383+ << std::endl;
384+ 
385+ return 1;
386+ }
387+}
388+ 
389+ 
390+int32_t main(int32_t argc, char *argv[])
391+{
392+ constexpr uint32_t totalLength =
393+ 8 * 2048;
394+ 
395+ constexpr float valueX = 1.2f;
396+ constexpr float valueY = 2.3f;
397+ 
398+ std::vector<float> x(
399+ totalLength,
400+ valueX);
401+ 
402+ std::vector<float> y(
403+ totalLength,
404+ valueY);
405+ 
406+ // 调用 Ascend C Kernel
407+ std::vector<float> output =
408+ kernel_mul(x, y);
409+ 
410+ std::vector<float> golden(
411+ totalLength,
412+ valueX * valueY);
413+ 
414+ return VerifyResult(
415+ output,
416+ golden);
417+}
A2026/CANN-Code-Detective/Challenge05-MulCustom/FineL1/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