{
"cells": [
{
"cell_type": "markdown",
"metadata": {},
"source": [
"# Conv2D V2 核函数开发\n",
"\n",
"本章以 **Conv2D V2** 算子为例，完整演示如何在 AscendC 中直接调用卷积核函数。内容涵盖：\n",
"\n",
"- 算子分析：卷积数学表达式、输入/输出规格、NPU 片上存储层次。\n",
"- 核函数开发：Tiling 结构体、片上缓冲区管理、im2col、Mmad、Fixpipe 等关键步骤。\n",
"- 主机侧调用：SetTilingData、CheckTilingSpace、ACL 内存管理与核函数启动。\n",
"\n",
"本章学习大纲如下：\n",
"\n",
"- 算子分析：明确 Conv2D 的数学表达式、输入输出规格及 NPU 存储层次。\n",
"- 核函数开发：逐步实现设备侧卷积核函数的各个模块。\n",
"- 主机侧调用代码：编写 Host 侧驱动程序，完成 Tiling 填充、内存管理与核函数启动。\n",
"- 编译与运行：使用 CMake + 毕昇编译器完成编译并执行。\n",
"\n",
"---\n",
"## 1. 环境准备\n",
"\n",
"正式开始学习之前，先要对 Jupyter 环境进行初始化。以下代码完成了初始化并将环境变量导入 Jupyter 环境，同时创建代码目录，保证能正常使用毕昇编译器完成算子的开发及编译。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"import os, subprocess\n",
"os.makedirs(\"Sources/01.03\", exist_ok=True)\n",
"env = subprocess.check_output(\"bash -l -c 'source $ASCEND_TOOLKIT_HOME/set_env.sh && env'\", shell=True, text=True)\n",
"for line in env.splitlines():\n",
" if \"=\" in line: os.environ.__setitem__(*line.split(\"=\", 1))\n",
"print(\"\\n Environment initialization process completed successfully!\")"
]
},