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

已合并

【代码侦探Challenge05】完成 zhaojianing2025 的 MulCustom 逐元素乘法算子 #3500

zhaojianing2025创建于 9月7日从zhaojianing2025/cann-outreach： challenge05-mulcustom-zhaojianing2025 合入到 cann/cann-outreach： master

【代码侦探Challenge05】完成 zhaojianing2025 的 MulCustom 逐元素乘法算子 #3500

已合并

zhaojianing2025创建于 9月7日

从zhaojianing2025/cann-outreach challenge05-mulcustom-zhaojianing2025合入到cann/cann-outreach master

共 3 个文件变更+205-0

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

@@ -0,0 +1,186 @@
	1	+// mul_custom.asc
	2	+// Ascend C 逐元素乘法算子：z[i] = x[i] * y[i]
	3	+
	4	+#include <algorithm>
	5	+#include <cstdint>
	6	+#include <iostream>
	7	+#include <iterator>
	8	+#include <vector>
	9	+#include "acl/acl.h"
	10	+#include "kernel_operator.h"
	11	+
	12	+constexpr uint32_t BUFFER_NUM = 2;
	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(GM_ADDR x, GM_ADDR y, GM_ADDR z,
	25	+ uint32_t totalLength, uint32_t tileNum)
	26	+ {
	27	+ blockLength = totalLength / AscendC::GetBlockNum();
	28	+ this->tileNum = tileNum;
	29	+ tileLength = blockLength / tileNum / BUFFER_NUM;
	30	+
	31	+ const uint32_t blockOffset = blockLength * AscendC::GetBlockIdx();
	32	+ xGm.SetGlobalBuffer((__gm__ float *)x + blockOffset, blockLength);
	33	+ yGm.SetGlobalBuffer((__gm__ float *)y + blockOffset, blockLength);
	34	+ zGm.SetGlobalBuffer((__gm__ float *)z + blockOffset, blockLength);
	35	+
	36	+ pipe.InitBuffer(inQueueX, BUFFER_NUM, tileLength * sizeof(float));
	37	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, tileLength * sizeof(float));
	38	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, tileLength * sizeof(float));
	39	+ }
	40	+
	41	+ __aicore__ inline void Process()
	42	+ {
	43	+ const int32_t loopCount = tileNum * BUFFER_NUM;
	44	+ for (int32_t i = 0; i < loopCount; ++i) {
	45	+ CopyIn(i);
	46	+ Compute();
	47	+ CopyOut(i);
	48	+ }
	49	+ }
	50	+
	51	+private:
	52	+ __aicore__ inline void CopyIn(int32_t progress)
	53	+ {
	54	+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
	55	+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
	56	+ const uint32_t offset = progress * tileLength;
	57	+ AscendC::DataCopy(xLocal, xGm[offset], tileLength);
	58	+ AscendC::DataCopy(yLocal, yGm[offset], tileLength);
	59	+ inQueueX.EnQue(xLocal);
	60	+ inQueueY.EnQue(yLocal);
	61	+ }
	62	+
	63	+ __aicore__ inline void Compute()
	64	+ {
	65	+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
	66	+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
	67	+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
	68	+
	69	+ AscendC::Mul(zLocal, xLocal, yLocal, tileLength);
	70	+
	71	+ outQueueZ.EnQue(zLocal);
	72	+ inQueueX.FreeTensor(xLocal);
	73	+ inQueueY.FreeTensor(yLocal);
	74	+ }
	75	+
	76	+ __aicore__ inline void CopyOut(int32_t progress)
	77	+ {
	78	+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
	79	+ AscendC::DataCopy(zGm[progress * tileLength], zLocal, tileLength);
	80	+ outQueueZ.FreeTensor(zLocal);
	81	+ }
	82	+
	83	+private:
	84	+ AscendC::TPipe pipe;
	85	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
	86	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;
	87	+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	88	+ AscendC::GlobalTensor<float> xGm;
	89	+ AscendC::GlobalTensor<float> yGm;
	90	+ AscendC::GlobalTensor<float> zGm;
	91	+ uint32_t blockLength;
	92	+ uint32_t tileNum;
	93	+ uint32_t tileLength;
	94	+};
	95	+
	96	+__global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z,
	97	+ MulCustomTilingData tiling)
	98	+{
	99	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	100	+ KernelMul op;
	101	+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
	102	+ op.Process();
	103	+}
	104	+
	105	+std::vector<float> kernel_mul(std::vector<float> &x, std::vector<float> &y)
	106	+{
	107	+ constexpr uint32_t blockDim = 8;
	108	+ constexpr uint32_t tileNum = 8;
	109	+ constexpr int32_t deviceId = 0;
	110	+ const uint32_t totalLength = x.size();
	111	+ const size_t totalByteSize = totalLength * sizeof(float);
	112	+ const MulCustomTilingData tiling = {totalLength, tileNum};
	113	+
	114	+ aclrtStream stream = nullptr;
	115	+ float *zHost = nullptr;
	116	+ uint8_t *xDevice = nullptr;
	117	+ uint8_t *yDevice = nullptr;
	118	+ uint8_t *zDevice = nullptr;
	119	+
	120	+ aclInit(nullptr);
	121	+ aclrtSetDevice(deviceId);
	122	+ aclrtCreateStream(&stream);
	123	+
	124	+ aclrtMallocHost(reinterpret_cast<void **>(&zHost), totalByteSize);
	125	+ aclrtMalloc(reinterpret_cast<void **>(&xDevice), totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	126	+ aclrtMalloc(reinterpret_cast<void **>(&yDevice), totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	127	+ aclrtMalloc(reinterpret_cast<void **>(&zDevice), totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	128	+
	129	+ aclrtMemcpy(xDevice, totalByteSize, x.data(), totalByteSize,
	130	+ ACL_MEMCPY_HOST_TO_DEVICE);
	131	+ aclrtMemcpy(yDevice, totalByteSize, y.data(), totalByteSize,
	132	+ ACL_MEMCPY_HOST_TO_DEVICE);
	133	+
	134	+ mul_custom<<<blockDim, nullptr, stream>>>((__gm__ uint8_t *)xDevice,
	135	+ (__gm__ uint8_t )yDevice, (__gm__ uint8_t )zDevice, tiling);
	136	+ aclrtSynchronizeStream(stream);
	137	+ aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize,
	138	+ ACL_MEMCPY_DEVICE_TO_HOST);
	139	+
	140	+ std::vector<float> output(zHost, zHost + totalLength);
	141	+
	142	+ aclrtFree(xDevice);
	143	+ aclrtFree(yDevice);
	144	+ aclrtFree(zDevice);
	145	+ aclrtFreeHost(zHost);
	146	+ aclrtDestroyStream(stream);
	147	+ aclrtResetDevice(deviceId);
	148	+ aclFinalize();
	149	+
	150	+ return output;
	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	+ }
	171	+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
	172	+ return 1;
	173	+}
	174	+
	175	+int32_t main(int32_t argc, char *argv[])
	176	+{
	177	+ constexpr uint32_t totalLength = 8 * 2048;
	178	+ constexpr float valueX = 1.2f;
	179	+ constexpr float valueY = 2.3f;
	180	+ std::vector<float> x(totalLength, valueX);
	181	+ std::vector<float> y(totalLength, valueY);
	182	+
	183	+ std::vector<float> output = kernel_mul(x, y);
	184	+ std::vector<float> golden(totalLength, valueX * valueY);
	185	+ return VerifyResult(output, golden);
	186	+}

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