Star268
2366
代码
代码
Issues203
Pull Requests768
Actions
项目讨论
Wiki
分析
Star268
2366
已开启
【代码侦探Challenge05】完成MulCustom逐元素乘法算子 #3609
JeffDing创建于 20 天前
【代码侦探Challenge05】完成MulCustom逐元素乘法算子 #3609
已开启
JeffDing创建于 20 天前
从JeffDing/cann-outreachChallenge_05合入到cann/cann-outreachmaster
共 3 个文件变更+198-0
A2026/CANN-Code-Detective/Challenge05-MulCustom/JeffDin/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/JeffDin/MulCustom/mul_custom.asc+179-0
@@ -0,0 +1,179 @@
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+ this->blockLength = totalLength / AscendC::GetBlockNum();
37+ this->tileNum = tileNum;
38+ this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
39+ xGm.SetGlobalBuffer((__gm__ float *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
40+ yGm.SetGlobalBuffer((__gm__ float *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
41+ zGm.SetGlobalBuffer((__gm__ float *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
42+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
43+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
44+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
45+ }
46+ __aicore__ inline void Process()
47+ {
48+ int32_t loopCount = this->tileNum * BUFFER_NUM;
49+ for (int32_t i = 0; i < loopCount; i++) {
50+ CopyIn(i);
51+ Compute(i);
52+ CopyOut(i);
53+ }
54+ }
55+ 
56+private:
57+ __aicore__ inline void CopyIn(int32_t progress)
58+ {
59+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
60+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
61+ AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
62+ AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], this->tileLength);
63+ inQueueX.EnQue(xLocal);
64+ inQueueY.EnQue(yLocal);
65+ }
66+ __aicore__ inline void Compute(int32_t progress)
67+ {
68+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
69+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
70+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
71+ AscendC::Mul(zLocal, xLocal, yLocal, this->tileLength);
72+ outQueueZ.EnQue(zLocal);
73+ inQueueX.FreeTensor(xLocal);
74+ inQueueY.FreeTensor(yLocal);
75+ }
76+ __aicore__ inline void CopyOut(int32_t progress)
77+ {
78+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
79+ AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, this->tileLength);
80+ outQueueZ.FreeTensor(zLocal);
81+ }
82+ 
83+private:
84+ AscendC::TPipe pipe;
85+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
86+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
87+ AscendC::GlobalTensor<float> xGm, yGm, zGm;
88+ uint32_t blockLength, tileNum, tileLength;
89+};
90+ 
91+// Kernel 入口函数
92+__global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, MulCustomTilingData tiling)
93+{
94+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
95+ KernelMul op;
96+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
97+ op.Process();
98+}
99+ 
100+// ----- Host 侧：kernel 直调封装 -----
101+ 
102+std::vector<float> kernel_mul(std::vector<float> &x, std::vector<float> &y)
103+{
104+ uint32_t blockDim = 8;
105+ uint32_t totalLength = x.size();
106+ MulCustomTilingData tiling = {totalLength, 8};
107+ aclInit(nullptr);
108+ int32_t deviceId = 0;
109+ aclrtSetDevice(deviceId);
110+ aclrtStream stream;
111+ aclrtCreateStream(&stream);
112+ size_t dataSize = totalLength * sizeof(float);
113+ void *xDevice, *yDevice, *zDevice;
114+ void *xHost, *yHost, *zHost;
115+ aclrtMallocHost(&xHost, dataSize);
116+ aclrtMallocHost(&yHost, dataSize);
117+ aclrtMallocHost(&zHost, dataSize);
118+ aclrtMalloc(&xDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
119+ aclrtMalloc(&yDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
120+ aclrtMalloc(&zDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
121+ aclrtMemcpy(xHost, dataSize, x.data(), dataSize, ACL_MEMCPY_HOST_TO_HOST);
122+ aclrtMemcpy(yHost, dataSize, y.data(), dataSize, ACL_MEMCPY_HOST_TO_HOST);
123+ aclrtMemcpy(xDevice, dataSize, xHost, dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
124+ aclrtMemcpy(yDevice, dataSize, yHost, dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
125+ mul_custom<<<blockDim, nullptr, stream>>>((__gm__ uint8_t*)xDevice, (__gm__ uint8_t*)yDevice, (__gm__ uint8_t*)zDevice, tiling);
126+ aclrtSynchronizeStream(stream);
127+ aclrtMemcpy(zHost, dataSize, zDevice, dataSize, ACL_MEMCPY_DEVICE_TO_HOST);
128+ std::vector<float> output(totalLength);
129+ std::copy((float *)zHost, (float *)zHost + totalLength, output.begin());
130+ aclrtFree(xDevice);
131+ aclrtFree(yDevice);
132+ aclrtFree(zDevice);
133+ aclrtFreeHost(xHost);
134+ aclrtFreeHost(yHost);
135+ aclrtFreeHost(zHost);
136+ aclrtDestroyStream(stream);
137+ aclrtResetDevice(deviceId);
138+ aclFinalize();
139+ return output;
140+}
141+ 
142+// ----- 验证与主函数 -----
143+ 
144+uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
145+{
146+ auto printTensor = [](std::vector<float> &tensor, const char *name) {
147+ constexpr size_t maxPrintSize = 20;
148+ std::cout << name << ": ";
149+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
150+ std::ostream_iterator<float>(std::cout, " "));
151+ if (tensor.size() > maxPrintSize) {
152+ std::cout << "...";
153+ }
154+ std::cout << std::endl;
155+ };
156+ printTensor(output, "Output");
157+ printTensor(golden, "Golden");
158+ if (std::equal(golden.begin(), golden.end(), output.begin())) {
159+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
160+ return 0;
161+ } else {
162+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
163+ return 1;
164+ }
165+}
166+ 
167+int32_t main(int32_t argc, char *argv[])
168+{
169+ constexpr uint32_t totalLength = 8 * 2048;
170+ constexpr float valueX = 1.2f;
171+ constexpr float valueY = 2.3f;
172+ std::vector<float> x(totalLength, valueX);
173+ std::vector<float> y(totalLength, valueY);
174+ 
175+ std::vector<float> output = kernel_mul(x, y);
176+ 
177+ std::vector<float> golden(totalLength, valueX * valueY);
178+ return VerifyResult(output, golden);
179+}
A2026/CANN-Code-Detective/Challenge05-MulCustom/JeffDin/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