Ascend C 多核 Tiling 切分策略全解析：以 Tiling Selector 矢量算子样例解读核间与核内数据切分-平芜编程栈

人工智能
深度学习
算子库
CANN
Ascend

【免费下载链接】asc-devkit

本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言，原生支持C和C++标准规范，主要由类库和语言扩展层构成，提供多层级API，满足多维场景算子开发诉求。

项目地址：https://gitcode.com/cann/asc-devkit

本文围绕 asc-devkit 仓库中的Tiling Selector 特性样例（02_tiling_selector）展开，以两个一维halfTensor 相加的矢量算子为例，系统讲解如何在 Host 侧根据输入数据量与参与计算的 AI Core 数量计算 Tiling 参数，并在"主块均分、尾块均分、尾核切分、尾核尾块切分"四种典型策略之间完成选择。读完本文，你将掌握多核矢量算子 Tiling 切分的完整实现套路：Tiling 结构体设计、Host 侧逐级切分算法、Device 侧按核偏移计算与分块循环处理，以及样例在 NPU/CPU/仿真模式下的编译运行与精度验证方法。

1. 为什么要做多核 Tiling：背景与核心概念

在昇腾 AI Core 上，核函数（Kernel）是 Device 侧执行的并行函数。为了提升算子执行效率，通常将输入数据切分后分配到多个逻辑核并行处理；同时，单个核的 Local Memory（如 Unified Buffer，即 UB）容量有限，无法一次容纳全部输入输出，必须分批搬运数据、分块计算。这种对输入数据进行切分、分块计算的过程称为Tiling，切分数据的算法即 Tiling 策略，而根据 shape 等信息计算切分参数（块大小、循环次数等）的程序称为Tiling 实现（Tiling 函数）。相关背景可参阅 multi_core_tiling/overview.md。

Tiling 切分遵循三条基本原则：

以 DataBlock 为最小切分单位：受硬件限制，UB 相关的数据搬运与矢量计算要求数据长度和起始地址 32 字节对齐。一个 DataBlock 为 32 字节，输入数据不满足对齐要求时需按 DataBlock 向上对齐。
充分利用 UB 空间：减少从 Global Memory 到 UB 的搬运次数，避免频繁搬运造成的性能瓶颈。
均衡利用多核计算能力：将计算尽量均衡地分配到各 AI Core，充分发挥多核并行优势。

围绕核间与核内两级切分，本文涉及四个关键概念：

主块：基于 Local Memory 可用空间确定的单个核一次处理的最大数据块，长度用mainTileLength表示，按 32 字节对齐。
尾块：单个核内最后一个长度小于mainTileLength的数据块，长度用lastTileLength表示。
整核：核间数据无法完全均分时，处理数据量较多的核。
尾核：核间数据无法完全均分时，处理数据量较少的核。

2. 样例总览与目录结构

02_tiling_selector 目录展示的是 Ascend C SIMDTiling Selector特性样例：根据输入数据量和 AI Core 数量计算 Tiling 参数，并选择核间（inter-core）与核内（intra-core）切分策略。其下唯一的算子开发样例为 vector_tiling_strategy，功能描述如下：

目录名称	功能描述
vector_tiling_strategy	以两个一维 Tensor 相加为例，展示矢量计算场景下多核 Tiling 切分策略的计算与选择流程。

vector_tiling_strategy样例的目录结构如下（详见 vector_tiling_strategy/README.md）：

├── vector_tiling_strategy │ ├── scripts │ │ ├── gen_data.py // 输入数据和真值数据生成脚本 │ │ └── verify_result.py // 输出结果验证脚本 │ ├── CMakeLists.txt // 编译工程文件 │ ├── vector_tiling_strategy.asc // Ascend C样例实现与调用样例 │ ├── data_utils.h // 数据读入写出函数 │ └── README.md // 样例说明文档

2.1 支持的产品及 CANN 软件版本

产品	CANN 软件版本
Ascend 950PR & 950DT 系列产品	≥ CANN 9.1.0
Atlas A3 系列产品	≥ CANN 9.0.0
Atlas A2 系列产品	≥ CANN 9.0.0

2.2 样例规格总览

本样例计算逻辑为z = x + y，两个输入与一个输出均为形状[1, N]、数据类型half、格式 ND 的一维 Tensor，核函数名为add_custom。样例固定使用8 个核，主块数据量为3200 个half元素（即 200 个 DataBlock），通过 CMake 编译参数SCENARIO_NUM选择不同数据量场景，共覆盖四种 Tiling 切分场景：

