{
"cells": [
{
"cell_type": "markdown",
"metadata": {},
"source": [
"# 2.2 算子执行流程\n",
"\n",
"本节详细介绍自定义算子的完整执行流程，包括 aclnn 单算子 API 调用和 GE 图模式两种主要调用方式的执行流程。\n",
"\n",
"---\n",
"\n",
"## 学习目标\n",
"\n",
"完成本节后，你将能够：\n",
"- 理解算子的执行流程架构\n",
"- 掌握 aclnn 单算子 API 调用流程\n",
"- 掌握 GE 图模式执行流程\n",
"- 了解 内核调用符<<<>>> kernel直调的执行方式\n",
"- 掌握各阶段的调试方法\n",
"\n",
"---"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"## 1. 环境准备\n",
"正式开始学习之前，先要对jupyter环境进行初始化。以下代码完成了初始化并将环境中的变量导入jupyter环境，同时完成了代码目录的创建。保证能正常导入代码以及使用bisheng编译器，完成算子的开发及编译。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"!mkdir -p Sources/02.02\n",
"!mkdir -p Sources/02.02/add\n",
"\n",
"import os, subprocess\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))"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 2. 算子执行通路介绍\n",
"\n",
"算子开发完成后，需要验证完整的通路，需对算子的执行过程有所了解，下图介绍了框架到算子的执行流程，共5个通路。\n",
"\n",
"<img src=\"./images/算子执行通路.png\" alt=\"算子执行通路\" width=\"700px\" >\n",
"\n",
"从上层框架到算子执行，一共有5条路径，开发过程中，重点关注的是路径3和路径5，也是本节课程重点介绍的路径。\n",
"\n",
"\n",
"路径3：即所谓的图模式。图模式是相对于单算子模式的，图模式下，通过有向无环图，描述网络的计算逻辑。图模式的Host调度，可以避免总是返回Python调用栈，避免冗余流程与数据结构转换，并且可以直接使用图编译阶段完成的Infer Shape与Tiling计算结果。\n",
"\n",
"路径5：即所谓的单算子模式或者eager模式（即时执行），由Host CPU逐个下发算子，一个算子的下发流程包含Python处理、Python到C++数据结构转换、Tiling计算、申请算子的Workspace内存和输出内存、Launch等Host操作。为了加速单算子Host调度，在Pytorch中，昇腾适配层采用了生产者-消费者双线程模式加速，生产者线程主要负责Launch之前的处理动作，消费者线程主要负责Launch算子，对标路径2。\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"Ascend C算子常见的调用方式有kernel直调、单算子调用、第三方框架中调用。工程化的算子通常以单算子调用和第三方框架中调用为主。具体调用方式和调用条件如下：\n",
"\n",
"<table >\n",
" <tr>\n",
" <th style=\"border: 1px solid #ddd; padding: 12px; text-align: left; font-weight: bold; min-width: 180px;\">调用方式</th>\n",
" <th style=\"border: 1px solid #ddd; padding: 12px; text-align: left; font-weight: bold; min-width: 180px;\">使用条件</th>\n",
" </tr>\n",
" <tr>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">单算子API</td>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">算子工程编译部署</td>\n",
" </tr>\n",
" <tr>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">单算子模型执行</td>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">算子入图开发，算子工程编译部署</td>\n",
" </tr>\n",
" <tr>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">IR构图</td>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">算子入图开发，算子工程编译部署</td>\n",
" </tr>\n",
" <tr>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">Pytorch框架调用</td>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">插件适配开发，算子工程编译部署</td>\n",
" </tr>\n",
" <tr>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">TensorFlow框架调用</td>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">插件适配开发，算子工程编译部署</td>\n",
" </tr>\n",
" <tr>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">Pybind调用</td>\n",
" <td style=\"border: 1px solid #ddd; padding: 10px; vertical-align: middle; text-align: left;\">算子工程编译部署</td>\n",
" </tr>\n",
"</table> "
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 3. 单算子调用执行流程\n",
"\n",
"下面我们以Add算子为例，基于torch框架，介绍单算子API的执行流程。\n",
"\n",
"Add算子的介绍参见CANN社区：https://gitcode.com/cann/ops-math/blob/master/math/add/README.md\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"\n",
"step1: 用户使用torch，调用Add算子。"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"import torch\n",
"import torch_npu\n",
"\n",
"x = torch.randint(-5, 5, (2, 2), dtype=torch.int32).npu()\n",
"y = torch.randint(-5, 5, (2, 2), dtype=torch.int32).npu()\n",
"\n",
"z = torch.add(x, y)"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"step2：torch层调用Add算子的aclnn接口，aclnn分为两段式接口，会先调用aclnnAddGetWorkspace分配内存，然后调用aclnnAdd接口。\n",
"\n",
"step3: aclnn根据注册的算子原型，找到算子执行的二进制，即kernel.o文件。\n",
"\n",
"step4: 进入Add算子的Tiling计算阶段，完成相关tilingData计算。\n",
"\n",
"step5: rts 进行kernelLaunch调度执行\n",
"\n",
"step6: Add算子kernel，从GM上CopyIn数据，并根据TilingData，在AICore上执行Compute计算，最后从UB上CopyOut到GM上。\n",
"\n",
"详细流程如图：\n",
"\n",
"<img src=\"./images/Add算子执行流程.png\" alt=\"Add算子执行接口流程\" width=\"700px\" >\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 4. 图模式调用执行流程\n",
"\n",
"\n",
"### 4.1 静态图执行流程\n",
"\n",
"静态图，即图上算子Shape在编译时是已知且明确的（shape为[10, 2]）。\n",
"\n",
"以Model粒度下发执行：\n",
"\n",
"- 模型加载时GE向rts申请模型流，按顺序下发所有算子的Kernel，此时模型不会立即执行\n",
"\n",
"- 模型执行开始时，由GE调用rts的rtModelExecute接口下发一个ModelExecute任务，触发流上的tasks启动，执行过程中没有host调用\n",
"\n",
"静态图执行示意：\n",
"\n",
"<img src=\"./images/静态图执行流程.png\" alt=\"静态图执行示意\" width=\"700px\" >\n",
"\n",
"\n",
"### 4.2 动态图执行流程\n",
"\n",
"动态图：图上算子Shape在编译时未知（shape中包含-1或-2），编译结果需要支持（range内）所有shape的执行。\n",
"\n",
"动态图分类：\n",
"\n",
"- 纯动态图 —— 网络输入时动态\n",
"\n",
"- 动静混合 —— 网络中部分输入动态，部分输入静态；或者网络输入为静态，但存在二三类算子导致inferShape推导出动态。\n",
"\n",
"动态图执行：\n",
"\n",
"走GE的动态图执行器（常说的runtime2.0）按照拓扑序调用每个算子的inferShape、Tiling，为每个算子分配输出内存和WorkSpace内存，组装args后调用rtKenelLaunch完成算子下发。\n",
"\n",
"\n",
"动态图执行示意：\n",
"\n",
"<img src=\"./images/动态图执行流程.png\" alt=\"动态图执行示意\" width=\"700px\" >\n"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 5、使用内核调用符<<<>>> kernel直调执行流程\n",
"\n",
"以Add算子为例子，演示内核调用符<<<>>> kernel直调执行流程。\n",
"\n",
"- 样例功能： \n",
" Add样例实现了两个数据相加，返回相加结果的功能。对应的数学表达式为： \n",
" ```\n",
" z = x + y\n",
" ```\n",
"\n",
"- 样例规格：\n",
" <table>\n",
" <tr><td rowspan=\"1\" align=\"center\">样例类型(OpType)</td><td colspan=\"4\" align=\"center\">Add</td></tr>\n",
" <tr><td rowspan=\"3\" align=\"center\">样例输入</td><td align=\"center\">name</td><td align=\"center\">shape</td><td align=\"center\">data type</td><td align=\"center\">format</td></tr>\n",
" <tr><td align=\"center\">x</td><td align=\"center\">[8, 2048]</td><td align=\"center\">float</td><td align=\"center\">ND</td></tr>\n",
" <tr><td align=\"center\">y</td><td align=\"center\">[8, 2048]</td><td align=\"center\">float</td><td align=\"center\">ND</td></tr>\n",
" <tr><td rowspan=\"1\" align=\"center\">样例输出</td><td align=\"center\">z</td><td align=\"center\">[8, 2048]</td><td align=\"center\">float</td><td align=\"center\">ND</td></tr>\n",
" <tr><td rowspan=\"1\" align=\"center\">核函数名</td><td colspan=\"4\" align=\"center\">add_custom</td></tr>\n",
" </table>\n",
"\n",
"- 样例实现：\n",
" - Kernel实现 \n",
" 本样例使用8个核完成计算，每个核处理2048个元素。计算偏移量为：\n",
" ```\n",
" block_idx * blockLength\n",
" ```\n",
"\n",
" Add样例的实现流程分为3个步骤：\n",
"\n",
" **第一步：搬运数据到UB（Unified Buffer）**\n",
" \n",
" 将GM（Global Memory）上的输入x和y搬运到UB（Unified Buffer）上的xLocal、yLocal中。\n",
" \n",
" **第二步：执行向量加法**\n",
" \n",
" 对xLocal、yLocal执行加法操作，计算结果存储在UB（Unified Buffer）上的zLocal中。\n",
" \n",
" **第三步：搬运结果到GM（Global Memory）**\n",
" \n",
" 将输出数据从zLocal搬运至GM（Global Memory）上的输出z中。\n",
"\n",
"- 调用实现 \n",
" 使用内核调用符<<<>>>调用核函数。"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"样例代码如下："
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile Sources/02.02/add/Add_kernel_call_test.cpp\n",
"\n",
"#include <cstdint>\n",
"#include <vector>\n",
"#include <algorithm>\n",
"#include <iterator>\n",
"\n",
"#include <iostream>\n",
"\n",
"#include \"acl/acl.h\"\n",
"\n",
"#include \"kernel_operator.h\"\n",
"\n",
"template <uint32_t blockLength>\n",
"__vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)\n",
"{\n",
" AscendC::InitSocState();\n",
"\n",
" AscendC::GlobalTensor<float> xGm, yGm, zGm;\n",
" xGm.SetGlobalBuffer(x + block_idx * blockLength, blockLength);\n",
" yGm.SetGlobalBuffer(y + block_idx * blockLength, blockLength);\n",
" zGm.SetGlobalBuffer(z + block_idx * blockLength, blockLength);\n",
"\n",
" AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;\n",
" AscendC::LocalTensor<float> xLocal = ubAllocator.Alloc<float, blockLength>();\n",
" AscendC::LocalTensor<float> yLocal = ubAllocator.Alloc<float, blockLength>();\n",
" AscendC::LocalTensor<float> zLocal = ubAllocator.Alloc<float, blockLength>();\n",
"\n",
" AscendC::DataCopy(xLocal, xGm, blockLength);\n",
" AscendC::DataCopy(yLocal, yGm, blockLength);\n",
" AscendC::PipeBarrier<PIPE_ALL>();\n",
"\n",
" AscendC::Add(zLocal, xLocal, yLocal, blockLength);\n",
" AscendC::PipeBarrier<PIPE_ALL>();\n",
"\n",
"#if 0\n",
" // Debug I/O. Set #if 1 to enable print.\n",
" AscendC::printf(\"%s\\n\", \"[ DumpTensor in xLocal]\");\n",
" AscendC::DumpTensor(xLocal, 1, 32);\n",
"\n",
" AscendC::printf(\"%s\\n\", \"[ DumpTensor in yLocal]\");\n",
" AscendC::DumpTensor(yLocal, 2, 32);\n",
"\n",
" AscendC::printf(\"%s\\n\", \"[ DumpTensor in zLocal]\");\n",
" AscendC::DumpTensor(zLocal, 3, 32);\n",
"#endif\n",
"\n",
" AscendC::DataCopy(zGm, zLocal, blockLength);\n",
" AscendC::PipeBarrier<PIPE_ALL>();\n",
"}\n",
"\n",
"std::vector<float> kernel_add(std::vector<float>& x, std::vector<float>& y)\n",
"{\n",
" constexpr uint32_t numBlocks = 8;\n",
" constexpr uint32_t blockLength = 2048;\n",
" uint32_t totalLength = x.size();\n",
" size_t totalByteSize = totalLength * sizeof(float);\n",
" int32_t deviceId = 0;\n",
" float* xDevice = nullptr;\n",
" float* yDevice = nullptr;\n",
" float* zDevice = nullptr;\n",
" uint8_t* zHost = nullptr;\n",
"\n",
" aclInit(nullptr);\n",
" aclrtSetDevice(deviceId);\n",
"\n",
" aclrtMalloc((void**)&xDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
" aclrtMalloc((void**)&yDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
" aclrtMalloc((void**)&zDevice, totalByteSize, ACL_MEM_MALLOC_HUGE_FIRST);\n",
" aclrtMallocHost((void**)&zHost, totalByteSize);\n",
"\n",
" aclrtMemcpy(xDevice, totalByteSize, x.data(), totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);\n",
" aclrtMemcpy(yDevice, totalByteSize, y.data(), totalByteSize, ACL_MEMCPY_HOST_TO_DEVICE);\n",
"\n",
" add_custom<blockLength><<<numBlocks, 0>>>(xDevice, yDevice, zDevice);\n",
" aclrtSynchronizeDevice();\n",
"\n",
" aclrtMemcpy(zHost, totalByteSize, zDevice, totalByteSize, ACL_MEMCPY_DEVICE_TO_HOST);\n",
" std::vector<float> z((float*)zHost, (float*)(zHost + totalByteSize));\n",
"\n",
" aclrtFree(xDevice);\n",
" aclrtFree(yDevice);\n",
" aclrtFree(zDevice);\n",
" aclrtFreeHost(zHost);\n",
"\n",
" aclrtResetDevice(deviceId);\n",
" aclFinalize();\n",
"\n",
" return z;\n",
"}\n",
"\n",
"uint32_t VerifyResult(std::vector<float>& output, std::vector<float>& golden)\n",
"{\n",
" auto printTensor = [](std::vector<float>& tensor, const char* name) {\n",
" constexpr size_t maxPrintSize = 20;\n",
" std::cout << name << \": \";\n",
" std::copy(\n",
" tensor.begin(), tensor.begin() + std::min(tensor.size(), maxPrintSize),\n",
" std::ostream_iterator<float>(std::cout, \" \"));\n",
" if (tensor.size() > maxPrintSize) {\n",
" std::cout << \"...\";\n",
" }\n",
" std::cout << std::endl;\n",
" };\n",
" printTensor(output, \"Output\");\n",
" printTensor(golden, \"Golden\");\n",
" if (std::equal(golden.begin(), golden.end(), output.begin())) {\n",
" std::cout << \"test pass!\" << std::endl;\n",
" return 0;\n",
" } else {\n",
" std::cout << \"test failed!\" << std::endl;\n",
" return 1;\n",
" }\n",
" return 0;\n",
"}\n",
"\n",
"int32_t main(int32_t argc, char* argv[])\n",
"{\n",
" constexpr uint32_t totalLength = 8 * 2048;\n",
" std::vector<float> x(totalLength);\n",
" std::vector<float> y(totalLength);\n",
"\n",
" for (uint32_t i = 0; i < totalLength; ++i) {\n",
" x[i] = i * 0.1f;\n",
" y[i] = i * 0.2f;\n",
" }\n",
"\n",
" std::vector<float> output = kernel_add(x, y);\n",
"\n",
" std::vector<float> golden(totalLength);\n",
" for (uint32_t i = 0; i < totalLength; ++i) {\n",
" golden[i] = x[i] + y[i];\n",
" }\n",
"\n",
" return VerifyResult(output, golden);\n",
"}"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"编写CMakeLists.txt"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"%%writefile Sources/02.02/add/CMakeLists.txt\n",
"\n",
"cmake_minimum_required(VERSION 3.16)\n",
"\n",
"set(CMAKE_ASC_ARCHITECTURES \"dav-2201\" CACHE STRING \"NPU architecture: dav-2201, dav-3510\")\n",
"\n",
"find_package(ASC REQUIRED)\n",
"\n",
"project(kernel_samples LANGUAGES ASC CXX)\n",
"\n",
"add_executable(demo\n",
" Add_kernel_call_test.cpp\n",
")\n",
"\n",
"target_compile_options(demo PRIVATE\n",
" $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>\n",
")"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"样例执行："
]
},
{
"cell_type": "code",
"execution_count": 45,
"metadata": {},
"outputs": [],
"source": [
"# 配置环境变量\n",
"!source $ASCEND_TOOLKIT_HOME/set_env.sh"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"!mkdir -p Sources/02.02/add/build # 创建并进入build目录\n",
"!cd Sources/02.02/add/build && cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j # 编译工程，默认npu模式\n",
"!./demo # 执行样例"
]
},
{
"cell_type": "markdown",
"metadata": {},
"source": [
"---\n",
"\n",
"## 课后练习\n",
"\n",
"完成以下题目，检验你对算子执行流程的理解：\n",
"\n",
"1. （判断题）aclnn 单算子 API 调用需要先完成算子工程的编译部署。 \n",
"\n",
"2. （判断题）静态图模式下，算子的 Shape 在编译时是已知且明确的，动态图模式下 Shape 在编译时未知。 \n",
"\n",
"3. （单选题）aclnn 接口采用两段式调用，先调用哪个接口分配内存？ \n",
" A. aclnnAdd \n",
" B. aclnnAddGetWorkspace \n",
" C. aclrtMalloc \n",
" D. aclrtMemcpy \n",
"\n",
"4. （单选题）内核调用符 <<<>>> 的语法中，threadsPerBlock 参数表示什么？ \n",
" A. 线程块的个数 \n",
" B. 每个线程块内的线程数量 \n",
" C. 动态申请内存大小 \n",
" D. 流同步标识 \n",
"\n",
"5. （单选题）动态图执行流程中，按照拓扑序调用每个算子的哪些操作？ \n",
" A. inferShape 和 Tiling \n",
" B. 编译和执行 \n",
" C. 数据生成和验证 \n",
" D. 内存分配和释放 \n",
"\n",
"**执行以下代码获取答案。**\n",
"\n",
"---\n",
"\n",
"上一节：[2.1 章节介绍](./02.01_chapter_intro.ipynb) | [下一节：2.3 算子目录结构](./02.03_operator_directory_structure.ipynb)"
]
},
{
"cell_type": "code",
"execution_count": null,
"metadata": {},
"outputs": [],
"source": [
"!cat ./answer/02.02_answer.txt"
]
}
],
"metadata": {
"kernelspec": {
"display_name": "Python 3",
"language": "python",
"name": "python3"
},
"language_info": {
"codemirror_mode": {
"name": "ipython",
"version": 3
},
"file_extension": ".py",
"mimetype": "text/x-python",
"name": "python",
"nbconvert_exporter": "python",
"pygments_lexer": "ipython3",
"version": "3.12.9"
}
},
"nbformat": 4,
"nbformat_minor": 2
}
