【代码侦探Challenge05】完成weixin_48384182的MulCustom逐元素乘法算子-cann-outreach-AtomGit

A2026/CANN-Code-Detective/Challenge02-DivBf16Custom/weixin_48384182/DivBf16Custom/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_bf16_test
	6	+ div_bf16_custom.asc
	7	+)
	8	+target_compile_options(div_bf16_test PRIVATE
	9	+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
	10	+)

A2026/CANN-Code-Detective/Challenge02-DivBf16Custom/weixin_48384182/DivBf16Custom/div_bf16_custom.asc+238-0

@@ -0,0 +1,238 @@
	1	+// div_bf16_custom.asc
	2	+
	3	+#include <cstdint>
	4	+#include <iostream>
	5	+#include <vector>
	6	+#include <algorithm>
	7	+#include <iterator>
	8	+#include <cstring>
	9	+#include "acl/acl.h"
	10	+#include "kernel_operator.h"
	11	+
	12	+class KernelDivBf16 {
	13	+public:
	14	+ __aicore__ inline KernelDivBf16() {}
	15	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength)
	16	+ {
	17	+ this->totalLength = totalLength;
	18	+ xGm.SetGlobalBuffer((__gm__ bfloat16_t *)x, totalLength);
	19	+ yGm.SetGlobalBuffer((__gm__ bfloat16_t *)y, totalLength);
	20	+ zGm.SetGlobalBuffer((__gm__ bfloat16_t *)z, totalLength);
	21	+
	22	+ // 初始化输入输出队列
	23	+ // 本题无需 Tiling，直接使用整个 totalLength 一次性计算
	24	+ pipe.InitBuffer(inQueueX, 1, totalLength * sizeof(bfloat16_t));
	25	+ pipe.InitBuffer(inQueueY, 1, totalLength * sizeof(bfloat16_t));
	26	+ pipe.InitBuffer(outQueueZ, 1, totalLength * sizeof(bfloat16_t));
	27	+
	28	+ // 完成 TBuf 初始化
	29	+ // 使用 InitBuffer 初始化两个 TBuf（tmpBuf0、tmpBuf1），
	30	+ // 大小为 totalLength * sizeof(float)，位置为 VECCALC
	31	+ pipe.InitBuffer(tmpBuf0, totalLength * sizeof(float));
	32	+ pipe.InitBuffer(tmpBuf1, totalLength * sizeof(float));
	33	+ }
	34	+
	35	+ __aicore__ inline void Process()
	36	+ {
	37	+ CopyIn();
	38	+ Compute();
	39	+ CopyOut();
	40	+ }
	41	+
	42	+private:
	43	+ __aicore__ inline void CopyIn()
	44	+ {
	45	+ // 完成 CopyIn 实现
	46	+ // 提示：本题无需 Tiling，直接使用整个 totalLength 进行 DataCopy
	47	+ // 1. AllocTensor 获取 xLocal、yLocal
	48	+ auto xLocal = inQueueX.AllocTensor<bfloat16_t>();
	49	+ auto yLocal = inQueueY.AllocTensor<bfloat16_t>();
	50	+ // 2. DataCopy 从 GM 搬运到 UB
	51	+ AscendC::DataCopy(xLocal, xGm[0], totalLength);
	52	+ AscendC::DataCopy(yLocal, yGm[0], totalLength);
	53	+ // 3. EnQue 入队
	54	+ inQueueX.EnQue(xLocal);
	55	+ inQueueY.EnQue(yLocal);
	56	+ }
	57	+
	58	+ __aicore__ inline void Compute()
	59	+ {
	60	+ // 完成 Compute 实现
	61	+ // 1. 从 inQueueX、inQueueY 中 DeQue 获取 xLocal、yLocal（bfloat16_t）
	62	+ auto xLocal = inQueueX.DeQue<bfloat16_t>();
	63	+ auto yLocal = inQueueY.DeQue<bfloat16_t>();
	64	+ // 2. 从 outQueueZ 中 AllocTensor 获取 zLocal（bfloat16_t）
	65	+ auto zLocal = outQueueZ.AllocTensor<bfloat16_t>();
	66	+ // 3. 从 tmpBuf0、tmpBuf1 中 Get<float> 获取临时 tensor
	67	+ auto tmpTensor0 = tmpBuf0.Get<float>();
	68	+ auto tmpTensor1 = tmpBuf1.Get<float>();
	69	+ // 4. 使用 Cast 将 xLocal、yLocal 从 bfloat16_t 转为 float，存入 tmpTensor
	70	+ // 提示：RoundMode 使用 CAST_NONE
	71	+ AscendC::Cast<float, bfloat16_t>(tmpTensor0, xLocal, AscendC::RoundMode::CAST_NONE, totalLength);
	72	+ AscendC::Cast<float, bfloat16_t>(tmpTensor1, yLocal, AscendC::RoundMode::CAST_NONE, totalLength);
	73	+ // 5. 使用 Div 执行 float 除法：tmpTensor0 = tmpTensor0 / tmpTensor1
	74	+ AscendC::Div<float>(tmpTensor0, tmpTensor0, tmpTensor1, totalLength);
	75	+ // 6. 使用 Cast 将结果从 float 转回 bfloat16_t，存入 zLocal
	76	+ // 提示：RoundMode 使用 CAST_RINT
	77	+ AscendC::Cast<bfloat16_t, float>(zLocal, tmpTensor0, AscendC::RoundMode::CAST_RINT, totalLength);
	78	+ // 7. EnQue(zLocal)，FreeTensor(xLocal)、FreeTensor(yLocal)
	79	+ outQueueZ.EnQue(zLocal);
	80	+ inQueueX.FreeTensor(xLocal);
	81	+ inQueueY.FreeTensor(yLocal);
	82	+ }
	83	+
	84	+ __aicore__ inline void CopyOut()
	85	+ {
	86	+ // 完成 CopyOut 实现
	87	+ // 提示：本题无需 Tiling，直接使用整个 totalLength 进行 DataCopy
	88	+ // 1. DeQue 获取 zLocal
	89	+ auto zLocal = outQueueZ.DeQue<bfloat16_t>();
	90	+ // 2. DataCopy 从 UB 搬运到 GM
	91	+ AscendC::DataCopy(zGm[0], zLocal, totalLength);
	92	+ // 3. FreeTensor
	93	+ outQueueZ.FreeTensor(zLocal);
	94	+ }
	95	+
	96	+private:
	97	+ uint32_t totalLength;
	98	+ AscendC::GlobalTensor<bfloat16_t> xGm;
	99	+ AscendC::GlobalTensor<bfloat16_t> yGm;
	100	+ AscendC::GlobalTensor<bfloat16_t> zGm;
	101	+ AscendC::TPipe pipe;
	102	+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueX;
	103	+ AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueY;
	104	+ AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueZ;
	105	+ // 完成 TBuf 声明
	106	+ // 声明两个 TBuf<AscendC::TPosition::VECCALC>：tmpBuf0、tmpBuf1
	107	+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuf0;
	108	+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuf1;
	109	+};
	110	+
	111	+__global__ __aicore__ void div_bf16_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength)
	112	+{
	113	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	114	+ KernelDivBf16 op;
	115	+ op.Init(x, y, z, totalLength);
	116	+ op.Process();
	117	+}
	118	+
	119	+// bfloat16 与 float 之间的转换辅助函数，实现为简化计算，不与kernel侧实现完全一致
	120	+inline uint16_t FloatToBfloat16Bits(float f)
	121	+{
	122	+ uint32_t bits;
	123	+ std::memcpy(&bits, &f, sizeof(float));
	124	+ return static_cast<uint16_t>(bits >> 16);
	125	+}
	126	+
	127	+inline float Bfloat16BitsToFloat(uint16_t bf16)
	128	+{
	129	+ uint32_t bits = static_cast<uint32_t>(bf16) << 16;
	130	+ float f;
	131	+ std::memcpy(&f, &bits, sizeof(float));
	132	+ return f;
	133	+}
	134	+
	135	+std::vector<float> kernel_div_bf16(std::vector<float> &x, std::vector<float> &y)
	136	+{
	137	+ uint32_t totalLength = static_cast<uint32_t>(x.size());
	138	+
	139	+ // 将 float 输入转为 bfloat16 (uint16_t) 比特表示
	140	+ std::vector<uint16_t> xBf16(totalLength);
	141	+ std::vector<uint16_t> yBf16(totalLength);
	142	+ for (uint32_t i = 0; i < totalLength; i++) {
	143	+ xBf16[i] = FloatToBfloat16Bits(x[i]);
	144	+ yBf16[i] = FloatToBfloat16Bits(y[i]);
	145	+ }
	146	+
	147	+ // 初始化 ACL
	148	+ aclrtStream stream = nullptr;
	149	+ uint8_t *xDevice = nullptr;
	150	+ uint8_t *yDevice = nullptr;
	151	+ uint8_t *zDevice = nullptr;
	152	+ std::vector<uint16_t> zHost(totalLength, 0);
	153	+
	154	+ aclInit(nullptr);
	155	+ aclrtSetDevice(0);
	156	+ aclrtCreateStream(&stream);
	157	+
	158	+ // 分配 device 内存并搬运数据
	159	+ uint32_t bufferSize = totalLength * sizeof(uint16_t);
	160	+ aclrtMalloc((void **)&xDevice, bufferSize, ACL_MEM_MALLOC_HUGE_FIRST);
	161	+ aclrtMalloc((void **)&yDevice, bufferSize, ACL_MEM_MALLOC_HUGE_FIRST);
	162	+ aclrtMalloc((void **)&zDevice, bufferSize, ACL_MEM_MALLOC_HUGE_FIRST);
	163	+ aclrtMemcpy(xDevice, bufferSize, xBf16.data(), bufferSize, ACL_MEMCPY_HOST_TO_DEVICE);
	164	+ aclrtMemcpy(yDevice, bufferSize, yBf16.data(), bufferSize, ACL_MEMCPY_HOST_TO_DEVICE);
	165	+
	166	+ // 启动 kernel：单核执行
	167	+ constexpr uint32_t blockDim = 1;
	168	+ div_bf16_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, totalLength);
	169	+ aclrtSynchronizeStream(stream);
	170	+
	171	+ // 回拷结果
	172	+ aclrtMemcpy(zHost.data(), bufferSize, zDevice, bufferSize, ACL_MEMCPY_DEVICE_TO_HOST);
	173	+
	174	+ // 将 bfloat16 比特转回 float
	175	+ std::vector<float> output(totalLength);
	176	+ for (uint32_t i = 0; i < totalLength; i++) {
	177	+ output[i] = Bfloat16BitsToFloat(zHost[i]);
	178	+ }
	179	+
	180	+ // 释放资源
	181	+ aclrtFree(xDevice);
	182	+ aclrtFree(yDevice);
	183	+ aclrtFree(zDevice);
	184	+ aclrtDestroyStream(stream);
	185	+ aclrtResetDevice(0);
	186	+ aclFinalize();
	187	+
	188	+ return output;
	189	+}
	190	+
	191	+uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
	192	+{
	193	+ auto printTensor = [](std::vector<float> &tensor, const char *name) {
	194	+ constexpr size_t maxPrintSize = 20;
	195	+ std::cout << name << ": ";
	196	+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
	197	+ std::ostream_iterator<float>(std::cout, " "));
	198	+ if (tensor.size() > maxPrintSize) {
	199	+ std::cout << "...";
	200	+ }
	201	+ std::cout << std::endl;
	202	+ };
	203	+ printTensor(output, "Output");
	204	+ printTensor(golden, "Golden");
	205	+ // bfloat16 精度较低，使用容差比较
	206	+ bool success = true;
	207	+ for (size_t i = 0; i < output.size(); i++) {
	208	+ if (std::abs(output[i] - golden[i]) > 1e-3) {
	209	+ success = false;
	210	+ break;
	211	+ }
	212	+ }
	213	+ if (success) {
	214	+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
	215	+ return 0;
	216	+ } else {
	217	+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
	218	+ return 1;
	219	+ }
	220	+}
	221	+
	222	+int32_t main(int32_t argc, char *argv[])
	223	+{
	224	+ constexpr uint32_t totalLength = 2048;
	225	+ constexpr float valueX = 1.2f;
	226	+ constexpr float valueY = 2.3f;
	227	+ std::vector<float> x(totalLength, valueX);
	228	+ std::vector<float> y(totalLength, valueY);
	229	+
	230	+ std::vector<float> output = kernel_div_bf16(x, y);
	231	+
	232	+ // golden 值需模拟 bfloat16 精度：float → bfloat16 → float → 除法 → bfloat16 → float
	233	+ float xBf16 = Bfloat16BitsToFloat(FloatToBfloat16Bits(valueX));
	234	+ float yBf16 = Bfloat16BitsToFloat(FloatToBfloat16Bits(valueY));
	235	+ float goldenValue = Bfloat16BitsToFloat(FloatToBfloat16Bits(xBf16 / yBf16));
	236	+ std::vector<float> golden(totalLength, goldenValue);
	237	+ return VerifyResult(output, golden);
	238	+}

