Add DivCustom sample code for Challenge01-cann-outreach-AtomGit

已开启

Add DivCustom sample code for Challenge01 #1796

西瓜WiFi创建于 7月31日从m0_62011685/cann-outreach_8370： master 合入到 cann/cann-outreach： master

Add DivCustom sample code for Challenge01 #1796

已开启

西瓜WiFi创建于 7月31日

从m0_62011685/cann-outreach_8370 master合入到cann/cann-outreach master

共 1 个文件变更+178-0

A2026/CANN-Code-Detective/Challenge01-DivCustom/SampleUse/DivCustom/div_custom.asc+178-0

@@ -0,0 +1,178 @@
	1	+// div_custom.asc
	2	+
	3	+#include <cstdint>
	4	+#include <iostream>
	5	+#include <vector>
	6	+#include <algorithm>
	7	+#include <iterator>
	8	+#include "acl/acl.h"
	9	+#include "kernel_operator.h"
	10	+
	11	+constexpr uint32_t BUFFER_NUM = 2; // tensor num for each queue
	12	+
	13	+struct DivCustomTilingData
	14	+{
	15	+ uint32_t totalLength;
	16	+ uint32_t tileNum;
	17	+};
	18	+
	19	+class KernelDiv {
	20	+public:
	21	+ __aicore__ inline KernelDiv(){}
	22	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t totalLength, uint32_t tileNum)
	23	+ {
	24	+ // 初始化 Global Memory 地址：将 GM_ADDR 绑定到 GlobalTensor
	25	+ xGm.SetGlobalBuffer((__gm__ float *)x, totalLength);
	26	+ yGm.SetGlobalBuffer((__gm__ float *)y, totalLength);
	27	+ zGm.SetGlobalBuffer((__gm__ float *)z, totalLength);
	28	+
	29	+ // 保存 Tiling 参数
	30	+ this->totalLength = totalLength;
	31	+ this->tileNum = tileNum;
	32	+ this->tileLength = totalLength / tileNum;
	atomgit-bot7月31日 🟡 Medium Priority 变更行：第 32 行 `this->tileLength = totalLength / tileNum;`。受影响的契约：当 `totalLength` 不能被 `tileNum` 整除时，整数除法舍去余数，每个 tile 只处理 `tileLength` 个元素，循环执行 `tileNum` 次总计处理 `tileNum * (totalLength / tileNum)` 个元素，末尾 `totalLength % tileNum` 个元素永远不会被处理。失败模式：在当前硬编码参数下（`totalLength=16384, tileNum=8`）恰好整除，不触发。但若任何人修改 `totalLength` 或 `tileNum` 使得不能整除（如 `totalLength=16385`），末尾元素将保持为设备内存中的未初始化数据，导致静默数据丢失和输出不正确。建议：在 Process() 循环中动态计算每个 tile 的实际长度：若为最后一个 tile，使用 totalLength - progress * tileLength 作为实际处理长度。同时在 Init 中可对 tileNum ≤ totalLength 做断言。不准确？
	33	+
	34	+ // 初始化管道和队列（Pipe 会自动分配 UB 空间）
	35	+ pipe.InitBuffer(inQueueX, BUFFER_NUM, tileLength * sizeof(float));
	36	+ pipe.InitBuffer(inQueueY, BUFFER_NUM, tileLength * sizeof(float));
	37	+ pipe.InitBuffer(outQueueZ, BUFFER_NUM, tileLength * sizeof(float));
	38	+ }
	39	+ __aicore__ inline void Process()
	40	+ {
	41	+ // 主处理循环：按 tile 迭代执行 CopyIn → Compute → CopyOut
	42	+ for (int32_t i = 0; i < tileNum; i++) {
	43	+ CopyIn(i);
	44	+ Compute(i);
	45	+ CopyOut(i);
	46	+ }
	47	+ }
	48	+
	49	+private:
	50	+ __aicore__ inline void CopyIn(int32_t progress)
	51	+ {
	52	+ // 从 Global Memory 搬数据到 Unified Buffer（UB）
	53	+ // 计算当前 tile 的偏移量
	54	+ uint32_t offset = progress * tileLength;
	55	+
	56	+ // 申请 UB 队列缓冲区
	57	+ LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
	58	+ LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
	59	+
	60	+ // 使用 DataCopy 将 GM 数据搬入 UB
	61	+ DataCopy(xLocal, xGm[offset], tileLength);
	62	+ DataCopy(yLocal, yGm[offset], tileLength);
	63	+
	64	+ // 入队，供 Compute 阶段使用
	65	+ inQueueX.EnQue(xLocal);
	66	+ inQueueY.EnQue(yLocal);
	67	+ }
	68	+ __aicore__ inline void Compute(int32_t progress)
	69	+ {
	70	+ // 出队，执行逐元素除法，入队
	71	+ LocalTensor<float> xLocal = inQueueX.DeQue<float>();
	72	+ LocalTensor<float> yLocal = inQueueY.DeQue<float>();
	73	+ LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
	74	+
	75	+ // 执行逐元素除法：zLocal = xLocal / yLocal
	76	+ Div(zLocal, xLocal, yLocal, tileLength);
	77	+
	78	+ // 入队，供 CopyOut 阶段使用
	79	+ outQueueZ.EnQue(zLocal);
	80	+
	81	+ // 释放已消费的输入缓冲区
	82	+ inQueueX.FreeTensor(xLocal);
	83	+ inQueueY.FreeTensor(yLocal);
	84	+ }
	85	+ __aicore__ inline void CopyOut(int32_t progress)
	86	+ {
	87	+ // 出队，将结果从 UB 搬回 Global Memory
	88	+ uint32_t offset = progress * tileLength;
	89	+ LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
	90	+
	91	+ // 将 UB 结果数据搬回 GM
	92	+ DataCopy(zGm[offset], zLocal, tileLength);
	93	+
	94	+ // 释放输出缓冲区
	95	+ outQueueZ.FreeTensor(zLocal);
	96	+ }
	97	+
	98	+private:
	99	+ // Global Tensor：管理 Global Memory 数据
	100	+ GlobalTensor<float> xGm;
	101	+ GlobalTensor<float> yGm;
	102	+ GlobalTensor<float> zGm;
	103	+
	104	+ // Tiling 参数
	105	+ uint32_t totalLength;
	106	+ uint32_t tileNum;
	107	+ uint32_t tileLength;
	108	+
	109	+ // TPipe：管理核内资源（管道）
	110	+ TPipe pipe;
	111	+
	112	+ // TQue：管理 UB 队列（Double Buffer 机制）
	113	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
	114	+ TQue<QuePosition::VECIN, BUFFER_NUM> inQueueY;
	115	+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
	116	+};
	117	+
	118	+__global__ __aicore__ void div_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, DivCustomTilingData tiling)
	119	+{
	120	+ // 实例化 KernelDiv 并调用 Init + Process
	121	+ KernelDiv kernel;
	122	+ kernel.Init(x, y, z, tiling.totalLength, tiling.tileNum);
	123	+ kernel.Process();
	124	+}
	125	+
	126	+std::vector<float> kernel_div(std::vector<float> &x, std::vector<float> &y)
	127	+{
	128	+ // 申请 GM 内存、搬运数据、启动 kernel、回拷结果
	129	+ constexpr uint32_t totalLength = 8 * 2048;
	130	+ constexpr uint32_t tileNum = 8;
	131	+ uint32_t size = totalLength * sizeof(float);
	132	+
	133	+ // 设置 Tiling 数据
	134	+ DivCustomTilingData tiling;
	135	+ tiling.totalLength = totalLength;
	136	+ tiling.tileNum = tileNum;
	137	+
	138	+ // 申请 device 侧 GM 内存
	139	+ float *deviceX = nullptr;
	140	+ float *deviceY = nullptr;
	141	+ float *deviceZ = nullptr;
	142	+ aclrtMalloc((void **)&deviceX, size, ACL_MEM_MALLOC_HUGE_FIRST);
	143	+ aclrtMalloc((void **)&deviceY, size, ACL_MEM_MALLOC_HUGE_FIRST);
	144	+ aclrtMalloc((void **)&deviceZ, size, ACL_MEM_MALLOC_HUGE_FIRST);
	atomgit-bot7月31日 🟠 High Priority 变更行：第 142-144 行三次 `aclrtMalloc` 调用均未检查返回值。受影响的契约：`aclrtMalloc` 返回 `aclError` 类型错误码，成功返回 `ACL_SUCCESS`。当 device 内存不足或运行时未初始化时分配会失败，返回的 `deviceX`/`deviceY`/`deviceZ` 保持为 `nullptr`。失败模式：若任一 `aclrtMalloc` 失败，后续第 147-148 行的 `aclrtMemcpy(deviceX, ...)` 将对空指针执行 device 端内存拷贝，导致运行时段错误或不可预期的行为。此外，若 `deviceX` 分配成功但 `deviceY` 失败，代码不会释放已成功的 `deviceX`（资源泄漏在崩溃时次要）。建议：对每个 aclrtMalloc 调用检查返回的 aclError，失败时执行清理并通过异常或错误码向上传播。不准确？
	145	+
	146	+ // 将 host 数据拷贝到 device
	147	+ aclrtMemcpy(deviceX, size, x.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	148	+ aclrtMemcpy(deviceY, size, y.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	atomgit-bot7月31日 🟠 High Priority 变更行：第 147-148 行和第 155 行，三次 `aclrtMemcpy` 调用均未检查返回值。受影响的契约：`aclrtMemcpy` 返回 `aclError`。若拷贝失败（如设备内存不可访问、权限不足），数据不会正确传输，但代码继续执行后续逻辑。失败模式： H2D 拷贝（第 147-148 行）失败 → device 内存中为未初始化数据，kernel 计算结果为垃圾值。 D2H 拷贝（第 155 行）失败 → `z` vector 中为未初始化数据，VerifyResult 将误报失败或意外通过。建议：对每个 aclrtMemcpy 调用检查返回值，失败时清理已分配资源并向上报告错误。不准确？
	149	+
	150	+ // 启动 kernel
	151	+ AscendC::LaunchKernel(div_custom, 1, deviceX, deviceY, deviceZ, tiling);
	152	+
	153	+ // 将 device 结果回拷到 host
	154	+ std::vector<float> z(totalLength);
	155	+ aclrtMemcpy(z.data(), size, deviceZ, size, ACL_MEMCPY_DEVICE_TO_HOST);
	156	+
	157	+ // 释放 device 内存
	158	+ aclrtFree(deviceX);
	159	+ aclrtFree(deviceY);
	160	+ aclrtFree(deviceZ);
	161	+
	162	+ return z;
	163	+}
	164	+
	165	+int32_t main(int32_t argc, char *argv[])
	166	+{
	167	+ constexpr uint32_t totalLength = 8 * 2048;
	168	+ constexpr float valueX = 1.2f;
	169	+ constexpr float valueY = 2.3f;
	170	+ std::vector<float> x(totalLength, valueX);
	171	+ std::vector<float> y(totalLength, valueY);
	172	+
	173	+ // 调用 kernel_div 函数执行算子计算
	174	+ std::vector<float> output = kernel_div(x, y);
	175	+
	176	+ std::vector<float> golden(totalLength, valueX / valueY);
	177	+ return VerifyResult(output, golden);
	atomgit-bot7月31日 🔴 Critical 变更行：第 177 行调用 `VerifyResult(output, golden)`。受影响的契约：`main()` 函数依赖 `VerifyResult` 来验证输出精度。该函数在模板文件（`Template/DivCustom/div_custom.asc` 第 60-81 行）中已定义，但在 SampleUse 实现中被遗漏。失败模式：编译时链接器报错 `undefined reference to VerifyResult(...)`，代码无法通过编译。建议：将模板中 VerifyResult 函数（第 60-81 行）完整拷贝到本文件第 164 行之前（kernel_div 函数之后、main 函数之前）。不准确？西瓜WiFi 7月31日评论： 7月31日评论：
	178	+}
	atomgit-bot7月31日 🟠 High Priority 变更行：`main()` 函数（第 165-178 行）和 `kernel_div()` 函数（第 126-163 行）。受影响的契约：CANN ACL 运行时要求在调用任何 `aclrt*` API（如 `aclrtMalloc`、`aclrtMemcpy`、`LaunchKernel`）之前必须先调用 `aclInit(nullptr)` 初始化，运行结束后调用 `aclFinalize()` 释放资源。失败模式：运行时 `aclrtMalloc` 等调用将返回 `ACL_ERROR_RT_NOT_INIT` 错误码（但因返回值未被检查而静默失败），后续对空指针的 `aclrtMemcpy` 导致段错误崩溃。建议：在 main() 开头添加 aclInit(nullptr)，在 return 之前添加 aclFinalize()，并检查两者的返回值。不准确？