
{
"cell_type": "markdown",
"metadata": {},
"outputs": [],
"source": [
"---\n\n",
"## 2. 算子分析\n\n",
"### 2.1 数学表达式\n\n",
"二维卷积（Conv2D）的数学表达式为：\n\n",
"$$Y[n, c_o, h_o, w_o] = \\sum_{c_i} \\sum_{k_h} \\sum_{k_w} X[n,\\, c_i,\\, h_o \\cdot s_H + k_h \\cdot d_H,\\, w_o \\cdot s_W + k_w \\cdot d_W] \\times W[c_o, c_i, k_h, k_w]$$\n\n",
"其中 $s_H, s_W$ 为步长，$d_H, d_W$ 为膨胀系数。\n\n",
"### 2.2 输入与输出\n\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>张量</strong></td>\n",
" <td align=\"center\"><strong>形状（NCHW）</strong></td>\n",
" <td align=\"center\" style=\"width: 180px\"><strong>取值范围</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">输入特征图 X </td>\n",
" <td align=\"left\">(BATCH, CI, HI, WI) </td>\n",
" <td align=\"left\">fp16</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">卷积核 W </td>\n",
" <td align=\"left\"> (CO, CI, KH, KW) </td>\n",
" <td align=\"left\">fp16</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">输出特征图 Y </td>\n",
" <td align=\"left\"> (BATCH, CO, HO, WO) </td>\n",
" <td align=\"left\">fp16</td>\n",
"</tr>\n",
"</table>\n",
"输出尺寸公式：\n\n$$HO = \\lfloor (HI + pad\\_top + pad\\_bottom - d_H(KH-1) - 1) / s_H \\rfloor + 1$$\n\n",
"### 2.3 NPU 上存储层次\n\nConv2D 的数据流经过以下存储层次（速度从快到慢）：\n\n",
"```\nGM (Global Memory) — 全局存储空间 (Ascend 950PR/Ascend 950DT)\n",
"L1 (L1 缓冲区) — 每核 512 MB\nL0A (矩阵 A 缓冲区) — 64 KB，送入 Mmad 左操作数\nL0B (矩阵 B 缓冲区) — 64 KB，送入 Mmad 右操作数\nL0C (累加器缓冲区) — 256 KB，存储 Mmad 结果（fp32）\n",
"```\n\n数据流示意：\n\n",
"```\nGM ──DataCopy──► L1 (LoadAL1 / LoadBL1)\n",
"L1 ──Load3D──► L0A (LoadAL0: im2col 在线展开)\n",
"L1 ──Load2D──► L0B (LoadBL0: 权重 NZ 布局)\n",
"L0A × L0B ──Mmad──► L0C\nL0C ──Fixpipe──► GM (fp32→fp16 类型转换 + NCHW 布局)\n",
"```\n\n### 2.4 本 Demo 的局限性与约束\n\n本 Demo 是一个**最小可运行实现**，**只支持Ascend 950PR/Ascend 950DT芯片型号**，目的是清晰展示 AscendC 卷积核函数的完整调用链。为了降低复杂度，做了以下简化，在实际生产代码中均需扩展：\n\n",
"#### 全载搬运（Full-Load）\n\n这是最核心的限制。`SetTilingData` 中令 `kL0 == kAL1 == kBL1 == KH × KW × CI`，即整个 K 维度在一次 L1→L0 搬运中全部完成，`ddr2l0LoopK` 始终为 1，K 循环只跑一次。\n\n",
"当输入通道 CI 或卷积核尺寸 KH/KW 增大时，全载所需的 L0A/L0B 空间会线性增长，很快超出 64 KB 的硬件上限。生产实现需要将 K 维度分块，分多次 Mmad 累加。\n\n",
"#### 单核，无多核分工\n\n`numBlocks = 1`，整个输出特征图由单核完成。没有按 batch、输出通道或空间维度分核的逻辑（无 `block_idx` 分片）。\n\n",
"#### 无流水 / 无双缓冲\n\nL1 队列深度为 1（`TQue<..., 1>`），L0 双缓冲关闭（`L0_SYBC_DB_CLOSE`）。数据搬运（MTE1）和矩阵计算（M）串行执行，没有流水重叠，硬件利用率低。\n\n",
"#### 数据类型固定为 fp16\n\n输入特征图、权重、输出均为 `half`（fp16），L0C 累加器为 `float`（fp32）。不支持 int8、bf16 等其他数据类型。\n\n",
"#### 数据布局固定为 NCHW\n\n`ConvFormat` 枚举只定义了 `NCHW = 1`，不支持 NHWC 等其他布局。\n\n",
"#### CI 和 CO 均固定为 16\n\n- **CI = 16**：`aL1SpaceSize = 8192` 硬编码，只够容纳 CI=16 的一个输入 patch；`cinAInCore / cinBInCore` 也直接由 `kAL1 / kernelHW` 推算，没有多 CI 块的循环。\n",
"- **CO = 16**：`nBL1 = 16`、`nL0 = 16` 硬编码，每次只处理 16 个输出通道，没有多 CO 块的循环。\n\n两者的根本原因相同：Demo 没有实现通道方向的分块循环，CI 和 CO 必须恰好等于 16（即一个 C0 块），而不是\"16 的倍数\"。\n\n",
"#### 只支持对称 Padding\n\n`Conv2DTilingData` 只存储 `padTop` 和 `padLeft`，没有独立的 `padBottom` 和 `padRight` 字段。`LoadAL1` 在运行时从输入/输出尺寸的几何关系推算 bottom/right padding，隐含 `padBottom = padTop`、`padRight = padLeft`。非对称 padding 无法通过当前结构体表达。\n\n",
"#### 空间尺寸上限（由全载搬运决定）\n\n`CheckTilingSpace` 会在启动前验证以下约束，超出则拒绝运行：\n\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>缓冲区</strong></td>\n",
" <td align=\"center\"><strong>硬件上限</strong></td>\n",
" <td align=\"center\" style=\"width: 300px\"><strong>全载用量公式</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">L1 </td>\n",
" <td align=\"left\">512 MB / 核</td>\n",
" <td align=\"left\">`(hiLoad × wiLoad × CI + nBL1 × kBL1) × 2` 字节</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">L0A </td>\n",
" <td align=\"left\"> 64 KB / 核 </td>\n",
" <td align=\"left\"> `AlignUp(HO×WO, 16) × (KH×KW×CI) × 2` 字节 </td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">L0B </td>\n",
" <td align=\"left\"> 64 KB / 核 </td>\n",
" <td align=\"left\"> `AlignUp(CO, 16) × (KH×KW×CI) × 2` 字节</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">L0C </td>\n",
" <td align=\"left\"> 256 KB / 核 </td>\n",
" <td align=\"left\"> `AlignUp(HO×WO, 16) × AlignUp(CO, 16) × 4` 字节 </td>\n",
"</tr>\n",
"</table>\n",
"其中 `hiLoad = (HO-1)×strideH + dilatedKH`，`wiLoad = (WO-1)×strideW + dilatedKW`。\n\n> **示例验证**：本 Demo 默认参数（CI=16, HI=WI=8, KH=KW=3, CO=16, 无 padding）下，输入全 1、权重全 1，每个输出点的计算值为：\n>\n> $$Y[n, c_o, h_o, w_o] = \\sum_{c_i=0}^{15} \\sum_{k_h=0}^{2} \\sum_{k_w=0}^{2} 1 \\times 1 = CI \\times KH \\times KW = 16 \\times 3 \\times 3 = 144$$\n>\n> 输出通道 $c_o$ 不影响每个点的数值，因为所有输出通道的权重均为 1，累加的元素数量（CI × KH × KW）相同。\n\n---\n\n## 3. 核函数开发\n\n基于对 Conv2D V2 算子的分析，我们开始逐步实现核函数代码。所有代码均写入 `Sources/01.03/minimal_demo.asc`，采用增量追加方式，便于逐节理解每个模块的作用。\n\n### 3.1 头文件与 Tiling 结构体\n\n首先引入必要的头文件，并定义 `Conv2DTilingData` 结构体。\n\n**Conv2DTilingData** 保存了核函数所需的全部参数：存储层次的分块尺寸以及 padding 偏移量。它由主机侧通过 `aclrtMemcpy` 传递给设备侧，再由 `CopyConv2dTiling` 从 GM 复制到片上内存后供核函数读取。\n\n**CopyConv2dTiling / GET_TILING_DATA**：核函数不能直接解引用 GM 指针读取结构体字段，必须先将数据复制到片上（L1 可访问）内存。`CopyConv2dTiling` 逐字完成这一复制，`GET_TILING_DATA` 宏将其封装为一行调用。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile Sources/01.03/minimal_demo.asc\n",
"/**\n",
" * Copyright (c) 2025 Huawei Technologies Co., Ltd.\n",
" */\n",
"\n",
"/*!\n",
" * \\file conv2d_v2_demo.asc\n",
" * \\brief Standalone demo for direct-calling Conv2D V2 kernel.\n",
" *\n",
" * ============================================================\n",
" * Tutorial overview\n",
" * ============================================================\n",
" * This file is the host-side driver for the Conv2D V2 kernel demo.\n",
" * It covers the full lifecycle of an AscendC kernel call:\n",
" *\n",
" * 1. Define conv shape constants (BATCH, CI, HI, WI, CO, KH, KW, …)\n",
" * 2. Fill in the tiling struct (SetTilingData) — tells the kernel how\n",
" * to partition work across the NPU memory hierarchy.\n",
" * 3. Validate that the tiling fits in on-chip SRAM (CheckTilingSpace).\n",
" * 4. Allocate device memory, copy inputs, launch the kernel.\n",
" * 5. Copy results back and dump to text files for inspection.\n",
" *\n",
" * NPU memory hierarchy (smallest / fastest → largest / slowest):\n",
" *\n",
" * GM (Global Memory) — off-chip (Ascend 950PR/Ascend 950DT)\n",
" * L1 (on-chip SRAM) — 512 MB per core\n",
" * L0A (Matrix A buffer) — 64 KB, feeds left operand of Mmad\n",
" * L0B (Matrix B buffer) — 64 KB, feeds right operand of Mmad\n",
" * L0C (Accumulator buffer) — 256 KB, stores Mmad result (fp32)\n",
" *\n",
" * Data flow for a single Conv2D tile:\n",
" *\n",
" * GM ──DataCopy──► L1 (LoadAL1 / LoadBL1)\n",
" * L1 ──Load3D──► L0A (LoadAL0: im2col on-the-fly)\n",
" * L1 ──Load2D──► L0B (LoadBL0: weight NZ layout)\n",
" * L0A × L0B ──Mmad──► L0C\n",
" * L0C ──Fixpipe──► GM (fp32→fp16 cast + NCHW layout)\n",
" * ============================================================\n",
" */\n",
"\n",
"#include <cstdint>\n",
"#include <cstring>\n",
"#include <iostream>\n",
"#include <fstream>\n",
"#include <cmath>\n",
"#include \"acl/acl.h\"\n",
"#include \"kernel_operator.h\"\n",
"\n",
"// ============================================================================\n",
"// Tiling data structure\n",
"// ============================================================================\n",
"//\n",
"// Conv2DTilingData holds all parameters the kernel needs: memory-hierarchy\n",
"// tiling values plus the two padding offsets.\n",
"// Passed from host to device via aclrtMemcpy, then copied from GM to\n",
"// on-chip memory by CopyConv2dTiling before the kernel reads any field.\n",
"//\n",
"// CopyConv2dTiling / GET_TILING_DATA: the kernel cannot dereference a GM\n",
"// pointer directly to read struct fields — the data must first be copied\n",
"// to on-chip (L1-accessible) memory. CopyConv2dTiling does that word-by-word.\n",
"\n",
"#pragma pack(1)\n",
"\n",
"struct Conv2DTilingData {\n",
" uint64_t orgHi = 0;\n",
" uint64_t orgWi = 0;\n",
" uint64_t orgHo = 0;\n",
" uint64_t orgWo = 0;\n",
" uint64_t singleCoreBatch = 0;\n",
" uint64_t singleCoreHo = 0;\n",
" uint64_t singleCoreWo = 0;\n",
" uint32_t orgCi = 0;\n",
" uint32_t orgCo = 0;\n",
" uint32_t singleCoreCi = 0;\n",
" uint32_t singleCoreCo = 0;\n",
" uint32_t kAL1 = 0;\n",
" uint32_t kBL1 = 0;\n",
" uint32_t nBL1 = 0;\n",
" uint32_t hoL0 = 0;\n",
" uint32_t woL0 = 0;\n",
" uint32_t kL0 = 0;\n",
" uint32_t nL0 = 0;\n",
" uint32_t orgHixWi = 0;\n",
" uint32_t kernelHxkernelW = 0;\n",
" uint32_t aL1SpaceSize = 0;\n",
" uint32_t cinAInCore = 0;\n",
" uint32_t cinBInCore = 0;\n",
" uint32_t mStep = 0;\n",
" uint32_t kStep = 0;\n",
" uint32_t fmapKStride = 0;\n",
" uint32_t coutOffsetBlock = 0;\n",
" uint32_t nL1DivBlockSize = 0;\n",
" uint32_t kernelH = 0;\n",
" uint32_t kernelW = 0;\n",
" uint32_t strideH = 0;\n",
" uint32_t strideW = 0;\n",
" uint32_t dilationH = 0;\n",
" uint32_t dilationW = 0;\n",
" int8_t offsetx = 0;\n",
" uint32_t padTop = 0;\n",
" uint32_t padLeft = 0;\n",
"};\n",
"\n",
"#pragma pack()\n",
"\n",
"__aicore__ inline void CopyConv2dTiling(Conv2DTilingData* dst, GM_ADDR tilingGM)\n",
"{\n",
" uint32_t* ptr = reinterpret_cast<uint32_t*>(dst);\n",
" auto src = reinterpret_cast<__gm__ uint32_t*>(tilingGM);\n",
" for (uint32_t i = 0; i < sizeof(Conv2DTilingData) / sizeof(uint32_t); i++, ptr++) {\n",
" *ptr = *(src + i);\n",
" }\n",
"}\n",
"\n",
"#define GET_TILING_DATA(tilingData, tilingArg) \\\n",
" Conv2DTilingData tilingData; \\\n",
" CopyConv2dTiling(&tilingData, tilingArg)\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.2 NPU存储层次与常量\n",
"\n",
"`#ifdef __CCE_AICORE__` 块内的代码只在设备侧编译。`namespace conv` 中定义了各个缓冲区的大小常量（L0A/L0B/L0C 均以字节为单位）、C0 块大小（32 字节）、padding 索引以及 Load3D/Load2D 所需的位域偏移量。\n",
"\n",
"这些常量直接对应 NPU 硬件规格，修改它们会导致越界访问，因此在教学中保持固定值。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"#ifdef __CCE_AICORE__\n",
"#include \"kernel_basic_intf.h\"\n",
"#include \"kernel_tiling/kernel_tiling.h\"\n",
"\n",
"using namespace AscendC;\n",
"\n",
"namespace conv {\n",
"\n",
"const static uint64_t L0A_SIZE = 65536;\n",
"const static uint64_t L0B_SIZE = 65536;\n",
"const static uint64_t L0C_SIZE = 262144;\n",
"const static uint64_t C0_SIZE = 32;\n",
"const static uint64_t PAD_IDX_T = 2;\n",
"const static uint64_t PAD_IDX_L = 0;\n",
"const static uint64_t PAD_IDX_R = 1;\n",
"const static uint64_t MAX_PAD_R = 255;\n",
"const static uint64_t MIN_HI_WI = 1;\n",
"const static uint32_t BLOCK_L0_N = 16;\n",
"const static uint32_t BLOCK_L0_M = 16;\n",
"const static uint32_t L0_SYBC_DB_CLOSE = 0x0;\n",
"\n",
"const static uint8_t MSTEP_OFFSET = 16;\n",
"const static uint8_t POSM_OFFSET = 48;\n",
"const static uint8_t POSK_OFFSET = 32;\n",
"const static uint8_t STRIDEH_OFFSET = 6;\n",
"const static uint8_t KERNELW_OFFSET = 12;\n",
"const static uint8_t KERNELH_OFFSET = 20;\n",
"const static uint8_t KERNELW_HIGHEST_BIT_OFFSET = 36;\n",
"const static uint8_t KERNELH_HIGHEST_BIT_OFFSET = 37;\n",
"const static uint8_t DILATIONW_OFFSET = 28;\n",
"const static uint8_t DILATIONH_OFFSET = 36;\n",
"const static uint8_t CIN_OFFSET = 48;\n",
"\n",
"const static uint64_t MASK_16 = 0xffff;\n",
"const static uint64_t MASK_8 = 0xff;\n",
"const static uint64_t MASK_6 = 0x3f;\n",
"const static uint64_t NINTH_BIT_MASK = 0x100;\n",
"\n",
"const static uint64_t DEQ_SCALAR_ONE = 1065353216;\n",
"static constexpr IsResetLoad3dConfig CONV_LOAD3DV2_DEFAULT_CONFIG = {false, false};\n",
"\n",
"#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 5102)\n",
"#define ASCEND_IS_AIC_CONV constexpr(true)\n",
"#define ASCEND_IS_AIV_CONV constexpr(true)\n",
"#else\n",
"#define ASCEND_IS_AIC_CONV ASCEND_IS_AIC\n",
"#define ASCEND_IS_AIV_CONV ASCEND_IS_AIV\n",
"#endif\n",
"\n",
"constexpr uint8_t UNIT_FLAG_ENABLE_ONLY = 2;\n",
"constexpr uint8_t UNIT_FLAG_ENABLE_WITH_FLIP = 3;\n",
"\n",
"static __aicore__ inline uint64_t AlignB(uint64_t a, uint64_t b)\n",
"{\n",
" return ((a + b - 1) / b) * b;\n",
"}\n",
"\n",
"static __aicore__ inline uint64_t CeilDiv(uint64_t a, uint64_t b)\n",
"{\n",
" return (a + b - 1) / b;\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.3 数据类型定义\n",
"\n",
"本节定义三个核心类型：\n",
"\n",
"- **ConvFormat**：枚举卷积数据布局，当前仅支持 NCHW。\n",
"- **ConvType**：将存储位置（TPosition）、数据布局（ConvFormat）和元素类型（TYPE）打包为一个模板类型标签，供 `Conv2dContext` 和 `Conv2dBase` 使用。\n",
"- **Conv2dContext**：保存单次卷积迭代所需的全部状态，包括片上缓冲区句柄（al0Buf/bl0Buf/l0cBuf）、队列（queueAL1/queueBL1）、GlobalTensor 指针以及从 Tiling 派生的循环控制变量。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"enum class ConvFormat : std::uint8_t {\n",
" NCHW = 1,\n",
"};\n",
"\n",
"template <TPosition POSITION, ConvFormat FORMAT, typename TYPE>\n",
"struct ConvType {\n",
" constexpr static TPosition pos = POSITION;\n",
" constexpr static ConvFormat format = FORMAT;\n",
" using T = TYPE;\n",
"};\n",
"\n",
"template <class FMAP_TYPE, class WEIGHT_TYPE, class OUTPUT_TYPE>\n",
"struct Conv2dContext {\n",
" using FmapT = typename FMAP_TYPE::T;\n",
" using WeightT = typename WEIGHT_TYPE::T;\n",
" using OutputT = typename OUTPUT_TYPE::T;\n",
" using L0cT = float;\n",
"\n",
" constexpr static uint64_t k0 = C0_SIZE / sizeof(WeightT);\n",
"\n",
" TPipe pipe;\n",
" GlobalTensor<FmapT> agm;\n",
" GlobalTensor<WeightT> bgm;\n",
" TBuf<TPosition::A2> al0Buf;\n",
" TBuf<TPosition::B2> bl0Buf;\n",
" TBuf<TPosition::CO1> l0cBuf;\n",
"\n",
" LocalTensor<FmapT> al1;\n",
" LocalTensor<WeightT> bl1;\n",
" LocalTensor<FmapT> wholeAl0Tensor;\n",
" LocalTensor<FmapT> al0;\n",
" LocalTensor<WeightT> wholeBl0Tensor;\n",
" LocalTensor<WeightT> bl0;\n",
" LocalTensor<L0cT> wholeCl0Tensor = LocalTensor<L0cT>(TPosition::CO1, 0, 0);\n",
" LocalTensor<L0cT> cl0 = LocalTensor<L0cT>(TPosition::CO1, 0, 0);\n",
"\n",
" TQue<QuePosition::A1, 1> queueAL1;\n",
" TQue<QuePosition::B1, 1> queueBL1;\n",
"\n",
" const Conv2DTilingData* __restrict convTiling = nullptr;\n",
"\n",
" uint64_t fmapOneBatchSize = 0;\n",
" uint64_t outputOneBatchSize = 0;\n",
" uint64_t dilatedKernelH = 0;\n",
" uint64_t dilatedKernelW = 0;\n",
" uint64_t orgCi = 0;\n",
" uint64_t orgCo = 0;\n",
" uint64_t orgHi = 0;\n",
" uint64_t orgWi = 0;\n",
" uint64_t orgHo = 0;\n",
" uint64_t orgWo = 0;\n",
" uint64_t kernelH = 0;\n",
" uint64_t kernelW = 0;\n",
" int64_t hiStartPos = -1;\n",
" int64_t wiStartPos = 0;\n",
" uint64_t singleCoreBatch = 0;\n",
" uint64_t singleCoreHo = 0;\n",
" uint64_t singleCoreWo = 0;\n",
" uint64_t singleCoreCi = 0;\n",
" uint64_t singleCoreCo = 0;\n",
"\n",
" uint64_t currentHoL0 = 0;\n",
" uint64_t currentWoL0 = 0;\n",
" uint64_t currentNBL1 = 0;\n",
" uint64_t currentML0Align = 0;\n",
" uint64_t currentNL0Align = 0;\n",
"\n",
" uint64_t ddr2l0LoopK = 0;\n",
" uint64_t maxKL0Iter = 0;\n",
" uint64_t kL0Tail = 0;\n",
" uint64_t kIter = 0;\n",
" uint64_t kAL0Iter = 0;\n",
" uint64_t kBL0Iter = 0;\n",
"\n",
" uint32_t bL1SpaceSize = 0;\n",
"};\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.4 上下文初始化 InitContext\n",
"\n",
"`InitContext` 从 Tiling 结构体中读取所有形状参数，初始化片上缓冲区（L0A/L0B/L0C）和 L1 队列，并预计算 K 维度的循环次数（`ddr2l0LoopK`）及尾块大小（`kL0Tail`）。\n",
"\n",
"关键逻辑：\n",
"- `dilatedKernelH/W`：考虑膨胀后的有效卷积核尺寸。\n",
"- `ddr2l0LoopK = CeilDiv(totalK, kL0)`：K 维度需要多少次 L0 迭代。\n",
"- `kL0Tail`：最后一次迭代的 K 块大小（若整除则等于 kL0）。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"template <class Ctx>\n",
"__aicore__ inline void InitContext(Ctx &ctx, const Conv2DTilingData *tiling)\n",
"{\n",
" ctx.convTiling = tiling;\n",
" ctx.singleCoreBatch = tiling->singleCoreBatch;\n",
"\n",
" if ASCEND_IS_AIC_CONV {\n",
" ctx.orgCi = tiling->orgCi;\n",
" ctx.orgHi = tiling->orgHi;\n",
" ctx.orgWi = tiling->orgWi;\n",
" ctx.orgCo = tiling->orgCo;\n",
" ctx.orgHo = tiling->orgHo;\n",
" ctx.orgWo = tiling->orgWo;\n",
" ctx.kernelH = tiling->kernelH;\n",
" ctx.kernelW = tiling->kernelW;\n",
" ctx.dilatedKernelH = 1 + (ctx.kernelH - 1) * tiling->dilationH;\n",
" ctx.dilatedKernelW = 1 + (ctx.kernelW - 1) * tiling->dilationW;\n",
" ctx.singleCoreCi = tiling->singleCoreCi;\n",
" ctx.singleCoreCo = tiling->singleCoreCo;\n",
" ctx.fmapOneBatchSize = tiling->orgCi * tiling->orgHi * tiling->orgWi;\n",
" ctx.outputOneBatchSize = tiling->orgCo * tiling->orgHo * tiling->orgWo;\n",
"\n",
" ctx.singleCoreHo = tiling->singleCoreHo;\n",
" ctx.singleCoreWo = tiling->singleCoreWo;\n",
" ctx.currentHoL0 = ctx.singleCoreHo;\n",
" ctx.currentWoL0 = ctx.singleCoreWo;\n",
" ctx.currentNBL1 = tiling->nBL1;\n",
"\n",
" ctx.pipe.InitBuffer(ctx.al0Buf, L0A_SIZE);\n",
" ctx.pipe.InitBuffer(ctx.bl0Buf, L0B_SIZE);\n",
" ctx.pipe.InitBuffer(ctx.l0cBuf, L0C_SIZE);\n",
" ctx.wholeAl0Tensor = ctx.al0Buf.template Get<typename Ctx::FmapT>();\n",
" ctx.wholeBl0Tensor = ctx.bl0Buf.template Get<typename Ctx::WeightT>();\n",
" ctx.wholeCl0Tensor = ctx.l0cBuf.template Get<typename Ctx::L0cT>();\n",
" ctx.bL1SpaceSize = tiling->nBL1 * tiling->kBL1;\n",
" ctx.pipe.InitBuffer(ctx.queueAL1, 1, tiling->aL1SpaceSize);\n",
" ctx.pipe.InitBuffer(ctx.queueBL1, 1, ctx.bL1SpaceSize * sizeof(typename Ctx::WeightT));\n",
"\n",
" uint64_t totalK = AlignB(ctx.singleCoreCi, Ctx::k0) * tiling->kernelHxkernelW;\n",
" ctx.ddr2l0LoopK = CeilDiv(totalK, tiling->kL0);\n",
" ctx.maxKL0Iter = ctx.ddr2l0LoopK - 1;\n",
" ctx.kL0Tail = totalK % tiling->kL0;\n",
" ctx.kL0Tail = ctx.kL0Tail == 0 ? tiling->kL0 : ctx.kL0Tail;\n",
"\n",
" ctx.currentML0Align = tiling->mStep;\n",
" ctx.currentNL0Align = tiling->nL0;\n",
" }\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.5 数据加载：LoadAL1 / LoadBL1\n",
"\n",
"**LoadAL1**：将特征图的一个输入 patch 从 GM 搬运到 L1（`queueAL1`）。\n",
"\n",
"核心步骤：\n",
"1. 根据当前输出 tile 的位置（`hiStartPos`/`wiStartPos`）计算需要加载的输入行列范围，并推导出四个方向的 padding 量。\n",
"2. 若整个 patch 都落在 padding 区域（`allPad`），则用 `InitConstValue` 填零，避免越界访问。\n",
"3. 否则，使用 `Dn2NzParams` 描述 NZ 布局的搬运参数，调用 `DataCopy` 完成实际数据搬运。\n",
"\n",
"**LoadBL1**：将权重从 GM 搬运到 L1（`queueBL1`）。\n",
"\n",
"- 1×1 卷积（`kernelHxkernelW == 1`）使用 `Nd2NzParams`（行优先到 NZ 的转换）。\n",
"- 非 1×1 卷积使用 `Dn2NzParams`，按 `coutOffsetBlock` 步长遍历输出通道块。\n",
"\n",
"#### DataCopy详解\n",
"\n",
"**功能定位**：DataCopy是通用数据搬运接口，支持多种数据搬运场景，并可在搬运过程中实现随路格式转换和量化激活等操作。该接口支持Local Memory与Global Memory之间的数据搬运，以及Local Memory内部的数据搬运，不涉及img2col变换。\n",
"\n",
"**核心参数**：\n",
"\n",
"```cpp\n",
"\n",
"DataCopyParams params;\n",
"params.blockCount = blockNum; // 搬运块数\n",
"params.blockLen = dataLength; // 每块数据长度\n",
"params.srcStride = srcStride; // 源步长\n",
"params.dstStride = dstStride; // 目的步长\n",
"// 带Padding的搬运\n",
"DataCopyPadParams padParams;\n",
"padParams.isPad = true;\n",
"padParams.leftPadding = leftPad;\n",
"padParams.rightPadding = rightPad;\n",
"\n",
"```\n",
"\n",
"**代码示例**：\n",
"\n",
"```cpp\n",
"\n",
"DataCopyParams dataCopyParams;\n",
"dataCopyParams.blockCount = cin1LoadL1;\n",
"dataCopyParams.blockLen = hiLoadL1 * orgWi;\n",
"dataCopyParams.srcStride = orgHixWi - blockLen;\n",
"DataCopy<FmapT>(al1[offset], agm[gmOffset], dataCopyParams);\n",
"// 带Padding搬运\n",
"DataCopyPadParams padParams(true, 0, rightPadding, 0);\n",
"DataCopyPad<ChannelWiseT>(tensorL1, tensorGm[offset], dataCopyParams, padParams);\n",
"\n",
"```\n",
"\n",
"**格式转换的必要性**：\n",
"\n",
"- **硬件适配**：Cube单元要求输入为分形矩阵格式（如NZ格式），而非NCHW/NHWC的连续存储格式\n",
"- **性能优化**：分形格式能够更好地利用矩阵乘法单元的计算能力，提高内存访问效率\n",
"- **计算流水**：格式转换与数据搬运融合，减少额外的转换开销\n",
"\n",
"**常用格式转换参数结构体**：\n",
"在卷积算子实现中，DataCopy不仅完成数据搬运，还承担着格式转换的关键角色。昇腾AI处理器的Cube单元要求输入数据为特定的分形格式（Fractal Format），而主流框架（如PyTorch、TensorFlow）通常采用NCHW或NHWC格式。因此，需要在数据从GM搬运至L1的过程中完成格式转换，以适配硬件计算单元的输入要求。为支持不同格式的数据转换，Ascend C提供了专门的参数结构体。以下介绍两种核心格式转换参数：Dn2NzParams和Nd2NzParams。\n",
"\n",
"#### Dn2NzParams详解\n",
"\n",
"**功能**：将NCHW格式转换为NZ格式。\n",
"\n",
"**Dn2NzParams结构体参数定义**：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>参数名称</strong></td>\n",
" <td align=\"center\"><strong>含义</strong></td>\n",
" <td align=\"center\" style=\"width: 180px\"><strong>取值范围</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dnNum</td>\n",
" <td align=\"left\">传输DN矩阵的数目。</td>\n",
" <td align=\"left\">[0, 4095]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">nValue</td>\n",
" <td align=\"left\">DN矩阵的行数（矩阵高度）。</td>\n",
" <td align=\"left\">[0, 16384]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dValue</td>\n",
" <td align=\"left\">DN矩阵的列数（矩阵宽度）。当dValue不满足32B对齐时，目的操作数中不足部分会被补齐为0。</td>\n",
" <td align=\"left\">[0, 2^32-1]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">srcDnMatrixStride</td>\n",
" <td align=\"left\">源操作数相邻DN矩阵起始地址间的偏移，单位：元素。</td>\n",
" <td align=\"left\">[0, 2^64-1]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">srcDValue</td>\n",
" <td align=\"left\">源操作数同一DN矩阵的相邻行起始地址间的偏移，单位：元素。</td>\n",
" <td align=\"left\">[1, 2^64-1]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstNzC0Stride</td>\n",
" <td align=\"left\">DN转换到NZ格式后，源操作数中的一列会转换为目的操作数的多行。dstNzC0Stride表示目的NZ矩阵中，来自源操作数同一列的多行数据相邻行起始地址间的偏移，单位：C0_SIZE（32B）。</td>\n",
" <td align=\"left\">[1, 65535]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstNzNStride</td>\n",
" <td align=\"left\">目的NZ矩阵中，Z型矩阵相邻行起始地址之间的偏移，单位：C0_SIZE（32B）。</td>\n",
" <td align=\"left\">[1, 65535]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstNzMatrixStride</td>\n",
" <td align=\"left\">目的NZ矩阵中，相邻NZ矩阵起始地址间的偏移，单位：元素。</td>\n",
" <td align=\"left\">[1, 2^32-1]</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"#### Nd2NzParams详解\n",
"\n",
"**功能**：将NHWC格式转换为NZ格式。\n",
"\n",
"**Nd2NzParams结构体参数定义**：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>参数名称</strong></td>\n",
" <td align=\"center\"><strong>含义</strong></td>\n",
" <td align=\"center\" style=\"width: 180px\"><strong>取值范围</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">ndNum</td>\n",
" <td align=\"left\">传输ND矩阵的数目。</td>\n",
" <td align=\"left\">[0, 4095]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">nValue</td>\n",
" <td align=\"left\">ND矩阵的行数（矩阵高度）。</td>\n",
" <td align=\"left\">[0, 16384]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dValue</td>\n",
" <td align=\"left\">ND矩阵的列数（矩阵宽度）。当dValue不满足32B对齐时，目的操作数中不足部分会被补齐为0。</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">srcNdMatrixStride</td>\n",
" <td align=\"left\">源操作数相邻ND矩阵起始地址间的偏移，单位：元素。</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">srcDValue</td>\n",
" <td align=\"left\">源操作数同一ND矩阵的相邻行起始地址间的偏移，单位：元素。</td>\n",
" <td align=\"left\">[1, 65535]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstNzC0Stride</td>\n",
" <td align=\"left\">ND转换到NZ格式后，源操作数中的一行会转换为目的操作数的多行。dstNzC0Stride表示目的NZ矩阵中，来自源操作数同一行的多行数据相邻行起始地址间的偏移，单位：C0_SIZE（32B）。</td>\n",
" <td align=\"left\">[1, 16384]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstNzNStride</td>\n",
" <td align=\"left\">目的NZ矩阵中，Z型矩阵相邻行起始地址之间的偏移，单位：C0_SIZE（32B）。</td>\n",
" <td align=\"left\">[1, 16384]</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstNzMatrixStride</td>\n",
" <td align=\"left\">目的NZ矩阵中，相邻NZ矩阵起始地址间的偏移，单位：元素。</td>\n",
" <td align=\"left\">[1, 65535]</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"详细的ND2NZ转换示意图请参考[昇腾社区官方文档](https://www.hiascend.com/document/detail/zh/canncommercial/850/API/ascendcopapi/atlasascendc_api_07_00127.html)中的\"图1 ND2NZ转换示意图（half数据类型）\"部分。该图清晰地展示了从ND格式到NZ格式的转换过程，包括源操作数和目的操作数的数据布局、各参数对应的偏移关系以及DataBlock的对齐方式。\n",
"\n",
"**约束说明**：\n",
"\n",
"- 针对Atlas 推理系列产品AI Core，使用Global Memory → Local Memory通路的ND2NZ搬运接口时，需要预留8K的UB空间，作为接口的临时数据存放区。\n",
"\n",
"#### 格式转换对比\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\"><strong>参数类型</strong></td>\n",
" <td align=\"center\"><strong>源格式</strong></td>\n",
" <td align=\"center\"><strong>目标格式</strong></td>\n",
" <td align=\"center\"><strong>适用框架</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">Dn2NzParams</td>\n",
" <td align=\"left\">NCHW</td>\n",
" <td align=\"left\">NZ</td>\n",
" <td align=\"left\">PyTorch、ONNX、Caffe</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">Nd2NzParams</td>\n",
" <td align=\"left\">NHWC</td>\n",
" <td align=\"left\">NZ</td>\n",
" <td align=\"left\">TensorFlow、推理引擎</td>\n",
"</tr>\n",
"</table>\n",
"\n"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"template <class Ctx>\n",
"__aicore__ inline void LoadAL1(Ctx &ctx)\n",
"{\n",
" ctx.al1 = ctx.queueAL1.template AllocTensor<typename Ctx::FmapT>();\n",
"\n",
" const auto *t = ctx.convTiling;\n",
" int64_t paddingHi = (ctx.singleCoreHo - 1) * t->strideH + ctx.dilatedKernelH;\n",
" int64_t paddingWi = (ctx.singleCoreWo - 1) * t->strideW + ctx.dilatedKernelW;\n",
"\n",
" int64_t hiTop = ctx.hiStartPos > 0 ? 0 : ctx.hiStartPos;\n",
" int64_t wiLeft = ctx.wiStartPos > 0 ? 0 : ctx.wiStartPos;\n",
" int64_t hiBottom = hiTop + paddingHi;\n",
" int64_t wiRight = wiLeft + paddingWi;\n",
"\n",
" uint64_t padTopL1 = hiTop < 0 ? uint64_t(-hiTop) : 0;\n",
" hiBottom = ctx.hiStartPos > 0 ? hiBottom + ctx.hiStartPos : hiBottom;\n",
" uint64_t padBottomL1 = hiBottom > (int64_t)ctx.orgHi ? uint64_t(hiBottom - ctx.orgHi) : 0;\n",
" uint64_t padLeftL1 = wiLeft < 0 ? uint64_t(-wiLeft) : 0;\n",
" wiRight = ctx.wiStartPos > 0 ? wiRight + ctx.wiStartPos : wiRight;\n",
" uint64_t padRightL1 = wiRight > (int64_t)ctx.orgWi ? uint64_t(wiRight - ctx.orgWi) : 0;\n",
"\n",
" bool allPad = (padTopL1 >= (uint64_t)paddingHi || padBottomL1 >= (uint64_t)paddingHi ||\n",
" padLeftL1 >= (uint64_t)paddingWi || padRightL1 >= (uint64_t)paddingWi);\n",
" uint64_t hiLoad = 0, wiLoad = 0;\n",
" if (!allPad) {\n",
" hiLoad = paddingHi - padTopL1 - padBottomL1;\n",
" hiLoad = hiLoad > ctx.orgHi ? ctx.orgHi : hiLoad;\n",
" wiLoad = paddingWi - padLeftL1 - padRightL1;\n",
" }\n",
"\n",
" uint8_t padList[PAD_SIZE] = {MAX_PAD_R, MAX_PAD_R, MAX_PAD_R, MAX_PAD_R};\n",
" if (unlikely(allPad)) {\n",
" Load3DSetFMatrixCal(MIN_HI_WI, MIN_HI_WI, padList);\n",
" } else {\n",
" padList[PAD_IDX_L] = padLeftL1;\n",
" padList[PAD_IDX_R] = padRightL1;\n",
" padList[PAD_IDX_T] = padTopL1;\n",
" Load3DSetFMatrixCal(hiLoad, wiLoad, padList);\n",
" }\n",
" Load3DSetPaddingCal(t->offsetx);\n",
"\n",
" if (allPad) {\n",
" InitConstValueParams<typename Ctx::FmapT> params(\n",
" 1, static_cast<uint16_t>(t->aL1SpaceSize / C0_SIZE), 0, 0);\n",
" InitConstValue<typename Ctx::FmapT>(ctx.al1, params);\n",
" ctx.queueAL1.EnQue(ctx.al1);\n",
" return;\n",
" }\n",
"\n",
" int64_t realHiTopGm = hiTop < 0 ? 0 : hiTop;\n",
" int64_t realWiTopGm = wiLeft < 0 ? 0 : wiLeft;\n",
" int64_t aL1GmOffset = realHiTopGm * ctx.orgWi + realWiTopGm;\n",
"\n",
" Dn2NzParams params;\n",
" if (likely(wiLoad == ctx.orgWi)) {\n",
" params.dnNum = 1;\n",
" params.nValue = hiLoad * wiLoad;\n",
" params.srcDnMatrixStride = 0;\n",
" params.dstNzMatrixStride = 0;\n",
" } else {\n",
" params.dnNum = hiLoad;\n",
" params.nValue = wiLoad;\n",
" params.srcDnMatrixStride = ctx.orgWi;\n",
" params.dstNzMatrixStride = wiLoad * Ctx::k0;\n",
" }\n",
" params.dValue = t->cinAInCore;\n",
" params.srcDValue = t->orgHixWi;\n",
" params.dstNzC0Stride = hiLoad * wiLoad;\n",
" params.dstNzNStride = 1;\n",
" DataCopy<typename Ctx::FmapT>(ctx.al1, ctx.agm[aL1GmOffset], params);\n",
" ctx.queueAL1.EnQue(ctx.al1);\n",
"}\n",
"\n",
"template <class Ctx>\n",
"__aicore__ inline void LoadBL1(Ctx &ctx)\n",
"{\n",
" ctx.bl1 = ctx.queueBL1.template AllocTensor<typename Ctx::WeightT>();\n",
"\n",
" const auto *t = ctx.convTiling;\n",
" bool oneByOne = (t->kernelHxkernelW == 1);\n",
" if (oneByOne) {\n",
" Nd2NzParams p;\n",
" p.ndNum = 1;\n",
" p.nValue = ctx.currentNBL1;\n",
" p.dValue = t->kBL1;\n",
" p.srcNdMatrixStride = 0;\n",
" p.srcDValue = ctx.orgCi;\n",
" p.dstNzC0Stride = t->nBL1;\n",
" p.dstNzNStride = 1;\n",
" p.dstNzMatrixStride = 0;\n",
" DataCopy<typename Ctx::WeightT>(ctx.bl1, ctx.bgm[0], p);\n",
" } else {\n",
" Dn2NzParams p;\n",
" p.dnNum = ctx.currentNBL1;\n",
" p.nValue = t->kernelHxkernelW;\n",
" p.dValue = t->cinBInCore;\n",
" p.srcDnMatrixStride = t->coutOffsetBlock;\n",
" p.srcDValue = t->kernelHxkernelW;\n",
" p.dstNzC0Stride = t->kernelHxkernelW * t->nBL1;\n",
" p.dstNzNStride = t->nBL1;\n",
" p.dstNzMatrixStride = Ctx::k0;\n",
" DataCopy<typename Ctx::WeightT>(ctx.bl1, ctx.bgm[0], p);\n",
" }\n",
" ctx.queueBL1.EnQue(ctx.bl1);\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.6 im2col：LoadAL0 / LoadBL0\n",
"\n",
"**LoadAL0**：将 L1 中的特征图 patch 通过 `Load3D` 指令在线展开为 im2col 矩阵，写入 L0A。\n",
"\n",
"`Load3D` 是 NPU 的专用硬件指令，能在数据搬运的同时完成 im2col 展开，无需额外的软件循环。第一次调用（`isFirst == true`）需要设置完整的卷积参数（步长、卷积核尺寸、膨胀系数、通道数），后续迭代只需更新 K 偏移量。\n",
"\n",
"#### Load3D详解\n",
"\n",
"**功能说明**：Load3D用于完成image to column操作，将多维feature map转为二维矩阵。支持如下数据通路：A1→A2; B1→B2。\n",
"\n",
"**Load3Dv2接口**：\n",
"\n",
"```cpp\n",
"\n",
"template <typename T, \n",
" const IsResetLoad3dConfig &defaultConfig = IS_RESER_LOAD3D_DEFAULT_CONFIG,\n",
" typename U = PrimT<T>,\n",
" typename Std::enable_if<Std::is_same<PrimT<T>, U>::value, bool>::type = true>\n",
"__aicore__ inline void LoadData(const LocalTensor<T>& dst, \n",
" const LocalTensor<T>& src, \n",
" const LoadData3DParamsV2<U>& loadDataParams)\n",
"\n",
"```\n",
"**参数说明**：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\"><strong>参数名称</strong></td>\n",
" <td align=\"center\"><strong>输入/输出</strong></td>\n",
" <td align=\"center\"><strong>含义</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dst</td>\n",
" <td align=\"left\">输出</td>\n",
" <td align=\"left\">目的操作数(LocalTensor)</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">src</td>\n",
" <td align=\"left\">输入</td>\n",
" <td align=\"left\">源操作数(LocalTensor/GlobalTensor)，类型需与dst保持一致</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">loadDataParams</td>\n",
" <td align=\"left\">输入</td>\n",
" <td align=\"left\">LoadData参数结构体(LoadData3DParamsV2)</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"**LoadData3DParamsV2 结构体参数**：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>参数名称</strong></td>\n",
" <td align=\"center\"><strong>含义</strong></td>\n",
" <td align=\"center\" style=\"width: 180px\"><strong>取值范围</strong></td>\n",
" <td align=\"center\" style=\"width: 80px\"><strong>默认值</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">padList</td>\n",
" <td align=\"left\">padding列表</td>\n",
" <td align=\"left\">[0,255]</td>\n",
" <td align=\"left\">{0,0,0,0}</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">l1H</td>\n",
" <td align=\"left\">源操作数height</td>\n",
" <td align=\"left\">[1, 32767]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">l1W</td>\n",
" <td align=\"left\">源操作数width</td>\n",
" <td align=\"left\">[1, 32767]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">channelSize</td>\n",
" <td align=\"left\">源操作数的通道数</td>\n",
" <td align=\"left\">[1, 63]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">kExtension</td>\n",
" <td align=\"left\">目的操作数width维度传输长度</td>\n",
" <td align=\"left\">[1, 65535]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">mExtension</td>\n",
" <td align=\"left\">目的操作数height维度传输长度</td>\n",
" <td align=\"left\">[1, 65535]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">kStartPt</td>\n",
" <td align=\"left\">目的操作数width维度起点</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
" <td align=\"left\">0</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">mStartPt</td>\n",
" <td align=\"left\">目的操作数height维度起点</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
" <td align=\"left\">0</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">strideW</td>\n",
" <td align=\"left\">卷积核w维度滑动步长</td>\n",
" <td align=\"left\">[1, 63]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">strideH</td>\n",
" <td align=\"left\">卷积核h维度滑动步长</td>\n",
" <td align=\"left\">[1, 63]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">filterW</td>\n",
" <td align=\"left\">卷积核width</td>\n",
" <td align=\"left\">[1, 255]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">filterH</td>\n",
" <td align=\"left\">卷积核height</td>\n",
" <td align=\"left\">[1, 255]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dilationFilterW</td>\n",
" <td align=\"left\">卷积核width膨胀系数</td>\n",
" <td align=\"left\">[1, 255]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dilationFilterH</td>\n",
" <td align=\"left\">卷积核height膨胀系数</td>\n",
" <td align=\"left\">[1, 255]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">enTranspose</td>\n",
" <td align=\"left\">是否启用转置功能(仅A2位置，half类型有效)</td>\n",
" <td align=\"left\">true/false</td>\n",
" <td align=\"left\">false</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">enSmallK</td>\n",
" <td align=\"left\">是否使能small k特性(已不再支持)</td>\n",
" <td align=\"left\">true/false</td>\n",
" <td align=\"left\">false</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">padValue</td>\n",
" <td align=\"left\">Pad填充值数值</td>\n",
" <td align=\"left\">与src数据类型一致</td>\n",
" <td align=\"left\">0</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">filterSizeW</td>\n",
" <td align=\"left\">是否在filterW基础上增加256个元素</td>\n",
" <td align=\"left\">true/false</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">filterSizeH</td>\n",
" <td align=\"left\">是否在filterH基础上增加256个元素</td>\n",
" <td align=\"left\">true/false</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">fMatrixCtrl</td>\n",
" <td align=\"left\">从左矩阵还是右矩阵获取FeatureMap属性描述(当前只支持false)</td>\n",
" <td align=\"left\">true/false</td>\n",
" <td align=\"left\">false</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"**LoadBL0**：将 L1 中的权重通过 `Load2D` 指令搬运到 L0B，转换为 NZ 布局。\n",
"\n",
"第一次调用需要设置 N/K 步长和源矩阵步长，后续迭代只需更新 K 起始位置。\n",
"\n",
"#### Load2D详解\n",
"\n",
"**功能定位**：Load2D用于将权重数据从L1或GM搬运至L0B，支持矩阵转置和分块搬运，适配Cube单元的输入格式要求。\n",
"\n",
"**Load2D接口**：\n",
"\n",
"```cpp\n",
"\n",
"template <typename T>\n",
"__aicore__ inline void LoadData(const LocalTensor<T>& dst, const LocalTensor<T>& src, const LoadData2DParams& loadDataParams)\n",
"\n",
"```\n",
"\n",
"**通用参数说明**：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>参数名称</strong></td>\n",
" <td align=\"center\" style=\"width: 100px\"><strong>输入/输出</strong></td>\n",
" <td align=\"center\"><strong>含义</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dst</td>\n",
" <td align=\"left\">输出</td>\n",
" <td align=\"left\">目的操作数，类型为LocalTensor。数据连续排列顺序由目的操作数所在TPosition决定：<br>- A2：ZZ格式，分形大小为16 * (32B / sizeof(T))<br>- B2：ZN格式，分形大小为(32B / sizeof(T)) * 16<br>- A1/B1：无格式要求，一般为NZ格式，分形大小为16 * (32B / sizeof(T))</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">src</td>\n",
" <td align=\"left\">输入</td>\n",
" <td align=\"left\">源操作数，类型为LocalTensor或GlobalTensor。数据类型需与dst保持一致。</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">loadDataParams</td>\n",
" <td align=\"left\">输入</td>\n",
" <td align=\"left\">LoadData参数结构体，类型为LoadData2DParams。</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"**LoadData2DParams结构体内参数说明**：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>参数名称</strong></td>\n",
" <td align=\"center\"><strong>含义</strong></td>\n",
" <td align=\"center\" style=\"width: 180px\"><strong>取值范围</strong></td>\n",
" <td align=\"center\" style=\"width: 80px\"><strong>默认值</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">startIndex</td>\n",
" <td align=\"left\">分形矩阵ID，说明搬运起始位置为源操作数中第几个分形（0为第1个分形矩阵）。单位：512B。</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
" <td align=\"left\">0</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">repeatTimes</td>\n",
" <td align=\"left\">迭代次数，每个迭代可以处理512B数据。</td>\n",
" <td align=\"left\">[1, 255]</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">srcStride</td>\n",
" <td align=\"left\">相邻迭代间，源操作数前一个分形与后一个分形起始地址的间隔。单位：512B。</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
" <td align=\"left\">0</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">sid</td>\n",
" <td align=\"left\">预留参数，配置为0即可。</td>\n",
" <td align=\"left\">-</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">dstGap</td>\n",
" <td align=\"left\">相邻迭代间，目的操作数前一个分形结束地址与后一个分形起始地址的间隔。单位：512B。<br><strong>注</strong>：Atlas 训练系列产品此参数不使能。</td>\n",
" <td align=\"left\">[0, 65535]</td>\n",
" <td align=\"left\">0</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">ifTranspose</td>\n",
" <td align=\"left\">是否启用转置功能，对每个分形矩阵进行转置。<br><strong>注意</strong>：只有A1->A2和B1->B2通路才能使能转置，使能转置时仅支持uint16_t/int16_t/half数据类型。</td>\n",
" <td align=\"left\">true/false</td>\n",
" <td align=\"left\">false</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">addrMode</td>\n",
" <td align=\"left\">预留参数，配置为0即可。</td>\n",
" <td align=\"left\">-</td>\n",
" <td align=\"left\">-</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"**约束说明**：\n",
"\n",
"- 操作数地址对齐要求请参见通用地址对齐约束。\n",
"- repeatTimes最大值为255，当搬运数据量超过255块时需分批搬运。\n",
"\n",
"**repeatTimes限制处理方案**：\n",
"\n",
"当搬运数据量超过255块时，需要分批搬运：\n",
"\n",
"```cpp\n",
"\n",
"if (repeatTimes > LOAD2D_MAX_REPEAT_TIMES) {\n",
" // 方案一：使用DataCopy替代\n",
" DataCopyParams dataCopyParams(1, blockLen, 0, 0);\n",
" DataCopy<WeightT>(bl1, bgm[offset], dataCopyParams);\n",
" \n",
" // 方案二：分批LoadData2D\n",
" for (i = 0; i < repeatTimes / 255; i++) {\n",
" LoadData<WeightT>(bl1[offset1], bgm[offset2], params255);\n",
" }\n",
" // 处理尾部\n",
" if (tail > 0) {\n",
" LoadData<WeightT>(bl1[offset1], bgm[offset2], paramsTail);\n",
" }\n",
"}\n",
"\n",
"```\n"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"template <class Ctx>\n",
"__aicore__ inline void LoadAL0(Ctx &ctx, bool isFirst)\n",
"{\n",
" if ASCEND_IS_AIV {\n",
" return;\n",
" }\n",
" const auto *t = ctx.convTiling;\n",
" uint64_t currentKL0 = (ctx.kIter == ctx.maxKL0Iter) ? ctx.kL0Tail : t->kL0;\n",
" uint64_t posK = ctx.kAL0Iter * t->kL0;\n",
" uint64_t channelSize = t->cinAInCore;\n",
" uint64_t currentML0 = ctx.currentHoL0 * ctx.currentWoL0;\n",
" uint16_t mExtension = static_cast<uint16_t>(currentML0);\n",
" uint16_t mStartPt = 0;\n",
"\n",
" Load3DBitModeParam param;\n",
" if (isFirst) {\n",
" uint64_t xmtmp = ((mExtension & MASK_16) << MSTEP_OFFSET) |\n",
" ((uint64_t)(mStartPt & MASK_16) << POSM_OFFSET);\n",
" uint64_t xt = ((uint64_t(t->strideW) & MASK_6) << 0) |\n",
" ((uint64_t(t->strideH) & MASK_6) << STRIDEH_OFFSET) |\n",
" ((uint64_t(ctx.kernelW) & MASK_8) << KERNELW_OFFSET) |\n",
" ((uint64_t(ctx.kernelW) & NINTH_BIT_MASK) << KERNELW_HIGHEST_BIT_OFFSET) |\n",
" ((uint64_t(ctx.kernelH) & MASK_8) << KERNELH_OFFSET) |\n",
" ((uint64_t(ctx.kernelH) & NINTH_BIT_MASK) << KERNELH_HIGHEST_BIT_OFFSET) |\n",
" ((uint64_t(t->dilationW) & MASK_8) << DILATIONW_OFFSET) |\n",
" ((uint64_t(t->dilationH) & MASK_8) << DILATIONH_OFFSET) |\n",
" ((uint64_t(channelSize) & MASK_16) << CIN_OFFSET);\n",
" param.SetConfig1(xt);\n",
" uint64_t xm = ((currentKL0 & MASK_16) << 0) | ((posK & MASK_16) << POSK_OFFSET) | xmtmp;\n",
" param.SetConfig0(xm);\n",
"\n",
" uint16_t dstStride =\n",
" (ctx.currentWoL0 == t->woL0 && ctx.currentHoL0 == t->hoL0)\n",
" ? static_cast<uint16_t>(t->fmapKStride)\n",
" : static_cast<uint16_t>(ctx.currentML0Align / BLOCK_L0_M);\n",
" LoadDataRepeatParam repeatParams = {0, 1, 0, dstStride};\n",
" SetLoadDataRepeat(repeatParams);\n",
" } else {\n",
" uint64_t xmtmp = ((mExtension & MASK_16) << MSTEP_OFFSET) |\n",
" ((uint64_t)(mStartPt & MASK_16) << POSM_OFFSET);\n",
" uint64_t xm = ((currentKL0 & MASK_16) << 0) | ((posK & MASK_16) << POSK_OFFSET) | xmtmp;\n",
" param.SetConfig0(xm);\n",
" }\n",
" LoadData<TPosition::A2, TPosition::A1, typename Ctx::FmapT>(ctx.al0, ctx.al1, param);\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.7 矩阵乘累加：Mad 与 CopyOut\n",
"\n",
"**Mad**：调用 `Mmad` 指令执行矩阵乘累加，将 L0A（特征图 im2col）× L0B（权重 NZ）的结果累加到 L0C（fp32 累加器）。\n",
"\n",
"关键参数：\n",
"- `cmatrixInitVal = (kIter == 0)`：第一次 K 迭代时清零 L0C，后续迭代累加。\n",
"- `unitFlag`：最后一次 K 迭代时设置 `UNIT_FLAG_ENABLE_WITH_FLIP`，触发 L0C 翻转（通知 Fixpipe 可以读取结果）。\n",
"\n",
"**CopyOut**：调用 `Fixpipe` 指令将 L0C 中的 fp32 结果转换为 fp16，并以 NCHW 布局写回 GM。\n",
"\n",
"`FixpipeParamsC310` 中的 `quantPre = F322F16` 指定了量化模式，`deqScalar = DEQ_SCALAR_ONE`（即浮点 1.0）表示不缩放。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"template <class Ctx>\n",
"__aicore__ inline void LoadBL0(Ctx &ctx, bool isFirst)\n",
"{\n",
" const auto *t = ctx.convTiling;\n",
" uint64_t kStep = (ctx.kIter != ctx.maxKL0Iter) ? t->kStep : (ctx.kL0Tail / Ctx::k0);\n",
" uint64_t ratioOfNToN0 = ctx.currentNL0Align / BLOCK_L0_N;\n",
"\n",
" Load2DBitModeParam param;\n",
" if (unlikely(isFirst)) {\n",
" param.SetMStartPosition(0);\n",
" param.SetKStartPosition(static_cast<uint32_t>(ctx.kBL0Iter * t->kStep));\n",
" param.SetMStep(static_cast<uint16_t>(ratioOfNToN0));\n",
" param.SetKStep(static_cast<uint16_t>(kStep));\n",
" param.SetSrcStride(static_cast<int32_t>(t->nL1DivBlockSize));\n",
" param.SetDstStride(static_cast<uint16_t>(ratioOfNToN0));\n",
" param.SetIfTranspose(false);\n",
" } else {\n",
" param.SetKStartPosition(static_cast<uint32_t>(ctx.kBL0Iter * t->kStep));\n",
" param.SetKStep(static_cast<uint16_t>(kStep));\n",
" }\n",
" LoadData<TPosition::B2, TPosition::B1, typename Ctx::WeightT>(ctx.bl0, ctx.bl1, param);\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.8 迭代控制：IterateOne 与 EndConv\n",
"\n",
"**IterateOne**：执行一次完整的卷积 tile 计算，包含以下流水步骤：\n",
"\n",
"1. 调用 `LoadAL1` / `LoadBL1` 将数据从 GM 搬运到 L1。\n",
"2. 进入 K 维度循环：每次迭代调用 `LoadAL0`（im2col）、`LoadBL0`（权重 NZ）、`Mad`（矩阵乘累加）。\n",
"3. 使用 `WaitFlag` / `SetFlag` 在 MTE1（数据搬运）和 M（矩阵计算）单元之间进行同步。\n",
"4. 调用 `CopyOut` 将 L0C 结果写回 GM。\n",
"\n",
"**EndConv**：释放 L1 队列中的所有事件资源，完成一次卷积的收尾工作。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"template <class Ctx>\n",
"__aicore__ inline void Mad(Ctx &ctx)\n",
"{\n",
" MmadParams p;\n",
" p.m = ctx.currentML0Align;\n",
" p.n = ctx.currentNL0Align;\n",
" p.k = (ctx.kIter == ctx.maxKL0Iter) ? ctx.kL0Tail : ctx.convTiling->kL0;\n",
" p.cmatrixInitVal = (ctx.kIter == 0);\n",
" p.cmatrixSource = false;\n",
" p.unitFlag = (ctx.kIter == ctx.maxKL0Iter) ? UNIT_FLAG_ENABLE_WITH_FLIP : UNIT_FLAG_ENABLE_ONLY;\n",
" Mmad<typename Ctx::L0cT, typename Ctx::FmapT, typename Ctx::WeightT>(\n",
" ctx.cl0, ctx.al0, ctx.bl0, p);\n",
"}\n",
"\n",
"template <class Ctx>\n",
"__aicore__ inline void CopyOut(Ctx &ctx, const GlobalTensor<typename Ctx::OutputT> &output)\n",
"{\n",
" uint64_t valueHoWo = ctx.orgHo * ctx.orgWo;\n",
" uint64_t offset = 0;\n",
"\n",
" FixpipeParamsC310<CO2Layout::COLUMN_MAJOR> p;\n",
" p.mSize = ctx.currentHoL0 * ctx.currentWoL0;\n",
" p.params.srcNzMatrixStride = 0;\n",
" p.params.dnNum = 1;\n",
" p.params.dstDnMatrixStride = 0;\n",
" p.nSize = ctx.currentNL0Align;\n",
" p.srcStride = AlignB(p.mSize, BLOCK_L0_M);\n",
" p.dstStride = valueHoWo;\n",
" p.params.srcNzC0Stride = 1;\n",
" p.quantPre = QuantMode_t::F322F16;\n",
" p.deqScalar = DEQ_SCALAR_ONE;\n",
" p.unitFlag = UNIT_FLAG_ENABLE_WITH_FLIP;\n",
"\n",
" Fixpipe<typename Ctx::OutputT, typename Ctx::L0cT, CFG_COLUMN_MAJOR>(\n",
" output[offset], ctx.cl0, p);\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 3.9 Conv2dBase 类与核函数入口\n",
"\n",
"**关闭 namespace conv**，然后定义三个 `constexpr` 变量指定特征图、权重和输出的数据布局（均为 NCHW）。\n",
"\n",
"**Conv2dBase**：模板类，封装了 `RunConv2dKernel` 方法。该方法：\n",
"1. 设置 GM 指针（特征图、权重、输出）。\n",
"2. 调用 `InitContext` 初始化上下文。\n",
"3. 调用 `IterateOne` 执行卷积计算。\n",
"4. 调用 `EndConv` 释放资源。\n",
"\n",
"**conv2d_v2_kernel**：全局核函数入口，使用 `GET_TILING_DATA` 宏从 GM 读取 Tiling 参数，实例化类型标签，创建 `Conv2dBase` 对象并调用 `RunConv2dKernel`。\n",
"\n",
"`#endif` 结束 `__CCE_AICORE__` 设备侧编译块。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"template <class Ctx>\n",
"__aicore__ inline void IterateOne(Ctx &ctx, const GlobalTensor<typename Ctx::OutputT> &output)\n",
"{\n",
" ctx.cl0 = ctx.wholeCl0Tensor;\n",
" ctx.kIter = 0;\n",
"\n",
" LoadAL1(ctx);\n",
" LoadBL1(ctx);\n",
" ctx.al1 = ctx.queueAL1.template DeQue<typename Ctx::FmapT>();\n",
" ctx.bl1 = ctx.queueBL1.template DeQue<typename Ctx::WeightT>();\n",
"\n",
" event_t ev = static_cast<event_t>(L0_SYBC_DB_CLOSE);\n",
" bool isFirst = true;\n",
" while (ctx.kIter < ctx.ddr2l0LoopK) {\n",
" AscendC::WaitFlag<AscendC::HardEvent::M_MTE1>(ev);\n",
" ctx.al0 = ctx.wholeAl0Tensor;\n",
" ctx.kAL0Iter = ctx.kIter;\n",
" LoadAL0(ctx, isFirst);\n",
" ctx.bl0 = ctx.wholeBl0Tensor;\n",
" ctx.kBL0Iter = ctx.kIter;\n",
" LoadBL0(ctx, isFirst);\n",
" AscendC::SetFlag<AscendC::HardEvent::MTE1_M>(ev);\n",
" AscendC::WaitFlag<AscendC::HardEvent::MTE1_M>(ev);\n",
" Mad(ctx);\n",
" AscendC::SetFlag<AscendC::HardEvent::M_MTE1>(ev);\n",
" ctx.kIter++;\n",
" isFirst = false;\n",
" }\n",
"\n",
" ctx.queueAL1.FreeTensor(ctx.al1);\n",
" ctx.queueBL1.FreeTensor(ctx.bl1);\n",
"\n",
" if ASCEND_IS_AIC_CONV {\n",
" CopyOut(ctx, output);\n",
" }\n",
"}\n",
"\n",
"template <class Ctx>\n",
"__aicore__ inline void EndConv(Ctx &ctx)\n",
"{\n",
" if ASCEND_IS_AIC_CONV {\n",
" ctx.queueAL1.FreeAllEvent();\n",
" ctx.queueBL1.FreeAllEvent();\n",
" }\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 4. Host侧调用代码\n",
"\n",
"Host侧代码负责：填充 Tiling 结构体、验证片上空间、分配设备内存、启动核函数、回收结果。\n",
"\n",
"### 4.1 卷积形状常量\n",
"\n",
"定义本 Demo 使用的卷积形状参数。所有后续的 Tiling 计算和内存分配都基于这些常量。\n",
"\n",
"`Fp16ToFloat` 是一个纯软件的 fp16 解码函数，用于在主机侧将输出结果转换为 float 以便打印。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"} // namespace conv\n",
"\n",
"using namespace conv;\n",
"\n",
"constexpr ConvFormat fmapFormat = ConvFormat::NCHW;\n",
"constexpr ConvFormat filterFormat = ConvFormat::NCHW;\n",
"constexpr ConvFormat outputFormat = ConvFormat::NCHW;\n",
"\n",
"template <class FMAP_TYPE, class WEIGHT_TYPE, class OUTPUT_TYPE>\n",
"class Conv2dBase {\n",
"public:\n",
" using FMAP_T = typename FMAP_TYPE::T;\n",
" using WEIGHT_T = typename WEIGHT_TYPE::T;\n",
" using OUTPUT_T = typename OUTPUT_TYPE::T;\n",
"\n",
" __aicore__ inline void RunConv2dKernel(GM_ADDR x, GM_ADDR filter, GM_ADDR y,\n",
" const Conv2DTilingData &tilingData)\n",
" {\n",
" ctx.hiStartPos = -static_cast<int64_t>(tilingData.padTop);\n",
" ctx.wiStartPos = -static_cast<int64_t>(tilingData.padLeft);\n",
" ctx.hiStartPos = ctx.hiStartPos > 0 ? 0 : ctx.hiStartPos;\n",
"\n",
" fmapGm.SetGlobalBuffer(reinterpret_cast<__gm__ FMAP_T *>(x));\n",
" filterGm.SetGlobalBuffer(reinterpret_cast<__gm__ WEIGHT_T *>(filter));\n",
" outputGm.SetGlobalBuffer(reinterpret_cast<__gm__ OUTPUT_T *>(y));\n",
"\n",
" InitContext(ctx, &tilingData);\n",
" ctx.hiStartPos = -static_cast<int64_t>(tilingData.padTop);\n",
" ctx.wiStartPos = -static_cast<int64_t>(tilingData.padLeft);\n",
"\n",
" ctx.agm = fmapGm;\n",
" ctx.bgm = filterGm;\n",
"\n",
" if ASCEND_IS_AIC_CONV {\n",
" IterateOne(ctx, outputGm);\n",
" }\n",
" EndConv(ctx);\n",
" }\n",
"\n",
"private:\n",
" conv::Conv2dContext<FMAP_TYPE, WEIGHT_TYPE, OUTPUT_TYPE> ctx;\n",
" GlobalTensor<FMAP_T> fmapGm;\n",
" GlobalTensor<WEIGHT_T> filterGm;\n",
" GlobalTensor<OUTPUT_T> outputGm;\n",
"};\n",
"#endif\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 4.2 卷积形状常量与辅助函数\n",
"\n",
"定义本 Demo 使用的卷积形状参数（BATCH、CI、HI、WI、CO、KH、KW 等）。所有后续的 Tiling 计算和内存分配都基于这些常量。\n",
"\n",
"`Fp16ToFloat` 是一个纯软件的 fp16 解码函数，用于在主机侧将输出结果转换为 float 以便打印。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"constexpr uint32_t BATCH = 1;\n",
"constexpr uint32_t CI = 16; // input channels\n",
"constexpr uint32_t HI = 8; // input height\n",
"constexpr uint32_t WI = 8; // input width\n",
"constexpr uint32_t CO = 16; // output channels (must be multiple of 16 for NZ layout)\n",
"constexpr uint32_t KH = 3; // kernel height\n",
"constexpr uint32_t KW = 3; // kernel width\n",
"constexpr uint32_t STRIDE_H = 1;\n",
"constexpr uint32_t STRIDE_W = 1;\n",
"constexpr uint32_t DILATION_H = 1;\n",
"constexpr uint32_t DILATION_W = 1;\n",
"constexpr uint32_t PAD_TOP = 0;\n",
"constexpr uint32_t PAD_BOTTOM = 0;\n",
"constexpr uint32_t PAD_LEFT = 0;\n",
"constexpr uint32_t PAD_RIGHT = 0;\n",
"// Standard conv output size formula\n",
"constexpr uint32_t HO = (HI + PAD_TOP + PAD_BOTTOM - DILATION_H * (KH - 1) - 1) / STRIDE_H + 1;\n",
"constexpr uint32_t WO = (WI + PAD_LEFT + PAD_RIGHT - DILATION_W * (KW - 1) - 1) / STRIDE_W + 1;\n",
"\n",
"__global__ __aicore__ void conv2d_v2_kernel(GM_ADDR x, GM_ADDR filter, GM_ADDR bias,\n",
" GM_ADDR offset_w, GM_ADDR y, __kfc_workspace__ GM_ADDR workspace, GM_ADDR tilingGm)\n",
"{\n",
"#ifdef __CCE_AICORE__\n",
" GET_TILING_DATA(tilingData, tilingGm);\n",
"\n",
" using fmapType = ConvType<TPosition::GM, fmapFormat, half>;\n",
" using weightType = ConvType<TPosition::GM, filterFormat, half>;\n",
" using outputType = ConvType<TPosition::GM, outputFormat, half>;\n",
"\n",
" Conv2dBase<fmapType, weightType, outputType> baseConv2d;\n",
" baseConv2d.RunConv2dKernel(x, filter, y, tilingData);\n",
"#endif\n",
"}\n",
"\n",
"constexpr uint8_t BLOCK_SIZE = 16;\n",
"\n",
"float Fp16ToFloat(uint16_t h)\n",
"{\n",
" uint32_t sign = (h >> 15) & 0x1;\n",
" uint32_t exp = (h >> 10) & 0x1F;\n",
" uint32_t frac = h & 0x3FF;\n",
" float result;\n",
" if (exp == 0) {\n",
" if (frac == 0) {\n",
" result = 0.0f;\n",
" } else {\n",
" result = std::ldexp(static_cast<float>(frac), -24);\n",
" }\n",
" } else if (exp == 31) {\n",
" result = (frac == 0) ? std::numeric_limits<float>::infinity() : std::numeric_limits<float>::quiet_NaN();\n",
" } else {\n",
" result = std::ldexp(1.0f + static_cast<float>(frac) / 1024.0f, static_cast<int>(exp) - 15);\n",
" }\n",
" return sign ? -result : result;\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 4.3 SetTilingData\n",
"\n",
"`SetTilingData` 根据卷积形状常量填充 `Conv2DTilingData` 结构体。\n",
"\n",
"关键字段说明：\n",
"- **kAL1 / kBL1**：L1 侧 K 维度大小 = KH × KW × CI（一次性加载全部 K）。\n",
"- **kL0**：L0 侧 K 维度大小，本 Demo 中等于 kAL1（单次 Mmad 完成全部 K 累加）。\n",
"- **mStep**：M 维度对齐到 16 的倍数（Mmad 要求 M 为 BLOCK_L0_M 的整数倍）。\n",
"- **cinAInCore / cinBInCore**：每核处理的输入通道数（= kAL1 / kernelHW）。\n",
"- **nL1DivBlockSize**：L1 中 N 方向的块数（= nBL1 / BLOCK_SIZE）。\n",
"\n",
"代码中标注了 `[QUIZ]` 的位置是本章课后练习的填空点。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"void SetTilingData(Conv2DTilingData* tiling)\n",
"{\n",
" tiling->padTop = PAD_TOP;\n",
" tiling->padLeft = PAD_LEFT;\n",
"\n",
" // ── ApiTiling: memory-hierarchy tiling parameters ──────────────────────\n",
" // These tell the kernel how large each on-chip buffer tile is.\n",
" // For this full-load demo the entire output tile fits in one L0 block,\n",
" // so hoL0/woL0 >= singleCoreHo/Wo and kL0 == kAL1 == kBL1 == total K.\n",
" tiling->orgHi = HI;\n",
" tiling->orgWi = WI;\n",
" tiling->orgHo = HO;\n",
" tiling->orgWo = WO;\n",
" tiling->singleCoreBatch = BATCH;\n",
" tiling->singleCoreHo = HO;\n",
" tiling->singleCoreWo = WO;\n",
" tiling->orgCi = CI;\n",
" tiling->orgCo = CO;\n",
" tiling->singleCoreCi = CI;\n",
" tiling->singleCoreCo = CO;\n",
" // [QUIZ 1/4] kAL1: total K dimension loaded into L1 for the fmap side.\n",
" // K = number of multiply-accumulate steps per output pixel\n",
" // = kernel_h × kernel_w × Cin\n",
" // Hint: use KH, KW, CI\n",
" /* BEGIN_TODO */\n",
" tiling->kAL1 = KH * KW * CI;\n",
" /* END_TODO */\n",
" // [QUIZ 2/4] kBL1: same K dimension for the weight side (must equal kAL1).\n",
" /* BEGIN_TODO */\n",
" tiling->kBL1 = KH * KW * CI;\n",
" /* END_TODO */\n",
" tiling->nBL1 = 16;\n",
" tiling->hoL0 = 16;\n",
" tiling->woL0 = WO;\n",
" // [QUIZ 3/4] kL0: K dimension actually loaded into L0A/L0B per Mmad call.\n",
" // In this full-load demo kL0 == kAL1 (one shot, no K-loop).\n",
" /* BEGIN_TODO */\n",
" tiling->kL0 = KH * KW * CI;\n",
" /* END_TODO */\n",
" tiling->nL0 = 16;\n",
" tiling->orgHixWi = HI * WI;\n",
" tiling->kernelHxkernelW = KH * KW;\n",
" tiling->aL1SpaceSize = 8192;\n",
"\n",
" // Derived fields — computed from the shape constants above\n",
" uint32_t kernelHW = KH * KW;\n",
" tiling->cinAInCore = tiling->kAL1 / kernelHW;\n",
" tiling->cinBInCore = tiling->kBL1 / kernelHW;\n",
" // [QUIZ 4/4] mStep: M dimension aligned up to the next multiple of 16.\n",
" // M = total output pixels per core = HO × WO\n",
" // Mmad requires M to be a multiple of BLOCK_L0_M (16).\n",
" // Hint: AlignUp(HO * WO, 16)\n",
" /* BEGIN_TODO */\n",
" tiling->mStep = ((HO * WO + 15) / 16) * 16;\n",
" /* END_TODO */\n",
" tiling->kStep = 1;\n",
" tiling->fmapKStride = 1;\n",
" tiling->coutOffsetBlock = CI * kernelHW;\n",
" tiling->nL1DivBlockSize = tiling->nBL1 / BLOCK_SIZE;\n",
" tiling->kernelH = KH;\n",
" tiling->kernelW = KW;\n",
" tiling->strideH = STRIDE_H;\n",
" tiling->strideW = STRIDE_W;\n",
" tiling->dilationH = DILATION_H;\n",
" tiling->dilationW = DILATION_W;\n",
" tiling->offsetx = 0;\n",
"}\n",
"\n",
"// ============================================================================\n",
"// Space validation\n",
"// ============================================================================\n",
"//\n",
"// Before launching the kernel, verify that the tiling parameters actually fit\n",
"// in the NPU's on-chip SRAM. If they don't, the kernel will silently corrupt\n",
"// memory or produce wrong results — there is no hardware bounds check.\n",
"//\n",
"// On-chip buffer sizes (Ascend 950PR/Ascend 950DT):\n",
"// L1 : 512 MB per core (fmap + weight staging area)\n",
"// L0A : 64 KB per core (matrix A — im2col output)\n",
"// L0B : 64 KB per core (matrix B — weight NZ)\n",
"// L0C : 256 KB per core (accumulator, fp32)\n",
"//\n",
"// Layout sizes:\n",
"// AL1 : hiLoad * wiLoad * CI * sizeof(fp16)\n",
"// where hiLoad = (HO-1)*strideH + dilatedKH\n",
"// wiLoad = (WO-1)*strideW + dilatedKW\n",
"// BL1 : nBL1 * kBL1 * sizeof(fp16)\n",
"// AL0 : AlignUp(HO*WO, 16) * kL0 * sizeof(fp16) [NZ: M-blocks × K]\n",
"// BL0 : AlignUp(CO, 16) * kL0 * sizeof(fp16) [NZ: N-blocks × K]\n",
"// CL0 : AlignUp(HO*WO, 16) * AlignUp(CO, 16) * sizeof(fp32)\n",
"//\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 4.4 CheckTilingSpace\n",
"\n",
"在启动核函数之前，必须验证 Tiling 参数是否真正适合 NPU 的片上 SRAM。若超出限制，核函数会静默地越界访问片上内存，产生错误结果或硬件异常——没有任何硬件边界检查。\n",
"\n",
"各缓冲区的大小限制（Ascend 950PR/Ascend 950DT）：\n",
"\n",
"<table style=\"text-align: left; margin-left: 0;\">\n",
"<tr style=\"background-color:#f0f0f0\">\n",
" <td align=\"center\" style=\"width: 150px\"><strong>缓冲区</strong></td>\n",
" <td align=\"center\"><strong>大小</strong></td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\"> L1 </td>\n",
" <td align=\"left\">512 MB / 核</td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">L0A </td>\n",
" <td align=\"left\"> 64 KB / 核 </td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\">L0B </td>\n",
" <td align=\"left\">64 KB / 核 </td>\n",
"</tr>\n",
"<tr>\n",
" <td align=\"left\"> L0C</td>\n",
" <td align=\"left\"> 256 KB / 核</td>\n",
"</tr>\n",
"</table>\n",
"\n",
"代码中标注了 `[QUIZ 1/5]` 到 `[QUIZ 5/5]` 的位置是本章课后练习的填空点。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"static bool CheckTilingSpace(const Conv2DTilingData* a)\n",
"{\n",
" const uint32_t fp16 = sizeof(uint16_t);\n",
" const uint32_t fp32 = sizeof(float);\n",
" const uint32_t align16 = 16;\n",
"\n",
" // L1 — fmap patch that load3d reads from\n",
" // [QUIZ 1/5] Effective kernel size after dilation.\n",
" // dilatedKH = 1 + (KH - 1) * DILATION_H\n",
" // (a 3×3 kernel with dilation=2 spans 5 rows in the input)\n",
" /* BEGIN_TODO */\n",
" uint32_t dilatedKH = 1 + (KH - 1) * DILATION_H;\n",
" uint32_t dilatedKW = 1 + (KW - 1) * DILATION_W;\n",
" /* END_TODO */\n",
" // [QUIZ 2/5] Input patch height/width that L1 must hold for one output tile.\n",
" // The last output row at index (HO-1) reads input row (HO-1)*strideH + (dilatedKH-1),\n",
" // so the patch spans: hiLoad = (HO-1)*STRIDE_H + dilatedKH rows.\n",
" /* BEGIN_TODO */\n",
" uint32_t hiLoad = (HO - 1) * STRIDE_H + dilatedKH;\n",
" uint32_t wiLoad = (WO - 1) * STRIDE_W + dilatedKW;\n",
" /* END_TODO */\n",
" // [QUIZ 3/5] AL1 bytes: the fmap patch is hiLoad × wiLoad × CI fp16 elements.\n",
" /* BEGIN_TODO */\n",
" uint64_t al1Bytes = (uint64_t)hiLoad * wiLoad * CI * fp16;\n",
" /* END_TODO */\n",
"\n",
" // L1 — weight block\n",
" uint64_t bl1Bytes = (uint64_t)a->nBL1 * a->kBL1 * fp16;\n",
"\n",
" // L1 total\n",
" const uint64_t L1_SIZE = 1024 * 1024ULL;\n",
" uint64_t l1Total = al1Bytes + bl1Bytes;\n",
"\n",
" // L0A — im2col result: M-blocks × K elements, fp16\n",
" // [QUIZ 4/5] AL0 bytes.\n",
" // After im2col, L0A holds mAlign × kL0 fp16 elements in NZ layout.\n",
" // mAlign = AlignUp(HO*WO, 16) — M must be a multiple of 16 for Mmad.\n",
" /* BEGIN_TODO */\n",
" uint32_t mAlign = ((HO * WO + align16 - 1) / align16) * align16;\n",
" uint64_t al0Bytes = (uint64_t)mAlign * a->kL0 * fp16;\n",
" /* END_TODO */\n",
"\n",
" // L0B — weight NZ: N-blocks × K elements, fp16\n",
" uint32_t nAlign = ((CO + align16 - 1) / align16) * align16;\n",
" uint64_t bl0Bytes = (uint64_t)nAlign * a->kL0 * fp16;\n",
"\n",
" // L0C — accumulator: M-blocks × N-blocks, fp32\n",
" // [QUIZ 5/5] CL0 bytes.\n",
" // L0C stores the Mmad result as fp32 (even though inputs are fp16).\n",
" // Size = mAlign × nAlign × sizeof(float).\n",
" // Note: fp32 = 4 bytes, NOT fp16 = 2 bytes — this is the most common mistake.\n",
" /* BEGIN_TODO */\n",
" uint64_t cl0Bytes = (uint64_t)mAlign * nAlign * fp32;\n",
" /* END_TODO */\n",
"\n",
" const uint64_t L0A_LIMIT = 64 * 1024ULL;\n",
" const uint64_t L0B_LIMIT = 64 * 1024ULL;\n",
" const uint64_t L0C_LIMIT = 256 * 1024ULL;\n",
"\n",
" bool ok = true;\n",
" if (l1Total > L1_SIZE) {\n",
" std::cerr << \"[TILING ERROR] L1 overflow: need \" << l1Total\n",
" << \" bytes (AL1=\" << al1Bytes << \" + BL1=\" << bl1Bytes\n",
" << \"), limit=\" << L1_SIZE << \"\\n\";\n",
" ok = false;\n",
" }\n",
" if (al0Bytes > L0A_LIMIT) {\n",
" std::cerr << \"[TILING ERROR] L0A overflow: need \" << al0Bytes\n",
" << \" bytes (M=\" << mAlign << \" K=\" << a->kL0\n",
" << \"), limit=\" << L0A_LIMIT << \"\\n\";\n",
" ok = false;\n",
" }\n",
" if (bl0Bytes > L0B_LIMIT) {\n",
" std::cerr << \"[TILING ERROR] L0B overflow: need \" << bl0Bytes\n",
" << \" bytes (N=\" << nAlign << \" K=\" << a->kL0\n",
" << \"), limit=\" << L0B_LIMIT << \"\\n\";\n",
" ok = false;\n",
" }\n",
" if (cl0Bytes > L0C_LIMIT) {\n",
" std::cerr << \"[TILING ERROR] L0C overflow: need \" << cl0Bytes\n",
" << \" bytes (M=\" << mAlign << \" N=\" << nAlign\n",
" << \"), limit=\" << L0C_LIMIT << \"\\n\";\n",
" ok = false;\n",
" }\n",
" if (ok) {\n",
" std::cout << \"[TILING OK] L1=\" << l1Total << \"/\" << L1_SIZE\n",
" << \" L0A=\" << al0Bytes << \"/\" << L0A_LIMIT\n",
" << \" L0B=\" << bl0Bytes << \"/\" << L0B_LIMIT\n",
" << \" L0C=\" << cl0Bytes << \"/\" << L0C_LIMIT << \"\\n\";\n",
" }\n",
" return ok;\n",
"}\n",
"\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"### 4.5 main 函数\n",
"\n",
"`main` 函数完整演示了 AscendC 核函数的调用生命周期：\n",
"\n",
"1. 调用 `SetTilingData` 填充 Tiling 结构体。\n",
"2. 调用 `CheckTilingSpace` 验证片上空间，若溢出则提前退出。\n",
"3. 初始化 ACL 运行环境（`aclInit` → `aclrtSetDevice` → `aclrtCreateStream`）。\n",
"4. 分配主机和设备内存，将输入数据（全 1.0 的 fp16）从主机复制到设备。\n",
"5. 通过 `<<<numBlocks, nullptr, stream>>>` 语法启动核函数，并用 `aclrtSynchronizeStream` 等待完成。\n",
"6. 将输出从设备复制回主机，并写入文本文件。\n",
"7. 释放所有内存和运行时资源。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile -a Sources/01.03/minimal_demo.asc\n",
"int32_t main(int32_t argc, char* argv[])\n",
"{\n",
" size_t inputSize = BATCH * CI * HI * WI * sizeof(uint16_t);\n",
" size_t filterSize = CO * CI * KH * KW * sizeof(uint16_t);\n",
" size_t outputSize = BATCH * CO * HO * WO * sizeof(uint16_t);\n",
" size_t tilingSize = sizeof(Conv2DTilingData);\n",
" size_t workspaceSize = 16 * 1024 * 1024;\n",
" uint32_t numBlocks = 1;\n",
"\n",
" Conv2DTilingData tilingBuf;\n",
" memset(&tilingBuf, 0, sizeof(Conv2DTilingData));\n",
" SetTilingData(&tilingBuf);\n",
"\n",
" // Validate that the tiling fits in on-chip SRAM before touching the device.\n",
" // An overflow here means the kernel would read/write out-of-bounds on-chip\n",
" // memory, producing silent wrong results or a hardware exception.\n",
" if (!CheckTilingSpace(&tilingBuf)) {\n",
" std::cerr << \"Tiling validation failed. Adjust shape constants or tiling params.\\n\";\n",
" return 1;\n",
" }\n",
"\n",
" aclInit(nullptr);\n",
" int32_t deviceId = 0;\n",
" aclrtSetDevice(deviceId);\n",
" aclrtStream stream = nullptr;\n",
" aclrtCreateStream(&stream);\n",
"\n",
" uint8_t* xHost;\n",
" uint8_t* xDevice;\n",
" aclrtMallocHost((void**)(&xHost), inputSize);\n",
" aclrtMalloc((void**)&xDevice, inputSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
" uint16_t* xHalf = reinterpret_cast<uint16_t*>(xHost);\n",
" for (uint32_t i = 0; i < BATCH * CI * HI * WI; ++i) {\n",
" xHalf[i] = 0x3C00; // fp16 1.0\n",
" }\n",
" aclrtMemcpy(xDevice, inputSize, xHost, inputSize, ACL_MEMCPY_HOST_TO_DEVICE);\n",
"\n",
" uint8_t* filterHost;\n",
" uint8_t* filterDevice;\n",
" aclrtMallocHost((void**)(&filterHost), filterSize);\n",
" aclrtMalloc((void**)&filterDevice, filterSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
" uint16_t* filterHalf = reinterpret_cast<uint16_t*>(filterHost);\n",
" for (uint32_t i = 0; i < CO * CI * KH * KW; ++i) {\n",
" filterHalf[i] = 0x3C00; // fp16 1.0\n",
" }\n",
" aclrtMemcpy(filterDevice, filterSize, filterHost, filterSize, ACL_MEMCPY_HOST_TO_DEVICE);\n",
"\n",
" uint8_t* yDevice;\n",
" uint8_t* yHost;\n",
" aclrtMallocHost((void**)(&yHost), outputSize);\n",
" aclrtMalloc((void**)&yDevice, outputSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
"\n",
" uint8_t* workspaceDevice;\n",
" aclrtMalloc((void**)&workspaceDevice, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
"\n",
" uint8_t* tilingHost;\n",
" uint8_t* tilingDevice;\n",
" aclrtMallocHost((void**)(&tilingHost), tilingSize);\n",
" aclrtMalloc((void**)&tilingDevice, tilingSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
" memcpy(tilingHost, &tilingBuf, tilingSize);\n",
" aclrtMemcpy(tilingDevice, tilingSize, tilingHost, tilingSize, ACL_MEMCPY_HOST_TO_DEVICE);\n",
"\n",
" conv2d_v2_kernel<<<numBlocks, nullptr, stream>>>(\n",
" xDevice, filterDevice, nullptr, nullptr, yDevice, workspaceDevice, tilingDevice);\n",
" aclrtSynchronizeStream(stream);\n",
"\n",
" aclrtMemcpy(yHost, outputSize, yDevice, outputSize, ACL_MEMCPY_DEVICE_TO_HOST);\n",
" std::cout << \"Conv2D V2 kernel execution completed.\" << std::endl;\n",
"\n",
" // Write input matrix to file\n",
" {\n",
" std::ofstream ofs(\"input_matrix.txt\");\n",
" uint16_t* xData = reinterpret_cast<uint16_t*>(xHost);\n",
" ofs << \"Input shape: [\" << BATCH << \", \" << CI << \", \" << HI << \", \" << WI << \"]\" << std::endl;\n",
" for (uint32_t n = 0; n < BATCH; ++n) {\n",
" for (uint32_t c = 0; c < CI; ++c) {\n",
" ofs << \"N=\" << n << \", C=\" << c << \":\" << std::endl;\n",
" for (uint32_t h = 0; h < HI; ++h) {\n",
" for (uint32_t w = 0; w < WI; ++w) {\n",
" uint32_t idx = ((n * CI + c) * HI + h) * WI + w;\n",
" ofs << Fp16ToFloat(xData[idx]) << \" \";\n",
" }\n",
" ofs << std::endl;\n",
" }\n",
" }\n",
" }\n",
" ofs.close();\n",
" std::cout << \"Input matrix written to input_matrix.txt\" << std::endl;\n",
" }\n",
"\n",
" // Write output matrix to file\n",
" {\n",
" std::ofstream ofs(\"output_matrix.txt\");\n",
" uint16_t* yData = reinterpret_cast<uint16_t*>(yHost);\n",
" ofs << \"Output shape: [\" << BATCH << \", \" << CO << \", \" << HO << \", \" << WO << \"]\" << std::endl;\n",
" for (uint32_t n = 0; n < BATCH; ++n) {\n",
" for (uint32_t c = 0; c < CO; ++c) {\n",
" ofs << \"N=\" << n << \", C=\" << c << \":\" << std::endl;\n",
" for (uint32_t h = 0; h < HO; ++h) {\n",
" for (uint32_t w = 0; w < WO; ++w) {\n",
" uint32_t idx = ((n * CO + c) * HO + h) * WO + w;\n",
" ofs << Fp16ToFloat(yData[idx]) << \" \";\n",
" }\n",
" ofs << std::endl;\n",
" }\n",
" }\n",
" }\n",
" ofs.close();\n",
" std::cout << \"Output matrix written to output_matrix.txt\" << std::endl;\n",
" }\n",
"\n",
" aclrtFree(xDevice);\n",
" aclrtFreeHost(xHost);\n",
" aclrtFree(filterDevice);\n",
" aclrtFreeHost(filterHost);\n",
" aclrtFree(yDevice);\n",
" aclrtFreeHost(yHost);\n",
" aclrtFree(workspaceDevice);\n",
" aclrtFree(tilingDevice);\n",
" aclrtFreeHost(tilingHost);\n",
"\n",
" aclrtDestroyStream(stream);\n",
" aclrtResetDevice(deviceId);\n",
" aclFinalize();\n",
"\n",
" return 0;\n",
"}\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 5. CMakeLists.txt\n",
"\n",
"以下是本项目的 CMake 构建配置。与 Add 算子示例相比，Conv2D V2 Demo 需要额外链接多个 CANN 运行时库（`tiling_api`、`ascendc_runtime`、`ascendcl` 等），并通过 `target_include_directories` 指定 AscendC 内部头文件路径。\n"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile Sources/01.03/CMakeLists.txt\n",
"cmake_minimum_required(VERSION 3.16)\n",
"\n",
"find_package(ASC REQUIRED)\n",
"\n",
"project(conv2d_v2_demo_minimal LANGUAGES ASC CXX)\n",
"\n",
"add_executable(minimal_demo\n",
" minimal_demo.asc\n",
")\n",
"\n",
"target_include_directories(minimal_demo PRIVATE\n",
" ${CMAKE_CURRENT_SOURCE_DIR}\n",
" ${ASCEND_CANN_PACKAGE_LINUX_PATH}/asc/impl\n",
" ${ASCEND_CANN_PACKAGE_LINUX_PATH}/asc/impl/basic_api\n",
" ${ASCEND_CANN_PACKAGE_LINUX_PATH}/asc/impl/basic_api/reg_compute\n",
" ${ASCEND_CANN_PACKAGE_LINUX_PATH}/asc/impl/simt_api\n",
" ${ASCEND_CANN_PACKAGE_LINUX_PATH}/include\n",
" ${ASCEND_CANN_PACKAGE_LINUX_PATH}/tikcpp/tikcfw\n",
")\n",
"\n",
"target_link_directories(minimal_demo PRIVATE\n",
" ${ASCEND_CANN_PACKAGE_PATH}/lib64\n",
")\n",
"\n",
"target_link_libraries(minimal_demo PRIVATE\n",
" tiling_api\n",
" register\n",
" platform\n",
" ascendc_runtime\n",
" runtime\n",
" ascendcl\n",
" error_manager\n",
" profapi\n",
" ge_common_base\n",
" unified_dlog\n",
" mmpa\n",
" ascend_dump\n",
" c_sec\n",
" m\n",
" dl\n",
")\n",
"\n",
"target_compile_options(minimal_demo PRIVATE\n",
" $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>\n",
")"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n\n## 6. 构建脚本 build.sh\n\n除了直接在 Jupyter 中运行 CMake 命令，`src/conv2d_op/build.sh` 提供了一个等价的一键构建脚本，适合在终端中使用。\n\n```bash\n#!/bin/bash\nset -e\n\nSCRIPT_DIR=$(cd \"$(dirname \"$0\")\" && pwd)\nBUILD_DIR=${SCRIPT_DIR}/build\n# TODO: Set ASCEND_HOME_PATH in your environment, or replace the fallback path below\n# with your actual CANN installation directory, e.g. /usr/local/Ascend/ascend-toolkit/latest\nCANN_PATH=${ASCEND_HOME_PATH:-/path/to/your/cann}\n\nrm -rf ${BUILD_DIR}\nmkdir -p ${BUILD_DIR} && cd ${BUILD_DIR}\n\ncmake .. -DASCEND_CANN_PACKAGE_PATH=${CANN_PATH}\nmake -j$(nproc)\n\necho \"Build success: ${BUILD_DIR}/minimal_demo\"\n```\n\n**使用方式**：\n\n```bash\n# 步骤一：设置环境变量\nexport ASCEND_HOME_PATH=/your/cann/install/path\nbash src/conv2d_op/build.sh\n\n# 步骤二：修改脚本中的 CANN_PATH 默认值为你的实际路径\n```\n\n> **注意**：脚本中 `CANN_PATH` 的默认值 `/path/to/your/cann` 是占位符，请替换为你的实际 CANN 安装目录，或在运行前设置 `ASCEND_HOME_PATH` 环境变量。"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 7. 编译与运行\n",
"\n",
"执行以下命令编译可执行文件："
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"!cd Sources/01.03 && mkdir -p build\n",
"!export ASC_DIR=$ASCEND_HOME_PATH/aarch64-linux/tikcpp/ascendc_kernel_cmake/ && \\\n",
"cd Sources/01.03/build/ && \\\n",
"cmake .. && \\\n",
"make"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"再执行以下代码，进行算子的实际运行："
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"!cd Sources/01.03/build/ && ./minimal_demo"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 课后实践\n",
"\n",
"本节共有 **9 处填空**，分布在 `SetTilingData`（4 处）和 `CheckTilingSpace`（5 处）中。\n",
"\n",
"请根据注释中的提示，在下面的代码框中补充完整的 `SetTilingData` 和 `CheckTilingSpace` 函数实现。\n",
"\n",
"**SetTilingData 填空（4 处）：**\n",
"\n",
"- `[QUIZ 1/4]` kAL1：L1 侧特征图的 K 维度大小 = KH × KW × CI\n",
"- `[QUIZ 2/4]` kBL1：L1 侧权重的 K 维度大小（与 kAL1 相等）\n",
"- `[QUIZ 3/4]` kL0：L0 侧每次 Mmad 的 K 维度大小（本 Demo 中等于 kAL1）\n",
"- `[QUIZ 4/4]` mStep：M 维度对齐到 16 的倍数，M = HO × WO\n",
"\n",
"**CheckTilingSpace 填空（5 处）：**\n",
"\n",
"- `[QUIZ 1/5]` dilatedKH/W：膨胀后的有效卷积核尺寸\n",
"- `[QUIZ 2/5]` hiLoad/wiLoad：L1 需要加载的输入 patch 尺寸\n",
"- `[QUIZ 3/5]` al1Bytes：AL1 缓冲区字节数\n",
"- `[QUIZ 4/5]` al0Bytes：AL0 缓冲区字节数（im2col 结果）\n",
"- `[QUIZ 5/5]` cl0Bytes：L0C 累加器字节数（注意是 fp32！）"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"# 在此处填写你的答案，然后运行本单元格验证\n",
"# 参考 SetTilingData 和 CheckTilingSpace 中的 [QUIZ] 注释\n",
"\n",
"# [QUIZ 1/4] kAL1 = ?\n",
"kAL1 = None # TODO: KH * KW * CI\n",
"\n",
"# [QUIZ 2/4] kBL1 = ?\n",
"kBL1 = None # TODO: KH * KW * CI\n",
"\n",
"# [QUIZ 3/4] kL0 = ?\n",
"kL0 = None # TODO: KH * KW * CI\n",
"\n",
"# [QUIZ 4/4] mStep = ?\n",
"HO = 6; WO = 6\n",
"mStep = None # TODO: AlignUp(HO * WO, 16)\n",
"\n",
"# [QUIZ 1/5] dilatedKH, dilatedKW = ?\n",
"KH = 3; KW = 3; DILATION_H = 1; DILATION_W = 1\n",
"dilatedKH = None # TODO: 1 + (KH - 1) * DILATION_H\n",
"dilatedKW = None # TODO: 1 + (KW - 1) * DILATION_W\n",
"\n",
"# [QUIZ 2/5] hiLoad, wiLoad = ?\n",
"STRIDE_H = 1; STRIDE_W = 1\n",
"hiLoad = None # TODO: (HO - 1) * STRIDE_H + dilatedKH\n",
"wiLoad = None # TODO: (WO - 1) * STRIDE_W + dilatedKW\n",
"\n",
"# [QUIZ 3/5] al1Bytes = ?\n",
"CI = 16\n",
"al1Bytes = None # TODO: hiLoad * wiLoad * CI * 2 (fp16 = 2 bytes)\n",
"\n",
"# [QUIZ 4/5] mAlign, al0Bytes = ?\n",
"mAlign = None # TODO: ((HO * WO + 15) // 16) * 16\n",
"al0Bytes = None # TODO: mAlign * kL0 * 2 (fp16 = 2 bytes)\n",
"\n",
"# [QUIZ 5/5] cl0Bytes = ?\n",
"CO = 16\n",
"nAlign = ((CO + 15) // 16) * 16\n",
"cl0Bytes = None # TODO: mAlign * nAlign * 4 (fp32 = 4 bytes, NOT 2!)\n",
"\n",
"print(\"Your answers:\")\n",
"print(f\" kAL1={kAL1}, kBL1={kBL1}, kL0={kL0}, mStep={mStep}\")\n",
"print(f\" dilatedKH={dilatedKH}, dilatedKW={dilatedKW}\")\n",
"print(f\" hiLoad={hiLoad}, wiLoad={wiLoad}\")\n",
"print(f\" al1Bytes={al1Bytes}, al0Bytes={al0Bytes}, cl0Bytes={cl0Bytes}\")"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"完成填写后，执行以下命令查看答案："
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"!cat ./answer/01.03_answer.txt"
]
}
],
"metadata": {
"kernelspec": {
"display_name": "Python 3",
"language": "python",
"name": "python3"
},
"language_info": {
"name": "python",
"version": "3.8.0"
}
},
"nbformat": 4,
"nbformat_minor": 5
}
