【代码侦探Challenge05】: complete Challenge05 MulCustom-cann-outreach-AtomGit

已合并

【代码侦探Challenge05】: complete Challenge05 MulCustom #3456

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

【代码侦探Challenge05】: complete Challenge05 MulCustom #3456

已合并

jczhaang创建于 9月6日

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

共 3 个文件变更+319-0

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

@@ -0,0 +1,300 @@
	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	+ // 请补充……
	37	+ // 1. 计算单核处理长度 blockLength = totalLength / GetBlockNum()
	38	+ // 2. 计算 tileLength = blockLength / tileNum / BUFFER_NUM
	39	+ // 3. 使用 SetGlobalBuffer 设置 xGm, yGm, zGm 的起始地址和长度
	40	+ // （注意用 GetBlockIdx() 计算每个核的偏移）
	41	+ // 4. 使用 pipe.InitBuffer 为 inQueueX, inQueueY, outQueueZ 分配内存
	42	+ // 每个队列分配 BUFFER_NUM 个块，每块大小 tileLength * sizeof(float)
	43	+ this->blockLength = totalLength / AscendC::GetBlockNum();
	44	+ this->tileNum = tileNum;
	45	+ this->tileLength = this->blockLength / this->tileNum / BUFFER_NUM;
	46	+
	47	+ xGm.SetGlobalBuffer(
	48	+ (__gm__ float )x + this->blockLength AscendC::GetBlockIdx(),
	49	+ this->blockLength);
	50	+ yGm.SetGlobalBuffer(
	51	+ (__gm__ float )y + this->blockLength AscendC::GetBlockIdx(),
	52	+ this->blockLength);
	53	+ zGm.SetGlobalBuffer(
	54	+ (__gm__ float )z + this->blockLength AscendC::GetBlockIdx(),
	55	+ this->blockLength);
	56	+
	57	+ pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
	58	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
	59	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileLength * sizeof(float));
	60	+ }
	61	+ __aicore__ inline void Process()
	62	+ {
	63	+ // 请补充……
	64	+ // 循环次数 = tileNum * BUFFER_NUM
	65	+ // 每次迭代调用 CopyIn(i) → Compute(i) → CopyOut(i)
	66	+ int32_t loopCount = this->tileNum * BUFFER_NUM;
	67	+
	68	+ for (int32_t i = 0; i < loopCount; i++) {
	69	+ CopyIn(i);
	70	+ Compute(i);
	71	+ CopyOut(i);
	72	+ }
	73	+ }
	74	+
	75	+private:
	76	+ __aicore__ inline void CopyIn(int32_t progress)
	77	+ {
	78	+ // 请补充……
	79	+ // 1. 从 inQueueX 中 AllocTensor<float> 获取 xLocal
	80	+ // 2. 从 inQueueY 中 AllocTensor<float> 获取 yLocal
	81	+ // 3. 使用 DataCopy 从 xGm[progress * tileLength] 搬入 tileLength 个元素到 xLocal
	82	+ // 4. 使用 DataCopy 从 yGm[progress * tileLength] 搬入 tileLength 个元素到 yLocal
	83	+ // 5. 将 xLocal, yLocal 入队 (EnQue)
	84	+ AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
	85	+ AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
	86	+
	87	+ AscendC::DataCopy(
	88	+ xLocal,
	89	+ xGm[progress * this->tileLength],
	90	+ this->tileLength);
	91	+
	92	+ AscendC::DataCopy(
	93	+ yLocal,
	94	+ yGm[progress * this->tileLength],
	95	+ this->tileLength);
	96	+
	97	+ inQueueX.EnQue(xLocal);
	98	+ inQueueY.EnQue(yLocal);
	99	+ }
	100	+ __aicore__ inline void Compute(int32_t progress)
	101	+ {
	102	+ // 请补充……
	103	+ // 1. 从 inQueueX 中 DeQue<float> 获取 xLocal
	104	+ // 2. 从 inQueueY 中 DeQue<float> 获取 yLocal
	105	+ // 3. 从 outQueueZ 中 AllocTensor<float> 获取 zLocal
	106	+ // 4. 使用 Mul(zLocal, xLocal, yLocal, tileLength) 计算逐元素乘法
	107	+ // 5. 将 zLocal 入队 (EnQue)
	108	+ // 6. 释放 xLocal, yLocal (FreeTensor)
	109	+ AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
	110	+ AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
	111	+ AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
	112	+
	113	+ AscendC::Mul(zLocal, xLocal, yLocal, this->tileLength);
	114	+
	115	+ outQueueZ.EnQue(zLocal);
	116	+
	117	+ inQueueX.FreeTensor(xLocal);
	118	+ inQueueY.FreeTensor(yLocal);
	119	+ }
	120	+ __aicore__ inline void CopyOut(int32_t progress)
	121	+ {
	122	+ // 请补充……
	123	+ // 1. 从 outQueueZ 中 DeQue<float> 获取 zLocal
	124	+ // 2. 使用 DataCopy 将 zLocal 搬出到 zGm[progress * tileLength]
	125	+ // 3. 释放 zLocal (FreeTensor)
	126	+ AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
	127	+
	128	+ AscendC::DataCopy(
	129	+ zGm[progress * this->tileLength],
	130	+ zLocal,
	131	+ this->tileLength);
	132	+
	133	+ outQueueZ.FreeTensor(zLocal);
	134	+ }
	135	+
	136	+private:
	137	+ // 请补充……
	138	+ // 声明以下成员变量：
	139	+ // TPipe pipe;
	140	+ // TQue<TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
	141	+ // TQue<TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	142	+ // GlobalTensor<float> xGm, yGm, zGm;
	143	+ // uint32_t blockLength, tileNum, tileLength;
	144	+
	145	+ AscendC::TPipe pipe;
	146	+
	147	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX;
	148	+ AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueY;
	149	+ AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
	150	+
	151	+ AscendC::GlobalTensor<float> xGm;
	152	+ AscendC::GlobalTensor<float> yGm;
	153	+ AscendC::GlobalTensor<float> zGm;
	154	+
	155	+ uint32_t blockLength;
	156	+ uint32_t tileNum;
	157	+ uint32_t tileLength;
	158	+};
	159	+
	160	+// Kernel 入口函数
	161	+__global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, MulCustomTilingData tiling)
	162	+{
	163	+ // 请补充……
	164	+ // 1. 添加 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)
	165	+ // 2. 实例化 KernelMul op
	166	+ // 3. 调用 op.Init(x, y, z, tiling.totalLength, tiling.tileNum)
	167	+ // 4. 调用 op.Process()
	168	+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
	169	+
	170	+ KernelMul op;
	171	+ op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
	172	+ op.Process();
	173	+}
	174	+
	175	+// ----- Host 侧：kernel 直调封装 -----
	176	+
	177	+std::vector<float> kernel_mul(std::vector<float> &x, std::vector<float> &y)
	178	+{
	179	+ // 请补充……
	180	+ // 1. 定义 blockDim = 8, totalLength = x.size()
	181	+ // 2. 定义 tiling: totalLength, tileNum = 8
	182	+ // 3. 初始化 ACL: aclInit, aclrtSetDevice, aclrtCreateStream
	183	+ // 4. 申请 Host/Device 内存 (aclrtMallocHost, aclrtMalloc)
	184	+ // 5. 将输入数据拷贝到 Device (aclrtMemcpy H2D)
	185	+ // 6. 调用 kernel: mul_custom<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling)
	186	+ // 7. 同步流 (aclrtSynchronizeStream)
	187	+ // 8. 将结果从 Device 拷贝回 Host (aclrtMemcpy D2H)
	188	+ // 9. 构造 vector<float> 返回
	189	+ // 10. 释放内存，销毁流，重置设备，aclFinalize
	190	+ constexpr uint32_t blockDim = 8;
	191	+ uint32_t totalLength = x.size();
	192	+ size_t totalByteSize = totalLength * sizeof(float);
	193	+ int32_t deviceId = 0;
	194	+ aclrtStream stream = nullptr;
	195	+
	196	+ MulCustomTilingData tiling = {/totalLength:/totalLength, /tileNum:/8};
	197	+
	198	+ uint8_t xHost = reinterpret_cast<uint8_t >(x.data());
	199	+ uint8_t yHost = reinterpret_cast<uint8_t >(y.data());
	200	+ uint8_t *zHost = nullptr;
	201	+
	202	+ uint8_t *xDevice = nullptr;
	203	+ uint8_t *yDevice = nullptr;
	204	+ uint8_t *zDevice = nullptr;
	205	+
	206	+ aclInit(nullptr);
	207	+ aclrtSetDevice(deviceId);
	208	+ aclrtCreateStream(&stream);
	209	+
	210	+ aclrtMallocHost((void **)(&zHost), totalByteSize);
	211	+
	212	+ aclrtMalloc((void **)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	213	+ aclrtMalloc((void **)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	214	+ aclrtMalloc((void **)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
	215	+
	216	+ aclrtMemcpy(
	217	+ xDevice,
	218	+ totalByteSize,
	219	+ xHost,
	220	+ totalByteSize,
	221	+ ACL_MEMCPY_HOST_TO_DEVICE);
	222	+
	223	+ aclrtMemcpy(
	224	+ yDevice,
	225	+ totalByteSize,
	226	+ yHost,
	227	+ totalByteSize,
	228	+ ACL_MEMCPY_HOST_TO_DEVICE);
	229	+
	230	+ mul_custom<<<blockDim, nullptr, stream>>>(
	231	+ xDevice,
	232	+ yDevice,
	233	+ zDevice,
	234	+ tiling);
	235	+
	236	+ aclrtSynchronizeStream(stream);
	237	+
	238	+ aclrtMemcpy(
	239	+ zHost,
	240	+ totalByteSize,
	241	+ zDevice,
	242	+ totalByteSize,
	243	+ ACL_MEMCPY_DEVICE_TO_HOST);
	244	+
	245	+ std::vector<float> z(
	246	+ (float *)zHost,
	247	+ (float *)(zHost + totalByteSize));
	248	+
	249	+ aclrtFree(xDevice);
	250	+ aclrtFree(yDevice);
	251	+ aclrtFree(zDevice);
	252	+ aclrtFreeHost(zHost);
	253	+
	254	+ aclrtDestroyStream(stream);
	255	+ aclrtResetDevice(deviceId);
	256	+ aclFinalize();
	257	+
	258	+ return z;
	259	+}
	260	+
	261	+// ----- 验证与主函数 -----
	262	+
	263	+uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
	264	+{
	265	+ auto printTensor = [](std::vector<float> &tensor, const char *name) {
	266	+ constexpr size_t maxPrintSize = 20;
	267	+ std::cout << name << ": ";
	268	+ std::copy(tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),
	269	+ std::ostream_iterator<float>(std::cout, " "));
	270	+ if (tensor.size() > maxPrintSize) {
	271	+ std::cout << "...";
	272	+ }
	273	+ std::cout << std::endl;
	274	+ };
	275	+ printTensor(output, "Output");
	276	+ printTensor(golden, "Golden");
	277	+ if (std::equal(golden.begin(), golden.end(), output.begin())) {
	278	+ std::cout << "[Success] Case accuracy is verification passed." << std::endl;
	279	+ return 0;
	280	+ } else {
	281	+ std::cout << "[Failed] Case accuracy is verification failed!" << std::endl;
	282	+ return 1;
	283	+ }
	284	+}
	285	+
	286	+int32_t main(int32_t argc, char *argv[])
	287	+{
	288	+ constexpr uint32_t totalLength = 8 * 2048;
	289	+ constexpr float valueX = 1.2f;
	290	+ constexpr float valueY = 2.3f;
	291	+ std::vector<float> x(totalLength, valueX);
	292	+ std::vector<float> y(totalLength, valueY);
	293	+
	294	+ // 请补充……
	295	+ // 调用 kernel_mul(x, y) 获取 output
	296	+ std::vector<float> output = kernel_mul(x, y);
	297	+
	298	+ std::vector<float> golden(totalLength, valueX * valueY);
	299	+ return VerifyResult(output, golden);
	300	+}

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