A2026/CANN-Code-Detective/Challenge02-DivBf16Custom/weixin_48384182/DivBf16Custom/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_bf16_test

A2026/CANN-Code-Detective/Challenge05-MulCustom/weixin_48384182/MulCustom/CMakeLists.txt+13-0

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

A2026/CANN-Code-Detective/Challenge05-MulCustom/weixin_48384182/MulCustom/mul_custom.asc+176-0

@@ -0,0 +1,176 @@
	1	+#include <cstdint>
	2	+#include <iostream>
	3	+#include <vector>
	4	+#include <algorithm>
	5	+#include <iterator>
	6	+#include "acl/acl.h"
	7	+#include "kernel_operator.h"
	8	+
	9	+constexpr uint32_t BUFFER_NUM = 2; // tensor num for each queue
	10	+
	11	+struct MulCustomTilingData
	12	+{
	13	+ uint32_t totalLength;
	14	+ uint32_t tileNum;
	15	+};
	16	+
	17	+class KernelMul {
	18	+public:
	19	+ __aicore__ inline KernelMul() {}
	20	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
	21	+ {
	22	+ // 请补充：多核数据切分，计算每个核处理的长度 blockLength
	23	+ this->blockLength = totalLength / AscendC::GetBlockNum();
	24	+ this->tileNum = tileNum;
	25	+ // 请补充：Double Buffer 下每段搬运长度 tileLength
	26	+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
	27	+ // 请补充：按核号偏移，设置每核的 GlobalBuffer 起始地址
	28	+ xGm.SetGlobalBuffer((__gm__ float )x + this->blockLength AscendC::GetBlockIdx(), this->blockLength);
	29	+ yGm.SetGlobalBuffer((__gm__ float )y + this->blockLength AscendC::GetBlockIdx(), this->blockLength);
	30	+ zGm.SetGlobalBuffer((__gm__ float )z + this->blockLength AscendC::GetBlockIdx(), this->blockLength);
	31	+ // 请补充：为输入/输出队列分配 UB 内存（Double Buffer）
	32	+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
	33	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
	34	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
	35	+ }
	36	+ __aicore__ inline void Process()
	37	+ {
	38	+ int32_t loopCount = this->tileNum * BUFFER_NUM;
	39	+ for (int32_t i = 0; i < loopCount; i++) {
	40	+ CopyIn(i);
	41	+ Compute(i);
	42	+ CopyOut(i);
	43	+ }
	44	+ }
	45	+
	46	+private:
	47	+ __aicore__ inline void CopyIn(int32_t progress)
	48	+ {
	49	+ // 请补充：申请局部张量，DataCopy 从 GM 搬入 UB，再 EnQue 入队
	50	+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
	51	+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
	52	+ AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
	53	+ AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
	54	+ inQueueX.EnQue(xLocal);
	55	+ inQueueY.EnQue(yLocal);
	56	+ }
	57	+ __aicore__ inline void Compute(int32_t progress)
	58	+ {
	59	+ // 请补充：出队输入，调用 AscendC::Mul 完成 z = x * y，结果入队并释放输入
	60	+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
	61	+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
	62	+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
	63	+ AscendC::Mul(zLocal, xLocal, yLocal, this->tileLength);
	64	+ outQueueZ.EnQue<float>(zLocal);
	65	+ inQueueX.FreeTensor(xLocal);
	66	+ inQueueY.FreeTensor(yLocal);
	67	+ }
	68	+ __aicore__ inline void CopyOut(int32_t progress)
	69	+ {
	70	+ // 请补充：出队结果，DataCopy 从 UB 搬回 GM，释放输出张量
	71	+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
	72	+ AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
	73	+ outQueueZ.FreeTensor(zLocal);
	74	+ }
	75	+
	76	+private:
	77	+ AscendC::TPipe pipe;
	78	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
	79	+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	80	+ AscendC::GlobalTensor<float> xGm;
	81	+ AscendC::GlobalTensor<float> yGm;
	82	+ AscendC::GlobalTensor<float> zGm;
	83	+ uint32_t blockLength;
	84	+ uint32_t tileNum;
	85	+ uint32_t tileLength;
	86	+};
	87	+
	88	+__global__ __vector__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, MulCustomTilingData tiling)
	89	+{
	90	+ KernelMul op;
	91	+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
	92	+ op.Process();
	93	+}
	94	+
	95	+std::vector<float> kernel_mul(std::vector<float> &x, std::vector<float> &y)
	96	+{
	97	+ constexpr uint32_t numBlocks = 8;
	98	+ uint32_t totalLength = x.size();
	99	+ size_t totalByteSize = totalLength * sizeof(float);
	100	+ int32_t deviceId = 0;
	101	+ aclrtStream stream = nullptr;
	102	+ MulCustomTilingData tiling = {/totalLength:/totalLength, /tileNum:/8};
	103	+ uint8_t xHost = reinterpret_cast<uint8_t >(x.data());
	104	+ uint8_t yHost = reinterpret_cast<uint8_t >(y.data());
	105	+ uint8_t *zHost = nullptr;
	106	+ uint8_t *xDevice = nullptr;
	107	+ uint8_t *yDevice = nullptr;
	108	+ uint8_t *zDevice = nullptr;
	109	+
	110	+ aclInit(nullptr);
	111	+ aclrtSetDevice(deviceId);
	112	+ aclrtCreateStream(&stream);
	113	+
	114	+ aclrtMallocHost((void **)(&zHost), totalByteSize);
	115	+ aclrtMalloc((void **)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	116	+ aclrtMalloc((void **)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	117	+ aclrtMalloc((void **)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	118	+
	119	+ aclrtMemcpy(xDevice, totalByteSize, xHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
	120	+ aclrtMemcpy(yDevice, totalByteSize, yHost, totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
	121	+
	122	+ mul_custom<<<numBlocks, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
	123	+ aclrtSynchronizeStream(stream);
	124	+
	125	+ aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
	126	+ std::vector<float> z((float )zHost, (float )(zHost + totalByteSize));
	127	+
	128	+ aclrtFree(xDevice);
	129	+ aclrtFree(yDevice);
	130	+ aclrtFree(zDevice);
	131	+ aclrtFreeHost(zHost);
	132	+
	133	+ aclrtDestroyStream(stream);
	134	+ aclrtResetDevice(deviceId);
	135	+ aclFinalize();
	136	+
	137	+ return z;
	138	+}
	139	+
	140	+uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
	141	+{
	142	+ auto printTensor = [](std::vector<float> &tensor, const char *name) {
	143	+ constexpr size_t maxPrintSize = 20;
	144	+ std::cout << name << ": ";
	145	+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
	146	+ std::ostream_iterator<float>(std::cout, " "));
	147	+ if (tensor.size() > maxPrintSize) {
	148	+ std::cout << "...";
	149	+ }
	150	+ std::cout << std::endl;
	151	+ };
	152	+ printTensor(output, "Output");
	153	+ printTensor(golden, "Golden");
	154	+ if (std::equal(golden.begin(), golden.end(), output.begin())) {
	155	+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
	156	+ return 0;
	157	+ } else {
	158	+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
	159	+ return 1;
	160	+ }
	161	+ return 0;
	162	+}
	163	+
	164	+int32_t main(int32_t argc, char *argv[])
	165	+{
	166	+ constexpr uint32_t totalLength = 8 * 2048;
	167	+ constexpr float valueX = 1.2f;
	168	+ constexpr float valueY = 2.3f;
	169	+ std::vector<float> x(totalLength, valueX);
	170	+ std::vector<float> y(totalLength, valueY);
	171	+
	172	+ std::vector<float> output = kernel_mul(x, y);
	173	+
	174	+ std::vector<float> golden(totalLength, valueX * valueY);
	175	+ return VerifyResult(output, golden);
	176	+}

A2026/CANN-Code-Detective/Challenge05-MulCustom/weixin_48384182/MulCustom/run.sh+18-0

@@ -0,0 +1,18 @@
	1	+#!/bin/bash
	2	+set -e
	3	+
	4	+# 激活 CANN 环境（CANNLab 中一般为 /home/developer/Ascend/cann-9.0.0 或 $ASCEND_TOOLKIT_HOME）
	5	+if [ -z "$ASCEND_TOOLKIT_HOME" ]; then
	6	+ for d in /home/developer/Ascend/ascend-toolkit/latest /usr/local/Ascend/ascend-toolkit/latest; do
	7	+ if [ -d "$d" ]; then export ASCEND_TOOLKIT_HOME="$d"; break; fi
	8	+ done
	9	+fi
	10	+source "$ASCEND_TOOLKIT_HOME/set_env.sh"
	11	+
	12	+rm -rf build
	13	+mkdir -p build && cd build
	14	+cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 .. && make -j
	15	+cd ..
	16	+
	17	+# 运行验证：输出 Output 与 Golden 均为 2.76 时打印 [Success] Case accuracy is verification passed.
	18	+./build/mul_test