diff --git a/OWNERS b/OWNERS index 0ebceeb59f2d3061240faa111907bcae38bad17a..c0e0590d6d61ece329a045509149c9e09ac765a4 100644 --- a/OWNERS +++ b/OWNERS @@ -22,6 +22,7 @@ approvers: - vicia - kong0808 - zhuliangying +- wang-xiaozhi reviewers: - wuzhaolinhuawei - horming diff --git a/atvc/README.md b/atvc/README.md index f3085fff3a644e4d04e2bd25fb70c3430f2ffe5f..ce4a1b871c216db6d824acd28afd10d273fb965c 100644 --- a/atvc/README.md +++ b/atvc/README.md @@ -1,15 +1,10 @@ # ATVC -ATVC(Ascend Template for Vector Compute)是一个用AscendC API搭建的C++模板头文件集合，旨在帮助开发快速实现AscendC Vector算子。它将AscendC Vector算子开发流程中的"动态部分"基于C++的模板解耦成可自定义的模块并提供一系列Vector算子基类以及Tiling算法等API。 +ATVC(AscendC Template for Vector Compute)是一个用AscendC API搭建的C++模板头文件集合，旨在帮助用户快速开发AscendC典型Vector算子。它将AscendC Vector算子开发流程中的计算实现解耦成可自定义的模块, 内部封装实现了kernel数据搬入搬出等底层通用操作及通用tiling计算，实现了高效的算子开发模式。 +相比传统AscendC算子开发方式，利用ATVC搭建的Vector算子可做到开发效率提升3-5倍。用户只需选择匹配的模板并完成核心计算逻辑就完成算子kernel侧开发，atvc还内置了每个模板库对应的通用tiling计算实现，可省去用户手写tiling的开发量就能达到不错的性能表现，极大提升算子开发效率。 -相比传统AscendC算子，利用ATVC搭建的Vector算子可做到开发代码量减少~50%， 用户只需搭建核心的计算逻辑便可灵活且快速完成Vector算子的编写，极大提升算子开发效率。 - -ATVC将Vector算子开发流程中的可定制化模块抽象出了Host层和Kernel层，两层的定义如下：
-- Host层：在CPU Host侧执行，它提供了一系列Tiling计算与策略选择的API，帮助用户计算出较优的数据搬运等运行态参数； -- Kernel层：在NPU侧调用，它是利用AscendC API搭建出的一系列Vector算子模板类，内置了算子开发中用户无需感知的数据搬入搬出以及资源申请等固定模块，并将核心计算模块开放给用户定义。 - -![architecture.png](./docs/data/architecture.png)
+![atvc_user_case.png](./docs/data/atvc_user_case.png)
请参阅[快速入门](./docs/1_quick_start.md)以快速了解ATVC的Add算子搭建流程。 @@ -70,16 +65,22 @@ Accuracy verification passed. ``` - +# 已支持的模版 +| Vector模版类型 | +| ------------------------------------------------------------ | +| Ele-wise模板 | +| Reduce模板 | +| Broadcast模板 | # 样例介绍 | 样例名 | 描述 | | ------------------------------------------------------------ | ------------------------------------------------------------ | -| [add](./examples/add/add.cpp) | ATVC实现的简单Add算子以及调用样例 | -| [add_with_scalar](./examples/add_with_scalar/add_with_scalar.cpp) | ATVC实现的通过标量控制计算逻辑的自定义EleWise算子以及调用样例 | -| [reduce_sum](./examples/reduce_sum/reduce_sum.cpp) | ATVC实现的自定义ReduceSum算子以及调用样例 | -| [sinh_custom](./examples/sinh_custom/sinh_custom.cpp) | ATVC实现的临时Tensor参与计算的Sinh自定义算子以及调用样例 | -| [broadcast_to](./examples/broadcast_to/broadcast_to.cpp) | ATVC实现的BroadcastTo自定义算子以及调用样例 | +| [add](./examples/add/add.cpp) | 使用ATVC的Ele-wise模板实现Add算子以及调用样例 | +| [sinh_custom](./examples/sinh_custom/sinh_custom.cpp) | 临时Tensor参与计算的自定义Ele-wise类算子以及调用样例 | +| [add_with_scalar](./examples/add_with_scalar/add_with_scalar.cpp) | 输入带标量的自定义Ele-wise类算子以及调用样例 | +| [reduce_sum](./examples/reduce_sum/reduce_sum.cpp) | 使用ATVC的Reduce模板实现自定义ReduceSum算子以及调用样例 | +| [broadcast_to](./examples/broadcast_to/broadcast_to.cpp) | 使用ATVC的Broadcast模板实现自定义BroadcastTo算子以及调用样例 | +更多算子类型介绍和如何选取模板参见参阅[快速入门](./docs/1_quick_start.md)。 @@ -87,7 +88,7 @@ Accuracy verification passed. | 算子模板 | 数据类型 | | ------------------------------------------------------------ | ------------------------------------------------------------ | -| EleWise | int32_t、float | +| Ele-wise | int32_t、float | | Reduce | int32_t、float | | Broadcast | int32_t、float | diff --git a/atvc/docs/1_quick_start.md b/atvc/docs/1_quick_start.md index 40625386d6167330a6d4fd4a8454f538c9ed1cef..3c5c469b49f1b6048ff53f5b6d4f9637a2273b22 100644 --- a/atvc/docs/1_quick_start.md +++ b/atvc/docs/1_quick_start.md @@ -133,4 +133,15 @@ $ bash run_test.sh broadcast_to # 执行broadcast样例 ``` ## 完整样例
-完整代码样例请参照[examples/add/add.cpp](../examples/add/add.cpp) \ No newline at end of file +完整代码样例请参照[examples/add/add.cpp](../examples/add/add.cpp) + +# 模板选择 +用户根据待开发的Vector算子定义特征，选择匹配的模板及其配套的tiling算法，若自定义算子不在当前模板库的范围内，建议使用基本AscnedC API 手写算子。 +## Ele-wise类算子 +Ele-wise类算子通常是指对张量进行元素级别的操作的函数或方法，包括但不限于加、减、乘、除及指数、对数、三角函数等数学函数。这类算子的特点是会逐元素进行计算操作，而不会改变输入数据的形状。常见的ELe-wise算子有Add、Sub、Exp、Log、Sin、Sqrt等。 +## Reduce类算子 +Reduce类算子通常是指对张量中的元素进行归约操作的算子，通常用来求和、求平均值等操作，可指定某几个维度进行归约计算，也可以将所有元素归约计算为一个标量。常见的Reduce类算子有ReduceSum(求和)、ReduceMean(求平均值)、ReduceProdcut(累乘)、ReduceMax(求最大值)、ReduceMin(求最小值)、ReduceAny(or操作)、ReduceAll(and操作)。 +## Broadcast +Broadcast算子是指完成广播操作，用于处理不同形状的张量间的运算。 +例如tensorA的shape是（1，5），tensorB的shape是（3，5），若要实现tensorC = tensorA + tensorB，实际上需要将TensorA广播为shape（3，5）,再进行tensorA和tensorB的相加操作。广播的过程，实际上就是将原数据在某个维度上进行复制扩展。 + diff --git a/atvc/docs/2_developer_guide.md b/atvc/docs/2_developer_guide.md index 625418a9be75509d42612fe13d157f105fa12bda..412e4ce5f3d19ca80d62dd7490023721d7ef0f94 100644 --- a/atvc/docs/2_developer_guide.md +++ b/atvc/docs/2_developer_guide.md @@ -62,7 +62,7 @@ using AddOpTraits = ATVC::OpTraits; // Add算 ### 1.1.2 Param -ATVC框架提供了`ATVC::EleWiseParam`、`ATVC::ReduceParam`、`ATVC::BroadCastParam` 三个结构体来描述算子内部调度的Tiling数据和其他资源变量。Param 作为Host侧Tiling API的输出，它将传入ATVC框架的Kernel层算子模板，并在运行时指导算子内部模块完成数据的循环搬运与调度计算。
+ATVC框架提供了`ATVC::EleWiseParam`、`ATVC::ReduceParam`、`ATVC::BroadcastParam` 三个结构体来描述算子内部调度的Tiling数据和其他资源变量。Param 作为Host侧Tiling API的输出，它将传入ATVC框架的Kernel层算子模板，并在运行时指导算子内部模块完成数据的循环搬运与调度计算。
以下为ElementWise类算子的`ATVC::EleWiseParam`参与计算的伪代码，详细使用流程请参考本文档的 2.1.5 Host层API： ```cpp @@ -81,7 +81,7 @@ EleWiseKernel<<>>(x, y, z, paramDevi ``` ### 1.1.3 Policy -编译态参数`Policy`(`ATVC::ReducePolicy`, `ATVC::BroadCastPolicy`)是ATVC框架里Kernel层对部分算子模板的拓展描述，它对应算子模板类在不同场景的实例化实现。它由Tiling API计算出，并在策略分派API(`ATVC::Host::ReduceAdapter`)里将运行态的Policy结果转化为模板参数并调用该场景下的最佳模板实现来完成高效的数据计算。
+编译态参数`Policy`(`ATVC::ReducePolicy`, `ATVC::BroadcastPolicy`)是ATVC框架里Kernel层对部分算子模板的拓展描述，它对应算子模板类在不同场景的实例化实现。它由Tiling API计算出，并在策略分派API(`ATVC::Host::ReduceAdapter`)里将运行态的Policy结果转化为模板参数并调用该场景下的最佳模板实现来完成高效的数据计算。
以下为Reduce算子开发场景中`ATVC::ReducePolicy`参与计算的伪代码，详细过程请参考2.2.5 Host层API： ```cpp @@ -172,22 +172,6 @@ public: // 完成变长参数的解析和数据调度计算 // } -private: - - __aicore__ inline void Init() { - // - // 初始化计算资源 - // - }; - - __aicore__ inline void Process() { - // - // 根据param_参数，循环搬运特定长度数据到Ub，调用compute_的仿函数完成计算后，再从UB搬出到GM - // - }; - - EleWiseCompute compute_; // 开发自定义的计算模板类 - __gm__ EleWiseParam* param_; // CalEleWiseTiling API计算出的运行态参数 } ``` @@ -235,15 +219,6 @@ __global__ __aicore__ void SinhCustom(GM_ADDR x, GM_ADDR y, GM_ADDR param) auto op = ATVC::Kernel::EleWiseOpTemplate>(); op.Run(x, y, param); // 按照输入、输出、param的顺序传入Run函数中；OpTraits内部的ATVC::OpTemps将由EleWiseOpTemplate内部申请资源，开发无需关注 } -// -// Host侧调用核函数样例 -// -int main() -{ - //... - SinhCustom<<>>(xDevice, yDevice, paramDevice); - // ... -} ```
@@ -262,15 +237,6 @@ extern "C" __global__ __aicore__ void SinhCustom(GM_ADDR x, GM_ADDR y, GM_ADDR p auto op = ATVC::Kernel::EleWiseOpTemplate>(); // 模板参数传入固定的SinhOpTraits op.Run(x, y, param); } -// -// Host侧调用核函数样例 -// -int main() -{ - //... - SinhCustom<<>>(xDevice, yDevice, paramDevice); - // ... -} ``` ### 2.1.5 Host层API @@ -590,20 +556,6 @@ public: // Reduce类算子Run接口按输入、输出、运行态参数param顺序传入 // } -private: - __aicore__ inline void Init() { - // - // 初始化计算资源 - // - }; - - __aicore__ inline void Process() { - // - // 根据param_参数，循环搬运特定长度数据到Ub，调用compute_的仿函数完成计算(其中包含了基块结果之间的更新、UB之间的数据结果更新等计算)后，再从UB搬出到GM - // - }; - ReduceCompute compute_; // 开发自定义的计算模板类 - __gm__ ReduceParam* param_; // CalcReduceTiling API计算出的运行态参数 } ``` @@ -646,16 +598,6 @@ __global__ __aicore__ void ReduceSumCustom(GM_ADDR x, GM_ADDR y, GM_ADDR param) auto op = ATVC::Kernel::ReduceOpTemplate, SelectPolicy>(); op.Run(x, y, param); } -// -// Host侧调用核函数样例 -// -int main() -{ - //... - static constexpr ATVC::ReducePolicy SelectPolicy { ATVC::AR_PATTERN::ARARA, ATVC::AR_COUNT::A3R0, 0 }; // 该policy由策略分派API给出 - ReduceSumCustom<<>>(xDevice, yDevice, paramDevice); - // ... -} ```
@@ -1069,7 +1011,7 @@ struct BroadcastPolicy { #include "broadcast/broadcast_utils/broadcast_buf_pool.h" namespace ATVC { namespace Kernel { -template +template class BroadcastOpTemplate { public: using DataType = typename BroadcastCompute::DataType; @@ -1087,17 +1029,6 @@ public: this->Process(); } -private: - __aicore__ inline void Init(GM_ADDR src, GM_ADDR dst, GM_ADDR broadcastParam) - { - ... // 完成参数、buffer、UB内存等资源的初始化 - } - - __aicore__ inline void Process() - { - ... // 循环完成多轮数据搬入、UB计算调度、数据搬出 - } - AscendC::GlobalTensor srcGlobal; AscendC::GlobalTensor dstGlobal; BroadcastCompute compute_; @@ -1151,22 +1082,6 @@ __global__ __aicore__ void BroadcastCustom(GM_ADDR x, GM_ADDR y, GM_ADDR broadca auto op = ATVC::Kernel::BroadcastOpTemplate, Policy>(); op.Run(x, y, broadcastParam); } - -// host侧调用示例 -int32_t main(int32_t argc, char* argv[]) -{ - // acl资源初始化 - ... - - // 调用kernel 核函数 - BroadcastCustom<<>>(xDevice, yDevice, paramDevice); - - // 释放Acl资源 - ... - - return 0; -} - ```
@@ -1220,7 +1135,7 @@ bool CalcBroadcastTiling(std::vector shapeIn, std::vector shap using DataType = typename ATVC::TypeListGet::Type; auto inputDtype = GetOriInputType(); BroadcastTilingInputParam opInput = {shapeIn, shapeOut, inputDtype}; - OpTiling::BroadCastOpTiling tiling(opInput, policy, param); + OpTiling::BroadcastOpTiling tiling(opInput, policy, param); if(!tiling.Run()) { printf("[ERROR] Tiling Error\n"); return false; diff --git a/atvc/docs/data/atvc_user_case.png b/atvc/docs/data/atvc_user_case.png new file mode 100644 index 0000000000000000000000000000000000000000..6bbdc693f2c2e95f5404d767dc69228c44bed556 Binary files /dev/null and b/atvc/docs/data/atvc_user_case.png differ diff --git a/atvc/include/broadcast/broadcast_host.h b/atvc/include/broadcast/broadcast_host.h index fbfb583a3b75559dfedafb7280b9b552e60e763c..e6035a42950ce803ec0ca6cbfcf2c3a59f5801bd 100644 --- a/atvc/include/broadcast/broadcast_host.h +++ b/atvc/include/broadcast/broadcast_host.h @@ -44,7 +44,7 @@ void PrintParam(BroadcastPolicy* policy, BroadcastParam* param) printf("[Broadcast] Tiling result: dstShape[%d] = %lu\n", i, param->tilingData.dstShape[i]); } printf("[Broadcast] Tiling result: policy.patternID = %d\n", policy->patternID); - printf("[Broadcast] Tiling result: workspaceSize = %d\n", param->workspaceSize); + printf("[Broadcast] Tiling result: workspaceSize = %u\n", param->workspaceSize); return; } @@ -63,7 +63,7 @@ bool CalcBroadcastTiling(std::vector shapeIn, std::vector shap using DataType = typename ATVC::TypeListGet::Type; auto inputDtype = GetOriInputType(); BroadcastTilingInputParam opInput = {shapeIn, shapeOut, inputDtype}; - OpTiling::BroadCastOpTiling tiling(opInput, policy, param); + OpTiling::BroadcastOpTiling tiling(opInput, policy, param); if(!tiling.Run()) { printf("[ERROR] Tiling Error\n"); return false; diff --git a/atvc/include/broadcast/broadcast_op_template.h b/atvc/include/broadcast/broadcast_op_template.h index b9a751818f065471f6a0c1bc421381de531ac0ca..8dff4f5d10cb14fd28c800911826a451febf30a7 100644 --- a/atvc/include/broadcast/broadcast_op_template.h +++ b/atvc/include/broadcast/broadcast_op_template.h @@ -31,7 +31,7 @@ struct BroadcastDataView }; namespace Kernel { -template +template class BroadcastOpTemplate { public: using DataType = typename BroadcastCompute::DataType; @@ -72,7 +72,7 @@ private: SyncDataQueue(); for (int i = 0; i < view.B1; i++) { uint32_t copyOutOffset; - if (SelectBroadCastPolicy.patternID == AB_PATTERN::ABA) { + if (SelectBroadcastPolicy.patternID == AB_PATTERN::ABA) { copyOutOffset = dimBCount * view.dimASize + dimACount * tilingData_->A2; } else { copyOutOffset = dimACount * tilingData_->A2 * view.dimBSize + dimBCount; @@ -112,7 +112,7 @@ private: AscendC::LocalTensor output; bufPool_.AllocTensor(output); SyncDataQueue(); - compute_.template Compute(input, inputOffset, output, + compute_.template Compute(input, inputOffset, output, OpsUtils::CeilAlign(tilingData_->A2, UB_ALIGN_COUNT), OpsUtils::CeilAlign(tilingData_->B2, UB_ALIGN_COUNT)); bufPool_.SetCopyOutSync(output); @@ -133,7 +133,7 @@ private: { uint32_t copyOutBaseOffset = 0; // 计算拷出偏移基址 - if (SelectBroadCastPolicy.patternID == AB_PATTERN::ABA) { + if (SelectBroadcastPolicy.patternID == AB_PATTERN::ABA) { if (tilingData_->A0 != 1) { // 核间A切分， 取部分A copyOutBaseOffset += view.dimAOffset; } @@ -153,7 +153,7 @@ private: __aicore__ inline void CalcView(BroadcastDataView &view) { - if (SelectBroadCastPolicy.patternID == AB_PATTERN::ABA) { + if (SelectBroadcastPolicy.patternID == AB_PATTERN::ABA) { view.dimASize = tilingData_->dstShape[1]; view.dimBSize = tilingData_->dstShape[0]; view.inShape[0] = 1; diff --git a/atvc/include/broadcast/tiling/broadcast_tiling.h b/atvc/include/broadcast/tiling/broadcast_tiling.h index a96ddd82907a1180a1a7aa3a5c9a1d8b1ae87b0f..5cc2a3c7e389fc0f36a8ddb13de696532e5f12c0 100644 --- a/atvc/include/broadcast/tiling/broadcast_tiling.h +++ b/atvc/include/broadcast/tiling/broadcast_tiling.h @@ -37,9 +37,9 @@ struct BroadcastTilingInputParam { namespace OpTiling { constexpr static int32_t BRC_BASIC_NUM = 4; // broadcast输入输出内存基本块分配个数 -class BroadCastOpTiling { +class BroadcastOpTiling { public: - BroadCastOpTiling(ATVC::BroadcastTilingInputParam& inputParam, + BroadcastOpTiling(ATVC::BroadcastTilingInputParam& inputParam, ATVC::BroadcastPolicy* policy, ATVC::BroadcastParam* param) : opInput_(inputParam), param_(param), policy_(policy) { diff --git a/atvc/include/reduce/reduce_host.h b/atvc/include/reduce/reduce_host.h index 9acc2f327e6a41f2193fb5945220dbecd474c674..fe99aca7dfe5349036a650edcf9370cfa36c9a9c 100644 --- a/atvc/include/reduce/reduce_host.h +++ b/atvc/include/reduce/reduce_host.h @@ -30,7 +30,7 @@ void PrintParam(ReducePolicy* policy, ReduceParam* param) printf("[Reduce] Tiling result: basicBlock = %zu\n", param->tilingData.basicBlock); printf("[Reduce] Tiling result: coreNum = %d\n", param->tilingData.coreNum); printf("[Reduce] Tiling result: nBufferNum = %d\n", param->nBufferNum); - printf("[Reduce] Tiling result: workspaceSize = %zu\n", param->workspaceSize); + printf("[Reduce] Tiling result: workspaceSize = %u\n", param->workspaceSize); printf("[Reduce] Tiling result: policy = (%d, %d, %d)\n", policy->patternID, policy->loopARCount, policy->loopInnerARCount); return;