【代码侦探Challenge05】实现逐元素乘法算子 MulCustom-cann-outreach-AtomGit

已开启

【代码侦探Challenge05】实现逐元素乘法算子 MulCustom #2431

aodebiao创建于 25 天前从gxhgxh2475698/cann-outreach_1729： challenge05-mulcustom 合入到 cann/cann-outreach： master

【代码侦探Challenge05】实现逐元素乘法算子 MulCustom #2431

已开启

aodebiao创建于 25 天前

从gxhgxh2475698/cann-outreach_1729 challenge05-mulcustom合入到cann/cann-outreach master

共 3 个文件变更+207-0

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

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

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