场景	`SCENARIO_NUM`	输入元素数	每核处理元素数	核内组成	切分特点
主块均分	0	256000	32000	10 个主块	各核数据量相同，核内只有主块
尾块均分	1	260096	32512	10 个主块 + 512 元素尾块	各核数据量相同，尾块等长
尾核切分	2	256064	前 4 核 32016 / 后 4 核 32000	整核：10 主块 + 16 元素尾块；尾核：10 主块	前 4 核为整核，尾核无尾块
尾核尾块切分	3	258112	前 4 核 32272 / 后 4 核 32256	整核：10 主块 + 272 元素尾块；尾核：10 主块 + 256 元素尾块	所有核都有尾块，整核多处理 1 个 DataBlock

主块（mainTileLength）是用户基于 UB 可用空间确定的一次可以处理的最大数据量，且满足 32 字节对齐。

3. 核心常量与 Tiling 结构体设计

3.1 常量定义

样例在 vector_tiling_strategy.asc 开头定义了全部切分基础常量：

constexpr uint32_t scenarioNum = SCENARIO_NUM; // 场景编号，由CMake编译参数注入 constexpr uint32_t numBlocks = 8; // 参与计算的逻辑核数 constexpr uint32_t dataBlockSize = 32; // 一个DataBlock的字节数 constexpr uint32_t halfSize = 2; // 一个half元素的字节数 constexpr uint32_t alignNum = dataBlockSize / halfSize; // 一个DataBlock包含的half元素数量，即16 constexpr uint32_t mainTileLength = 3200; // 一个主块包含的half元素数量

其中alignNum = 16是全部分支判断的核心粒度：任何剩余数据都以"一个 DataBlock 含 16 个half"为单位进行整除与取余。

3.2 四种 Tiling 结构体

不同场景需要传递的切分信息不同，样例定义了四种 Tiling 结构体（见 vector_tiling_strategy.asc）：

主块均分场景：各核数据量相同，仅需每核长度与主块个数。

