昇腾C数学计算API使用样例-北京尧图网络

更多样例【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit样例一下面的样例展示了数学库kernel侧API和Tiling API GetXxxMaxMinTmpSize的配套使用方法具体流程如下Host侧调用Tiling接口获取所需临时空间的大小并将其写入tiling data中kernel侧再读取tiling data获取相应的临时空间大小并根据此分配临时空间。Host侧Tiling API使用样例:#include vector #include register/op_def_registry.h #include register/tilingdata_base.h #include tiling/tiling_api.h namespace optiling { BEGIN_TILING_DATA_DEF(AsinCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, srcSize); TILING_DATA_FIELD_DEF(uint32_t, tmpBufferSize); END_TILING_DATA_DEF; static ge::graphStatus TilingFunc(gert::TilingContext* context) { // Input source shapes. auto shape_input context-GetInputTensor(0)-GetOriginShape(); std::vectorint64_t srcDims {shape_input.GetDim(0), shape_input.GetDim(1)}; uint32_t srcSize 1; for (auto dim : srcDims) { srcSize * dim; } uint32_t typeSize 2; ge::Shape shape(srcDims); uint32_t minValue 0; uint32_t maxValue 0; AscendC::GetAsinMaxMinTmpSize(shape, typeSize, false, maxValue, minValue); auto platformInfo context-GetPlatformInfo(); auto ascendcPlatform platform_ascendc::PlatformAscendC(platformInfo); uint64_t tailSize 0; // ub剩余空间大小 ascendcPlatform.GetCoreMemSize( platform_ascendc::CoreMemType::UB, tailSize); // 本样例中使用完整的ub空间实际情况下tailSize需要减掉用户已使用的ub空间 auto tmpSize tailSize maxValue ? maxValue : tailSize; AsinCustomTilingData tiling; tiling.set_srcSize(srcSize); tiling.set_tmpBufferSize(tmpSize); context-SetBlockDim(1); tiling.SaveToBuffer(context-GetRawTilingData()-GetData(), context-GetRawTilingData()-GetCapacity()); context-GetRawTilingData()-SetDataSize(tiling.GetDataSize()); context-SetTilingKey(1); return ge::GRAPH_SUCCESS; } } // namespace optilingkernel侧读取tiling data获取相应的临时空间大小并根据此分配临时空间#include kernel_operator.h template typename srcType class KernelAsin { public: __aicore__ inline KernelAsin() {} __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize, uint32_t tmpBufferSize) { srcGlobal.SetGlobalBuffer(reinterpret_cast__gm__ srcType*(srcGm), srcSize); dstGlobal.SetGlobalBuffer(reinterpret_cast__gm__ srcType*(dstGm), srcSize); pipe.InitBuffer(inQueue, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(outQueue, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(tmpBuf, tmpBufferSize); bufferSize srcSize; } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensorsrcType srcLocal inQueue.AllocTensorsrcType(); AscendC::DataCopy(srcLocal, srcGlobal, bufferSize); inQueue.EnQue(srcLocal); } __aicore__ inline void Compute() { AscendC::LocalTensorsrcType dstLocal outQueue.AllocTensorsrcType(); AscendC::LocalTensorsrcType srcLocal inQueue.DeQuesrcType(); AscendC::LocalTensoruint8_t sharedTmpBuffer tmpBuf.Getuint8_t(); AscendC::AsinsrcType, false(dstLocal, srcLocal, sharedTmpBuffer, bufferSize); outQueue.EnQuesrcType(dstLocal); inQueue.FreeTensor(srcLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensorsrcType dstLocal outQueue.DeQuesrcType(); AscendC::DataCopy(dstGlobal, dstLocal, bufferSize); outQueue.FreeTensor(dstLocal); } private: AscendC::GlobalTensorsrcType srcGlobal; AscendC::GlobalTensorsrcType dstGlobal; AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueue; AscendC::TBufAscendC::TPosition::VECCALC tmpBuf; AscendC::TQueAscendC::TPosition::VECOUT, 1 outQueue; uint32_t bufferSize 0; }; extern C __global__ __aicore__ void asin_custom(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelAsinhalf op; op.Init(srcGm, dstGm, tilingData.srcSize, tilingData.tmpBufferSize); if (TILING_KEY_IS(1)) { op.Process(); } }样例二下面的样例展示了数学库Kernel侧API和PlatformAscendC::ReserveLocalMemory的配合使用方法流程如下Host侧调用ReserveLocalMemory接口预留Unified Buffer内存空间并通过GetCoreMemSize接口获取实际可用的Unified Buffer内存大小。基于实际可用的内存大小计算能够支持的最大Shape最大数据规模。这种方式可以避免多次调用GetXXXTmpMaxMinSize接口来获取合适的临时空间大小。Host侧Tiling API使用样例#include vector #include register/op_def_registry.h #include register/tilingdata_base.h #include tiling/tiling_api.h namespace optiling { BEGIN_TILING_DATA_DEF(MathCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, calSize); END_TILING_DATA_DEF; static ge::graphStatus TilingFunc(gert::TilingContext* context) { auto platformInfo context-GetPlatformInfo(); auto ascendcPlatform platform_ascendc::PlatformAscendC(platformInfo); uint64_t tailSize 0; // Unified Buffer剩余空间大小 ascendcPlatform.ReserveLocalMemory(platform_ascendc::ReservedSize::RESERVED_SIZE_8K); ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, tailSize); uint64_t calSize tailSize / (3 * 4); // src dst tmp float type MathCustomTilingData tiling; tiling.set_calSize(calSize); context-SetBlockDim(1); tiling.SaveToBuffer(context-GetRawTilingData()-GetData(), context-GetRawTilingData()-GetCapacity()); context-GetRawTilingData()-SetDataSize(tiling.GetDataSize()); context-SetTilingKey(1); return ge::GRAPH_SUCCESS; } } // namespace optilingKernel侧通过读取TilingData获取相应的计算规模参数。随后Kernel调用数学库提供的不带临时空间的API执行多种数学运算#include kernel_operator.h template typename srcType class KernelMath { public: __aicore__ inline KernelMath() {} __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize) { srcGlobal.SetGlobalBuffer(reinterpret_cast__gm__ srcType*(srcGm), srcSize); dstGlobal.SetGlobalBuffer(reinterpret_cast__gm__ srcType*(dstGm), srcSize); pipe.InitBuffer(inQueue, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(outQueue, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(tmpBuf, srcSize * sizeof(srcType)); bufferSize srcSize; } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensorsrcType srcLocal inQueue.AllocTensorsrcType(); AscendC::DataCopy(srcLocal, srcGlobal, bufferSize); inQueue.EnQue(srcLocal); } __aicore__ inline void Compute() { AscendC::LocalTensorsrcType dstLocal outQueue.AllocTensorsrcType(); AscendC::LocalTensorsrcType srcLocal inQueue.DeQuesrcType(); AscendC::LocalTensorsrcType tmp tmpBuf.GetsrcType(); AscendC::AsinsrcType, false(tmp, srcLocal, bufferSize); AscendC::PipeBarrierPIPE_V(); AscendC::SinsrcType, false(dstLocal, tmp, bufferSize); outQueue.EnQuesrcType(dstLocal); inQueue.FreeTensor(srcLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensorsrcType dstLocal outQueue.DeQuesrcType(); AscendC::DataCopy(dstGlobal, dstLocal, bufferSize); outQueue.FreeTensor(dstLocal); } private: AscendC::GlobalTensorsrcType srcGlobal; AscendC::GlobalTensorsrcType dstGlobal; AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueue; AscendC::TBufAscendC::TPosition::VECCALC tmpBuf; AscendC::TQueAscendC::TPosition::VECOUT, 1 outQueue; uint32_t bufferSize 0; }; extern C __global__ __aicore__ void math_custom(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelMathfloat op; op.Init(srcGm, dstGm, tilingData.calSize); if (TILING_KEY_IS(1)) { op.Process(); } }样例三下面的样例展示了数学库kernel侧API和Tiling API GetXxxTmpBufferFactorSize的配套使用方法具体流程如下Host侧调用Tiling接口获取maxLiveNodeCount和extraBuf并推算算子单次最大计算元素数量将其写入tiling data中kernel侧再读取tiling data获取该值基于该值分配临时空间。Host侧Tiling API使用样例:#include vector #include cassert #include register/op_def_registry.h #include register/tilingdata_base.h #include tiling/tiling_api.h namespace optiling { BEGIN_TILING_DATA_DEF(AsinCustomTilingData) TILING_DATA_FIELD_DEF(uint32_t, srcSize); TILING_DATA_FIELD_DEF(uint32_t, tmpBufferSize); END_TILING_DATA_DEF; static ge::graphStatus TilingFunc(gert::TilingContext* context) { // Input source shapes. auto shape_input context-GetInputTensor(0)-GetOriginShape(); std::vectorint64_t srcDims {shape_input.GetDim(0), shape_input.GetDim(1)}; uint32_t srcSize 1; uint32_t srcCurSize 1; for (auto dim : srcDims) { srcSize * dim; } uint32_t typeSize 2; auto platformInfo context-GetPlatformInfo(); auto ascendcPlatform platform_ascendc::PlatformAscendC(platformInfo); uint64_t tailSize 0; // ub剩余空间大小 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, tailSize); uint32_t asinMaxLiveNodeCount 0; uint32_t asinExtraBuf 0; uint32_t acosMaxLiveNodeCount 0; uint32_t acosExtraBuf 0; AscendC::GetAsinTmpBufferFactorSize(typeSize, asinMaxLiveNodeCount, asinExtraBuf); AscendC::GetAcosTmpBufferFactorSize(typeSize, acosMaxLiveNodeCount, acosExtraBuf); // tmp的大小需要减去UB上调用api接口输入和输出占用的大小 // 该示例中包括Asin接口的输入输出以及Acos的输入输出其中Asin接口的输出作为Acos的输入因此一共需要3份src的空间大小 auto tmpSize tailSize - srcSize * typeSize * 3; assert(tmpSize asinExtraBuf); assert(tmpSize acosExtraBuf); // 计算Asin算子单次最大计算元素数量 if (asinMaxLiveNodeCount ! 0) { srcAsinCurSize (tmpSize - asinExtraBuf) / asinMaxLiveNodeCount / typeSize; } else { srcAsinCurSize srcSize; } // 计算Acos算子单次最大计算元素数量 if (acosMaxLiveNodeCount ! 0) { srcAcosCurSize (tmpSize - acosExtraBuf) / acosMaxLiveNodeCount / typeSize; } else { srcAcosCurSize srcSize; } srcCurSize std::min(srcAsinCurSize, srcAcosCurSize); AsinCustomTilingData tiling; tiling.set_srcSize(srcSize); tiling.set_srcCurSize(srcCurSize); tiling.set_tmpBufferSize(tmpSize); context-SetBlockDim(1); tiling.SaveToBuffer(context-GetRawTilingData()-GetData(), context-GetRawTilingData()-GetCapacity()); context-GetRawTilingData()-SetDataSize(tiling.GetDataSize()); context-SetTilingKey(1); return ge::GRAPH_SUCCESS; } } // namespace optilingkernel侧样例:#include kernel_operator.h template typename srcType class KernelAsin { public: __aicore__ inline KernelAsin() {} __aicore__ inline void Init( GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSizeIn, uint32_t srcCurSizeIn, uint32_t tmpBufferSize) { srcSize srcSizeIn; srcCurSize srcCurSizeIn; srcGlobal.SetGlobalBuffer(reinterpret_cast__gm__ srcType*(srcGm), srcSize); dstGlobal.SetGlobalBuffer(reinterpret_cast__gm__ srcType*(dstGm), srcSize); pipe.InitBuffer(inQueue, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(outQueue, 1, srcSize * sizeof(srcType)); pipe.InitBuffer(tmpBuf1, 1, srcCurSize * sizeof(srcType)); pipe.InitBuffer(tmpBuf, 1, tmpBufferSize); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { AscendC::LocalTensorsrcType srcLocal inQueue.AllocTensorsrcType(); AscendC::DataCopy(srcLocal, srcGlobal, srcSize); inQueue.EnQue(srcLocal); } __aicore__ inline void Compute() { AscendC::LocalTensorsrcType dstLocal outQueue.AllocTensorsrcType(); AscendC::LocalTensorsrcType srcLocal inQueue.DeQuesrcType(); AscendC::LocalTensoruint8_t sharedTmpBuffer tmpBuf.Getuint8_t(); AscendC::LocalTensorsrcType tmpresBuffer tmpBuf1.GetsrcType(); for (int32_t offset 0; offset srcSize; offset srcCurSize) { AscendC::AsinsrcType, false(tmpresBuffer, srcLocal[offset], sharedTmpBuffer, srcCurSize); AscendC::PipeBarrierPIPE_V(); AscendC::AcossrcType, false(dstLocal[offset], tmpresBuffer, sharedTmpBuffer, srcCurSize); AscendC::PipeBarrierPIPE_V(); } outQueue.EnQuesrcType(dstLocal); inQueue.FreeTensor(srcLocal); } __aicore__ inline void CopyOut() { AscendC::LocalTensorsrcType dstLocal outQueue.DeQuesrcType(); AscendC::DataCopy(dstGlobal, dstLocal, srcSize); outQueue.FreeTensor(dstLocal); } private: AscendC::GlobalTensorsrcType srcGlobal; AscendC::GlobalTensorsrcType dstGlobal; AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueue; AscendC::TBufAscendC::TPosition::VECCALC, 1 tmpBuf; AscendC::TBufAscendC::TPosition::VECCALC, 1 tmpBuf1; AscendC::TQueAscendC::TPosition::VECOUT, 1 outQueue; uint32_t srcSize, srcCurSize; }; extern C __global__ __aicore__ void kernel_asin_operator( GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR workspace, GM_ADDR tiling) { GET_TILING_DATA(tilingData, tiling); KernelAsinhalf op; op.Init(srcGm, dstGm, tilingData.srcSize, tilingData.srcCurSize, tilingData.tmpBufferSize); if (TILING_KEY_IS(1)) { op.Process(); } }样例四下面以Exp接口的关键调用代码为例辅以调用前后数据打印结果展示模板参数isReuseSource的使用及其相关影响。模板参数isReuseSource为bool类型若isReuseSource为false则接口内部计算时不复用源操作数的内存空间若isReuseSource为true接口内部计算时会复用源操作数的内存空间存放一些中间结果节省内存空间开发者需要注意接口执行完成后源操作数的内存空间不再是原始值。isReuseSource false// 调用Exp前dstLocal、srcLocal数值 // dstLocal: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0] // srcLocal: [-7.5, -6.5, -5.5, -4.5, -3.5, -2.5, -1.5, -0.5, 0.5, 1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5] AscendC::Expfloat, 15, false(dstLocal, srcLocal, 16); // 调用Exp后dstLocal、srcLocal数值 // dstLocal: [0.000553084, 0.00150344, 0.00408677, 0.011109, 0.0301974, 0.082085, 0.22313, // 0.606531, 1.64872, 4.48169, 12.1825, 33.1155, 90.0171, 244.692, 665.142, 1808.04] srcLocal: [-7.5, -6.5, -5.5, -4.5, // -3.5, -2.5, -1.5, -0.5, 0.5, 1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5]isReuseSource true// 调用Exp前dstLocal、srcLocal数值 // dstLocal: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0] // srcLocal: [-7.5, -6.5, -5.5, -4.5, -3.5, -2.5, -1.5, -0.5, 0.5, 1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5] AscendC::Expfloat, 15, true(dstLocal, srcLocal, 16); // 调用Exp后dstLocal、srcLocal数值 // dstLocal: [0.000553084, 0.00150344, 0.00408677, 0.011109, 0.0301974, 0.082085, 0.22313, // 0.606531, 1.64872, 4.48169, 12.1825, 33.1155, 90.0171, 244.692, 665.142, 1808.04] srcLocal: [0.5, 0.5, 0.5, 0.5, 0.5, // 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5]【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

昇腾C数学计算API使用样例

相关推荐

昇腾C数据过滤临时空间API

5分钟部署Infisical：开源密钥管理平台实战指南

SQL Server错误代码解析与故障排除指南

更多资讯

大夹子电工作业用途考前押题怎么准备才能稳过实操考试

电工证求职网站有哪些？工地太忙根本没时间复习考试，内部消息告诉你怎么应对

考电工证包就业怎么网上查询？学历不高也能报！

鸡西架子工特种作业证真实经历：怕被中介坑？老手教你避坑

架子工特种作业证书多久拿证？怕考不过白交培训费的看过来

低压电工证审查收费吗？学历不高怕报不上名？上岗必备全攻略来了

今日更新

高压电工证哪里可以挂多久拿证全攻略：实操考试心里没底怕挂科别慌

电工证的理论题目在哪考，外省证怎么转过来

工地太忙根本没时间复习考试，职业学校的电工证值不值得考？

本周精选

安监局电工证分数查询怎么查才能知道含金量高不高

新疆电工作业干活累吗良心建议助你避开考不过的坑

中牟县电工证怎么报名外省证能转吗

本月热门

面向复杂光照与角度的宠物面部识别鲁棒性增强策略

基于DRV8818与STM32的双极步进电机驱动方案实战解析

DRV8818PWPR+PIC18F45K80双极步进电机控制方案实战

读完文章，想聊聊您的网站？