【代码侦探Challenge05】新增 MulCustom 逐元素乘法算子代码提交-cann-outreach-AtomGit

已开启

【代码侦探Challenge05】新增 MulCustom 逐元素乘法算子代码提交 #2840

gcw_JQJBklQH创建于 23 天前从gcw_JQJBklQH/cann-outreach： five 合入到 cann/cann-outreach： master

【代码侦探Challenge05】新增 MulCustom 逐元素乘法算子代码提交 #2840

已开启

gcw_JQJBklQH创建于 23 天前

从gcw_JQJBklQH/cann-outreach five合入到cann/cann-outreach master

共 3 个文件变更+298-0

A2026/CANN-Code-Detective/Challenge05-MulCustom/gcw_JQJBklQH/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/gcw_JQJBklQH/mul_custom.asc+279-0

@@ -0,0 +1,279 @@
	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	+ // 多核数据切分：每个 AI Core 处理一段连续数据，核内再按双缓冲 tile 切分。
	37	+ this->blockLength = totalLength / AscendC::GetBlockNum();
	38	+ this->tileNum = tileNum;
	39	+ this->tileLength = this->blockLength / this->tileNum / BUFFER_NUM;
	40	+
	41	+ // 根据当前核号计算本核在 GM 中的起始偏移，并设置三个 GlobalTensor 的起始地址和长度。
	42	+ uint32_t blockOffset = this->blockLength * AscendC::GetBlockIdx();
	43	+ xGm.SetGlobalBuffer((__gm__ float *)x + blockOffset, this->blockLength);
	44	+ yGm.SetGlobalBuffer((__gm__ float *)y + blockOffset, this->blockLength);
	45	+ zGm.SetGlobalBuffer((__gm__ float *)z + blockOffset, this->blockLength);
	46	+
	47	+ // 为输入/输出队列分配内存：每个队列 BUFFER_NUM 个块，每块 tileLength * sizeof(float)。
	48	+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
	49	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
	50	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
	51	+ }
	52	+ __aicore__ inline void Process()
	53	+ {
	54	+ // 循环次数 = tileNum * BUFFER_NUM，每次迭代执行 CopyIn → Compute → CopyOut。
	55	+ int32_t loopCount = this->tileNum * BUFFER_NUM;
	56	+ for (int32_t i = 0; i < loopCount; i++) {
	57	+ CopyIn(i);
	58	+ Compute(i);
	59	+ CopyOut(i);
	60	+ }
	61	+ }
	62	+
	63	+private:
	64	+ __aicore__ inline void CopyIn(int32_t progress)
	65	+ {
	66	+ // 从输入队列申请 UB 缓冲区。
	67	+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
	68	+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
	69	+
	70	+ // 从 GM 搬运 tileLength 个元素到 UB。
	71	+ uint32_t offset = progress * this->tileLength;
	72	+ AscendC::DataCopy(xLocal, xGm[offset], this->tileLength);
	73	+ AscendC::DataCopy(yLocal, yGm[offset], this->tileLength);
	74	+
	75	+ // 入队，标识数据搬运完成，可供 Compute 阶段 DeQue 使用。
	76	+ inQueueX.EnQue(xLocal);
	77	+ inQueueY.EnQue(yLocal);
	78	+ }
	79	+ __aicore__ inline void Compute(int32_t progress)
	80	+ {
	81	+ // 从输入队列出队取回 xLocal、yLocal，并向输出队列申请 zLocal。
	82	+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
	83	+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
	84	+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
	85	+
	86	+ // 逐元素乘法：zLocal[i] = xLocal[i] * yLocal[i]。
	87	+ AscendC::Mul(zLocal, xLocal, yLocal, this->tileLength);
	88	+
	89	+ // 计算结果入队供 CopyOut 使用，并释放输入缓冲区。
	90	+ outQueueZ.EnQue<float>(zLocal);
	91	+ inQueueX.FreeTensor(xLocal);
	92	+ inQueueY.FreeTensor(yLocal);
	93	+ }
	94	+ __aicore__ inline void CopyOut(int32_t progress)
	95	+ {
	96	+ // 从输出队列出队取回计算结果 zLocal。
	97	+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
	98	+
	99	+ // 将结果从 UB 搬回 GM。
	100	+ uint32_t offset = progress * this->tileLength;
	101	+ AscendC::DataCopy(zGm[offset], zLocal, this->tileLength);
	102	+
	103	+ // 释放输出缓冲区。
	104	+ outQueueZ.FreeTensor(zLocal);
	105	+ }
	106	+
	107	+private:
	108	+ AscendC::TPipe pipe;
	109	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
	110	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;
	111	+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	112	+ AscendC::GlobalTensor<float> xGm;
	113	+ AscendC::GlobalTensor<float> yGm;
	114	+ AscendC::GlobalTensor<float> zGm;
	115	+ uint32_t blockLength;
	116	+ uint32_t tileNum;
	117	+ uint32_t tileLength;
	118	+};
	119	+
	120	+// Kernel 入口函数
	121	+__global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, MulCustomTilingData tiling)
	122	+{
	123	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	124	+ KernelMul op;
	125	+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
	126	+ op.Process();
	127	+}
	128	+
	129	+// ----- Host 侧：kernel 直调封装 -----
	130	+
	131	+// ACL 调用返回值检查，失败时打印错误信息并返回 false。
	132	+static bool CheckAclResult(aclError result, const char *operation)
	133	+{
	134	+ if (result == ACL_SUCCESS) {
	135	+ return true;
	136	+ }
	137	+ std::cerr << "[Error] " << operation << " failed, error code: " << result << std::endl;
	138	+ return false;
	139	+}
	140	+
	141	+std::vector<float> kernel_mul(std::vector<float> &x, std::vector<float> &y)
	142	+{
	143	+ constexpr uint32_t blockDim = 8;
	144	+ constexpr uint32_t tileNum = 8;
	145	+ const uint32_t totalLength = static_cast<uint32_t>(x.size());
	146	+
	147	+ MulCustomTilingData tiling{totalLength, tileNum};
	148	+
	149	+ int32_t deviceId = 0;
	150	+ aclrtStream stream = nullptr;
	151	+ uint8_t *xDevice = nullptr;
	152	+ uint8_t *yDevice = nullptr;
	153	+ uint8_t *zDevice = nullptr;
	154	+ bool deviceSet = false;
	155	+ bool aclInitialized = false;
	156	+
	157	+ // 资源释放 lambda：无论成功或失败，统一释放设备内存、销毁流、复位设备。
	158	+ auto cleanup = [&]() {
	159	+ if (xDevice != nullptr) {
	160	+ aclrtFree(xDevice);
	161	+ }
	162	+ if (yDevice != nullptr) {
	163	+ aclrtFree(yDevice);
	164	+ }
	165	+ if (zDevice != nullptr) {
	166	+ aclrtFree(zDevice);
	167	+ }
	168	+ if (stream != nullptr) {
	169	+ aclrtDestroyStream(stream);
	170	+ }
	171	+ if (deviceSet) {
	172	+ aclrtResetDevice(deviceId);
	173	+ }
	174	+ if (aclInitialized) {
	175	+ aclFinalize();
	176	+ }
	177	+ };
	178	+
	179	+ // 初始化 ACL。
	180	+ if (!CheckAclResult(aclInit(nullptr), "aclInit")) {
	181	+ return {};
	182	+ }
	183	+ aclInitialized = true;
	184	+
	185	+ if (!CheckAclResult(aclrtSetDevice(deviceId), "aclrtSetDevice")) {
	186	+ cleanup();
	187	+ return {};
	188	+ }
	189	+ deviceSet = true;
	190	+
	191	+ if (!CheckAclResult(aclrtCreateStream(&stream), "aclrtCreateStream")) {
	192	+ cleanup();
	193	+ return {};
	194	+ }
	195	+
	196	+ // 申请 Device 内存。
	197	+ size_t byteSize = static_cast<size_t>(totalLength) * sizeof(float);
	198	+ if (!CheckAclResult(
	199	+ aclrtMalloc(reinterpret_cast<void **>(&xDevice), byteSize, ACL_MEM_MALLOC_HUGE_FIRST),
	200	+ "aclrtMalloc(xDevice)") \|\|
	201	+ !CheckAclResult(
	202	+ aclrtMalloc(reinterpret_cast<void **>(&yDevice), byteSize, ACL_MEM_MALLOC_HUGE_FIRST),
	203	+ "aclrtMalloc(yDevice)") \|\|
	204	+ !CheckAclResult(
	205	+ aclrtMalloc(reinterpret_cast<void **>(&zDevice), byteSize, ACL_MEM_MALLOC_HUGE_FIRST),
	206	+ "aclrtMalloc(zDevice)")) {
	207	+ cleanup();
	208	+ return {};
	209	+ }
	210	+
	211	+ // 将输入数据从 Host 拷贝到 Device（H2D）。
	212	+ if (!CheckAclResult(
	213	+ aclrtMemcpy(xDevice, byteSize, x.data(), byteSize, ACL_MEMCPY_HOST_TO_DEVICE),
	214	+ "aclrtMemcpy(xHostToDevice)") \|\|
	215	+ !CheckAclResult(
	216	+ aclrtMemcpy(yDevice, byteSize, y.data(), byteSize, ACL_MEMCPY_HOST_TO_DEVICE),
	217	+ "aclrtMemcpy(yHostToDevice)")) {
	218	+ cleanup();
	219	+ return {};
	220	+ }
	221	+
	222	+ // 调用 kernel。
	223	+ mul_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);
	224	+ if (!CheckAclResult(aclrtSynchronizeStream(stream), "aclrtSynchronizeStream")) {
	225	+ cleanup();
	226	+ return {};
	227	+ }
	228	+
	229	+ // 将结果从 Device 拷贝回 Host（D2H）。
	230	+ std::vector<float> z(totalLength, 0.0f);
	231	+ if (!CheckAclResult(
	232	+ aclrtMemcpy(z.data(), byteSize, zDevice, byteSize, ACL_MEMCPY_DEVICE_TO_HOST),
	233	+ "aclrtMemcpy(zDeviceToHost)")) {
	234	+ cleanup();
	235	+ return {};
	236	+ }
	237	+
	238	+ cleanup();
	239	+ return z;
	240	+}
	241	+
	242	+// ----- 验证与主函数 -----
	243	+
	244	+uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
	245	+{
	246	+ auto printTensor = [](std::vector<float> &tensor, const char *name) {
	247	+ constexpr size_t maxPrintSize = 20;
	248	+ std::cout << name << ": ";
	249	+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
	250	+ std::ostream_iterator<float>(std::cout, " "));
	251	+ if (tensor.size() > maxPrintSize) {
	252	+ std::cout << "...";
	253	+ }
	254	+ std::cout << std::endl;
	255	+ };
	256	+ printTensor(output, "Output");
	257	+ printTensor(golden, "Golden");
	258	+ if (output.size() == golden.size() && std::equal(golden.begin(), golden.end(), output.begin())) {
	259	+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
	260	+ return 0;
	261	+ } else {
	262	+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
	263	+ return 1;
	264	+ }
	265	+}
	266	+
	267	+int32_t main(int32_t argc, char *argv[])
	268	+{
	269	+ constexpr uint32_t totalLength = 8 * 2048;
	270	+ constexpr float valueX = 1.2f;
	271	+ constexpr float valueY = 2.3f;
	272	+ std::vector<float> x(totalLength, valueX);
	273	+ std::vector<float> y(totalLength, valueY);
	274	+
	275	+ std::vector<float> output = kernel_mul(x, y);
	276	+
	277	+ std::vector<float> golden(totalLength, valueX * valueY);
	278	+ return VerifyResult(output, golden);
	279	+}

A2026/CANN-Code-Detective/Challenge05-MulCustom/gcw_JQJBklQH/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