struct MainTileOnlyTiling { uint32_t blockLength; // 每个核处理的元素数量 uint32_t tileNum; // 主块数量 };

尾块均分场景：在上一结构体基础上新增尾块长度。

struct MainTileWithTailBlockTiling { uint32_t blockLength; // 每个核处理的元素数量 uint32_t tileNum; // 主块数量 uint32_t lastTileLength; // 尾块元素数量 };

尾核切分场景：整核与尾核参数不同，用former*描述整核、tail*描述尾核；尾核只含主块，无需tailLastTileLength。

struct MainTileWithTailCoreTiling { uint32_t formerNum; // 整核数量 uint32_t formerLength; // 整核处理的元素数量 uint32_t formerTileNum; // 整核中主块的数量 uint32_t formerLastTileLength; // 整核中尾块的元素数量 uint32_t tailLength; // 尾核处理的元素数量 uint32_t tailTileNum; // 尾核中主块的数量 };

尾核尾块切分场景：尾核也存在尾块，补充tailLastTileLength。

struct TailBlockAndTailCoreTiling { uint32_t formerNum; // 整核数量 uint32_t formerLength; // 整核处理的元素数量 uint32_t formerTileNum; // 整核中主块的数量 uint32_t formerLastTileLength; // 整核中尾块的元素数量 uint32_t tailLength; // 尾核处理的元素数量 uint32_t tailTileNum; // 尾核中主块的数量 uint32_t tailLastTileLength; // 尾核中尾块的元素数量 };

4. 四种切分场景详解

4.1 场景 0：主块均分（Main Tile Equal Split）

输入：x、y均为[1, 256000]个half元素；输出：z为[1, 256000]个half元素。
Tiling 切分：8 个核各处理 32000 个元素；每个核包含 10 个主块，每个主块 3200 个元素。
说明：所有核处理数据量相同，核内只包含主块。

数据量校验：256000 / (8 × 3200) = 10，整除无余数，因此每个核恰好分到 10 个主块。对应的MainTileOnlyTiling结构体取值为{32000, 10}。

4.2 场景 1：尾块均分（Tail Block Equal Split）

输入：x、y均为[1, 260096]个half元素；输出：z为[1, 260096]个half元素。
Tiling 切分：8 个核各处理 32512 个元素；每个核包含 10 个主块和 512 个元素的尾块。
说明：所有核处理数据量相同，每个核在主块后都有等长尾块。

数据量校验：260096 = 8 × 32512，其中32512 = 10 × 3200 + 512，512 = 32 × 16，即每核多分 32 个 DataBlock 作为等长尾块。MainTileWithTailBlockTiling结构体取值为{32512, 10, 512}。

4.3 场景 2：尾核切分（Tail Core Split）

输入：x、y均为[1, 256064]个half元素；输出：z为[1, 256064]个half元素。
Tiling 切分：前 4 个整核各处理 32016 个元素，后 4 个尾核各处理 32000 个元素；整核包含 10 个主块和 16 个元素的尾块，尾核包含 10 个主块。
说明：前 4 个整核比尾核多处理 1 个 DataBlock，尾核无尾块。

数据量校验：256064 = 8 × 32000 + 64，剩余 64 个元素即 4 个 DataBlock，只能分给前 4 个核（每核多 1 个 DataBlock，即 16 个元素），因此formerNum = 4。MainTileWithTailCoreTiling结构体取值为{4, 32016, 10, 16, 32000, 10}。

4.4 场景 3：尾核尾块切分（Tail Core Split with Tail Blocks）

输入：x、y均为[1, 258112]个half元素；输出：z为[1, 258112]个half元素。
Tiling 切分：前 4 个整核各处理 32272 个元素，后 4 个尾核各处理 32256 个元素；整核包含 10 个主块和 272 个元素的尾块，尾核包含 10 个主块和 256 个元素的尾块。
说明：所有核都有尾块，前 4 个整核再多处理 1 个 DataBlock。

数据量校验：258112 = 8 × 32000 + 2112；2112 = 8 × 256 + 64，即每核先分 16 个 DataBlock 基础尾块（256 元素），剩余 64 元素即 4 个 DataBlock 再分给前 4 核。TailBlockAndTailCoreTiling结构体取值为{4, 32272, 10, 272, 32256, 10, 256}。

四种场景的输入输出规格表（名称 / shape / 数据类型 / 格式）汇总如下，均满足"输入输出同为 ND 格式的half，核函数名add_custom"：

场景	输入 x / y shape	输出 z shape	数据类型	格式	核函数名
0	[1, 256000]	[1, 256000]	half	ND	add_custom
1	[1, 260096]	[1, 260096]	half	ND	add_custom
2	[1, 256064]	[1, 256064]	half	ND	add_custom
3	[1, 258112]	[1, 258112]	half	ND	add_custom

5. Host 侧 Tiling 计算流程

Host 侧输入为totalLength，逐级计算主块数量、尾块数量、尾核数量，最终输出 Tiling 切分结果。完整实现位于 vector_tiling_strategy.asc 的LaunchAddByTiling函数中。

第 0 步：按 DataBlock 对齐总数据量。

uint32_t totalLengthAligned = AlignUp(totalLength, alignNum);

其中AlignUp实现为((value + align - 1U) / align) * align。本样例四种场景的totalLength均为alignNum=16的整数倍，因此对齐前后数值不变。

第 1 步：在所有核上均分相同数量的主块。

uint32_t mainTileNum = totalLengthAligned / (numBlocks * mainTileLength); uint32_t mainTileRemainder = totalLengthAligned % (numBlocks * mainTileLength);

此时每个核至少处理mainTileNum * mainTileLength个元素，mainTileRemainder是主块分配后的剩余数据量。当mainTileRemainder == 0时进入主块均分场景并结束：

if (mainTileRemainder == 0U) { MainTileOnlyTiling tiling{baseLength, baseTileNum}; printf("Tiling mode: main_tile_only\n"); add_custom<MainTileOnlyTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); return; }

第 2 步：mainTileRemainder不为 0 时，按 DataBlock 给所有核补齐等长尾块。

uint32_t tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); uint32_t remainingTailLength = mainTileRemainder % (numBlocks * alignNum); uint32_t baseLength = mainTileNum * mainTileLength + tailBlockNumEachCore * alignNum;

baseLength是每个核处理的基础数据量；remainingTailLength是每核补齐等长 DataBlock 尾块后尚未分配的元素数量，其值一定是alignNum的整数倍。

第 3 步：判断尾块均分场景。remainingTailLength == 0表示所有核均分到等长尾块，进入尾块均分场景：

if (remainingTailLength == 0U) { MainTileWithTailBlockTiling tiling{baseLength, baseTileNum, baseLastTileLength}; printf("Tiling mode: main_tile_with_tail_block\n"); add_custom<MainTileWithTailBlockTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); return; }

第 4 步：remainingTailLength不为 0 时，将最后剩余的 DataBlock 分给前若干个整核，并计算整核与尾核的数据长度。

uint32_t formerNum = remainingTailLength / alignNum; uint32_t formerLength = baseLength + alignNum; // 前formerNum个核都多处理一个DataBlock的数据 uint32_t tailLength = baseLength; // 尾核处理的数据量

第 5 步：判断尾核切分 vs 尾核尾块切分。tailBlockNumEachCore == 0时各核未分配到基础尾块，尾核内只包含主块，进入尾核切分场景；否则进入尾核尾块切分场景：

if (tailBlockNumEachCore == 0U) { uint32_t formerLastTileLength = alignNum; MainTileWithTailCoreTiling tiling{formerNum, formerLength, formerTileNum, formerLastTileLength, tailLength, tailTileNum}; printf("Tiling mode: main_tile_with_tail_core\n"); add_custom<MainTileWithTailCoreTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); } else { uint32_t formerLastTileLength = baseLastTileLength + alignNum; if (formerLastTileLength == mainTileLength) { // 若整核尾块恰好凑满一个主块，则合并为主块 formerTileNum += 1U; formerLastTileLength = 0U; } uint32_t tailLastTileLength = baseLastTileLength; TailBlockAndTailCoreTiling tiling{formerNum, formerLength, formerTileNum, formerLastTileLength, tailLength, tailTileNum, tailLastTileLength}; printf("Tiling mode: tail_block_and_tail_core\n"); add_custom<TailBlockAndTailCoreTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); }

整核尾块长度的合并处理是一个容易被忽略的细节：当baseLastTileLength + alignNum恰好等于mainTileLength时，整核的"主块 + 尾块"会退化为多一个完整主块，此时应把formerLastTileLength置 0、formerTileNum加 1，保证核内循环按主块粒度处理。

以四个场景代入上述流程，可完整还原结构体取值：

场景	关键中间量	命中分支	结构体取值
0	`mainTileRemainder = 0`	main_tile_only	`{32000, 10}`
1	`mainTileRemainder = 4096`，`remainingTailLength = 0`	main_tile_with_tail_block	`{32512, 10, 512}`
2	`remainingTailLength = 64`，`tailBlockNumEachCore = 0`	main_tile_with_tail_core	`{4, 32016, 10, 16, 32000, 10}`
3	`remainingTailLength = 64`，`tailBlockNumEachCore = 16`	tail_block_and_tail_core	`{4, 32272, 10, 272, 32256, 10, 256}`

6. Device 侧核函数实现

6.1 核函数入口与模板化 Tiling 传递

核函数add_custom是模板函数，TilingData类型即第 3 节定义的四种结构体之一。Host 侧调用add_custom<TilingData><<<numBlocks, 0, stream>>>时，Tiling 结构体按值作为核函数参数传入：

template <typename TilingData> __vector__ __global__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, TilingData tiling) { AscendC::InitSocState(); KernelAdd op; op.Init(x, y, z, tiling); op.Process(); AscendC::PipeBarrier<PIPE_ALL>(); }

6.2 Init：按核索引计算 Global Memory 偏移

每个核访问的 Global Memory 地址范围不同，Init通过 GetBlockIdx() 获取当前逻辑 AI Core 索引，并结合 Tiling 参数计算起始偏移。GetBlockIdx 的返回值范围为[0, numBlocks)，numBlocks由<<<>>>调用配置。

主块均分 / 尾块均分（各核数据量相同）：偏移为blockIdx × blockLength，尾块均分场景额外保存lastTileLength。

__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, MainTileWithTailBlockTiling tiling) { uint32_t offset = AscendC::GetBlockIdx() * tiling.blockLength; InitGm(x, y, z, offset, tiling.blockLength); this->tileNum = tiling.tileNum; this->lastTileLength = tiling.lastTileLength; }

