已合并
【docs】更新多核tiling 算子实践 #4157
Chen Ning创建于 7月8日
【docs】更新多核tiling 算子实践 #4157
已合并
共 24 个文件变更+702-847
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:5c7d3b244bf8a48c90c6b4853286475db328eb001564c7c0691f68ac993fbfd8 | ||
| 3 | +size 11083 | ||
Rdocs/zh/guide/figures/多核及Tiling示意图.png→docs/zh/guide/figures/operator_implementation_components.png+2-2
| @@ -1,3 +1,3 @@ | |||
| 1 | version https://git-lfs.github.com/spec/v1 | 1 | version https://git-lfs.github.com/spec/v1 |
| 2 | -oid sha256:0cb62b9830229bae4ce03acda937a1c79c4675941337a9de1e4fed3ae5530dc8 | 2 | +oid sha256:681bd210b51a623f4f22784385d4aebea8830114688f08fb6ac9b890a8613a3e |
| 3 | -size 9923 | 3 | +size 6943 |
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:a8a8028d39c3b0ab8a8f217e2257fff21dfefe880693e5ffd4ffa0f70e42f005 | ||
| 3 | +size 28373 | ||
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:bf71a885875cb6a8f72f58e79963ff15816139658f2182af7a7f1568ba429cc8 | ||
| 3 | +size 19101 | ||
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:7aec122245c269e69acaf7bf844d7a8d20f599d062fa7ac5dc934cdfed5ea75e | ||
| 3 | +size 22892 | ||
| @@ -0,0 +1,3 @@ | |||
| 1 | +version https://git-lfs.github.com/spec/v1 | ||
| 2 | +oid sha256:88b99a303d1c5097363433f433e7108b1c69cf635646b1e7f35210abd2fc2bce | ||
| 3 | +size 23576 | ||
| @@ -199,12 +199,12 @@ | |||
| 199 | - [概述](算子实践参考/SIMD算子实现/矢量编程/概述.md) | 199 | - [概述](算子实践参考/SIMD算子实现/矢量编程/概述.md) |
| 200 | - [基础矢量算子](算子实践参考/SIMD算子实现/矢量编程/基础矢量算子.md) | 200 | - [基础矢量算子](算子实践参考/SIMD算子实现/矢量编程/基础矢量算子.md) |
| 201 | - [TBuf的使用](算子实践参考/SIMD算子实现/矢量编程/TBuf的使用.md) | 201 | - [TBuf的使用](算子实践参考/SIMD算子实现/矢量编程/TBuf的使用.md) |
| 202 | - - [多核&Tiling切分](算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/概述.md) | 202 | + - [多核Tiling切分](算子实践参考/SIMD算子实现/矢量编程/多核Tiling切分/概述.md) |
| 203 | - - [概述](算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/概述.md) | 203 | + - [概述](算子实践参考/SIMD算子实现/矢量编程/多核Tiling切分/概述.md) |
| 204 | - - [多核Tiling](算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/多核Tiling.md) | 204 | + - [主块均分](算子实践参考/SIMD算子实现/矢量编程/多核Tiling切分/主块均分.md) |
| 205 | - - [尾块Tiling](算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/尾块Tiling.md) | 205 | + - [尾块均分](算子实践参考/SIMD算子实现/矢量编程/多核Tiling切分/尾块均分.md) |
| 206 | - - [尾核Tiling](算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/尾核Tiling.md) | 206 | + - [尾核切分](算子实践参考/SIMD算子实现/矢量编程/多核Tiling切分/尾核切分.md) |
| 207 | - - [尾核&尾块](算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/尾核-尾块.md) | 207 | + - [尾核尾块切分](算子实践参考/SIMD算子实现/矢量编程/多核Tiling切分/尾核尾块切分.md) |
| 208 | 208 | ||
| 209 | - [DoubleBuffer场景](算子实践参考/SIMD算子实现/矢量编程/DoubleBuffer场景.md) | 209 | - [DoubleBuffer场景](算子实践参考/SIMD算子实现/矢量编程/DoubleBuffer场景.md) |
| 210 | - [Broadcast场景](算子实践参考/SIMD算子实现/矢量编程/Broadcast场景.md) | 210 | - [Broadcast场景](算子实践参考/SIMD算子实现/矢量编程/Broadcast场景.md) |
| @@ -20,7 +20,7 @@ | |||
| 20 | totalLength : ((totalLength + alignNum - 1) / alignNum) * alignNum; | 20 | totalLength : ((totalLength + alignNum - 1) / alignNum) * alignNum; |
| 21 | ``` | 21 | ``` |
| 22 | 22 | ||
| 23 | -2. 根据totalLengthAligned，计算每个核的计算数据长度blockLength，分核策略可参照[尾核Tiling](多核-Tiling切分/尾核Tiling.md)。 | 23 | +2. 根据totalLengthAligned，计算每个核的计算数据长度blockLength，分核策略可参照[尾核切分](多核Tiling切分/尾核切分.md)。 |
| 24 | 3. 计算其余Tiling参数。 | 24 | 3. 计算其余Tiling参数。 |
| 25 | 25 | ||
| 26 | 对当前Unified Buffer可用空间以32字节为粒度，进行切分，计算出数据块个数UB\_BLOCK\_NUM。根据是否开启DoubleBuffer计算出当前可用的最大数据块个数，记作MAX\_AVAILABLE\_UB\_BLOCK\_NUM。最后，以MAX\_AVAILABLE\_UB\_BLOCK\_NUM为粒度，对blockLength进行切分。为方便演示，如下代码直接给出UB\_BLOCK\_NUM，作为当前Unified Buffer可用空间包含的block（32字节）数。 | 26 | 对当前Unified Buffer可用空间以32字节为粒度，进行切分，计算出数据块个数UB\_BLOCK\_NUM。根据是否开启DoubleBuffer计算出当前可用的最大数据块个数，记作MAX\_AVAILABLE\_UB\_BLOCK\_NUM。最后，以MAX\_AVAILABLE\_UB\_BLOCK\_NUM为粒度，对blockLength进行切分。为方便演示，如下代码直接给出UB\_BLOCK\_NUM，作为当前Unified Buffer可用空间包含的block（32字节）数。 |
| @@ -115,7 +115,7 @@ __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t | |||
| 115 | } | 115 | } |
| 116 | ``` | 116 | ``` |
| 117 | 117 | ||
| 118 | -由于开启DoubleBuffer后，切分后的主块数据块个数翻倍，在Process函数中，需要将BUFFER\_NUM带入计算循环次数；尾块独立计算，不开启DoubleBuffer。后续主尾块在CopyIn、Compute、CopyOut阶段中的处理，与[尾块tiling处理](多核-Tiling切分/尾块Tiling.md)相同。 | 118 | +由于开启DoubleBuffer后，切分后的主块数据块个数翻倍，在Process函数中，需要将BUFFER\_NUM带入计算循环次数；尾块独立计算，不开启DoubleBuffer。后续主尾块在CopyIn、Compute、CopyOut函数中的处理，与[尾块均分](多核Tiling切分/尾块均分.md)相同。 |
| 119 | 119 | ||
| 120 | ``` | 120 | ``` |
| 121 | __aicore__ inline void Process() | 121 | __aicore__ inline void Process() |
Ddocs/zh/guide/算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/多核Tiling.md+0-219
| @@ -1,219 +0,0 @@ | |||
| 1 | -# 多核Tiling<a name="ZH-CN_TOPIC_0000002532228159"></a> | ||
| 2 | - | ||
| 3 | -基于Ascend C方式实现带有Tiling的算子的开发流程如下图所示。 | ||
| 4 | - | ||
| 5 | -**图1** 算子开发流程<a name="zh-cn_topic_0000002236197677_fig18330201774715"></a> | ||
| 6 | - | ||
| 7 | - | ||
| 8 | -## 算子分析<a name="zh-cn_topic_0000002236197677_section179012044886"></a> | ||
| 9 | - | ||
| 10 | -本样例为输入数据在核间均分、核内均分场景。本样例的Tiling策略为：数据整体长度TOTAL\_LENGTH为8 \* 2048，数据平均分配到8个核上运行，每个核上计算的数据长度BLOCK\_LENGTH为2048，将单核上的数据切分成16块（此处切分成16块仅用来作为Tiling的样例，并不代表性能最佳，仅供参考），每块数据的长度TILE\_LENGTH为128。数据切分示意如下图所示： | ||
| 11 | - | ||
| 12 | -**图2** 数据切分示意图<a name="zh-cn_topic_0000002236197677_fig1986021174914"></a> | ||
| 13 | - | ||
| 14 | - | ||
| 15 | -通过以上分析，得到Ascend C Add算子的设计规格如下： | ||
| 16 | - | ||
| 17 | -- 算子类型（OpType）：Add | ||
| 18 | -- 算子输入输出： | ||
| 19 | - | ||
| 20 | - **表1** Add算子输入输出规格 | ||
| 21 | - | ||
| 22 | - <a name="table4934296305"></a> | ||
| 23 | - <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p> | ||
| 24 | - </th> | ||
| 25 | - <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p> | ||
| 26 | - </th> | ||
| 27 | - <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.3"><p id="p1950391883014"><a name="p1950391883014"></a><a name="p1950391883014"></a><strong id="b2503111803020"><a name="b2503111803020"></a><a name="b2503111803020"></a>data type</strong></p> | ||
| 28 | - </th> | ||
| 29 | - <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.4"><p id="p14503218133015"><a name="p14503218133015"></a><a name="p14503218133015"></a><strong id="b8503141818301"><a name="b8503141818301"></a><a name="b8503141818301"></a>format</strong></p> | ||
| 30 | - </th> | ||
| 31 | - </tr> | ||
| 32 | - </thead> | ||
| 33 | - <tbody><tr id="row393589203016"><td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.1 "><p id="p1950331810308"><a name="p1950331810308"></a><a name="p1950331810308"></a>x（输入）</p> | ||
| 34 | - </td> | ||
| 35 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.2 "><p id="p12328141215344"><a name="p12328141215344"></a><a name="p12328141215344"></a>(8, 2048)</p> | ||
| 36 | - </td> | ||
| 37 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.3 "><p id="p135031118173010"><a name="p135031118173010"></a><a name="p135031118173010"></a>half</p> | ||
| 38 | - </td> | ||
| 39 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.4 "><p id="p19503131815305"><a name="p19503131815305"></a><a name="p19503131815305"></a>ND</p> | ||
| 40 | - </td> | ||
| 41 | - </tr> | ||
| 42 | - <tr id="row6935119173013"><td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.1 "><p id="p75031182305"><a name="p75031182305"></a><a name="p75031182305"></a>y（输入）</p> | ||
| 43 | - </td> | ||
| 44 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.2 "><p id="p532871211344"><a name="p532871211344"></a><a name="p532871211344"></a>(8, 2048)</p> | ||
| 45 | - </td> | ||
| 46 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.3 "><p id="p45031818103018"><a name="p45031818103018"></a><a name="p45031818103018"></a>half</p> | ||
| 47 | - </td> | ||
| 48 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.4 "><p id="p1503918103012"><a name="p1503918103012"></a><a name="p1503918103012"></a>ND</p> | ||
| 49 | - </td> | ||
| 50 | - </tr> | ||
| 51 | - <tr id="row59354943016"><td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.1 "><p id="p1450316186305"><a name="p1450316186305"></a><a name="p1450316186305"></a>z（输出）</p> | ||
| 52 | - </td> | ||
| 53 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.2 "><p id="p632831214347"><a name="p632831214347"></a><a name="p632831214347"></a>(8, 2048)</p> | ||
| 54 | - </td> | ||
| 55 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.3 "><p id="p15503918173017"><a name="p15503918173017"></a><a name="p15503918173017"></a>half</p> | ||
| 56 | - </td> | ||
| 57 | - <td class="cellrowborder" valign="top" width="25%" headers="mcps1.2.5.1.4 "><p id="p1503101813012"><a name="p1503101813012"></a><a name="p1503101813012"></a>ND</p> | ||
| 58 | - </td> | ||
| 59 | - </tr> | ||
| 60 | - </tbody> | ||
| 61 | - </table> | ||
| 62 | - | ||
| 63 | -- 核函数名称：tiling\_strategy\_custom | ||
| 64 | -- 使用的主要接口： | ||
| 65 | - - DataCopy：数据搬移接口 | ||
| 66 | - - Add：矢量基础算术接口 | ||
| 67 | - - EnQue、DeQue等接口：Queue队列管理接口 | ||
| 68 | - | ||
| 69 | -- 算子实现文件名称：tiling\_strategy.asc | ||
| 70 | - | ||
| 71 | -## Tiling实现<a name="zh-cn_topic_0000002236197677_section480741815522"></a> | ||
| 72 | - | ||
| 73 | -前述场景中算子的输入和输出均为固定shape，然而在实际的算子开发场景中，这些信息是支持动态变化的，场景会更加灵活和复杂。动态shape场景下，输入的shape是未知的。一些与输入shape相关的变量（比如每次搬运的块大小等），需要通过Tiling计算出来，然后传递到kernel侧，kernel侧使用该参数进行后续的计算。 | ||
| 74 | - | ||
| 75 | -具体实现方式为：分析设计Tiling参数、定义Tiling结构体，在Host侧通过上下文获取输入输出的shape信息，根据shape信息，计算Tiling参数并设置到对应的Tiling结构体中；通过核函数入口参数将Tiling信息传入核函数，在核函数内通过解析Tiling结构体，获取并使用相关参数来实现核函数内部逻辑，详细介绍请参考[Host侧tiling实现](../../../../编程指南/高级编程/Aclnn算子工程化开发/设计与实现/Host侧Tiling实现.md)。本节将以上述分析中的切分策略为例，说明如何实现Tiling。 | ||
| 76 | - | ||
| 77 | -基于本节的切分策略，Tiling需要定义如下参数： | ||
| 78 | - | ||
| 79 | -- blockLength：每个核的计算数据长度； | ||
| 80 | -- tileNum：每个核需要计算的数据块个数； | ||
| 81 | -- tileLength：每个核内每个数据块的长度。 | ||
| 82 | - | ||
| 83 | -根据确定的Tiling参数，使用C++语法定义TilingData结构体，代码如下。 | ||
| 84 | - | ||
| 85 | -``` | ||
| 86 | -struct AddCustomTilingData { | ||
| 87 | - uint32_t blockLength; | ||
| 88 | - uint32_t tileNum; | ||
| 89 | - uint32_t tileLength; | ||
| 90 | - ... | ||
| 91 | -} | ||
| 92 | -``` | ||
| 93 | - | ||
| 94 | -接下来完成Tiling参数的计算。由于每个核内数据被切分为16块，根据使用的核数和核内切分数，计算Tiling参数，并写入到Tiling结构体内。代码示例如下： | ||
| 95 | - | ||
| 96 | -``` | ||
| 97 | -constexpr int32_t NUM_BLOCKS = 8; // 使用的核数 | ||
| 98 | -constexpr int32_t TILE_NUM = 16; // 核内切分数量 | ||
| 99 | -void GenerateTilingData(uint8_t* tilingBuf, uint32_t numBlocks) | ||
| 100 | -{ | ||
| 101 | - uint32_t totalLength; | ||
| 102 | - // 此处省略如何获取数据总长TOTAL_LENGTH，可以根据具体情况实现。本章节仅介绍Tiling相关内容。 | ||
| 103 | - AddCustomTilingData* tiling = reinterpret_cast<AddCustomTilingData *>(tilingBuf); | ||
| 104 | - uint32_t blockLength = TOTAL_LENGTH / numBlocks; | ||
| 105 | - uint32_t tileNum = TILE_NUM; | ||
| 106 | - uint32_t tileLength = blockLength / tileNum; | ||
| 107 | - | ||
| 108 | - tiling->blockLength = blockLength; | ||
| 109 | - tiling->tileNum = tileNum; | ||
| 110 | - tiling->tileLength = tileLength; | ||
| 111 | -} | ||
| 112 | -``` | ||
| 113 | - | ||
| 114 | -最后，在Host侧调用程序中，调用上述Tiling参数计算函数，计算出相关参数，然后传递到Kernel侧核函数。 | ||
| 115 | - | ||
| 116 | -``` | ||
| 117 | - constexpr int32_t NUM_BLOCKS = 8; | ||
| 118 | - ... | ||
| 119 | - uint8_t *tiling = nullptr; | ||
| 120 | - size_t tilingSize = sizeof(AddCustomTilingData); | ||
| 121 | - GenerateTilingData(tiling, NUM_BLOCKS); // 调用tiling参数计算函数 | ||
| 122 | - .... | ||
| 123 | - | ||
| 124 | - tiling_strategy_custom<<<NUM_BLOCKS, 0, stream>>>(xDevice, yDevice, zDevice, | ||
| 125 | - *reinterpret_cast<AddCustomTilingData*>(tiling)); | ||
| 126 | - .... | ||
| 127 | -``` | ||
| 128 | - | ||
| 129 | -## 算子类实现<a name="zh-cn_topic_0000002236197677_section849945172010"></a> | ||
| 130 | - | ||
| 131 | -Kernel侧算子实现仍遵循[矢量算子核函数实现流程](../基础矢量算子.md#zh-cn_topic_0000002201157438_fig16061570280)，接下来重点介绍本场景中算子类实现的不同点。 | ||
| 132 | - | ||
| 133 | -- 设置输入输出Global Tensor的Global Memory内存地址。 | ||
| 134 | - | ||
| 135 | - 由于本样例中将数据分配到了多个核上进行处理，每个核处理不同的数据，因此不同核要处理的数据在Global Memory上的地址不同，在初始化函数Init中，需要获取单核所需处理的输入输出在Global Memory上的内存偏移地址，并将该偏移地址设置到GlobalTensor中。 | ||
| 136 | - | ||
| 137 | - 以获取输入x在Global Memory上的内存偏移地址为例，数据整体长度TOTAL\_LENGTH为8 \* 2048，平均分配到8个核上运行，每个核上处理的数据长度blockLength为2048，调用[GetBlockIdx](../../../../../api/SIMD-API/basic_api/tool_interface/system_resources_and_variables/GetBlockIdx.md)接口获取当前核的index，x + blockLength \* GetBlockIdx\(\)即为单核处理程序中x在Global Memory上的内存偏移地址，获取偏移地址后，使用GlobalTensor类的[SetGlobalBuffer](../../../../../api/SIMD-API/basic_api/data_structures/GlobalTensor/SetGlobalBuffer.md)接口设定该核上Global Memory的起始地址以及长度，具体示意图请参考[图3](#zh-cn_topic_0000002236197677_fig398721711313)。代码如下所示： | ||
| 138 | - | ||
| 139 | - ``` | ||
| 140 | - xGm.SetGlobalBuffer((__gm__ half *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength); | ||
| 141 | - ``` | ||
| 142 | - | ||
| 143 | - **图3** 多核并行处理示意图<a name="zh-cn_topic_0000002236197677_fig398721711313"></a> | ||
| 144 | -  | ||
| 145 | - | ||
| 146 | -- 通过Pipe内存管理对象为输入输出Queue分配内存。 | ||
| 147 | - | ||
| 148 | - 对于单核上的处理数据，可以进行数据切块（Tiling），在本示例中，仅作为参考，将单核上的数据（2048个数）切分成16块（并不意味着16块就是性能最优），每块tileLength（128）个数据。数据切分示意图如[图4](#zh-cn_topic_0000002236197677_fig1319211154719)所示。 | ||
| 149 | - | ||
| 150 | - **图4** 单核数据切分示意图<a name="zh-cn_topic_0000002236197677_fig1319211154719"></a> | ||
| 151 | -  | ||
| 152 | - | ||
| 153 | - 与[基础矢量算子](../基础矢量算子.md)相比，在通过Pipe内存管理对象为输入输出Queue分配内存时，需使用单核内每个数据块的长度tileLength作为分配内存的长度。比如，为输入x的Queue分配内存，可以通过如下代码段实现，Pipe为inQueueX分配了一块大小为tileLength \* sizeof\(half\)个字节的内存块，每个内存块能容纳tileLength（128）个half类型数据。 | ||
| 154 | - | ||
| 155 | - ``` | ||
| 156 | - pipe->InitBuffer(inQueueX, 1, this->tileLength * sizeof(half)) | ||
| 157 | - ``` | ||
| 158 | - | ||
| 159 | -具体的初始化函数代码如下： | ||
| 160 | - | ||
| 161 | -``` | ||
| 162 | -__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, AddCustomTilingData tiling, AscendC::TPipe* pipeIn) | ||
| 163 | -{ | ||
| 164 | - pipe = pipeIn; | ||
| 165 | - this->blockLength = tiling.blockLength; | ||
| 166 | - this->tileNum = tiling.tileNum; | ||
| 167 | - this->tileLength = tiling.tileLength; | ||
| 168 | - // 计算每个核上的地址偏移 | ||
| 169 | - xGm.SetGlobalBuffer((__gm__ half *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength); | ||
| 170 | - yGm.SetGlobalBuffer((__gm__ half *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength); | ||
| 171 | - zGm.SetGlobalBuffer((__gm__ half *)z + this->blockLength * AscendC::GetBlockIdx(), this->blockLength); | ||
| 172 | - // pipe alloc memory to queue, the unit is Bytes | ||
| 173 | - pipe->InitBuffer(inQueueX, 1, this->tileLength * sizeof(half)); | ||
| 174 | - pipe->InitBuffer(inQueueY, 1, this->tileLength * sizeof(half)); | ||
| 175 | - pipe->InitBuffer(outQueueZ, 1, this->tileLength * sizeof(half)); | ||
| 176 | -} | ||
| 177 | -``` | ||
| 178 | - | ||
| 179 | -每个核需要对tileNum个数据块分别进行搬入、计算、搬出处理，因此Process函数内将tileNum作为循环上限。 | ||
| 180 | - | ||
| 181 | -``` | ||
| 182 | -__aicore__ inline void Process() | ||
| 183 | -{ | ||
| 184 | - int32_t loopCount = this->tileNum; | ||
| 185 | - // tiling strategy, pipeline parallel | ||
| 186 | - for (int32_t i = 0; i < loopCount; i++) { | ||
| 187 | - CopyIn(i, this->tileLength); | ||
| 188 | - Compute(i, this->tileLength); | ||
| 189 | - CopyOut(i, this->tileLength); | ||
| 190 | - } | ||
| 191 | -} | ||
| 192 | -``` | ||
| 193 | - | ||
| 194 | -对应的，每个核内搬入、搬出每个数据块时，需定位到每个数据块所在Global Memory上的内存偏移地址，因此在CopyIn和CopyOut函数内部使用DataCopy接口时，需增加每个数据块的地址偏移。Compute函数没有变化，与[基础矢量算子](../基础矢量算子.md)相同。 | ||
| 195 | - | ||
| 196 | -CopyIn函数实现代码如下： | ||
| 197 | - | ||
| 198 | -``` | ||
| 199 | -__aicore__ inline void CopyIn(int32_t progress, uint32_t tileLength) | ||
| 200 | -{ | ||
| 201 | - ... | ||
| 202 | - // copy progress_th tile from global tensor to local tensor | ||
| 203 | - AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], tileLength); | ||
| 204 | - AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], tileLength); | ||
| 205 | - ... | ||
| 206 | -} | ||
| 207 | -``` | ||
| 208 | - | ||
| 209 | -CopyOut函数实现代码如下： | ||
| 210 | - | ||
| 211 | -``` | ||
| 212 | - __aicore__ inline void CopyOut(int32_t progress, uint32_t tileLength) | ||
| 213 | -{ | ||
| 214 | - ... | ||
| 215 | - // copy progress_th tile from local tensor to global tensor | ||
| 216 | - AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, tileLength); | ||
| 217 | - ... | ||
| 218 | -} | ||
| 219 | -``` | ||
Ddocs/zh/guide/算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/尾块Tiling.md+0-207
| @@ -1,207 +0,0 @@ | |||
| 1 | -# 尾块Tiling<a name="ZH-CN_TOPIC_0000002500548090"></a> | ||
| 2 | - | ||
| 3 | -如下图中的示例，算子的输入shape为（1，2048），支持的数据类型为half类型，输入数据可以对齐到一个datablock的大小（32字节），输入数据为2048 \* 2 / 32 = 128个datablock，因此可以平均分配到每个核上（假设使用8个核），每个核上处理256个数，16个datablock。此时不需要进行尾块处理。 | ||
| 4 | - | ||
| 5 | -**图1** shape对齐场景<a name="zh-cn_topic_0000002201157442_fig55115212126"></a> | ||
| 6 | - | ||
| 7 | - | ||
| 8 | -针对一些shape，比如算子的输入shape为（1，1904），支持的数据类型为half类型，输入数据可以对齐到一个datablock的大小（32字节），可以平均分配到每个核上（假设使用8个核），每个核上处理238个数，238个数无法均分到datablock上，分满14个datablock后，剩余14个数（28字节），多核切分后需要进行尾块处理。 | ||
| 9 | - | ||
| 10 | -对于不同shape的输入进行数据切分时，可能会发生Tiling后的数据平均分配到多核上，但每个核内的数据无法均分的情况。针对此种场景，在Tiling参数中增加变量lastTileLength，用来表示最后一个分块，即尾块的大小。因此，在定义算子的Tiling结构体时包含以下四个成员： | ||
| 11 | - | ||
| 12 | -- blockLength：每个核上计算的数据长度； | ||
| 13 | -- tileNum：每个核上切分的主块数据块的个数； | ||
| 14 | -- tileLength：每个核上主块数据块的长度； | ||
| 15 | -- lastTileLength：每个核上尾块的长度。 | ||
| 16 | - | ||
| 17 | -**图2** 多核Tiling尾块示意图<a name="zh-cn_topic_0000002201157442_fig48081401102"></a> | ||
| 18 | - | ||
| 19 | - | ||
| 20 | -## Tiling实现<a name="zh-cn_topic_0000002201157442_section469034017306"></a> | ||
| 21 | - | ||
| 22 | -算子的Tiling结构体定义如下： | ||
| 23 | - | ||
| 24 | -``` | ||
| 25 | -struct AddCustomTilingData { | ||
| 26 | - uint32_t blockLength; | ||
| 27 | - uint32_t tileNum; | ||
| 28 | - uint32_t tileLength; | ||
| 29 | - uint32_t lastTileLength; | ||
| 30 | - ... | ||
| 31 | -}; | ||
| 32 | -``` | ||
| 33 | - | ||
| 34 | -Host侧Tiling实现的主要内容为计算以上四个成员变量。步骤如下： | ||
| 35 | - | ||
| 36 | -1. 判断数据总长度totalLength是否满足32字节对齐，如不满足，则计算totalLength向上32字节对齐后的长度totalLengthAligned。 | ||
| 37 | - | ||
| 38 | - ``` | ||
| 39 | - constexpr uint32_t BLOCK_SIZE = 32; | ||
| 40 | - // 为方便计算，这里根据数据类型定义变量alignNum作为对齐数 | ||
| 41 | - uint32_t alignNum = BLOCK_SIZE / dataTypeSize; | ||
| 42 | - // totalLength为数据总量 | ||
| 43 | - totalLengthAligned = (totalLength % alignNum == 0U) ? | ||
| 44 | - static_cast<uint32_t>(totalLength) : | ||
| 45 | - ((static_cast<uint32_t>(totalLength) + alignNum - 1) / alignNum) * alignNum; | ||
| 46 | - ``` | ||
| 47 | - | ||
| 48 | -2. 判断totalLengthAligned是否能被使用的核数NumBlocks均分，如果可以，则计算每个核上计算数据长度blockLength。 | ||
| 49 | - | ||
| 50 | - ``` | ||
| 51 | - constexpr uint32_t NUM_BLOCKS = 8; | ||
| 52 | - constexpr uint32_t UB_BLOCK_NUM = 100; // 此处为方便验证，使用UB_BLOCK_NUM作为Unified Buffer可用的Block数量，因此可得出可用UB空间的大小为UB_BLOCK_NUM * BLOCK_SIZE | ||
| 53 | - uint32_t blockLength, tileNum; | ||
| 54 | - if ((totalLengthAligned / alignNum) % NUM_BLOCKS == 0U) { | ||
| 55 | - blockLength = totalLengthAligned / NUM_BLOCKS; | ||
| 56 | - } | ||
| 57 | - ``` | ||
| 58 | - | ||
| 59 | -3. 计算tileNum。为了减少数据搬运开销，应尽量使用核内的Unified Buffer空间。基于每个核上的计算量以及可用Unified Buffer空间的大小，计算tileNum。 | ||
| 60 | - | ||
| 61 | - ``` | ||
| 62 | - tileNum = blockLength / (alignNum * UB_BLOCK_NUM); | ||
| 63 | - ``` | ||
| 64 | - | ||
| 65 | -4. 根据计算出的tileNum，计算tileLength和lastTileLength。 | ||
| 66 | - | ||
| 67 | - 如果每个核的计算量能够被当前可用Unified Buffer空间均分，则按照无尾块场景处理。 | ||
| 68 | - | ||
| 69 | - ``` | ||
| 70 | - if (static_cast<uint32_t>(blockLength / alignNum) % UB_BLOCK_NUM == 0U) { | ||
| 71 | - // 单核的计算量能被当前可用UB空间均分，仅有主块，无尾块 | ||
| 72 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 73 | - lastTileLength = 0U; | ||
| 74 | - } | ||
| 75 | - ``` | ||
| 76 | - | ||
| 77 | - 反之，按照尾块场景处理，尾块长度为单核计算数据长度 - tileNum \* tileLength。 | ||
| 78 | - | ||
| 79 | - ``` | ||
| 80 | - if (tileNum == 0U) { | ||
| 81 | - // 单核需要计算的长度小于UB可用空间，按照仅有尾块处理 | ||
| 82 | - tileLength = 0U; | ||
| 83 | - lastTileLength = static_cast<uint32_t>(((blockLength + alignNum - 1) / alignNum) * alignNum); | ||
| 84 | - } else { | ||
| 85 | - // 同时有主块和尾块 | ||
| 86 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 87 | - lastTileLength = static_cast<uint32_t>(blockLength - tileNum * tileLength); | ||
| 88 | - } | ||
| 89 | - ``` | ||
| 90 | - | ||
| 91 | -Host侧Tiling实现的代码如下： | ||
| 92 | - | ||
| 93 | -``` | ||
| 94 | -constexpr uint32_t BLOCK_SIZE = 32; | ||
| 95 | -constexpr uint32_t NUM_BLOCKS = 8; | ||
| 96 | -constexpr uint32_t UB_BLOCK_NUM = 100; // 此处为方便验证，使用UB_BLOCK_NUM作为UB可用的Block数量，因此可得出可用UB空间的大小为UB_BLOCK_NUM * BLOCK_SIZE | ||
| 97 | -... | ||
| 98 | - | ||
| 99 | -uint32_t alignNum = BLOCK_SIZE / dataTypeSize; // 为方便计算，这里根据数据类型定义变量alignNum作为对齐数，dataTypeSize为运算数据的数据类型对应的字节数 | ||
| 100 | -// totalLength为数据总量 | ||
| 101 | -totalLengthAligned = (totalLength % alignNum == 0U) ? | ||
| 102 | - static_cast<uint32_t>(totalLength) : | ||
| 103 | - ((static_cast<uint32_t>(totalLength) + alignNum - 1) / alignNum) * alignNum; | ||
| 104 | -uint32_t blockLength, tileNum; | ||
| 105 | -if ((totalLengthAligned / alignNum) % NUM_BLOCKS == 0U) { | ||
| 106 | - blockLength = totalLengthAligned / NUM_BLOCKS; | ||
| 107 | - tileNum = blockLength / alignNum / UB_BLOCK_NUM; | ||
| 108 | - | ||
| 109 | - if (tileNum == 0) { | ||
| 110 | - // 单核需要计算的长度小于UB可用空间，按照仅有尾块处理 | ||
| 111 | - tileLength = 0; | ||
| 112 | - lastTileLength = ((blockLength + alignNum - 1) / alignNum) * alignNum; | ||
| 113 | - } else if ((blockLength / alignNum) % UB_BLOCK_NUM == 0) { | ||
| 114 | - // 单核的计算量能被当前可用UB空间均分，仅有主块，无尾块 | ||
| 115 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 116 | - lastTileLength = 0; | ||
| 117 | - } else { | ||
| 118 | - // 同时有主块和尾块 | ||
| 119 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 120 | - lastTileLength = blockLength - tileNum * tileLength; | ||
| 121 | - } | ||
| 122 | - ... | ||
| 123 | -} | ||
| 124 | -``` | ||
| 125 | - | ||
| 126 | -\(1，1904\)形状的输入数据计算后，tiling结构体内各个变量的值如下： | ||
| 127 | - | ||
| 128 | -``` | ||
| 129 | -struct AddCustomTilingData { | ||
| 130 | - uint32_t blockLength = 238; // 每个核计算238个half，8个核共计算1904个half | ||
| 131 | - uint32_t tileNum = 0; // 可用的UB空间足够，为仅有尾块的场景 | ||
| 132 | - uint32_t tileLength = 0; // 没有主块，主块长度为0 | ||
| 133 | - uint32_t lastTileLength = 240; // 238个half未32B对齐，对齐到240个half搬运 | ||
| 134 | - ... | ||
| 135 | -}; | ||
| 136 | -``` | ||
| 137 | - | ||
| 138 | -## 算子类实现<a name="zh-cn_topic_0000002201157442_section2027413973119"></a> | ||
| 139 | - | ||
| 140 | -与[多核Tiling](多核Tiling.md)相比，在Init函数中通过Pipe内存管理对象为输入输出Queue分配内存时，取tileLength与lastTileLength中的最大值作为分配内存的长度。例如，当单核需要计算的长度小于UB可用空间时，按照仅有尾块处理，此时tileLength为0，而lastTileLength为数据块长度。因此，需要取两者中的较大值来分配内存。 | ||
| 141 | - | ||
| 142 | -``` | ||
| 143 | -uint32_t initBufferLength = AscendC::Std::max(this->tileLength, this->lastTileLength); | ||
| 144 | -pipe->InitBuffer(inQueueX, 1, this->initBufferLength * sizeof(dataType)); | ||
| 145 | -``` | ||
| 146 | - | ||
| 147 | -由于尾块长度为lastTileLength，与主块数据块的长度不同，因此在CopyIn函数、Compute函数、CopyOut函数中传入本次循环待处理的数据块长度参数tileLength，即待处理的主块或尾块的数据长度。 | ||
| 148 | - | ||
| 149 | -Process函数实现代码如下： | ||
| 150 | - | ||
| 151 | -``` | ||
| 152 | -__aicore__ inline void Process() | ||
| 153 | -{ | ||
| 154 | - // 计算主块数据，对应数据块长度为tileLength | ||
| 155 | - for (uint32_t i = 0; i < this->tileNum; i++) { | ||
| 156 | - CopyIn(i, this->tileLength); | ||
| 157 | - Compute(i, this->tileLength); | ||
| 158 | - CopyOut(i, this->tileLength); | ||
| 159 | - } | ||
| 160 | - // 计算尾块数据，对应数据块长度为lastTileLength | ||
| 161 | - if (this->lastTileLength > 0) { | ||
| 162 | - CopyIn(this->tileNum, this->lastTileLength); | ||
| 163 | - Compute(this->tileNum, this->lastTileLength); | ||
| 164 | - CopyOut(this->tileNum, this->lastTileLength); | ||
| 165 | - } | ||
| 166 | -} | ||
| 167 | -``` | ||
| 168 | - | ||
| 169 | -CopyIn函数实现代码如下： | ||
| 170 | - | ||
| 171 | -``` | ||
| 172 | -__aicore__ inline void CopyIn(int32_t progress, uint32_t tileLength) | ||
| 173 | -{ | ||
| 174 | - AscendC::LocalTensor<T> xLocal = inQueueX.AllocTensor<T>(); | ||
| 175 | - AscendC::LocalTensor<T> yLocal = inQueueY.AllocTensor<T>(); | ||
| 176 | - AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], tileLength); | ||
| 177 | - AscendC::DataCopy(yLocal, yGm[progress * this->tileLength], tileLength); | ||
| 178 | - inQueueX.EnQue(xLocal); | ||
| 179 | - inQueueY.EnQue(yLocal); | ||
| 180 | -} | ||
| 181 | -``` | ||
| 182 | - | ||
| 183 | -Compute函数实现代码如下： | ||
| 184 | - | ||
| 185 | -``` | ||
| 186 | -__aicore__ inline void Compute(int32_t progress, uint32_t tileLength) | ||
| 187 | -{ | ||
| 188 | - AscendC::LocalTensor<T> xLocal = inQueueX.DeQue<T>(); | ||
| 189 | - AscendC::LocalTensor<T> yLocal = inQueueY.DeQue<T>(); | ||
| 190 | - AscendC::LocalTensor<T> zLocal = outQueueZ.AllocTensor<T>(); | ||
| 191 | - AscendC::Add(zLocal, xLocal, yLocal, tileLength); | ||
| 192 | - outQueueZ.EnQue<T>(zLocal); | ||
| 193 | - inQueueX.FreeTensor(xLocal); | ||
| 194 | - inQueueY.FreeTensor(yLocal); | ||
| 195 | -} | ||
| 196 | -``` | ||
| 197 | - | ||
| 198 | -CopyOut函数实现代码如下： | ||
| 199 | - | ||
| 200 | -``` | ||
| 201 | -__aicore__ inline void CopyOut(int32_t progress, uint32_t tileLength) | ||
| 202 | -{ | ||
| 203 | - AscendC::LocalTensor<T> zLocal = outQueueZ.DeQue<T>(); | ||
| 204 | - AscendC::DataCopy(zGm[progress * this->tileLength], zLocal, tileLength); | ||
| 205 | - outQueueZ.FreeTensor(zLocal); | ||
| 206 | -} | ||
| 207 | -``` | ||
Ddocs/zh/guide/算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/尾核-尾块.md+0-117
| @@ -1,117 +0,0 @@ | |||
| 1 | -# 尾核&尾块<a name="ZH-CN_TOPIC_0000002500468234"></a> | ||
| 2 | - | ||
| 3 | -对于不同shape的输入进行数据切分时，可能会发生数据无法平均分配到多个核、同时每个核内的数据无法均分的情况。参考[核间均分场景下的尾块处理](尾块Tiling.md)与[核间不均分场景下的尾核处理](尾核Tiling.md)的处理方式，将两者结合起来考虑整核的尾块、尾核的尾块的处理方式。 | ||
| 4 | - | ||
| 5 | -## Tiling实现<a name="zh-cn_topic_0000002236277533_section132981926144714"></a> | ||
| 6 | - | ||
| 7 | -由于本场景中核间、核内的数据均无法均分，在[核间不均分场景下的尾核处理](尾核Tiling.md)定义的Tiling结构体的基础上增加两个成员变量： | ||
| 8 | - | ||
| 9 | -- formerLastTileLength：数据量多的核最后一个分块大小，即整核的尾块大小。 | ||
| 10 | - | ||
| 11 | - 计算时，先按[尾核Tiling](尾核Tiling.md)中提到的分核策略，切分数据量多的核。 | ||
| 12 | - | ||
| 13 | - ``` | ||
| 14 | - // shape需要对齐到的datablock | ||
| 15 | - uint32_t totalLengthAligned = (totalLength % alignNum == 0U) ? | ||
| 16 | - static_cast<uint32_t>(totalLength) : | ||
| 17 | - ((static_cast<uint32_t>(totalLength) + alignNum - 1) / alignNum) * alignNum; | ||
| 18 | - // 计算整核数量 | ||
| 19 | - uint32_t formerNum = (totalLengthAligned / alignNum) % numBlocks; | ||
| 20 | - // 计算整核的数据量 | ||
| 21 | - uint32_t formerLength = static_cast<uint32_t>(((totalLengthAligned + numBlocks - 1) / numBlocks + alignNum - 1) / alignNum) * alignNum; | ||
| 22 | - ``` | ||
| 23 | - | ||
| 24 | - 再按[尾块Tiling](尾块Tiling.md)中的切分策略，计算尾块长度。 | ||
| 25 | - | ||
| 26 | - ``` | ||
| 27 | - TilingParamsCalc(formerLength, alignNum, formerTileNum, formerTileLength, formerLastTileLength); | ||
| 28 | - | ||
| 29 | - void TilingParamsCalc(uint32_t length, uint32_t alignNum, uint32_t& tileNum, uint32_t& tileLength, | ||
| 30 | - uint32_t& lastTileLength) | ||
| 31 | - { | ||
| 32 | - tileNum = length / (alignNum * UB_BLOCK_NUM); | ||
| 33 | - if (tileNum == 0U) { | ||
| 34 | - tileLength = 0U; | ||
| 35 | - lastTileLength = static_cast<uint32_t>(((length + alignNum - 1) / alignNum) * alignNum); | ||
| 36 | - } else if (static_cast<uint32_t>(length / alignNum) % UB_BLOCK_NUM == 0U) { | ||
| 37 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 38 | - lastTileLength = 0U; | ||
| 39 | - } else { | ||
| 40 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 41 | - lastTileLength = static_cast<uint32_t>(length - tileNum * tileLength); | ||
| 42 | - } | ||
| 43 | - } | ||
| 44 | - ``` | ||
| 45 | - | ||
| 46 | -- tailLastTileLength：数据量少的核最后一个分块大小，即尾核的尾块大小。 | ||
| 47 | - | ||
| 48 | - 计算时，先按[尾核Tiling](尾核Tiling.md)中提到的分核策略，切分数据量少的核。 | ||
| 49 | - | ||
| 50 | - ``` | ||
| 51 | - // 计算尾核数量 | ||
| 52 | - uint32_t tailNum = numBlocks - formerNum; | ||
| 53 | - // 计算尾核的数据量 | ||
| 54 | - uint32_t tailLength = (totalLengthAligned / numBlocks / alignNum) * alignNum; | ||
| 55 | - ``` | ||
| 56 | - | ||
| 57 | - 再按[尾块Tiling](尾块Tiling.md)中的切分策略，计算尾块长度。 | ||
| 58 | - | ||
| 59 | - ``` | ||
| 60 | - TilingParamsCalc(tailLength, alignNum, tailTileNum, tailTileLength, tailLastTileLength); | ||
| 61 | - | ||
| 62 | - void TilingParamsCalc(uint32_t length, uint32_t alignNum, uint32_t& tileNum, uint32_t& tileLength, | ||
| 63 | - uint32_t& lastTileLength) | ||
| 64 | - { | ||
| 65 | - tileNum = length / (alignNum * UB_BLOCK_NUM); | ||
| 66 | - if (tileNum == 0U) { | ||
| 67 | - tileLength = 0U; | ||
| 68 | - lastTileLength = static_cast<uint32_t>(((length + alignNum - 1) / alignNum) * alignNum); | ||
| 69 | - } else if (static_cast<uint32_t>(length / alignNum) % UB_BLOCK_NUM == 0U) { | ||
| 70 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 71 | - lastTileLength = 0U; | ||
| 72 | - } else { | ||
| 73 | - tileLength = UB_BLOCK_NUM * alignNum; | ||
| 74 | - lastTileLength = static_cast<uint32_t>(length - tileNum * tileLength); | ||
| 75 | - } | ||
| 76 | - } | ||
| 77 | - ``` | ||
| 78 | - | ||
| 79 | -## 算子类实现<a name="zh-cn_topic_0000002236277533_section182075114713"></a> | ||
| 80 | - | ||
| 81 | -Kernel侧Init函数和Process函数的实现需将[核间均分场景下的尾块处理](尾块Tiling.md)与[核间不均分场景下的尾核处理](尾核Tiling.md)的实现结合起来。 | ||
| 82 | - | ||
| 83 | -Init函数中由于整核和尾核对应的tileLength和lastTileLength不同。因此需按照[核间不均分场景下的尾核处理](尾核Tiling.md)中提到的分别处理整核和尾核。后续对主块和尾块的CopyIn、Compute、CopyOut函数的处理方式与[核间均分场景下的处理](尾块Tiling.md)方式相同。 | ||
| 84 | - | ||
| 85 | -Init函数实现代码如下： | ||
| 86 | - | ||
| 87 | -``` | ||
| 88 | -__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, AddCustomTilingData tiling, AscendC::TPipe* pipeIn) | ||
| 89 | -{ | ||
| 90 | - pipe = pipeIn; | ||
| 91 | - if (AscendC::GetBlockIdx() < tiling.formerNum) { | ||
| 92 | - this->tileNum = tiling.formerTileNum; | ||
| 93 | - this->tileLength = tiling.formerTileLength; | ||
| 94 | - this->lastTileLength = tiling.formerLastTileLength; | ||
| 95 | - uint64_t offset = tiling.formerLength * AscendC::GetBlockIdx(); | ||
| 96 | - xGm.SetGlobalBuffer((__gm__ half *)x + offset, tiling.formerLength); | ||
| 97 | - yGm.SetGlobalBuffer((__gm__ half *)y + offset, tiling.formerLength); | ||
| 98 | - zGm.SetGlobalBuffer((__gm__ half *)z + offset, tiling.formerLength); | ||
| 99 | - } else { | ||
| 100 | - this->tileNum = tiling.tailTileNum; | ||
| 101 | - this->tileLength = tiling.tailTileLength; | ||
| 102 | - this->lastTileLength = tiling.tailLastTileLength; | ||
| 103 | - uint64_t offset = tiling.formerLength * tiling.formerNum | ||
| 104 | - + tiling.tailLength * (AscendC::GetBlockIdx() - tiling.formerNum); | ||
| 105 | - xGm.SetGlobalBuffer((__gm__ half *)x + offset, tiling.tailLength); | ||
| 106 | - yGm.SetGlobalBuffer((__gm__ half *)y + offset, tiling.tailLength); | ||
| 107 | - zGm.SetGlobalBuffer((__gm__ half *)z + offset, tiling.tailLength); | ||
| 108 | - } | ||
| 109 | - | ||
| 110 | - // 只有尾块的场景下，tileLength为0，因此取tileLength和lastTileLength的最大值来初始化 | ||
| 111 | - uint32_t initBufferLength = AscendC::Std::max(this->tileLength, this->lastTileLength); | ||
| 112 | - pipe->InitBuffer(inQueueX, 1, this->initBufferLength * sizeof(half)); | ||
| 113 | - pipe->InitBuffer(inQueueY, 1, this->initBufferLength * sizeof(half)); | ||
| 114 | - pipe->InitBuffer(outQueueZ, 1, this->initBufferLength * sizeof(half)); | ||
| 115 | -} | ||
| 116 | -``` | ||
| 117 | - | ||
Ddocs/zh/guide/算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/尾核Tiling.md+0-113
| @@ -1,113 +0,0 @@ | |||
| 1 | -# 尾核Tiling<a name="ZH-CN_TOPIC_0000002532388125"></a> | ||
| 2 | - | ||
| 3 | -对于不同shape的输入进行数据切分时，可能会发生数据无法平均分配到多个核的情况。例如当算子的输入shape为\[1, 1999\]，使用核数为8，数据类型为half时，需要计算的数据总量为1 \* 1999 \* sizeof\(half\) = 3998字节，3998字节既不满足32字节对齐，也无法平均分配到8个核上。因此该场景下，对数据进行多核切分后，每个核的计算数据量不同。此种情况下，应该尽可能均匀的分配数据，所有核上的计算数据量有两种情况，将计算量较多的核称为整核，计算量较少的核称为尾核。 | ||
| 4 | - | ||
| 5 | -**图1** 数据对齐示意图<a name="zh-cn_topic_0000002201317270_fig5261155043319"></a> | ||
| 6 | - | ||
| 7 | - | ||
| 8 | -## Tiling实现<a name="zh-cn_topic_0000002201317270_section209650864618"></a> | ||
| 9 | - | ||
| 10 | -- 因为AI处理器在进行数据搬运和Vector计算时，对于搬运的数据长度和Unified Buffer首地址都有必须32字节对齐的要求，**首先待处理数据需要先保证向上对齐到32字节的大小**。该场景下后续搬运和计算的处理细节请参考[非对齐场景](../非对齐场景.md)。如下代码片段展示了将数据对齐到datablock大小的示例： | ||
| 11 | - | ||
| 12 | - ``` | ||
| 13 | - constexpr uint32_t SIZE_OF_HALF = 2; | ||
| 14 | - constexpr uint32_t BLOCK_SIZE = 32; | ||
| 15 | - constexpr uint32_t NUM_BLOCKS = 8; | ||
| 16 | - constexpr uint32_t ALIGN_NUM = BLOCK_SIZE / SIZE_OF_HALF; | ||
| 17 | - // shape需要对齐到的32字节，假设原totalLength为1999，向上满足32字节对齐后为2000 | ||
| 18 | - uint32_t totalLengthAligned = (totalLength % ALIGN_NUM == 0U) ? | ||
| 19 | - static_cast<uint32_t>(totalLength) : | ||
| 20 | - ((static_cast<uint32_t>(totalLength) + ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM; | ||
| 21 | - ``` | ||
| 22 | - | ||
| 23 | -- 满足32字节对齐后的数据，应尽可能的均分到每个核上。如果无法均分，那么先将可以均分的部分平均分配，剩余的部分分配给部分核，会有部分核多算一个datablock。为了保证切分后的数据仍是满足32字节对齐的，以ALIGN\_NUM（ALIGN\_NUM个数据为32字节）为粒度，将数据分配到所有核上。在本样例中，数据类型为half，ALIGN\_NUM = BLOCK\_SIZE / sizeof\(half\) = 16。将对齐后的数据总量按ALIGN\_NUM为粒度分成x个数据块，x = 2000 / 16 = 125。 | ||
| 24 | - | ||
| 25 | - AI处理器的核数NUM\_BLOCKS为8，无法将125个数据块均分到8个核上。按照以下步骤将数据块尽可能的均分到每个核上： | ||
| 26 | - | ||
| 27 | - 1. 计算x / NUM\_BLOCKS = 15； | ||
| 28 | - 2. 计算x % NUM\_BLOCKS = 5。 | ||
| 29 | - | ||
| 30 | - 根据上述步骤得出，如果每个核上分配15个数据块，那么将有5个数据块剩余。将这5个剩余的数据块分配到5个核上，这样可以得到5个计算16个数据块的整核和3个计算15个数据块的尾核。下图展示了数据无法均分时多核切分的示例。 | ||
| 31 | - | ||
| 32 | - **图2** 无法均分到每个核上的示例<a name="zh-cn_topic_0000002201317270_fig68506184496"></a> | ||
| 33 | -  | ||
| 34 | - | ||
| 35 | -基于上文，设计如下的算子Tiling结构体成员： | ||
| 36 | - | ||
| 37 | -- formerNum：分配到数据量较多的核数，即整核的核数。 | ||
| 38 | -- tailNum：分配到数据量较少的核数，即尾核的核数。 | ||
| 39 | -- formerLength：整核计算的数据长度。 | ||
| 40 | -- tailLength：尾核计算的数据长度。 | ||
| 41 | - | ||
| 42 | -Tiling参数的计算代码如下： | ||
| 43 | - | ||
| 44 | -``` | ||
| 45 | -constexpr uint32_t NUM_BLOCKS = 8; | ||
| 46 | -constexpr uint32_t SIZE_OF_HALF = 2; | ||
| 47 | -constexpr uint32_t BLOCK_SIZE = 32; | ||
| 48 | -// shape需要对齐到的最小单位 | ||
| 49 | -constexpr uint32_t ALIGN_NUM = BLOCK_SIZE / SIZE_OF_HALF; | ||
| 50 | -... | ||
| 51 | -void GenerateTilingData(uint8_t* tilingBuf, uint32_t numBlocks) | ||
| 52 | -{ | ||
| 53 | - // shape需要对齐到的datablock,假设原totalLength为1999，向上满足32字节对齐后为2000 | ||
| 54 | - uint32_t totalLengthAligned = (totalLength % ALIGN_NUM == 0U) ? | ||
| 55 | - static_cast<uint32_t>(totalLength) : | ||
| 56 | - ((static_cast<uint32_t>(totalLength) + ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM; | ||
| 57 | - // 核心数为8，一个datablock包含16个数，那么：datablock的总数：2000 / 16 = 125 | ||
| 58 | - // 有5个核会分到16个datablock：125 % 8 =5，可以称之为整核 | ||
| 59 | - // 有3个核会分到15个datablock：8 - 5 = 3，可以称之为尾核 | ||
| 60 | - uint32_t formerNum = (totalLengthAligned / ALIGN_NUM) % numBlocks; | ||
| 61 | - uint32_t tailNum = numBlocks - formerNum; | ||
| 62 | - // 整核计算的数据长度：totalLengthAligned / NUM_BLOCKS为每个核上计算的元素个数，formerLength为上述元素个数向上32字节对齐的结果 | ||
| 63 | - uint32_t formerLength = | ||
| 64 | - static_cast<uint32_t>(((totalLengthAligned + numBlocks - 1) / numBlocks + ALIGN_NUM - 1) / ALIGN_NUM) * ALIGN_NUM; | ||
| 65 | - // 尾核计算的数据长度：totalLengthAligned / NUM_BLOCKS为每个核上计算的元素个数，tailLength为上述元素个数向下32字节对齐的结果 | ||
| 66 | - uint32_t tailLength = (totalLengthAligned / numBlocks / ALIGN_NUM) * ALIGN_NUM; | ||
| 67 | - ... | ||
| 68 | -} | ||
| 69 | -``` | ||
| 70 | - | ||
| 71 | -## 算子类实现<a name="zh-cn_topic_0000002201317270_section11477133219466"></a> | ||
| 72 | - | ||
| 73 | -在Kernel侧的Init函数中，计算输入在Global Memory上的内存偏移地址时，应对整核和尾核加以区分。 | ||
| 74 | - | ||
| 75 | -整核上，输入的内存偏移地址计算代码如下： | ||
| 76 | - | ||
| 77 | -``` | ||
| 78 | -xGm.SetGlobalBuffer((__gm__ T*)x + formerLength * AscendC::GetBlockIdx(), formerLength); | ||
| 79 | -``` | ||
| 80 | - | ||
| 81 | -尾核上，计算输入的内存偏移地址时，需在全部整核的数据长度基础上加上尾核的偏移量，代码如下： | ||
| 82 | - | ||
| 83 | -``` | ||
| 84 | -xGm.SetGlobalBuffer((__gm__ T*)x + formerLength * formerNum + tailLength * (AscendC::GetBlockIdx() - formerNum), tailLength); | ||
| 85 | -``` | ||
| 86 | - | ||
| 87 | -完整的Init函数实现代码如下： | ||
| 88 | - | ||
| 89 | -``` | ||
| 90 | -__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, AddCustomTilingData tiling, AscendC::TPipe* pipeIn) | ||
| 91 | -{ | ||
| 92 | - pipe = pipeIn; | ||
| 93 | - if (AscendC::GetBlockIdx() < tiling.formerNum) { | ||
| 94 | - this->tileLength = tiling.formerLength; | ||
| 95 | - uint64_t offset = tiling.formerLength * AscendC::GetBlockIdx(); | ||
| 96 | - xGm.SetGlobalBuffer((__gm__ half *)x + offset, tiling.formerLength); | ||
| 97 | - yGm.SetGlobalBuffer((__gm__ half *)y + offset, tiling.formerLength); | ||
| 98 | - zGm.SetGlobalBuffer((__gm__ half *)z + offset, tiling.formerLength); | ||
| 99 | - } else { | ||
| 100 | - this->tileLength = tiling.tailLength; | ||
| 101 | - uint64_t offset = tiling.formerLength * tiling.formerNum | ||
| 102 | - + tiling.tailLength * (AscendC::GetBlockIdx() - tiling.formerNum); | ||
| 103 | - xGm.SetGlobalBuffer((__gm__ half *)x +offset, tiling.tailLength); | ||
| 104 | - yGm.SetGlobalBuffer((__gm__ half *)y + offset, tiling.tailLength); | ||
| 105 | - zGm.SetGlobalBuffer((__gm__ half *)z + offset, tiling.tailLength); | ||
| 106 | - } | ||
| 107 | - pipe->InitBuffer(inQueueX, 1, this->tileLength * sizeof(half)); | ||
| 108 | - pipe->InitBuffer(inQueueY, 1, this->tileLength * sizeof(half)); | ||
| 109 | - pipe->InitBuffer(outQueueZ, 1, this->tileLength * sizeof(half)); | ||
| 110 | -} | ||
| 111 | -``` | ||
| 112 | - | ||
| 113 | -其余实现与[多核Tiling](多核Tiling.md)中的实现一致，这里不重复进行说明。 | ||
Ddocs/zh/guide/算子实践参考/SIMD算子实现/矢量编程/多核-Tiling切分/概述.md+0-34
| @@ -1,34 +0,0 @@ | |||
| 1 | -# 概述<a name="ZH-CN_TOPIC_0000002500468232"></a> | ||
| 2 | - | ||
| 3 | -Ascend C核函数是运行在一个核上的处理函数，上述介绍的[基础矢量算子](../基础矢量算子.md)与[TBuf的使用](../TBuf的使用.md)样例均为在单核上运行的算子，不涉及Host侧Tiling实现。矢量算子实现的组成如下图所示。 | ||
| 4 | - | ||
| 5 | -为了提高算子的执行效率，通常在算子中实现多核并行计算，即对输入数据进行切分，并将不同的数据块分配到不同的核上处理。此外，由于单个核上内部存储Local Memory大小有限，存在无法一次完整地容纳算子的输入和输出数据的场景，因此需要每次搬运一部分输入进行计算然后搬出，再搬运下一部分输入进行计算，直到获得最终的完整结果，这个数据切分、分块计算的过程称之为**Tiling**。切分数据的算法称为Tiling算法或者Tiling策略。根据算子的shape等信息来确定数据切分算法相关参数（比如每次搬运的块大小，以及总共循环多少次）的计算程序，称之为**Tiling实现**，也叫Tiling函数（Tiling Function）。由于Tiling实现中完成的均为标量计算，AI Core并不擅长，所以我们将其独立出来放在Host侧CPU上执行。核函数内部通过解析Host侧传入的Tiling结构体获取Tiling信息，根据Tiling信息控制数据搬入、搬出Local Memory的流程；通过调用计算、数据搬运、内存管理、任务同步API，实现算子逻辑。 | ||
| 6 | - | ||
| 7 | -**图1** 算子实现组成<a name="zh-cn_topic_0000002256238457_fig2052955220511"></a> | ||
| 8 | - | ||
| 9 | - | ||
| 10 | -由于硬件限制，在对输入数据进行数据切分时应遵循以下几个原则： | ||
| 11 | - | ||
| 12 | -1. 由于AI Core中Unified Buffer上的物理限制，要求Unified Buffer上的数据存储空间必须保持32字节对齐。 | ||
| 13 | - - 输入数据不满足32字节对齐时，需要取输入数据长度向上对齐到32字节的长度作为输入数据总长度。 | ||
| 14 | - - 进行Tiling有关计算时，以32字节为最小单位进行计算。 | ||
| 15 | - | ||
| 16 | -2. 尽可能最大利用Unified Buffer空间。 | ||
| 17 | - | ||
| 18 | - AI Core与外部存储交互时会产生性能开销，频繁的进行数据搬运会导致性能瓶颈，因此应尽可能充分利用Unified Buffer空间，减少从Global Memory上搬运数据的次数。 | ||
| 19 | - | ||
| 20 | -3. AI处理器包含多个AI Core，应该充分均衡利用多核计算能力，将计算均衡分配到多个AI Core上。 | ||
| 21 | - | ||
| 22 | -本章将基于以上原则对几种典型场景进行说明。 | ||
| 23 | - | ||
| 24 | -**图2** 多核及Tiling示意图<a name="zh-cn_topic_0000002256238457_fig188456591607"></a> | ||
| 25 | - | ||
| 26 | - | ||
| 27 | -数据切分示意如上图所示，将长度为TOTAL\_LENGTH的算子输入分配到多个核上进行计算，每个核上计算的数据长度为BLOCK\_LENGTH。对于每个核的计算数据，基于Local Memory的大小进一步切分，切分数据块的个数为TILE\_NUM，得到的每个数据块的长度为TILE\_LENGTH。 | ||
| 28 | - | ||
| 29 | -根据每个核计算的数据量是否相同、核内每个数据块的数据量是否相同，切分策略可能会存在以下几种场景： | ||
| 30 | - | ||
| 31 | -1. 核间均分，核内均分：每个核处理的数据量相同，核内每个数据块的数据量相同。在此场景中，通过[多核Tiling](多核Tiling.md)将数据均匀分配到各个核上执行，每个核上每次计算的数据长度相同。 | ||
| 32 | -2. 核间均分，核内不均分：每个核处理的数据量相同，核内各数据块的数据量不完全相同。此场景基于多核Tiling，核内数据不能切分为多个数据量相同且32字节对齐的数据块，需要通过[尾块Tiling](尾块Tiling.md)处理尾块数据的计算。 | ||
| 33 | -3. 核间不均分，核内均分：每个核处理的数据量不同，核内每个数据块的数据量相同。在此场景中，通过[尾核Tiling](尾核Tiling.md)的处理解决数据无法在各核间均匀分配的问题。 | ||
| 34 | -4. 核间不均分，核内不均分：每个核处理的数据量不同，核内各数据块的数据量不完全相同。该场景下需要同时考虑[尾核&尾块](尾核-尾块.md)，处理多核间及核内数据的合理切分。 | ||
| @@ -0,0 +1,127 @@ | |||
| 1 | +# 主块均分 | ||
| 2 | + | ||
| 3 | +长度为`totalLength`的输入数据可以按参与计算的核数和主块长度`mainTileLength`均匀切分，主块均分后无剩余数据。各核处理的数据量相同，每个核处理`tileNum`个主块，核内只包含主块。 | ||
| 4 | + | ||
| 5 | +**图 1** 数据切分示意图 | ||
| 6 | + | ||
| 7 | + | ||
| 8 | + | ||
| 9 | +本场景样例的完整代码请参考[场景0：主块均分](../../../../../../../examples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/README.md#scenario0-main-tile-only)。样例中固定使用8个核，主块长度`mainTileLength`为3200个`half`元素，数据整体长度为256000个`half`元素。数据平均分配到8个核上处理，每个核计算32000个元素；每个核内共切分成10个主块，不涉及尾块和尾核处理。 | ||
| 10 | + | ||
| 11 | +## Tiling实现 | ||
| 12 | + | ||
| 13 | +基于本节的切分策略，Tiling需要定义如下参数： | ||
| 14 | + | ||
| 15 | +- blockLength：每个核处理的数据长度，单位为元素。 | ||
| 16 | +- tileNum：每个核需要计算的主块个数。 | ||
| 17 | + | ||
| 18 | +据此定义`MainTileOnlyTiling`结构体，代码如下： | ||
| 19 | + | ||
| 20 | +```cpp | ||
| 21 | +struct MainTileOnlyTiling { | ||
| 22 | + uint32_t blockLength; // 每个核处理的元素数量 | ||
| 23 | + uint32_t tileNum; // 主块数量 | ||
| 24 | +}; | ||
| 25 | +``` | ||
| 26 | + | ||
| 27 | +接下来完成Tiling参数的计算。涉及的主要变量包括：逻辑核数`numBlocks`、一个DataBlock的字节数`dataBlockSize`、一个`half`元素的字节数`halfSize`、一个DataBlock包含的`half`元素数量`alignNum`、一个主块包含的`half`元素数量`mainTileLength`。步骤如下： | ||
| 28 | + | ||
| 29 | +1. 根据`dataBlockSize`和`halfSize`计算`alignNum`，本样例中`alignNum`为16。`totalLength`为输入数据的元素总数，将其按`alignNum`向上对齐后得到对齐后的元素总数`totalLengthAligned`。 | ||
| 30 | + | ||
| 31 | + ```cpp | ||
| 32 | + constexpr uint32_t dataBlockSize = 32; | ||
| 33 | + constexpr uint32_t halfSize = 2; | ||
| 34 | + constexpr uint32_t alignNum = dataBlockSize / halfSize; | ||
| 35 | + | ||
| 36 | + uint32_t totalLengthAligned = AlignUp(totalLength, alignNum); | ||
| 37 | + ``` | ||
| 38 | + | ||
| 39 | +2. 用`totalLengthAligned`除以`numBlocks * mainTileLength`得到每个核均分到的主块数量`mainTileNum`，取余得到主块均分后的剩余元素数量`mainTileRemainder`；再将`mainTileNum`乘以`mainTileLength`得到每个核已分配的元素数量`baseLength`，`baseTileNum`直接取`mainTileNum`。 | ||
| 40 | + | ||
| 41 | + ```cpp | ||
| 42 | + constexpr uint32_t numBlocks = 8; | ||
| 43 | + constexpr uint32_t mainTileLength = 3200; | ||
| 44 | + | ||
| 45 | + uint32_t mainTileNum = totalLengthAligned / (numBlocks * mainTileLength); | ||
| 46 | + uint32_t mainTileRemainder = totalLengthAligned % (numBlocks * mainTileLength); | ||
| 47 | + uint32_t baseLength = mainTileNum * mainTileLength; | ||
| 48 | + uint32_t baseTileNum = mainTileNum; | ||
| 49 | + ``` | ||
| 50 | + | ||
| 51 | +3. 当`mainTileRemainder`为0时，表示所有数据都能按主块均匀分配到各核，此时使用`MainTileOnlyTiling`结构体传递Tiling切分信息。 | ||
| 52 | + | ||
| 53 | + ```cpp | ||
| 54 | + if (mainTileRemainder == 0U) { | ||
| 55 | + MainTileOnlyTiling tiling{baseLength, baseTileNum}; | ||
| 56 | + add_custom<MainTileOnlyTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 57 | + return; | ||
| 58 | + } | ||
| 59 | + ``` | ||
| 60 | + | ||
| 61 | +## 算子类实现 | ||
| 62 | + | ||
| 63 | +Kernel侧算子采用[静态Tensor编程](../../../../编程指南/编程模型/AI-Core-SIMD编程/基于Tensor的CPP编程/静态Tensor编程.md)方式实现，通过`LocalTensor`构造函数直接指定Local Memory地址。静态Tensor编程中的内存管理和同步管理请参考对应章节，本节重点说明各核GM地址的计算以及主块数据的循环处理流程。 | ||
| 64 | + | ||
| 65 | +- Kernel入口接收Host侧传入的Tiling结构体，并依次调用算子类的`Init`与`Process`函数完成初始化与计算。主块均分场景中，传入的结构体类型为`MainTileOnlyTiling`。 | ||
| 66 | + | ||
| 67 | + ```cpp | ||
| 68 | + template <typename TilingData> | ||
| 69 | + __vector__ __global__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, TilingData tiling) | ||
| 70 | + { | ||
| 71 | + AscendC::InitSocState(); | ||
| 72 | + KernelAdd op; | ||
| 73 | + op.Init(x, y, z, tiling); | ||
| 74 | + op.Process(); | ||
| 75 | + AscendC::PipeBarrier<PIPE_ALL>(); | ||
| 76 | + } | ||
| 77 | + ``` | ||
| 78 | + | ||
| 79 | +- 由于数据被分配到多个核上处理，每个核访问的Global Memory地址范围各不相同。`Init`函数通过`tiling.blockLength`获取每个核处理的元素数量，并结合[GetBlockIdx](../../../../../api/SIMD-API/basic_api/tool_interface/system_resources_and_variables/GetBlockIdx.md)计算当前核所处理数据在Global Memory上的起始偏移`offset`，随后将偏移与长度传入`InitGm`完成GM地址设置；同时将`tiling.tileNum`保存到成员变量，供`Process`函数使用。 | ||
| 80 | + | ||
| 81 | + ```cpp | ||
| 82 | + __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, | ||
| 83 | + MainTileOnlyTiling tiling) | ||
| 84 | + { | ||
| 85 | + uint32_t offset = AscendC::GetBlockIdx() * tiling.blockLength; | ||
| 86 | + InitGm(x, y, z, offset, tiling.blockLength); | ||
| 87 | + this->tileNum = tiling.tileNum; | ||
| 88 | + this->lastTileLength = 0U; | ||
| 89 | + } | ||
| 90 | + ``` | ||
| 91 | + | ||
| 92 | +- `InitGm`函数通过[SetGlobalBuffer](../../../../../api/SIMD-API/basic_api/data_structures/GlobalTensor/SetGlobalBuffer.md)为输入`xGm`、`yGm`和输出`zGm`设置Global Memory地址。以输入x为例，`(__gm__ half*)x + offset`即为当前核处理数据在Global Memory上的起始地址，`length`为该核处理的元素数量。 | ||
| 93 | + | ||
| 94 | + ```cpp | ||
| 95 | + __aicore__ inline void InitGm( | ||
| 96 | + __gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, uint32_t offset, uint32_t length) | ||
| 97 | + { | ||
| 98 | + xGm.SetGlobalBuffer((__gm__ half*)x + offset, length); | ||
| 99 | + yGm.SetGlobalBuffer((__gm__ half*)y + offset, length); | ||
| 100 | + zGm.SetGlobalBuffer((__gm__ half*)z + offset, length); | ||
| 101 | + } | ||
| 102 | + ``` | ||
| 103 | + | ||
| 104 | +- `Process`函数基于成员变量`tileNum`循环处理主块数据。`xAddr`、`yAddr`、`zAddr`为静态Tensor预先分配的Local Memory地址。循环次数采用`tileNum + (lastTileLength > 0U ? 1U : 0U)`的通用形式计算，主块均分场景下`lastTileLength`为0，循环次数即等于`tileNum`。每轮循环以`mainTileLength`作为当前主块的元素数量，通过`loopIdx * mainTileLength`计算核内偏移`startElement`，依次完成输入搬入、Add计算和输出搬出。 | ||
| 105 | + | ||
| 106 | + ```cpp | ||
| 107 | + __aicore__ inline void Process() | ||
| 108 | + { | ||
| 109 | + AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, mainTileLength); | ||
| 110 | + AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, mainTileLength); | ||
| 111 | + AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, mainTileLength); | ||
| 112 | + | ||
| 113 | + uint32_t totalLoop = this->tileNum + (this->lastTileLength > 0U ? 1U : 0U); | ||
| 114 | + for (uint32_t loopIdx = 0; loopIdx < totalLoop; loopIdx++) { | ||
| 115 | + uint32_t curLen = (loopIdx < this->tileNum) ? mainTileLength : this->lastTileLength; | ||
| 116 | + uint32_t startElement = loopIdx * mainTileLength; | ||
| 117 | + // ... 同步管理代码省略，请参考完整样例代码 ... | ||
| 118 | + AscendC::DataCopy(xLocal, xGm[startElement], curLen); | ||
| 119 | + AscendC::DataCopy(yLocal, yGm[startElement], curLen); | ||
| 120 | + // ... 同步管理代码省略 ... | ||
| 121 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 122 | + // ... 同步管理代码省略 ... | ||
| 123 | + AscendC::DataCopy(zGm[startElement], zLocal, curLen); | ||
| 124 | + // ... 同步管理代码省略 ... | ||
| 125 | + } | ||
| 126 | + } | ||
| 127 | + ``` | ||
| @@ -0,0 +1,94 @@ | |||
| 1 | +# 尾块均分 | ||
| 2 | + | ||
| 3 | +该场景中，完成主块分配后仍有剩余数据，但剩余数据无法继续按主块均匀分配到每个核。为了均衡利用多核计算能力，继续以DataBlock为最小单位将主块均分后剩余数据`mainTileRemainder`均衡分配到每个核，每个核分配到`tailBlockNumEachCore`个DataBlock的数据，且按DataBlock均分后无剩余。最终各核处理的数据量相同，核内除主块外还存在等长尾块。 | ||
| 4 | + | ||
| 5 | +本场景样例的完整代码请参考[场景1：尾块均分](../../../../../../../examples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/README.md#scenario1-main-tile-tail-block)。样例中固定使用8个核，主块长度`mainTileLength`为3200个`half`元素，数据整体长度为260096个`half`元素。数据平均分配到8个核上处理，每个核计算32512个元素；每个核内包含10个主块和512个元素的尾块。数据切分示意如下图所示。 | ||
| 6 | + | ||
| 7 | +**图 1** 数据切分示意图 | ||
| 8 | + | ||
| 9 | + | ||
| 10 | + | ||
| 11 | +与主块均分相比，本场景每个核内除主块外还存在尾块，因此Tiling参数在`blockLength`和`tileNum`基础上新增`lastTileLength`，用于记录尾块的元素数量。 | ||
| 12 | + | ||
| 13 | +## Tiling实现 | ||
| 14 | + | ||
| 15 | +算子的Tiling结构体定义如下： | ||
| 16 | + | ||
| 17 | +```cpp | ||
| 18 | +struct MainTileWithTailBlockTiling { | ||
| 19 | + uint32_t blockLength; // 每个核处理的元素数量 | ||
| 20 | + uint32_t tileNum; // 主块数量 | ||
| 21 | + uint32_t lastTileLength; // 尾块元素数量 | ||
| 22 | +}; | ||
| 23 | +``` | ||
| 24 | + | ||
| 25 | +Host侧Tiling沿用[主块均分](主块均分.md)中的DataBlock对齐和主块分配计算，得到：主块均分后的剩余元素数量`mainTileRemainder`、每个核已分配的基础元素数量`baseLength`、以及对应的主块数量`baseTileNum`。当`mainTileRemainder`不为0时，主块均分分支不成立，继续执行以下步骤： | ||
| 26 | + | ||
| 27 | +1. 以DataBlock为最小单位，将主块均分后的剩余数据`mainTileRemainder`均匀分配到每个核上。一个DataBlock包含`alignNum`个`half`元素，因此代码中以`alignNum`作为均分的单位。用`mainTileRemainder`除以`numBlocks * alignNum`得到每个核分配到的DataBlock数量`tailBlockNumEachCore`，取余得到按DataBlock均分后仍无法均分到每个核的剩余元素数量`remainingTailLength`；再将`tailBlockNumEachCore`乘以`alignNum`得到每个核基础尾块的元素数量`baseLastTileLength`，并累加到`baseLength`上得到每个核的基础元素数量。 | ||
| 28 | + | ||
| 29 | + ```cpp | ||
| 30 | + uint32_t tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); | ||
 | |||
| 31 | + uint32_t remainingTailLength = mainTileRemainder % (numBlocks * alignNum); | ||
| 32 | + uint32_t baseLastTileLength = tailBlockNumEachCore * alignNum; | ||
| 33 | + baseLength += baseLastTileLength; | ||
| 34 | + ``` | ||
| 35 | + | ||
| 36 | +2. 判断`remainingTailLength`是否为0。等于0表示`mainTileRemainder`能以DataBlock为最小单位均匀分配到每个核，此时进入尾块均分场景，使用`MainTileWithTailBlockTiling`结构体传递Tiling切分信息。 | ||
| 37 | + | ||
| 38 | + ```cpp | ||
| 39 | + if (remainingTailLength == 0U) { | ||
| 40 | + MainTileWithTailBlockTiling tiling{baseLength, baseTileNum, baseLastTileLength}; | ||
| 41 | + add_custom<MainTileWithTailBlockTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 42 | + ... | ||
| 43 | + } | ||
| 44 | + ``` | ||
| 45 | + | ||
| 46 | +对形状为`(1, 260096)`的输入数据计算后，Tiling结构体内各个变量的值如下： | ||
| 47 | + | ||
| 48 | +```cpp | ||
| 49 | +MainTileWithTailBlockTiling tiling{ | ||
| 50 | + 32512, // blockLength：每个核计算32512个half，8个核共计算260096个half | ||
| 51 | + 10, // tileNum：每个核包含10个主块 | ||
| 52 | + 512 // lastTileLength：每个核包含512个half的尾块 | ||
| 53 | +}; | ||
| 54 | +``` | ||
| 55 | + | ||
| 56 | +## 算子类实现 | ||
| 57 | + | ||
| 58 | +Kernel侧算子仍采用[静态Tensor编程](../../../../编程指南/编程模型/AI-Core-SIMD编程/基于Tensor的CPP编程/静态Tensor编程.md)方式实现。与[主块均分](主块均分.md)相比，本场景中每个核处理的数据长度仍相同，因此Global Memory偏移仍由`tiling.blockLength`和`GetBlockIdx()`计算；差异在于`Init`函数需要额外保存`tiling.lastTileLength`，供`Process`函数判断最后一次循环是否处理尾块。 | ||
| 59 | + | ||
| 60 | +```cpp | ||
| 61 | +__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, | ||
| 62 | + MainTileWithTailBlockTiling tiling) | ||
| 63 | +{ | ||
| 64 | + uint32_t offset = AscendC::GetBlockIdx() * tiling.blockLength; | ||
| 65 | + InitGm(x, y, z, offset, tiling.blockLength); | ||
| 66 | + this->tileNum = tiling.tileNum; | ||
| 67 | + this->lastTileLength = tiling.lastTileLength; | ||
| 68 | +} | ||
| 69 | +``` | ||
| 70 | + | ||
| 71 | +由于尾块长度`lastTileLength`与主块长度`mainTileLength`不同，`Process`函数需要根据当前循环位置确定本次待处理的数据长度。`totalLoop`为当前核需要处理的循环次数：前`tileNum`轮处理主块，最后一轮处理尾块；`curLen`为本轮处理的元素数量，`startElement`为本轮在核内的元素偏移。 | ||
| 72 | + | ||
| 73 | +```cpp | ||
| 74 | +__aicore__ inline void Process() | ||
| 75 | +{ | ||
| 76 | + AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, mainTileLength); | ||
| 77 | + AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, mainTileLength); | ||
| 78 | + AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, mainTileLength); | ||
| 79 | + | ||
| 80 | + uint32_t totalLoop = this->tileNum + (this->lastTileLength > 0U ? 1U : 0U); | ||
| 81 | + for (uint32_t loopIdx = 0; loopIdx < totalLoop; loopIdx++) { | ||
| 82 | + uint32_t curLen = (loopIdx < this->tileNum) ? mainTileLength : this->lastTileLength; | ||
| 83 | + uint32_t startElement = loopIdx * mainTileLength; | ||
| 84 | + // ... 同步管理代码省略，请参考完整样例代码 ... | ||
| 85 | + AscendC::DataCopy(xLocal, xGm[startElement], curLen); | ||
| 86 | + AscendC::DataCopy(yLocal, yGm[startElement], curLen); | ||
| 87 | + // ... 同步管理代码省略 ... | ||
| 88 | + AscendC::Add(zLocal, xLocal, yLocal, curLen); | ||
| 89 | + // ... 同步管理代码省略 ... | ||
| 90 | + AscendC::DataCopy(zGm[startElement], zLocal, curLen); | ||
| 91 | + // ... 同步管理代码省略 ... | ||
| 92 | + } | ||
| 93 | +} | ||
| 94 | +``` | ||
| @@ -0,0 +1,121 @@ | |||
| 1 | +# 尾核切分 | ||
| 2 | + | ||
| 3 | +在实际算子开发中，受输入数据总量与切分单位的制约，并非所有场景都能将计算完全均匀地分配到每个核上。该场景中，主块均分后存在剩余数据`mainTileRemainder`，但剩余数据无法以DataBlock为最小单位均匀分配到每个核（`tailBlockNumEachCore`为0，`remainingTailLength`不为0），只能继续分配给部分核。获得额外DataBlock的核为整核，其余为尾核；此时核间负载不均衡，需要分别处理整核和尾核。由于`tailBlockNumEachCore`为0，各核未分配到基础尾块，因此尾核内只包含主块，整核比尾核多处理一个DataBlock的尾块。数据切分后，前`formerNum`个核为整核，各处理`formerLength`个元素（含主块和尾块）；其余核为尾核，各处理`tailLength`个元素（仅含主块）。 | ||
| 4 | + | ||
| 5 | +本场景样例的完整代码请参考[场景2：尾核切分](../../../../../../../examples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/README.md#scenario2-main-tile-tail-core)。样例中固定使用8个核，主块长度`mainTileLength`为3200个`half`元素，数据整体长度为256064个`half`元素。数据按DataBlock对齐后，前4个核各处理32016个元素，后4个核各处理32000个元素；整核包含10个主块和16个元素的尾块，尾核只包含10个主块。数据切分示意如下图所示。 | ||
| 6 | + | ||
| 7 | +**图 1** 数据切分示意图 | ||
| 8 | + | ||
| 9 | + | ||
| 10 | + | ||
| 11 | +本场景中整核和尾核处理的元素数量、主块数量及尾块情况不同，无法仅用一组`blockLength`、`tileNum`和`lastTileLength`描述所有核。因此Tiling参数新增`formerNum`记录整核数量，并分别使用`former*`和`tail*`参数描述记录整核和尾核相关Tiling信息的参数；由于尾核内只包含主块，无需定义`tailLastTileLength`。Tiling结构体包含以下成员： | ||
| 12 | + | ||
| 13 | +- formerNum：分配到数据量较多的核数，即整核的核数。 | ||
| 14 | +- formerLength：整核计算的元素数量。 | ||
| 15 | +- formerTileNum：整核中主块的数量。 | ||
| 16 | +- formerLastTileLength：整核中尾块的元素数量。 | ||
| 17 | +- tailLength：尾核计算的元素数量。 | ||
| 18 | +- tailTileNum：尾核中主块的数量。 | ||
| 19 | + | ||
| 20 | +## Tiling实现 | ||
| 21 | + | ||
| 22 | +Tiling结构体定义如下： | ||
| 23 | + | ||
| 24 | +```cpp | ||
| 25 | +struct MainTileWithTailCoreTiling { | ||
| 26 | + uint32_t formerNum; // 整核数量 | ||
| 27 | + uint32_t formerLength; // 整核处理的元素数量 | ||
| 28 | + uint32_t formerTileNum; // 整核中主块的数量 | ||
| 29 | + uint32_t formerLastTileLength; // 整核中尾块的元素数量 | ||
| 30 | + uint32_t tailLength; // 尾核处理的元素数量 | ||
| 31 | + uint32_t tailTileNum; // 尾核中主块的数量 | ||
| 32 | +}; | ||
| 33 | +``` | ||
| 34 | + | ||
| 35 | +Host侧Tiling沿用[尾块均分](尾块均分.md)中的DataBlock对齐、主块分配和基础尾块计算，得到：按DataBlock均分后仍无法均分到每个核的剩余元素数量`remainingTailLength`、每个核分配到的DataBlock数量`tailBlockNumEachCore`、每个核基础尾块的元素数量`baseLastTileLength`、每个核已分配的基础元素数量`baseLength`，以及对应的主块数量`baseTileNum`。当`remainingTailLength`不为0时，尾块均分分支不成立，继续执行以下步骤： | ||
| 36 | + | ||
| 37 | +1. `remainingTailLength`对应的剩余数据无法均分到每个核，只能分配给前`formerNum`个整核。用`remainingTailLength`除以`alignNum`得到整核数量`formerNum`；整核比尾核多处理一个DataBlock（含`alignNum`个`half`元素），因此`formerLength`由`baseLength`加上`alignNum`得到，而`tailLength`直接取`baseLength`。 | ||
| 38 | + | ||
| 39 | + ```cpp | ||
| 40 | + if (remainingTailLength == 0U) { | ||
| 41 | + ... | ||
| 42 | + return; | ||
| 43 | + } | ||
| 44 | + | ||
| 45 | + uint32_t formerNum = remainingTailLength / alignNum; | ||
| 46 | + uint32_t formerLength = baseLength + alignNum; | ||
| 47 | + uint32_t tailLength = baseLength; | ||
| 48 | + ``` | ||
| 49 | + | ||
| 50 | +2. 整核和尾核的主块数量相同，均等于`baseTileNum`。 | ||
| 51 | + | ||
| 52 | + ```cpp | ||
| 53 | + uint32_t formerTileNum = baseTileNum; | ||
| 54 | + uint32_t tailTileNum = baseTileNum; | ||
| 55 | + ``` | ||
| 56 | + | ||
| 57 | +3. 进一步判断`tailBlockNumEachCore`是否为0。等于0表示各核未分配到基础尾块（`baseLastTileLength`为0），因此尾核内只包含主块；整核比尾核多处理一个DataBlock，因此整核尾块长度`formerLastTileLength`直接设置为`alignNum`。此时进入尾核切分场景，使用`MainTileWithTailCoreTiling`结构体传递Tiling切分信息。 | ||
| 58 | + | ||
| 59 | + ```cpp | ||
| 60 | + if (tailBlockNumEachCore == 0U) { | ||
| 61 | + uint32_t formerLastTileLength = alignNum; | ||
| 62 | + MainTileWithTailCoreTiling tiling{ | ||
| 63 | + formerNum, formerLength, formerTileNum, formerLastTileLength, tailLength, tailTileNum}; | ||
| 64 | + add_custom<MainTileWithTailCoreTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 65 | + } | ||
| 66 | + ``` | ||
| 67 | + | ||
| 68 | +对形状为`(1, 256064)`的输入数据计算后，Tiling结构体内各个变量的值如下： | ||
| 69 | + | ||
| 70 | +```cpp | ||
| 71 | +MainTileWithTailCoreTiling tiling{ | ||
| 72 | + 4, // formerNum：前4个核为整核 | ||
| 73 | + 32016, // formerLength：整核计算32016个half | ||
| 74 | + 10, // formerTileNum：整核包含10个主块 | ||
| 75 | + 16, // formerLastTileLength：整核包含16个half的尾块 | ||
| 76 | + 32000, // tailLength：尾核计算32000个half | ||
| 77 | + 10 // tailTileNum：尾核包含10个主块 | ||
| 78 | +}; | ||
| 79 | +``` | ||
| 80 | + | ||
| 81 | +## 算子类实现 | ||
| 82 | + | ||
| 83 | +Kernel侧算子仍采用[静态Tensor编程](../../../../编程指南/编程模型/AI-Core-SIMD编程/基于Tensor的CPP编程/静态Tensor编程.md)方式实现。与[尾块均分](尾块均分.md)相比，本场景中整核和尾核处理的数据长度不同，因此`Init`函数需要根据`GetBlockIdx()`判断当前核是整核还是尾核，并分别设置Global Memory偏移和循环参数。`blockIdx`表示当前逻辑核的索引。 | ||
| 84 | + | ||
| 85 | +整核使用`formerLength`、`formerTileNum`和`formerLastTileLength`，计算如下： | ||
| 86 | + | ||
| 87 | +```cpp | ||
| 88 | +uint32_t offset = tiling.formerLength * blockIdx; | ||
| 89 | +InitGm(x, y, z, offset, tiling.formerLength); | ||
| 90 | +``` | ||
| 91 | + | ||
| 92 | +尾核使用`tailLength`和`tailTileNum`，且尾核内只包含主块，因此`lastTileLength`设置为0。尾核的Global Memory偏移需要在全部整核数据之后继续计算，代码如下： | ||
| 93 | + | ||
| 94 | +```cpp | ||
| 95 | +uint32_t offset = tiling.formerLength * tiling.formerNum + tiling.tailLength * (blockIdx - tiling.formerNum); | ||
| 96 | +InitGm(x, y, z, offset, tiling.tailLength); | ||
| 97 | +``` | ||
| 98 | + | ||
| 99 | +完整的`Init`函数实现代码如下： | ||
| 100 | + | ||
| 101 | +```cpp | ||
| 102 | +__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, | ||
| 103 | + MainTileWithTailCoreTiling tiling) | ||
| 104 | +{ | ||
| 105 | + uint32_t blockIdx = AscendC::GetBlockIdx(); | ||
| 106 | + if (blockIdx < tiling.formerNum) { | ||
| 107 | + uint32_t offset = tiling.formerLength * blockIdx; | ||
| 108 | + InitGm(x, y, z, offset, tiling.formerLength); | ||
| 109 | + this->tileNum = tiling.formerTileNum; | ||
| 110 | + this->lastTileLength = tiling.formerLastTileLength; | ||
| 111 | + } else { | ||
| 112 | + uint32_t offset = | ||
| 113 | + tiling.formerLength * tiling.formerNum + tiling.tailLength * (blockIdx - tiling.formerNum); | ||
| 114 | + InitGm(x, y, z, offset, tiling.tailLength); | ||
| 115 | + this->tileNum = tiling.tailTileNum; | ||
| 116 | + this->lastTileLength = 0U; | ||
| 117 | + } | ||
| 118 | +} | ||
| 119 | +``` | ||
| 120 | + | ||
| 121 | +完成`Init`后，`Process`函数仍按照`tileNum`和`lastTileLength`进行循环处理。整核需要处理主块和尾块，尾核只处理主块；循环逻辑与[尾块均分](尾块均分.md)中的`Process`函数一致，此处不再展开。 | ||
| @@ -0,0 +1,105 @@ | |||
| 1 | +# 尾核尾块切分 | ||
| 2 | + | ||
| 3 | +该场景中，主块均分后存在剩余数据`mainTileRemainder`，剩余数据部分能以DataBlock为最小单位均匀分配到每个核，但仍有少量DataBlock无法均分（`remainingTailLength`不为0），只能继续分配给部分核。获得额外DataBlock的核为整核，其余为尾核；整核和尾核都包含尾块，且整核处理的数据量更多。由于`tailBlockNumEachCore`不为0，各核都分配到基础尾块，因此尾核内也存在尾块；整核在尾核基础上多处理一个DataBlock。数据切分后，前`formerNum`个核为整核，各处理`formerLength`个元素（含主块和较长的尾块）；其余核为尾核，各处理`tailLength`个元素（含主块和较短的尾块）。 | ||
| 4 | + | ||
| 5 | +本场景样例的完整代码请参考[场景3：尾核尾块切分](../../../../../../../examples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/README.md#scenario3-tail-block-tail-core)。样例中固定使用8个核，主块长度`mainTileLength`为3200个`half`元素，数据整体长度为258112个`half`元素。数据按DataBlock对齐后，前4个核各处理32272个元素，后4个核各处理32256个元素；整核包含10个主块和272个元素的尾块，尾核包含10个主块和256个元素的尾块。数据切分示意如下图所示。 | ||
| 6 | + | ||
| 7 | +**图 1** 数据切分示意图 | ||
| 8 | + | ||
| 9 | + | ||
| 10 | + | ||
| 11 | +本场景中整核和尾核处理的元素数量、主块数量及尾块长度均不同，需要分别记录两类核的Tiling参数。与[尾核切分](尾核切分.md)相比，由于尾核内也存在尾块，Tiling参数在`former*`和`tail*`基础上新增`tailLastTileLength`记录尾核的尾块元素数量。Tiling结构体包含以下成员： | ||
| 12 | + | ||
| 13 | +- formerNum：分配到数据量较多的核数，即整核的核数。 | ||
| 14 | +- formerLength：整核计算的元素数量。 | ||
| 15 | +- formerTileNum：整核中主块的数量。 | ||
| 16 | +- formerLastTileLength：整核中尾块的元素数量。 | ||
| 17 | +- tailLength：尾核计算的元素数量。 | ||
| 18 | +- tailTileNum：尾核中主块的数量。 | ||
| 19 | +- tailLastTileLength：尾核中尾块的元素数量。 | ||
| 20 | + | ||
| 21 | +## Tiling实现 | ||
| 22 | + | ||
| 23 | +Tiling结构体定义如下： | ||
| 24 | + | ||
| 25 | +```cpp | ||
| 26 | +struct TailBlockAndTailCoreTiling { | ||
| 27 | + uint32_t formerNum; // 整核数量 | ||
| 28 | + uint32_t formerLength; // 整核处理的元素数量 | ||
| 29 | + uint32_t formerTileNum; // 整核中主块的数量 | ||
| 30 | + uint32_t formerLastTileLength; // 整核中尾块的元素数量 | ||
| 31 | + uint32_t tailLength; // 尾核处理的元素数量 | ||
| 32 | + uint32_t tailTileNum; // 尾核中主块的数量 | ||
| 33 | + uint32_t tailLastTileLength; // 尾核中尾块的元素数量 | ||
| 34 | +}; | ||
| 35 | +``` | ||
| 36 | + | ||
| 37 | +Host侧Tiling沿用[尾核切分](尾核切分.md)中的DataBlock对齐、主块分配、基础尾块计算及整核/尾核划分，得到：整核数量`formerNum`、每个整核处理的元素数量`formerLength`、每个尾核处理的元素数量`tailLength`、每个核基础尾块的元素数量`baseLastTileLength`、整核中主块的数量`formerTileNum`，以及尾核中主块的数量`tailTileNum`。当`remainingTailLength`和`tailBlockNumEachCore`均不为0时，尾核切分分支不成立，进入尾核尾块切分场景。此时整核比尾核多处理一个DataBlock，整核尾块长度`formerLastTileLength`为基础尾块长度`baseLastTileLength`与`alignNum`（一个DataBlock包含的`half`元素数量）之和；若两者之和等于`mainTileLength`，则将其合并为一个主块（`formerTileNum`加1，`formerLastTileLength`置0）。尾核尾块长度`tailLastTileLength`直接取`baseLastTileLength`。 | ||
| 38 | + | ||
| 39 | +```cpp | ||
| 40 | +uint32_t formerLastTileLength = baseLastTileLength + alignNum; | ||
| 41 | +if (formerLastTileLength == mainTileLength) { | ||
| 42 | + formerTileNum += 1U; | ||
| 43 | + formerLastTileLength = 0U; | ||
| 44 | +} | ||
| 45 | +uint32_t tailLastTileLength = baseLastTileLength; | ||
| 46 | +TailBlockAndTailCoreTiling tiling{ | ||
| 47 | + formerNum, formerLength, formerTileNum, formerLastTileLength, | ||
| 48 | + tailLength, tailTileNum, tailLastTileLength}; | ||
| 49 | +add_custom<TailBlockAndTailCoreTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 50 | +``` | ||
| 51 | + | ||
| 52 | +对形状为`(1, 258112)`的输入数据计算后，Tiling结构体内各个变量的值如下： | ||
| 53 | + | ||
| 54 | +```cpp | ||
| 55 | +TailBlockAndTailCoreTiling tiling{ | ||
| 56 | + 4, // formerNum：前4个核为整核 | ||
| 57 | + 32272, // formerLength：整核计算32272个half | ||
| 58 | + 10, // formerTileNum：整核包含10个主块 | ||
| 59 | + 272, // formerLastTileLength：整核包含272个half的尾块 | ||
| 60 | + 32256, // tailLength：尾核计算32256个half | ||
| 61 | + 10, // tailTileNum：尾核包含10个主块 | ||
| 62 | + 256 // tailLastTileLength：尾核包含256个half的尾块 | ||
| 63 | +}; | ||
| 64 | +``` | ||
| 65 | + | ||
| 66 | +## 算子类实现 | ||
| 67 | + | ||
| 68 | +Kernel侧算子仍采用[静态Tensor编程](../../../../编程指南/编程模型/AI-Core-SIMD编程/基于Tensor的CPP编程/静态Tensor编程.md)方式实现。与[尾核切分](尾核切分.md)相比，本场景中整核和尾核都包含尾块，因此`Init`函数需要根据`GetBlockIdx()`判断当前核是整核还是尾核，并分别设置Global Memory偏移、`tileNum`和`lastTileLength`。 | ||
| 69 | +整核使用`formerLength`、`formerTileNum`和`formerLastTileLength`，计算如下： | ||
| 70 | + | ||
| 71 | +```cpp | ||
| 72 | +uint32_t offset = tiling.formerLength * blockIdx; | ||
| 73 | +InitGm(x, y, z, offset, tiling.formerLength); | ||
| 74 | +``` | ||
| 75 | + | ||
| 76 | +尾核使用`tailLength`、`tailTileNum`和`tailLastTileLength`。尾核的Global Memory偏移需要在全部整核数据之后继续计算，代码如下： | ||
| 77 | + | ||
| 78 | +```cpp | ||
| 79 | +uint32_t offset = tiling.formerLength * tiling.formerNum + tiling.tailLength * (blockIdx - tiling.formerNum); | ||
| 80 | +InitGm(x, y, z, offset, tiling.tailLength); | ||
| 81 | +``` | ||
| 82 | + | ||
| 83 | +完整的`Init`函数实现代码如下： | ||
| 84 | + | ||
| 85 | +```cpp | ||
| 86 | +__aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, | ||
| 87 | + TailBlockAndTailCoreTiling tiling) | ||
| 88 | +{ | ||
| 89 | + uint32_t blockIdx = AscendC::GetBlockIdx(); | ||
| 90 | + if (blockIdx < tiling.formerNum) { | ||
| 91 | + uint32_t offset = tiling.formerLength * blockIdx; | ||
| 92 | + InitGm(x, y, z, offset, tiling.formerLength); | ||
| 93 | + this->tileNum = tiling.formerTileNum; | ||
| 94 | + this->lastTileLength = tiling.formerLastTileLength; | ||
| 95 | + } else { | ||
| 96 | + uint32_t offset = | ||
| 97 | + tiling.formerLength * tiling.formerNum + tiling.tailLength * (blockIdx - tiling.formerNum); | ||
| 98 | + InitGm(x, y, z, offset, tiling.tailLength); | ||
| 99 | + this->tileNum = tiling.tailTileNum; | ||
| 100 | + this->lastTileLength = tiling.tailLastTileLength; | ||
| 101 | + } | ||
| 102 | +} | ||
| 103 | +``` | ||
| 104 | + | ||
| 105 | +完成`Init`后，`Process`函数仍按照`tileNum`和`lastTileLength`进行循环处理。整核和尾核都需要处理主块和尾块，循环逻辑与[尾块均分](尾块均分.md)中的`Process`函数一致，此处不再展开。 | ||
| @@ -0,0 +1,45 @@ | |||
| 1 | +# 概述 | ||
| 2 | + | ||
| 3 | +Ascend C核函数是在Device侧执行的并行函数。为了提高算子的执行效率，通常采用多核并行计算，将输入数据切分后分配到不同的逻辑核上处理；同时，由于单个核的Local Memory容量有限，无法一次完整容纳算子的输入和输出数据，需要分批搬运数据进行计算。这种对输入数据进行切分、分块计算的过程称为**Tiling**。 | ||
| 4 | + | ||
| 5 | +切分数据的算法称为Tiling算法或Tiling策略；根据算子shape等信息计算切分相关参数（如每次搬运的块大小、循环次数等）的程序，称为**Tiling实现**，也叫Tiling函数（Tiling Function）。[基础矢量算子](../基础矢量算子.md)与[TBuf的使用](../TBuf的使用.md)样例均在单核上运行，不涉及Host侧Tiling实现。 | ||
| 6 | + | ||
| 7 | +由于Tiling实现完成的均为标量计算，而AI Core并不擅长这类计算，因此通常将其独立出来由Host侧执行；对于动态shape算子，切分参数无法在编译期确定，也需要在Host侧计算。 | ||
| 8 | + | ||
| 9 | +开发者需要分析并设计Tiling参数、定义Tiling结构体，在Host侧通过上下文获取输入输出shape信息，根据shape信息计算Tiling参数并设置到对应的Tiling结构体中；随后通过核函数入口参数将Tiling信息传入核函数。核函数内部解析Host侧传入的Tiling结构体，根据Tiling信息控制数据搬入、搬出Local Memory的流程，并完成计算逻辑。 | ||
| 10 | + | ||
| 11 | +**图 1** 算子实现组成 | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | +如上图所示，算子实现由Device侧的Kernel实现与Host侧的Tiling实现两部分组成：Kernel实现负责在Device侧完成数据搬运、矢量计算、内存管理和任务同步等逻辑；Tiling实现用于在Kernel执行前，根据输入shape等信息在Host侧计算Kernel所需的切分参数。对于输入和输出shape固定、切分策略固定的简单算子，Host侧Tiling实现可以省略；对于动态shape或多核并行场景，每个核处理的数据量、核内分块数量等参数通常需要在Host侧计算后传递给Kernel使用。 | ||
| 16 | + | ||
| 17 | +对输入数据进行切分时应遵循以下原则： | ||
| 18 | + | ||
| 19 | +1. 以DataBlock为最小切分单位。受硬件限制，进行Unified Buffer相关的数据搬运和矢量计算时，搬运的数据长度和操作数的起始地址都需要保证32字节对齐，即以DataBlock为单位进行操作，一个DataBlock大小为32字节；输入数据不满足DataBlock大小对齐要求时，需按DataBlock大小向上对齐。 | ||
| 20 | + | ||
| 21 | +2. 充分利用Unified Buffer空间。AI Core与外部存储交互时会产生性能开销，频繁地进行数据搬运会导致性能瓶颈，因此应尽可能充分利用Unified Buffer空间，减少从Global Memory搬运数据的次数。 | ||
| 22 | + | ||
| 23 | +3. 均衡利用多核计算能力。AI处理器包含多个AI Core，应将计算均衡分配到各核上，充分发挥多核并行优势。 | ||
| 24 | + | ||
| 25 | +本章基于[矢量计算多核Tiling切分策略样例](../../../../../../../examples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/README.md)，以两个一维Tensor相加的Add矢量算子为例，对多核Tiling切分的典型场景进行说明。样例实现文件为[vector_tiling_strategy.asc](../../../../../../../examples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/vector_tiling_strategy.asc)。 | ||
| 26 | + | ||
| 27 | +**图 2** 多核Tiling示意图<a name="fig-2"></a> | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | +如上图所示，将长度为`totalLength`的算子输入分配到多个核上进行计算。Host侧Tiling先将输入数据量按DataBlock对齐，再按主块长度`mainTileLength`在各核上分配相同数量的主块，随后把剩余数据以DataBlock为最小单位尽量均衡地分配到各核。当剩余数据能均分到每个核时，各核处理的数据量相同；无法均分时，各核处理的数据量分为两类——处理数据量较多的整核和处理数据量较少的尾核，其中整核数量记为`formerNum`。对于单个核的计算数据，核内先处理`tileNum`个长度为`mainTileLength`的主块；如果最后剩余的数据长度不足一个主块，则通过`lastTileLength`记录尾块长度。Tiling切分完成后，上述参数被封装到Tiling结构体并通过核函数入口传入Kernel侧；Kernel侧根据`GetBlockIdx()`和Tiling参数计算当前核需处理的输入数据的Global Memory偏移，并按`tileNum`和`lastTileLength`循环完成数据搬入、计算和搬出。 | ||
| 32 | + | ||
| 33 | +本章涉及如下概念： | ||
| 34 | + | ||
| 35 | +- 主块：基于Local Memory可用空间确定的单个核一次处理的最大数据块。主块长度使用`mainTileLength`表示，且需要按32字节对齐。 | ||
| 36 | +- 尾块：单个核内最后一个长度小于`mainTileLength`的数据块。尾块长度使用`lastTileLength`表示。 | ||
| 37 | +- 整核：核间数据无法完全均分时，处理数据量较多的核。 | ||
| 38 | +- 尾核：核间数据无法完全均分时，处理数据量较少的核。 | ||
| 39 | + | ||
| 40 | +根据数据切分的情况，有如下几种场景： | ||
| 41 | + | ||
| 42 | +1. [主块均分](主块均分.md)：输入数据可以按参与计算的核数和主块长度均匀切分。完成主块分配后没有剩余数据，各核处理的数据量相同，且核内只包含主块。 | ||
| 43 | +2. [尾块均分](尾块均分.md)：完成主块分配后仍有剩余数据，但剩余数据能以DataBlock为最小单位均匀分配到每个核。各核处理的数据量相同，核内除主块外还存在等长尾块。 | ||
| 44 | +3. [尾核切分](尾核切分.md)：完成主块分配后仍有剩余数据，但剩余数据不足以为每个核各分配一个DataBlock，只能以DataBlock为最小单位分配给部分核。获得额外DataBlock的核为整核，其余为尾核；尾核内只包含主块。 | ||
| 45 | +4. [尾核尾块切分](尾核尾块切分.md)：完成主块分配后仍有剩余数据，每个核都能分配到若干DataBlock，但仍有少量DataBlock只能分配给部分核。此时既存在处理数据量更多的整核，尾核也需要处理尾块。 | ||
| @@ -4,10 +4,11 @@ | |||
| 4 | 4 | ||
| 5 | - [基础矢量算子](基础矢量算子.md)：开发一个简单的Add矢量算子。 | 5 | - [基础矢量算子](基础矢量算子.md)：开发一个简单的Add矢量算子。 |
| 6 | - [TBuf的使用](TBuf的使用.md)：在算子计算过程中使用临时空间存储运算的中间结果。 | 6 | - [TBuf的使用](TBuf的使用.md)：在算子计算过程中使用临时空间存储运算的中间结果。 |
| 7 | -- [多核Tiling](多核-Tiling切分/多核Tiling.md)：算子在AI处理器的多个核上运行，所有核的计算数据量相等且32字节对齐。 | 7 | +- [多核Tiling切分](多核Tiling切分/概述.md)：当算子需要多核并行计算或支持动态shape时，对输入数据进行切分、分块计算，并将不同数据块分配到多个AI Core上处理。 |
| 8 | -- [尾块Tiling](多核-Tiling切分/尾块Tiling.md)：算子在AI处理器的多个核上运行，所有核的计算数据量相等，每个核上除最后一个数据块（尾块）外，其余数据块的数据量相等，每个核都需要处理尾块数据的计算。 | 8 | + - [主块均分](多核Tiling切分/主块均分.md)：算子在AI处理器的多个核上运行，所有核的计算数据量相等且32字节对齐。 |
| 9 | -- [尾核Tiling](多核-Tiling切分/尾核Tiling.md)：算子在AI处理器的多个核上运行，数据无法平均分配到每个核。将所有核分为多个整核和多个尾核，整核的计算数据量相等，尾核的计算数据量相等。 | 9 | + - [尾块均分](多核Tiling切分/尾块均分.md)：算子在AI处理器的多个核上运行，所有核的计算数据量相等，每个核上除最后一个数据块（尾块）外，其余数据块的数据量相等，每个核都需要处理尾块数据的计算。 |
| 10 | -- [尾核&尾块](多核-Tiling切分/尾核-尾块.md)：算子在AI处理器的多个核上运行，数据无法平均分配到每个核，同时每个核内的数据无法均分，除最后一个数据块（尾块）外，其余数据块的数据量相等，每个核都需要单独处理尾块数据的计算。 | 10 | + - [尾核切分](多核Tiling切分/尾核切分.md)：算子在AI处理器的多个核上运行，数据无法平均分配到每个核。将所有核分为多个整核和多个尾核，整核的计算数据量相等，尾核的计算数据量相等。 |
| 11 | + - [尾核尾块切分](多核Tiling切分/尾核尾块切分.md)：算子在AI处理器的多个核上运行，数据无法平均分配到每个核，同时每个核内的数据无法均分，除最后一个数据块（尾块）外，其余数据块的数据量相等，每个核都需要单独处理尾块数据的计算。 | ||
| 11 | - [DoubleBuffer场景](DoubleBuffer场景.md)：开启double buffer，算子中的多条流水并行执行。 | 12 | - [DoubleBuffer场景](DoubleBuffer场景.md)：开启double buffer，算子中的多条流水并行执行。 |
| 12 | - [Broadcast场景](Broadcast场景.md)：算子中两个输入的shape（形状）不相等，需要将一个输入的shape进行Broadcast（广播）后，再执行计算。 | 13 | - [Broadcast场景](Broadcast场景.md)：算子中两个输入的shape（形状）不相等，需要将一个输入的shape进行Broadcast（广播）后，再执行计算。 |
| 13 | - [非对齐场景](非对齐场景.md)：更多数据非32字节对齐场景的处理方案。 | 14 | - [非对齐场景](非对齐场景.md)：更多数据非32字节对齐场景的处理方案。 |
| @@ -4,6 +4,8 @@ | |||
| 4 | 4 | ||
| 5 | 【描述】AI处理器的物理核数是固定的，当L2 Cache切分之后，可能发生部分核有计算拖尾的情况，即每次所有核计算量除以每个核处理的数据量不能被核数整除，导致最后需要部分尾核来计算尾块数据。而在尾核计算时，部分核始终处于空闲状态，从而使得算子的整体性能变差。如[图1](#fig26395321518)，假设总的数据量为TotalSize，L2 Cache切分之后分为两份TotalSize / 2，每个核每次的计算量为TotalSize / 2 / 25，即需要25个核进行处理，由于AI处理器的核数为20，因此每次计算时，1到5核的每个核需要多算一份数据，导致发生拖尾的情况。 | 5 | 【描述】AI处理器的物理核数是固定的，当L2 Cache切分之后，可能发生部分核有计算拖尾的情况，即每次所有核计算量除以每个核处理的数据量不能被核数整除，导致最后需要部分尾核来计算尾块数据。而在尾核计算时，部分核始终处于空闲状态，从而使得算子的整体性能变差。如[图1](#fig26395321518)，假设总的数据量为TotalSize，L2 Cache切分之后分为两份TotalSize / 2，每个核每次的计算量为TotalSize / 2 / 25，即需要25个核进行处理，由于AI处理器的核数为20，因此每次计算时，1到5核的每个核需要多算一份数据，导致发生拖尾的情况。 |
| 6 | 6 | ||
| 7 | +更详细的多核Tiling切分过程请参见[多核Tiling切分](../../SIMD算子实现/矢量编程/多核Tiling切分/概述.md)。 | ||
| 8 | + | ||
| 7 | 【反例】 | 9 | 【反例】 |
| 8 | 10 | ||
| 9 | **图1** 计算拖尾示意图<a name="fig26395321518"></a> | 11 | **图1** 计算拖尾示意图<a name="fig26395321518"></a> |
| @@ -15,10 +15,10 @@ set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, | |||
| 15 | 15 | ||
| 16 | # 场景参数（通过命令行传入） | 16 | # 场景参数（通过命令行传入） |
| 17 | # SCENARIO_NUM: Tiling切分场景编号 | 17 | # SCENARIO_NUM: Tiling切分场景编号 |
| 18 | -# 0 - 仅主块 | 18 | +# 0 - 主块均分 |
| 19 | -# 1 - 主块+尾块 | 19 | +# 1 - 尾块均分 |
| 20 | -# 2 - 主块+尾核 | 20 | +# 2 - 尾核切分 |
| 21 | -# 3 - 尾块+尾核 | 21 | +# 3 - 尾核尾块切分 |
| 22 | set(SCENARIO_NUM "0" CACHE STRING "Tiling scenario number: 0-3") | 22 | set(SCENARIO_NUM "0" CACHE STRING "Tiling scenario number: 0-3") |
| 23 | 23 | ||
| 24 | message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}") | 24 | message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}") |
| @@ -4,7 +4,7 @@ | |||
| 4 | 4 | ||
| 5 | 本样例以两个一维Tensor相加为例，介绍矢量计算多核Tiling切分策略。该策略根据输入数据量计算不同的核间和核内切分参数，使数据在多个AI Core之间尽量均衡分配，从而充分利用多核计算能力实现高效计算。 | 5 | 本样例以两个一维Tensor相加为例，介绍矢量计算多核Tiling切分策略。该策略根据输入数据量计算不同的核间和核内切分参数，使数据在多个AI Core之间尽量均衡分配，从而充分利用多核计算能力实现高效计算。 |
| 6 | 6 | ||
| 7 | -样例固定使用8个AI Core，核内主块数据量为3200个`half`元素（即200个DataBlock），通过CMake编译参数`SCENARIO_NUM`选择不同数据量场景。 | 7 | +样例固定使用8个核，核内主块数据量为3200个`half`元素（即200个DataBlock），通过CMake编译参数`SCENARIO_NUM`选择不同数据量场景。 |
| 8 | 8 | ||
| 9 | ## 本样例支持的产品及CANN软件版本 | 9 | ## 本样例支持的产品及CANN软件版本 |
| 10 | 10 | ||
| @@ -31,30 +31,38 @@ | |||
| 31 | 31 | ||
| 32 | 本样例实现两个一维Tensor相加，计算逻辑为`z = x + y`。样例通过CMake编译参数`SCENARIO_NUM`选择不同数据量场景，对应不同的Tiling切分场景，所有场景数据格式为ND，输入输出均为`half`类型，核函数名为`add_custom`。 | 32 | 本样例实现两个一维Tensor相加，计算逻辑为`z = x + y`。样例通过CMake编译参数`SCENARIO_NUM`选择不同数据量场景，对应不同的Tiling切分场景，所有场景数据格式为ND，输入输出均为`half`类型，核函数名为`add_custom`。 |
| 33 | 33 | ||
| 34 | -样例固定使用8个AI Core，主块数据量为3200个`half`元素，其中“主块”（`MAIN_TILE_LENGTH`）是用户基于UB可用空间确定的一次可以处理的最大数据量，且满足32字节对齐。 | 34 | +样例固定使用8个核，主块数据量为3200个`half`元素，其中“主块”（`mainTileLength`）是用户基于UB可用空间确定的一次可以处理的最大数据量，且满足32字节对齐。 |
| 35 | 35 | ||
| 36 | -**场景0：仅主块** | 36 | +<a name="scenario0-main-tile-only"></a> |
| 37 | + | ||
| 38 | +### 场景0：主块均分 | ||
| 37 | 39 | ||
| 38 | - 输入：`x`、`y`均为[1, 256000]个`half`元素 | 40 | - 输入：`x`、`y`均为[1, 256000]个`half`元素 |
| 39 | - 输出：`z`为[1, 256000]个`half`元素 | 41 | - 输出：`z`为[1, 256000]个`half`元素 |
| 40 | - Tiling切分：8个核各处理32000个元素；每个核包含10个主块，每个主块3200个元素 | 42 | - Tiling切分：8个核各处理32000个元素；每个核包含10个主块，每个主块3200个元素 |
| 41 | -- 说明：所有核处数据量相同，核内只包含主块 | 43 | +- 说明：所有核处理数据量相同，核内只包含主块 |
| 42 | 44 | ||
| 43 | -**场景1：主块+尾块** | 45 | +<a name="scenario1-main-tile-tail-block"></a> |
| 46 | + | ||
| 47 | +### 场景1：尾块均分 | ||
| 44 | 48 | ||
| 45 | - 输入：`x`、`y`均为[1, 260096]个`half`元素 | 49 | - 输入：`x`、`y`均为[1, 260096]个`half`元素 |
| 46 | - 输出：`z`为[1, 260096]个`half`元素 | 50 | - 输出：`z`为[1, 260096]个`half`元素 |
| 47 | - Tiling切分：8个核各处理32512个元素；每个核包含10个主块和512个元素的尾块 | 51 | - Tiling切分：8个核各处理32512个元素；每个核包含10个主块和512个元素的尾块 |
| 48 | - 说明：所有核处理数据量相同，每个核在主块后都有等长尾块 | 52 | - 说明：所有核处理数据量相同，每个核在主块后都有等长尾块 |
| 49 | 53 | ||
| 50 | -**场景2：主块+尾核** | 54 | +<a name="scenario2-main-tile-tail-core"></a> |
| 55 | + | ||
| 56 | +### 场景2：尾核切分 | ||
| 51 | 57 | ||
| 52 | - 输入：`x`、`y`均为[1, 256064]个`half`元素 | 58 | - 输入：`x`、`y`均为[1, 256064]个`half`元素 |
| 53 | - 输出：`z`为[1, 256064]个`half`元素 | 59 | - 输出：`z`为[1, 256064]个`half`元素 |
| 54 | - Tiling切分：前4个整核各处理32016个元素，后4个尾核各处理32000个元素；整核包含10个主块和16个元素的尾块，尾核包含10个主块 | 60 | - Tiling切分：前4个整核各处理32016个元素，后4个尾核各处理32000个元素；整核包含10个主块和16个元素的尾块，尾核包含10个主块 |
| 55 | - 说明：前4个整核比尾核多处理1个DataBlock，尾核无尾块 | 61 | - 说明：前4个整核比尾核多处理1个DataBlock，尾核无尾块 |
| 56 | 62 | ||
| 57 | -**场景3：尾块+尾核** | 63 | +<a name="scenario3-tail-block-tail-core"></a> |
| 64 | + | ||
| 65 | +### 场景3：尾核尾块切分 | ||
| 58 | 66 | ||
| 59 | - 输入：`x`、`y`均为[1, 258112]个`half`元素 | 67 | - 输入：`x`、`y`均为[1, 258112]个`half`元素 |
| 60 | - 输出：`z`为[1, 258112]个`half`元素 | 68 | - 输出：`z`为[1, 258112]个`half`元素 |
| @@ -113,42 +121,58 @@ | |||
| 113 | totalLengthAligned = AlignUp(totalLength, alignNum); | 121 | totalLengthAligned = AlignUp(totalLength, alignNum); |
| 114 | ``` | 122 | ``` |
| 115 | 123 | ||
| 116 | -2. 先在所有核上分配相同数量的主块`MAIN_TILE_LENGTH`。 | 124 | +2. 先在所有核上分配相同数量的主块`mainTileLength`。 |
| 117 | 125 | ||
| 118 | ```cpp | 126 | ```cpp |
| 119 | - mainTileNum = totalLengthAligned / (numBlocks * MAIN_TILE_LENGTH); | 127 | + mainTileNum = totalLengthAligned / (numBlocks * mainTileLength); |
| 120 | - mainTileRemainder = totalLengthAligned % (numBlocks * MAIN_TILE_LENGTH); | 128 | + mainTileRemainder = totalLengthAligned % (numBlocks * mainTileLength); |
| 121 | ``` | 129 | ``` |
| 122 | 130 | ||
| 123 | - 此时每个核至少处理`mainTileNum * MAIN_TILE_LENGTH`个元素，`mainTileRemainder`表示完成主块分配后的剩余数据量。 | 131 | + 此时每个核至少处理`mainTileNum * mainTileLength`个元素，`mainTileRemainder`表示完成主块分配后的剩余数据量。当`mainTileRemainder`为0时，进入主块均分场景并结束Tiling计算。 |
| 124 | - | ||
| 125 | -3. `mainTileRemainder`再按DataBlock给所有核补齐等长尾块。 | ||
| 126 | 132 | ||
| 127 | ```cpp | 133 | ```cpp |
| 128 | - tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); | 134 | + if (mainTileRemainder == 0) { |
| 129 | - formerCoreRemainder = mainTileRemainder % (numBlocks * alignNum); | 135 | + // 使用MainTileOnlyTiling。 |
| 130 | - baseLength = mainTileNum * MAIN_TILE_LENGTH + tailBlockNumEachCore * alignNum; | 136 | + return; |
| 131 | - ``` | ||
| 132 | - | ||
| 133 | - `baseLength`表示每个核处理的基础数据量。`formerCoreRemainder`表示每核补齐等长DataBlock尾块后，剩余需要分给前若干个核的数据量。 | ||
| 134 | - | ||
| 135 | -4. 最后剩余的DataBlock分给前若干个整核。 | ||
| 136 | - | ||
| 137 | - ```cpp | ||
| 138 | - if (formerCoreRemainder == 0) { | ||
| 139 | - formerNum = numBlocks; | ||
| 140 | - formerLength = baseLength; | ||
| 141 | - tailNum = 0; | ||
| 142 | - tailLength = 0; | ||
| 143 | - } else { | ||
| 144 | - formerNum = formerCoreRemainder / alignNum; | ||
| 145 | - formerLength = baseLength + alignNum; | ||
| 146 | - tailNum = numBlocks - formerNum; | ||
| 147 | - tailLength = baseLength; | ||
| 148 | } | 137 | } |
| 149 | ``` | 138 | ``` |
| 150 | 139 | ||
| 151 | - `formerCoreRemainder == 0`时不存在尾核，所有核都处理`baseLength`个元素。`formerCoreRemainder != 0`时，`formerNum`表示整核数量，前`formerNum`个核处理`formerLength`个元素；`tailNum`表示尾核数量，剩余核处理`tailLength`个元素。 | 140 | +3. `mainTileRemainder`不为0时，再按DataBlock给所有核补齐等长尾块。 |
| 141 | + | ||
| 142 | + ```cpp | ||
| 143 | + tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); | ||
| 144 | + remainingTailLength = mainTileRemainder % (numBlocks * alignNum); | ||
| 145 | + baseLength = mainTileNum * mainTileLength + tailBlockNumEachCore * alignNum; | ||
| 146 | + ``` | ||
| 147 | + | ||
| 148 | + `baseLength`表示每个核处理的基础数据量。`remainingTailLength`表示每核补齐等长DataBlock尾块后尚未分配的元素数量，其值为`alignNum`的整数倍。 | ||
| 149 | + | ||
| 150 | +4. 判断尾块均分场景。当`remainingTailLength`为0时，表示所有核均分到等长尾块，进入尾块均分场景。 | ||
| 151 | + | ||
| 152 | + ```cpp | ||
| 153 | + if (remainingTailLength == 0) { | ||
| 154 | + // 使用MainTileWithTailBlockTiling。 | ||
| 155 | + return; | ||
| 156 | + } | ||
| 157 | + ``` | ||
| 158 | + | ||
| 159 | +5. `remainingTailLength`不为0时，将最后剩余的DataBlock分给前若干个整核，并计算整核和尾核的数据长度。 | ||
| 160 | + | ||
| 161 | + ```cpp | ||
| 162 | + formerNum = remainingTailLength / alignNum; | ||
| 163 | + formerLength = baseLength + alignNum; | ||
| 164 | + tailLength = baseLength; | ||
| 165 | + ``` | ||
| 166 | + | ||
| 167 | +6. 判断尾核切分场景。`tailBlockNumEachCore`为0时，尾核内只包含主块，进入尾核切分场景；否则，进入尾核尾块切分场景。 | ||
| 168 | + | ||
| 169 | + ```cpp | ||
| 170 | + if (tailBlockNumEachCore == 0) { | ||
| 171 | + // 使用MainTileWithTailCoreTiling。 | ||
| 172 | + } else { | ||
| 173 | + // 使用TailBlockAndTailCoreTiling。 | ||
| 174 | + } | ||
| 175 | + ``` | ||
| 152 | 176 | ||
| 153 | ## 编译运行 | 177 | ## 编译运行 |
| 154 | 178 | ||
| @@ -193,7 +217,7 @@ | |||
| 193 | |------|--------|------| | 217 | |------|--------|------| |
| 194 | | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 | | 218 | | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 | |
| 195 | | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和 Atlas A3 训练系列产品/Atlas A3 推理系列产品，dav-3510 对应 Ascend 950PR/Ascend 950DT | | 219 | | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和 Atlas A3 训练系列产品/Atlas A3 推理系列产品，dav-3510 对应 Ascend 950PR/Ascend 950DT | |
| 196 | - | `SCENARIO_NUM` | `0`（默认）、`1`、`2`、`3` | 场景编号：0（仅主块）、1（主块+尾块）、2（主块+尾核）、3（尾块+尾核） | | 220 | + | `SCENARIO_NUM` | `0`（默认）、`1`、`2`、`3` | 场景编号：0（主块均分）、1（尾块均分）、2（尾核切分）、3（尾核尾块切分） | |
| 197 | 221 | ||
| 198 | - 执行结果 | 222 | - 执行结果 |
| 199 | 223 | ||
| @@ -4,7 +4,7 @@ | |||
| 4 | 4 | ||
| 5 | This sample uses two one-dimensional Tensors as inputs for addition. It introduces a multi-core Tiling split strategy for vector computing. The strategy calculates different inter-core and intra-core split parameters based on the input data volume. It distributes data across multiple AI Cores as evenly as possible to improve multi-core computing efficiency. | 5 | This sample uses two one-dimensional Tensors as inputs for addition. It introduces a multi-core Tiling split strategy for vector computing. The strategy calculates different inter-core and intra-core split parameters based on the input data volume. It distributes data across multiple AI Cores as evenly as possible to improve multi-core computing efficiency. |
| 6 | 6 | ||
| 7 | -The sample uses 8 AI Cores. The main tile data volume in each core is 3200 `half` elements, which equals 200 DataBlocks. Use the CMake build parameter `SCENARIO_NUM` to select a data volume scenario. | 7 | +The sample uses 8 Cores. The main tile data volume in each core is 3200 `half` elements, which equals 200 DataBlocks. Use the CMake build parameter `SCENARIO_NUM` to select a data volume scenario. |
| 8 | 8 | ||
| 9 | ## Supported Products and CANN Software Versions | 9 | ## Supported Products and CANN Software Versions |
| 10 | 10 | ||
| @@ -31,30 +31,30 @@ The sample uses 8 AI Cores. The main tile data volume in each core is 3200 `half | |||
| 31 | 31 | ||
| 32 | This sample adds two one-dimensional Tensors. The compute logic is `z = x + y`. Use the CMake build parameter `SCENARIO_NUM` to select a data volume scenario. Each scenario maps to a different Tiling split scenario. All scenarios use the ND data format. The input and output data type is `half`. The kernel function name is `add_custom`. | 32 | This sample adds two one-dimensional Tensors. The compute logic is `z = x + y`. Use the CMake build parameter `SCENARIO_NUM` to select a data volume scenario. Each scenario maps to a different Tiling split scenario. All scenarios use the ND data format. The input and output data type is `half`. The kernel function name is `add_custom`. |
| 33 | 33 | ||
| 34 | -The sample uses 8 AI Cores. The main tile data volume is 3200 `half` elements. The main tile (`MAIN_TILE_LENGTH`) is the maximum data volume that one operation can process based on the available Unified Buffer (UB) space. It meets 32-byte alignment. | 34 | +The sample uses 8 Cores. The main tile data volume is 3200 `half` elements. The main tile (`mainTileLength`) is the maximum data volume that one operation can process based on the available Unified Buffer (UB) space. It meets 32-byte alignment. |
| 35 | 35 | ||
| 36 | -**Scenario 0: Main Tiles Only** | 36 | +**Scenario 0: Main Tile Equal Split** |
| 37 | 37 | ||
| 38 | - Input: `x` and `y` each contain [1, 256000] `half` elements. | 38 | - Input: `x` and `y` each contain [1, 256000] `half` elements. |
| 39 | - Output: `z` contains [1, 256000] `half` elements. | 39 | - Output: `z` contains [1, 256000] `half` elements. |
| 40 | - Tiling split: Each of the 8 cores processes 32000 elements. Each core contains 10 main tiles, and each main tile contains 3200 elements. | 40 | - Tiling split: Each of the 8 cores processes 32000 elements. Each core contains 10 main tiles, and each main tile contains 3200 elements. |
| 41 | - Description: All cores process the same data volume. Each core contains only main tiles. | 41 | - Description: All cores process the same data volume. Each core contains only main tiles. |
| 42 | 42 | ||
| 43 | -**Scenario 1: Main Tiles and Tail Block** | 43 | +**Scenario 1: Tail Block Equal Split** |
| 44 | 44 | ||
| 45 | - Input: `x` and `y` each contain [1, 260096] `half` elements. | 45 | - Input: `x` and `y` each contain [1, 260096] `half` elements. |
| 46 | - Output: `z` contains [1, 260096] `half` elements. | 46 | - Output: `z` contains [1, 260096] `half` elements. |
| 47 | - Tiling split: Each of the 8 cores processes 32512 elements. Each core contains 10 main tiles and a tail block with 512 elements. | 47 | - Tiling split: Each of the 8 cores processes 32512 elements. Each core contains 10 main tiles and a tail block with 512 elements. |
| 48 | - Description: All cores process the same data volume. Each core has an equal-length tail block after the main tiles. | 48 | - Description: All cores process the same data volume. Each core has an equal-length tail block after the main tiles. |
| 49 | 49 | ||
| 50 | -**Scenario 2: Main Tiles and Tail Core** | 50 | +**Scenario 2: Tail Core Split** |
| 51 | 51 | ||
| 52 | - Input: `x` and `y` each contain [1, 256064] `half` elements. | 52 | - Input: `x` and `y` each contain [1, 256064] `half` elements. |
| 53 | - Output: `z` contains [1, 256064] `half` elements. | 53 | - Output: `z` contains [1, 256064] `half` elements. |
| 54 | - Tiling split: The first 4 former cores each process 32016 elements, and the remaining 4 tail cores each process 32000 elements. Each former core contains 10 main tiles and a tail block with 16 elements. Each tail core contains 10 main tiles. | 54 | - Tiling split: The first 4 former cores each process 32016 elements, and the remaining 4 tail cores each process 32000 elements. Each former core contains 10 main tiles and a tail block with 16 elements. Each tail core contains 10 main tiles. |
| 55 | - Description: The first 4 former cores each process one more DataBlock than the tail cores. The tail cores do not have tail blocks. | 55 | - Description: The first 4 former cores each process one more DataBlock than the tail cores. The tail cores do not have tail blocks. |
| 56 | 56 | ||
| 57 | -**Scenario 3: Tail Block and Tail Core** | 57 | +**Scenario 3: Tail Core Split with Tail Blocks** |
| 58 | 58 | ||
| 59 | - Input: `x` and `y` each contain [1, 258112] `half` elements. | 59 | - Input: `x` and `y` each contain [1, 258112] `half` elements. |
| 60 | - Output: `z` contains [1, 258112] `half` elements. | 60 | - Output: `z` contains [1, 258112] `half` elements. |
| @@ -113,42 +113,58 @@ The Tiling calculation process of this sample is as follows. The host side recei | |||
| 113 | totalLengthAligned = AlignUp(totalLength, alignNum); | 113 | totalLengthAligned = AlignUp(totalLength, alignNum); |
| 114 | ``` | 114 | ``` |
| 115 | 115 | ||
| 116 | -2. Allocate the same number of main tiles, `MAIN_TILE_LENGTH`, to all cores. | 116 | +2. Allocate the same number of main tiles, `mainTileLength`, to all cores. |
| 117 | 117 | ||
| 118 | ```cpp | 118 | ```cpp |
| 119 | - mainTileNum = totalLengthAligned / (numBlocks * MAIN_TILE_LENGTH); | 119 | + mainTileNum = totalLengthAligned / (numBlocks * mainTileLength); |
| 120 | - mainTileRemainder = totalLengthAligned % (numBlocks * MAIN_TILE_LENGTH); | 120 | + mainTileRemainder = totalLengthAligned % (numBlocks * mainTileLength); |
| 121 | ``` | 121 | ``` |
| 122 | 122 | ||
| 123 | - At this point, each core processes at least `mainTileNum * MAIN_TILE_LENGTH` elements. `mainTileRemainder` indicates the remaining data volume after main tile allocation. | 123 | + At this point, each core processes at least `mainTileNum * mainTileLength` elements. `mainTileRemainder` indicates the remaining data volume after main tile allocation. If `mainTileRemainder` is 0, select the main-tiles-only scenario and finish the Tiling calculation. |
| 124 | - | ||
| 125 | -3. Allocate `mainTileRemainder` to all cores as equal-length DataBlock tail blocks. | ||
| 126 | 124 | ||
| 127 | ```cpp | 125 | ```cpp |
| 128 | - tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); | 126 | + if (mainTileRemainder == 0) { |
| 129 | - formerCoreRemainder = mainTileRemainder % (numBlocks * alignNum); | 127 | + // Use MainTileOnlyTiling. |
| 130 | - baseLength = mainTileNum * MAIN_TILE_LENGTH + tailBlockNumEachCore * alignNum; | 128 | + return; |
| 131 | - ``` | ||
| 132 | - | ||
| 133 | - `baseLength` indicates the base data volume processed by each core. `formerCoreRemainder` indicates the remaining data volume that must be allocated to the first several cores after each core receives equal-length DataBlock tail blocks. | ||
| 134 | - | ||
| 135 | -4. Allocate the remaining DataBlocks to the first several former cores. | ||
| 136 | - | ||
| 137 | - ```cpp | ||
| 138 | - if (formerCoreRemainder == 0) { | ||
| 139 | - formerNum = numBlocks; | ||
| 140 | - formerLength = baseLength; | ||
| 141 | - tailNum = 0; | ||
| 142 | - tailLength = 0; | ||
| 143 | - } else { | ||
| 144 | - formerNum = formerCoreRemainder / alignNum; | ||
| 145 | - formerLength = baseLength + alignNum; | ||
| 146 | - tailNum = numBlocks - formerNum; | ||
| 147 | - tailLength = baseLength; | ||
| 148 | } | 129 | } |
| 149 | ``` | 130 | ``` |
| 150 | 131 | ||
| 151 | - If `formerCoreRemainder == 0`, no tail core exists, and all cores process `baseLength` elements. If `formerCoreRemainder != 0`, `formerNum` indicates the number of former cores. The first `formerNum` cores process `formerLength` elements. `tailNum` indicates the number of tail cores, and the remaining cores process `tailLength` elements. | 132 | +3. If `mainTileRemainder` is not 0, allocate it to all cores as equal-length DataBlock tail blocks. |
| 133 | + | ||
| 134 | + ```cpp | ||
| 135 | + tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); | ||
| 136 | + remainingTailLength = mainTileRemainder % (numBlocks * alignNum); | ||
| 137 | + baseLength = mainTileNum * mainTileLength + tailBlockNumEachCore * alignNum; | ||
| 138 | + ``` | ||
| 139 | + | ||
| 140 | + `baseLength` indicates the base data volume processed by each core. `remainingTailLength` indicates the number of elements that remain unallocated after each core receives equal-length DataBlock tail blocks. Its value is an integer multiple of `alignNum`. | ||
| 141 | + | ||
| 142 | +4. Check for the tail block equal split scenario. If `remainingTailLength` is 0, all cores receive equal-length tail blocks, and this scenario is selected. | ||
| 143 | + | ||
| 144 | + ```cpp | ||
| 145 | + if (remainingTailLength == 0) { | ||
| 146 | + // Use MainTileWithTailBlockTiling. | ||
| 147 | + return; | ||
| 148 | + } | ||
| 149 | + ``` | ||
| 150 | + | ||
| 151 | +5. If `remainingTailLength` is not 0, allocate the remaining DataBlocks to the first several former cores and calculate the data lengths of former and tail cores. | ||
| 152 | + | ||
| 153 | + ```cpp | ||
| 154 | + formerNum = remainingTailLength / alignNum; | ||
| 155 | + formerLength = baseLength + alignNum; | ||
| 156 | + tailLength = baseLength; | ||
| 157 | + ``` | ||
| 158 | + | ||
| 159 | +6. Check for the tail core split scenario. If `tailBlockNumEachCore` is 0, tail cores contain only main tiles, and this scenario is selected. Otherwise, select the tail core split with tail blocks scenario. | ||
| 160 | + | ||
| 161 | + ```cpp | ||
| 162 | + if (tailBlockNumEachCore == 0) { | ||
| 163 | + // Use MainTileWithTailCoreTiling. | ||
| 164 | + } else { | ||
| 165 | + // Use TailBlockAndTailCoreTiling. | ||
| 166 | + } | ||
| 167 | + ``` | ||
| 152 | 168 | ||
| 153 | ## Build and Run | 169 | ## Build and Run |
| 154 | 170 | ||
| @@ -193,7 +209,7 @@ In the sample root directory, perform the following steps to build and run the s | |||
| 193 | |------|--------|------| | 209 | |------|--------|------| |
| 194 | | `CMAKE_ASC_RUN_MODE` | `npu` (default), `cpu`, `sim` | Run mode: NPU execution, CPU debugging, and NPU simulation | | 210 | | `CMAKE_ASC_RUN_MODE` | `npu` (default), `cpu`, `sim` | Run mode: NPU execution, CPU debugging, and NPU simulation | |
| 195 | | `CMAKE_ASC_ARCHITECTURES` | `dav-2201` (default), `dav-3510` | NPU architecture: `dav-2201` maps to Atlas A2 training series products, Atlas A2 inference series products, Atlas A3 training series products, and Atlas A3 inference series products. `dav-3510` maps to Ascend 950PR and Ascend 950DT | | 211 | | `CMAKE_ASC_ARCHITECTURES` | `dav-2201` (default), `dav-3510` | NPU architecture: `dav-2201` maps to Atlas A2 training series products, Atlas A2 inference series products, Atlas A3 training series products, and Atlas A3 inference series products. `dav-3510` maps to Ascend 950PR and Ascend 950DT | |
| 196 | - | `SCENARIO_NUM` | `0` (default), `1`, `2`, `3` | Scenario number: 0 (main tiles only), 1 (main tiles and tail block), 2 (main tiles and tail core), and 3 (tail block and tail core) | | 212 | + | `SCENARIO_NUM` | `0` (default), `1`, `2`, `3` | Scenario number: 0 (main tile equal split), 1 (tail block equal split), 2 (tail core split), and 3 (tail core split with tail blocks) | |
| 197 | 213 | ||
| 198 | - Execution result. | 214 | - Execution result. |
| 199 | 215 | ||
Mexamples/01_simd_cpp_api/02_features/02_tiling_selector/vector_tiling_strategy/vector_tiling_strategy.asc+65-70
| @@ -12,7 +12,7 @@ | |||
| 12 | * \file vector_tiling_strategy.asc | 12 | * \file vector_tiling_strategy.asc |
| 13 | * \brief | 13 | * \brief |
| 14 | * 以两个一维Tensor相加为例，演示如何根据数据总量与核数的关系计算Tiling参数。 | 14 | * 以两个一维Tensor相加为例，演示如何根据数据总量与核数的关系计算Tiling参数。 |
| 15 | - * 样例覆盖四种典型场景：仅主块、主块+尾块、主块+尾核、尾块+尾核。 | 15 | + * 样例覆盖四种典型场景：主块均分、尾块均分、尾核切分、尾核尾块切分。 |
| 16 | */ | 16 | */ |
| 17 | 17 | ||
| 18 | #include <cstdint> | 18 | #include <cstdint> |
| @@ -27,57 +27,44 @@ | |||
| 27 | #endif | 27 | #endif |
| 28 | 28 | ||
| 29 | constexpr uint32_t scenarioNum = SCENARIO_NUM; | 29 | constexpr uint32_t scenarioNum = SCENARIO_NUM; |
| 30 | -constexpr uint32_t NUM_BLOCKS = 8; | 30 | +constexpr uint32_t numBlocks = 8; // 参与计算的逻辑核数 |
| 31 | -constexpr uint32_t BLOCK_SIZE = 32; | 31 | +constexpr uint32_t dataBlockSize = 32; // 一个DataBlock的字节数 |
| 32 | -constexpr uint32_t HALF_SIZE = 2; | 32 | +constexpr uint32_t halfSize = 2; // 一个half元素的字节数 |
| 33 | -constexpr uint32_t ALIGN_NUM = BLOCK_SIZE / HALF_SIZE; | 33 | +constexpr uint32_t alignNum = dataBlockSize / halfSize; // 一个DataBlock包含的half元素数量 |
| 34 | -constexpr uint32_t MAIN_TILE_LENGTH = 3200; // 主块数据量 | 34 | +constexpr uint32_t mainTileLength = 3200; // 一个主块包含的half元素数量 |
| 35 | 35 | ||
| 36 | struct MainTileOnlyTiling { | 36 | struct MainTileOnlyTiling { |
| 37 | - uint32_t blockLength; // 每个核处理的数据量 | 37 | + uint32_t blockLength; // 每个核处理的元素数量 |
| 38 | uint32_t tileNum; // 主块数量 | 38 | uint32_t tileNum; // 主块数量 |
| 39 | }; | 39 | }; |
| 40 | 40 | ||
| 41 | struct MainTileWithTailBlockTiling { | 41 | struct MainTileWithTailBlockTiling { |
| 42 | - uint32_t blockLength; // 每个核处理的数据量 | 42 | + uint32_t blockLength; // 每个核处理的元素数量 |
| 43 | uint32_t tileNum; // 主块数量 | 43 | uint32_t tileNum; // 主块数量 |
| 44 | - uint32_t lastTileLength; // 尾块数据量 | 44 | + uint32_t lastTileLength; // 尾块元素数量 |
| 45 | }; | 45 | }; |
| 46 | 46 | ||
| 47 | struct MainTileWithTailCoreTiling { | 47 | struct MainTileWithTailCoreTiling { |
| 48 | uint32_t formerNum; // 整核数量 | 48 | uint32_t formerNum; // 整核数量 |
| 49 | - uint32_t formerLength; // 整核处理的数据量 | 49 | + uint32_t formerLength; // 整核处理的元素数量 |
| 50 | uint32_t formerTileNum; // 整核中主块的数量 | 50 | uint32_t formerTileNum; // 整核中主块的数量 |
| 51 | - uint32_t formerLastTileLength; // 整核中尾块的数据量 | 51 | + uint32_t formerLastTileLength; // 整核中尾块的元素数量 |
| 52 | - uint32_t tailLength; // 尾核处理的数据量 | 52 | + uint32_t tailLength; // 尾核处理的元素数量 |
| 53 | uint32_t tailTileNum; // 尾核中主块的数量 | 53 | uint32_t tailTileNum; // 尾核中主块的数量 |
| 54 | }; | 54 | }; |
| 55 | 55 | ||
| 56 | struct TailBlockAndTailCoreTiling { | 56 | struct TailBlockAndTailCoreTiling { |
| 57 | uint32_t formerNum; // 整核数量 | 57 | uint32_t formerNum; // 整核数量 |
| 58 | - uint32_t formerLength; // 整核处理的数据量 | 58 | + uint32_t formerLength; // 整核处理的元素数量 |
| 59 | uint32_t formerTileNum; // 整核中主块的数量 | 59 | uint32_t formerTileNum; // 整核中主块的数量 |
| 60 | - uint32_t formerLastTileLength; // 整核中尾块的数据量 | 60 | + uint32_t formerLastTileLength; // 整核中尾块的元素数量 |
| 61 | - uint32_t tailLength; // 尾核处理的数据量 | 61 | + uint32_t tailLength; // 尾核处理的元素数量 |
| 62 | uint32_t tailTileNum; // 尾核中主块的数量 | 62 | uint32_t tailTileNum; // 尾核中主块的数量 |
| 63 | - uint32_t tailLastTileLength; // 尾核中尾块的数据量 | 63 | + uint32_t tailLastTileLength; // 尾核中尾块的元素数量 |
| 64 | -}; | ||
| 65 | - | ||
| 66 | -struct TileInfo { | ||
| 67 | - uint32_t tileNum; | ||
| 68 | - uint32_t lastTileLength; | ||
| 69 | }; | 64 | }; |
| 70 | 65 | ||
| 71 | uint32_t AlignUp(uint32_t value, uint32_t align) { return ((value + align - 1U) / align) * align; } | 66 | uint32_t AlignUp(uint32_t value, uint32_t align) { return ((value + align - 1U) / align) * align; } |
| 72 | 67 | ||
| 73 | -TileInfo CalcTileInfo(uint32_t length) | ||
| 74 | -{ | ||
| 75 | - TileInfo info; | ||
| 76 | - info.tileNum = length / MAIN_TILE_LENGTH; | ||
| 77 | - info.lastTileLength = length - info.tileNum * MAIN_TILE_LENGTH; | ||
| 78 | - return info; | ||
| 79 | -} | ||
| 80 | - | ||
| 81 | class KernelAdd { | 68 | class KernelAdd { |
| 82 | public: | 69 | public: |
| 83 | __aicore__ inline KernelAdd() = default; | 70 | __aicore__ inline KernelAdd() = default; |
| @@ -137,14 +124,14 @@ public: | |||
| 137 | 124 | ||
| 138 | __aicore__ inline void Process() | 125 | __aicore__ inline void Process() |
| 139 | { | 126 | { |
| 140 | - AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, MAIN_TILE_LENGTH); | 127 | + AscendC::LocalTensor<half> xLocal(AscendC::TPosition::VECCALC, xAddr, mainTileLength); |
| 141 | - AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, MAIN_TILE_LENGTH); | 128 | + AscendC::LocalTensor<half> yLocal(AscendC::TPosition::VECCALC, yAddr, mainTileLength); |
| 142 | - AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, MAIN_TILE_LENGTH); | 129 | + AscendC::LocalTensor<half> zLocal(AscendC::TPosition::VECCALC, zAddr, mainTileLength); |
| 143 | 130 | ||
| 144 | uint32_t totalLoop = this->tileNum + (this->lastTileLength > 0U ? 1U : 0U); | 131 | uint32_t totalLoop = this->tileNum + (this->lastTileLength > 0U ? 1U : 0U); |
| 145 | for (uint32_t loopIdx = 0; loopIdx < totalLoop; loopIdx++) { | 132 | for (uint32_t loopIdx = 0; loopIdx < totalLoop; loopIdx++) { |
| 146 | - uint32_t curLen = (loopIdx < this->tileNum) ? MAIN_TILE_LENGTH : this->lastTileLength; | 133 | + uint32_t curLen = (loopIdx < this->tileNum) ? mainTileLength : this->lastTileLength; |
| 147 | - uint32_t startElement = loopIdx * MAIN_TILE_LENGTH; | 134 | + uint32_t startElement = loopIdx * mainTileLength; |
| 148 | 135 | ||
| 149 | if (loopIdx != 0U) { | 136 | if (loopIdx != 0U) { |
| 150 | AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); | 137 | AscendC::WaitFlag<AscendC::HardEvent::V_MTE2>(EVENT_ID0); |
| @@ -187,8 +174,8 @@ private: | |||
| 187 | 174 | ||
| 188 | private: | 175 | private: |
| 189 | static constexpr uint32_t xAddr = 0; | 176 | static constexpr uint32_t xAddr = 0; |
| 190 | - static constexpr uint32_t yAddr = MAIN_TILE_LENGTH * HALF_SIZE; | 177 | + static constexpr uint32_t yAddr = mainTileLength * halfSize; |
| 191 | - static constexpr uint32_t zAddr = yAddr + MAIN_TILE_LENGTH * HALF_SIZE; | 178 | + static constexpr uint32_t zAddr = yAddr + mainTileLength * halfSize; |
| 192 | 179 | ||
| 193 | AscendC::GlobalTensor<half> xGm; | 180 | AscendC::GlobalTensor<half> xGm; |
| 194 | AscendC::GlobalTensor<half> yGm; | 181 | AscendC::GlobalTensor<half> yGm; |
| @@ -211,51 +198,59 @@ void LaunchAddByTiling(uint8_t* xDevice, uint8_t* yDevice, uint8_t* zDevice, uin | |||
| 211 | { | 198 | { |
| 212 | // Host侧Tiling逐级计算主块数量、尾块数量、尾核数量。 | 199 | // Host侧Tiling逐级计算主块数量、尾块数量、尾核数量。 |
| 213 | // 第零步，将数据量按32字节对齐 | 200 | // 第零步，将数据量按32字节对齐 |
| 214 | - uint32_t totalLengthAligned = AlignUp(totalLength, ALIGN_NUM); | 201 | + uint32_t totalLengthAligned = AlignUp(totalLength, alignNum); |
| 215 | 202 | ||
| 216 | // 第一步，所有核均分相同数量的主块。 | 203 | // 第一步，所有核均分相同数量的主块。 |
| 217 | - uint32_t mainTileNum = totalLengthAligned / (NUM_BLOCKS * MAIN_TILE_LENGTH); | 204 | + uint32_t mainTileNum = totalLengthAligned / (numBlocks * mainTileLength); |
| 218 | - uint32_t mainTileRemainder = totalLengthAligned % (NUM_BLOCKS * MAIN_TILE_LENGTH); | 205 | + uint32_t mainTileRemainder = totalLengthAligned % (numBlocks * mainTileLength); |
| 206 | + uint32_t baseLength = mainTileNum * mainTileLength; | ||
| 207 | + uint32_t baseTileNum = mainTileNum; | ||
| 208 | + | ||
| 209 | + if (mainTileRemainder == 0U) { | ||
| 210 | + MainTileOnlyTiling tiling{baseLength, baseTileNum}; | ||
| 211 | + printf("Tiling mode: main_tile_only\n"); | ||
| 212 | + add_custom<MainTileOnlyTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 213 | + return; | ||
| 214 | + } | ||
| 219 | 215 | ||
| 220 | // 第二步，mainTileRemainder按DataBlock给每个核补齐等长尾块。 | 216 | // 第二步，mainTileRemainder按DataBlock给每个核补齐等长尾块。 |
| 221 | - uint32_t tailBlockNumEachCore = mainTileRemainder / (NUM_BLOCKS * ALIGN_NUM); | 217 | + uint32_t tailBlockNumEachCore = mainTileRemainder / (numBlocks * alignNum); |
| 222 | - uint32_t formerCoreRemainder = mainTileRemainder % (NUM_BLOCKS * ALIGN_NUM); | 218 | + uint32_t remainingTailLength = mainTileRemainder % (numBlocks * alignNum); |
| 223 | - uint32_t baseLength = mainTileNum * MAIN_TILE_LENGTH + tailBlockNumEachCore * ALIGN_NUM; | 219 | + uint32_t baseLastTileLength = tailBlockNumEachCore * alignNum; |
| 224 | - TileInfo baseInfo = CalcTileInfo(baseLength); | 220 | + baseLength += baseLastTileLength; |
| 225 | 221 | ||
| 226 | - // 如果没有剩余DataBlock，所有核都处理baseLength，只需要区分是否存在尾块。 | 222 | + // 如果没有剩余DataBlock，所有核都处理baseLength，每个核都有等长尾块。 |
| 227 | - if (formerCoreRemainder == 0U) { | 223 | + if (remainingTailLength == 0U) { |
| 228 | - if (baseInfo.lastTileLength == 0U) { | 224 | + MainTileWithTailBlockTiling tiling{baseLength, baseTileNum, baseLastTileLength}; |
| 229 | - MainTileOnlyTiling tiling{baseLength, baseInfo.tileNum}; | 225 | + printf("Tiling mode: main_tile_with_tail_block\n"); |
| 230 | - printf("Tiling mode: main_tile_only\n"); | 226 | + add_custom<MainTileWithTailBlockTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); |
| 231 | - add_custom<MainTileOnlyTiling><<<NUM_BLOCKS, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 232 | - } else { | ||
| 233 | - MainTileWithTailBlockTiling tiling{baseLength, baseInfo.tileNum, baseInfo.lastTileLength}; | ||
| 234 | - printf("Tiling mode: main_tile_with_tail_block\n"); | ||
| 235 | - add_custom<MainTileWithTailBlockTiling><<<NUM_BLOCKS, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | ||
| 236 | - } | ||
| 237 | return; | 227 | return; |
| 238 | } | 228 | } |
| 239 | 229 | ||
| 240 | // 第三步，最后剩余的DataBlock分给前formerNum个整核。 | 230 | // 第三步，最后剩余的DataBlock分给前formerNum个整核。 |
| 241 | - uint32_t formerNum = formerCoreRemainder / ALIGN_NUM; | 231 | + uint32_t formerNum = remainingTailLength / alignNum; |
| 242 | - uint32_t formerLength = baseLength + ALIGN_NUM; // 前formerNum个核都多处理一个DataBlock的数据 | 232 | + uint32_t formerLength = baseLength + alignNum; // 前formerNum个核都多处理一个DataBlock的数据 |
| 243 | - uint32_t tailLength = baseLength; | 233 | + uint32_t tailLength = baseLength; // 尾核处理的数据量 |
| 244 | - TileInfo formerInfo = CalcTileInfo(formerLength); | 234 | + uint32_t formerTileNum = baseTileNum; |
| 245 | - TileInfo tailInfo = baseInfo; | 235 | + uint32_t tailTileNum = baseTileNum; |
| 246 | - bool formerOnlyOneDataBlockTail = (formerInfo.lastTileLength == ALIGN_NUM); | 236 | + // 未分配基础尾块时进入尾核切分场景，否则进入尾核尾块切分场景。 |
| 247 | - bool tailHasNoTailTile = (tailInfo.lastTileLength == 0U); | 237 | + if (tailBlockNumEachCore == 0U) { |
| 248 | - if (formerOnlyOneDataBlockTail && tailHasNoTailTile) { | 238 | + uint32_t formerLastTileLength = alignNum; |
| 249 | - MainTileWithTailCoreTiling tiling{formerNum, formerLength, formerInfo.tileNum, formerInfo.lastTileLength, | 239 | + MainTileWithTailCoreTiling tiling{formerNum, formerLength, formerTileNum, formerLastTileLength, |
| 250 | - tailLength, tailInfo.tileNum}; | 240 | + tailLength, tailTileNum}; |
| 251 | printf("Tiling mode: main_tile_with_tail_core\n"); | 241 | printf("Tiling mode: main_tile_with_tail_core\n"); |
| 252 | - add_custom<MainTileWithTailCoreTiling><<<NUM_BLOCKS, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | 242 | + add_custom<MainTileWithTailCoreTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); |
| 253 | } else { | 243 | } else { |
| 254 | - TailBlockAndTailCoreTiling tiling{ | 244 | + uint32_t formerLastTileLength = baseLastTileLength + alignNum; |
| 255 | - formerNum, formerLength, formerInfo.tileNum, formerInfo.lastTileLength, | 245 | + if (formerLastTileLength == mainTileLength) { |
| 256 | - tailLength, tailInfo.tileNum, tailInfo.lastTileLength}; | 246 | + formerTileNum += 1U; |
| 247 | + formerLastTileLength = 0U; | ||
| 248 | + } | ||
| 249 | + uint32_t tailLastTileLength = baseLastTileLength; | ||
| 250 | + TailBlockAndTailCoreTiling tiling{formerNum, formerLength, formerTileNum, formerLastTileLength, | ||
| 251 | + tailLength, tailTileNum, tailLastTileLength}; | ||
| 257 | printf("Tiling mode: tail_block_and_tail_core\n"); | 252 | printf("Tiling mode: tail_block_and_tail_core\n"); |
| 258 | - add_custom<TailBlockAndTailCoreTiling><<<NUM_BLOCKS, 0, stream>>>(xDevice, yDevice, zDevice, tiling); | 253 | + add_custom<TailBlockAndTailCoreTiling><<<numBlocks, 0, stream>>>(xDevice, yDevice, zDevice, tiling); |
| 259 | } | 254 | } |
| 260 | } | 255 | } |
| 261 | 256 | ||
| @@ -265,7 +260,7 @@ int32_t main(int32_t argc, char* argv[]) | |||
| 265 | (scenarioNum == 1) ? 260096 : | 260 | (scenarioNum == 1) ? 260096 : |
| 266 | (scenarioNum == 2) ? 256064 : | 261 | (scenarioNum == 2) ? 256064 : |
| 267 | 258112; | 262 | 258112; |
| 268 | - size_t dataByteSize = static_cast<size_t>(dataLen) * HALF_SIZE; | 263 | + size_t dataByteSize = static_cast<size_t>(dataLen) * halfSize; |
| 269 | 264 | ||
| 270 | printf("Usage: %s (compile-time SCENARIO_NUM=%u)\n", argv[0], scenarioNum); | 265 | printf("Usage: %s (compile-time SCENARIO_NUM=%u)\n", argv[0], scenarioNum); |
| 271 | printf("Running scenario %u, totalLength=%u\n", scenarioNum, dataLen); | 266 | printf("Running scenario %u, totalLength=%u\n", scenarioNum, dataLen); |
增加align num和datablock的关系