尾核切分 / 尾核尾块切分（整核与尾核数据量不同）：先判断当前核是否为整核（blockIdx < tiling.formerNum），整核偏移为formerLength × blockIdx；尾核偏移需跳过全部整核数据，为formerLength × formerNum + tailLength × (blockIdx - formerNum)：

__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, TailBlockAndTailCoreTiling tiling) { uint32_t blockIdx = AscendC::GetBlockIdx(); if (blockIdx < tiling.formerNum) { // 整核 uint32_t offset = tiling.formerLength * blockIdx; InitGm(x, y, z, offset, tiling.formerLength); this->tileNum = tiling.formerTileNum; this->lastTileLength = tiling.formerLastTileLength; } else { // 尾核 uint32_t offset = tiling.formerLength * tiling.formerNum + tiling.tailLength * (blockIdx - tiling.formerNum); InitGm(x, y, z, offset, tiling.tailLength); this->tileNum = tiling.tailTileNum; this->lastTileLength = tiling.tailLastTileLength; } }

InitGm通过 SetGlobalBuffer 为输入xGm、yGm和输出zGm设置 Global Memory 地址：(__gm__ half*)x + offset为当前核数据的起始地址，length为该核处理的元素数量。

6.3 Process：分块循环搬运、计算与同步

Process采用静态 Tensor 编程方式实现（参考 static_tensor_programming.md）：通过LocalTensor构造函数直接指定 VECCALC 位置上的固定地址，避免 TPipe 框架初始化开销。xAddr、yAddr、zAddr依次线性排布在 UB 上：

static constexpr uint32_t xAddr = 0; static constexpr uint32_t yAddr = mainTileLength * halfSize; static constexpr uint32_t zAddr = yAddr + mainTileLength * halfSize;

循环次数采用通用形式tileNum + (lastTileLength > 0 ? 1 : 0)计算：前tileNum轮处理主块，最后一轮（若存在）处理尾块；curLen为当前轮元素数，startElement = loopIdx * mainTileLength为核内偏移。由于静态 Tensor 编程需要开发者自行管理流水同步，Process内通过SetFlag/WaitFlag维护 MT E2 搬入、Vector 计算、MTE3 搬出三条流水之间的依赖：

__aicore__ inline void Process() { AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, mainTileLength); AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, mainTileLength); AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, mainTileLength); uint32_t totalLoop = this->tileNum + (this->lastTileLength > 0U ? 1U : 0U); for (uint32_t loopIdx = 0; loopIdx < totalLoop; loopIdx++) { uint32_t curLen = (loopIdx < this->tileNum) ? mainTileLength : this->lastTileLength; uint32_t startElement = loopIdx * mainTileLength; if (loopIdx != 0U) { AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); // 反向同步：上一次计算完成后再搬入 } AscendC::DataCopy(xLocal, xGm[startElement], curLen); // 搬入输入x AscendC::DataCopy(yLocal, yGm[startElement], curLen); // 搬入输入y AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0); // 正向同步：搬入完成后再计算 if (loopIdx != 0U) { AscendC::WaitFlag<AscendC::HardEvent::MTE3_V>(EVENT_ID0); // 反向同步：上一次搬出完成后再计算 } AscendC::Add(zLocal, xLocal, yLocal, curLen); // 矢量加计算 if (loopIdx != totalLoop - 1U) { AscendC::SetFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); } AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0); // 正向同步：计算完成后再搬出 AscendC::DataCopy(zGm[startElement], zLocal, curLen); // 搬出输出z if (loopIdx != totalLoop - 1U) { AscendC::SetFlag<AscendC::HardEvent::MTE3_V>(EVENT_ID0); } } }

这里的同步模式与多核 Tiling 概述文档中的描述一致：正向同步（MTE2_V、V_MTE3）保证"先搬入后计算、先计算后搬出"，反向同步（V_MTE2、MTE3_V）防止上一轮数据被覆盖。四种场景共用同一个Process函数，差异全部由tileNum与lastTileLength两个成员变量体现，这正是 Tiling 结构体设计的价值所在。

7. 编译运行与结果验证

7.1 环境准备

根据当前环境上 CANN 开发套件包的安装方式配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

7.2 样例执行

在样例根目录（vector_tiling_strategy）下执行：

SCENARIO_NUM=0 # 设置场景编号 mkdir -p build && cd build; # 创建并进入build目录 cmake .. -DCMAKE_ASC_ARCHITECTURES=dav-2201 -DSCENARIO_NUM=$SCENARIO_NUM;make -j; # 编译工程 python3 ../scripts/gen_data.py -scenarioNum=$SCENARIO_NUM # 生成测试输入数据 ./demo # 执行编译生成的可执行程序，执行样例 python3 ../scripts/verify_result.py ./output/output.bin ./output/golden.bin # 验证输出结果是否正确

使用 CPU 调试或 NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=cpu或-DCMAKE_ASC_RUN_MODE=sim参数：

cmake .. -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 -DSCENARIO_NUM=$SCENARIO_NUM;make -j; # CPU调试模式 cmake .. -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 -DSCENARIO_NUM=$SCENARIO_NUM;make -j; # NPU仿真模式

注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行rm CMakeCache.txt后重新 cmake。

7.3 编译选项说明

编译选项的默认值与取值校验定义在 CMakeLists.txt 中，SCENARIO_NUM通过target_compile_definitions以宏定义形式注入源码（constexpr uint32_t scenarioNum = SCENARIO_NUM;），且 CMake 侧做了取值校验（SCENARIO_NUM必须匹配^[0-3]$，否则FATAL_ERROR）：

选项	可选值	说明
`CMAKE_ASC_RUN_MODE`	`npu`（默认）、`cpu`、`sim`	运行模式：NPU 运行、CPU 调试、NPU 仿真
`CMAKE_ASC_ARCHITECTURES`	`dav-2201`（默认）、`dav-3510`	NPU 架构：`dav-2201`对应 Atlas A2 系列产品和 Atlas A3 系列产品，`dav-3510`对应 Ascend 950PR & 950DT 系列产品
`SCENARIO_NUM`	`0`（默认）、`1`、`2`、`3`	场景编号：0（主块均分）、1（尾块均分）、2（尾核切分）、3（尾核尾块切分）

7.4 数据生成与精度验证

数据生成脚本 gen_data.py 内置了四种场景的数据长度映射SCENARIO_LENGTHS = {0: 256000, 1: 260096, 2: 256064, 3: 258112}，以固定随机种子（2026 + scenario）在[-8, 8]均匀分布中采样float16数据，生成input_x.bin、input_y.bin输入与golden.bin真值。

验证脚本 verify_result.py 读取output.bin与golden.bin，使用np.isclose按相对容差1e-4、绝对容差1e-5逐元素比对，统计错误率并允许ERROR_TOL = 1e-4以内的误差（half 浮点运算的舍入误差）。执行结果如下即说明精度对比成功：

test pass!

7.5 主函数执行流程

vector_tiling_strategy.asc 中的main函数展示了完整的 ACL 运行时调用链：aclInit→aclrtSetDevice→aclrtCreateStream→aclrtMallocHost/aclrtMalloc申请内存 → data_utils.h 中的ReadFile读取输入 →aclrtMemcpy拷贝至 Device →LaunchAddByTiling发起计算 →aclrtSynchronizeStream等待完成 → 拷回 Host 并WriteFile落盘 → 释放资源并aclFinalize。运行时会打印当前场景与数据量：

Usage: ./demo (compile-time SCENARIO_NUM=0) Running scenario 0, totalLength=256000 Tiling mode: main_tile_only Scenario 0 completed.

8. 策略选择准则与扩展思考

综合第 5 节的逐级判断，四种 Tiling 切分策略的选择完全由两个取余结果驱动，可归纳为如下决策树：

mainTileRemainder == 0→主块均分（MainTileOnlyTiling），核内仅主块，各核完全均衡；
remainingTailLength == 0→尾块均分（MainTileWithTailBlockTiling），各核均衡，核内含等长尾块；
tailBlockNumEachCore == 0→尾核切分（MainTileWithTailCoreTiling），出现整核/尾核，尾核无尾块；
其余情况 →尾核尾块切分（TailBlockAndTailCoreTiling），整核与尾核都含尾块。

该决策逻辑可以直接移植到真实算子开发中：Tiling 结构体按需扩展（例如整核/尾核参数分组）、Host 侧以 DataBlock（alignNum）为粒度做整除取余、Device 侧用GetBlockIdx()分支计算偏移，这三步是所有多核矢量算子 Tiling 的通用范式。进一步扩展时，可参考仓库中 multi_core_tiling 的四个分册（主块均分、尾块均分、尾核切分、尾核尾块切分）了解每个场景更细粒度的实现推导；矢量计算 API（如 Add）与静态 Tensor 编程的同步管理细节可分别查阅对应文档。若要在自有工程中落地，可将样例的LaunchAddByTiling泛化为"输入 shape → 切分参数 → 核函数调用"的工具函数，并根据实际 UB 容量重新计算mainTileLength。