入门功能落地【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit本节使用基础API进行矩阵乘法的编程指导。算子中数据搬运和计算流水线通过事件同步进行协调降低编程复杂度。算子概述Matmul计算公式 $$ C_{M, N} A_{M, K} * B_{K, N} $$A、B 为源操作数A 为左矩阵形状为[M, K]B 为右矩阵形状为[K, N]C 为目的操作数存放矩阵乘结果的矩阵形状为[M, N]算子实现开发流程基于Ascend C方式实现矩阵算子的流程如下所示算子分析分析算子的输入、输出以及计算逻辑的实现明确需要调用的Ascend C接口。核函数定义定义Ascend C算子入口函数。算子实现完成核函数的内部实现调用私有成员函数CopyInA、CopyInB、DataLoadA、DataLoadB、Compute、CopyOut完成矩阵算子的四级流水操作。算子分析在开发算子代码之前需要分析算子的输入、输出以及计算逻辑的实现明确需要调用的Ascend C接口。表 1Ascend C Matmul算子设计规格算子类型OpTypeMatmul算子输入nameshapedata typeformata(m, k) (512, 512)halfNDb(k, n) (512, 1024)halfND算子输出c(m, n) (512, 1024)halfND核函数名称mmad_custom使用的主要接口DataCopyGM到L1 Buffer数据搬运 ND转NZ格式转换接口LoadDataL1到L0数据搬运 NZ转ZZ/ZN格式转换接口Mmad矩阵乘计算接口FixpipeL0C到GM数据搬出 NZ转ND格式转换 精度转换接口算子实现文件名称matmul.asc核函数定义本样例中函数名为mmad_custom核函数名称可自定义有3个参数abc其中ab都为输入内存c为输出内存。使用__global__函数类型限定符来标识它是一个核函数可以被调用使用__cube__函数类型限定符来标识该核函数在设备端aicore上的Cube核执行。算子类的Init函数完成内存初始化相关工作Process函数完成算子实现的核心逻辑。核函数在开始时调用AscendC::InitSocState()初始化硬件状态结束时调用AscendC::PipeBarrierPIPE_ALL()等待所有流水线完成。template uint32_t M, uint32_t K, uint32_t N, uint32_t singleCoreM, uint32_t singleCoreK, uint32_t singleCoreN, uint32_t baseM, uint32_t baseK, uint32_t baseN __global__ __cube__ void mmad_custom(GM_ADDR a, GM_ADDR b, GM_ADDR c) { AscendC::InitSocState(); KernelMatmulM, K, N, singleCoreM, singleCoreK, singleCoreN, baseM, baseK, baseN op; op.Init(a, b, c); op.Process(); AscendC::PipeBarrierPIPE_ALL(); }[!NOTE]说明 核函数使用模板参数传入矩阵的分块信息这样可以在编译期确定循环次数有利于编译器优化。其中M/K/N为总矩阵维度singleCoreM/singleCoreK/singleCoreN为单核处理的分块大小baseM/baseK/baseN为Cube单元单次计算的子块大小。算子类实现本节介绍一个Matmul矩阵算子的初步实现主要包括负责初始化的Init函数和负责核心数据处理的Process函数。Matmul矩阵算子的整体实现逻辑如下Stage1CopyIn任务GM -- L1使用DataCopy接口将Global Memory 中的数据搬运到L1 BufferL1A/L1B同时利用随路格式转换功能将ND格式转换为NZ格式。Stage2Split任务L1 -- L0使用LoadData接口将数据从L1A/L1B中搬运到L0A/L0B同时完成NZ到ZZ/ZN的格式转换。Stage3Compute任务L0上计算使用Mmad接口在L0上完成矩阵乘计算结果存放在CO1L0C Buffer。Stage4CopyOut任务L0C -- GM使用Fixpipe接口将结果从CO1直接搬运到Global Memory同时利用Fixpipe的随路格式转换功能将NZ格式转回ND格式。下文将以Matmul算子为例对上述步骤进行详细介绍。完整样例代码请参见matmul样例。实现代码说明KernelMatmul算子类的定义如下class KernelMatmul { public: __aicore__ inline KernelMatmul() {} __aicore__ inline void Init(__gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* c) { aGM.SetGlobalBuffer((__gm__ half*)a); bGM.SetGlobalBuffer((__gm__ half*)b); cGM.SetGlobalBuffer((__gm__ half*)c); InitGMOffsets(); } __aicore__ inline void Process() { AscendC::LocalTensorhalf a1Local(AscendC::TPosition::A1, a1Addr, baseM * baseK); AscendC::LocalTensorhalf a2Local(AscendC::TPosition::A2, a2Addr, baseM * baseK); AscendC::LocalTensorhalf b1Local(AscendC::TPosition::B1, b1Addr, baseK * baseN); AscendC::LocalTensorhalf b2Local(AscendC::TPosition::B2, b2Addr, baseK * baseN); AscendC::LocalTensorfloat cLocal(AscendC::TPosition::CO1, cAddr, baseM * baseN); constexpr uint32_t kLoopCount singleCoreK / baseK; constexpr uint32_t mLoopCount singleCoreM / baseM; constexpr uint32_t nLoopCount singleCoreN / baseN; for (uint32_t mIndex 0; mIndex mLoopCount; mIndex) { for (uint32_t nIndex 0; nIndex nLoopCount; nIndex) { for (uint32_t kIndex 0; kIndex kLoopCount; kIndex) { CopyInA(a1Local, kIndex, mIndex); CopyInB(b1Local, kIndex, nIndex); AscendC::PipeBarrierPIPE_ALL(); DataLoadA(a1Local, a2Local, mIndex); DataLoadB(b1Local, b2Local, nIndex); AscendC::PipeBarrierPIPE_ALL(); Compute(cLocal, a2Local, b2Local, kIndex, mIndex, nIndex); AscendC::PipeBarrierPIPE_ALL(); } AscendC::PipeBarrierPIPE_ALL(); CopyOut(cLocal, mIndex, nIndex); } } } private: __aicore__ inline void CopyInA(...) __aicore__ inline void CopyInB(..) __aicore__ inline void DataLoadA(...) __aicore__ inline void DataLoadB(...) __aicore__ inline void Compute(...) __aicore__ inline void CopyOut(...) private: // ... };下面逐个函数说明。InitGMOffsets函数实现InitGMOffsets函数负责根据当前核的编号通过GetBlockIdx接口获取计算出该核在全局矩阵中负责的数据偏移量。CopyInA / CopyInB函数实现Stage1GM -- L11. CopyInA函数按当前K方向和M方向的循环索引计算矩阵A对应子块在Global Memory中的起始地址通过DataCopy接口将其搬运到L1 Buffer。后续计算依赖NZ分形格式搬运过程中同步完成ND→NZ的格式转换。2. CopyInB函数按当前K方向和N方向的循环索引计算矩阵B对应子块在Global Memory中的起始地址通过DataCopy接口将其搬运到L1 Buffer同步完成ND→NZ格式转换。DataLoadA / DataLoadB函数实现Stage2: L1 -- L01. DataLoadA 函数将矩阵A从L1 Buffer搬运到L0 BufferL0A同时通过LoadData接口完成NZ到ZZ格式的转换。Mmad指令要求左矩阵以ZZ格式行优先的Z分形布局输入DataLoadA通过这次搬运一步完成数据拆分 格式转换。2. DataLoadB 函数将矩阵B从L1 Buffer搬运到L0 BufferL0B同时通过LoadData接口完成NZ到ZN格式的转换。Mmad指令要求右矩阵以ZN格式列优先的N分形布局输入DataLoadB通过转置搬运满足这一要求。Compute函数实现Stage3矩阵乘计算Compute函数调用Mmad完成单个子块baseM × baseN × baseK的矩阵乘法。当K方向为第一次计算时需初始化结果矩阵后续K分片计算结果直接累加最终将完整的子块累加结果存放在L0C Buffer中。CopyOut函数实现Stage4L0C -- GMCopyOut函数通过Fixpipe接口将计算结果从CO1L0C Buffer直接搬运到Global Memory同时完成NZ到ND的格式转换和精度转换float32 - float16。流水同步机制说明CopyIn、DataLoad、Compute和CopyOut四个阶段的操作分别在不同的硬件流水线上执行。为保证数据依赖正确——即每个阶段读取的数据必须由前一阶段写入完毕需要在各阶段之间设置同步屏障。本样例使用全局流水同步PipeBarrier.md)PIPE_ALL统一协调所有流水线的完成状态确保各阶段严格串行执行。功能调试printf该接口提供CPU域或NPU域调试场景下的格式化输出功能。在算子kernel侧实现代码中需要输出日志信息的地方调用printf接口打印相关内容。示例如下AscendC::printf(matmul blockIdx%d\n, AscendC::GetBlockIdx());[!CAUTION]注意 printfPRINTF接口打印功能会对算子实际运行的性能带来一定影响通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP0的方式关闭打印功能。DumpTensor基于算子工程开发的算子可以使用该接口Dump指定Tensor的内容。同时支持打印自定义的附加信息仅支持uint32_t数据类型的信息比如打印当前行号等。在算子kernel侧实现代码中需要打印Tensor数据的地方调用DumpTensor接口打印相关内容。样例如下AscendC::DumpTensor(cLocal, baseM * baseN);[!CAUTION]注意 DumpTensor接口打印功能会对算子实际运行的性能带来一定影响通常在调测阶段使用。开发者可以按需通过设置ASCENDC_DUMP0来关闭打印功能。性能调试msProf工具介绍msProf工具是单算子性能分析工具。包含msprof op和msprof op simulator两种使用方式。该工具协助用户定位算子内存、算子代码以及算子指令的异常实现全方位的算子调优。当前支持基于不同运行模式上板或仿真和不同文件形式可执行文件或算子二进制.o文件进行性能数据的采集和自动解析。上板性能采集通过上板性能采集可以直接测定算子在昇腾AI处理器上的运行时间。该方式适合在板环境中快速定位算子性能问题。基于可执行文件demo通过msprof op执行算子调优msprof op ./demo性能数据说明命令完成后会在默认目录下生成以“OPPROF_{timestamp}_XXX”命名的文件夹,性能数据文件夹结构示例如下├──dump # 原始的性能数据用户无需关注 ├──ArithmeticUtilization.csv # cube/vector指令cycle占比 ├──L2Cache.csv # L2 Cache命中率影响MTE2建议合理规划数据搬运逻辑增加命中率 ├──Memory.csv # UBL1和主存储器读写带宽速率 ├──MemoryL0.csv # L0AL0B和L0C读写带宽速率 ├──MemoryUB.csv # Vector和Scalar到UB的读写带宽速率 ├──OpBasicInfo.csv # 算子基础信息 ├──PipeUtilization.csv # 采集计算单元和搬运单元耗时和占比 ├──ResourceConflictRatio.csv # UB上的 bank group、bank conflict和资源冲突率在所有指令中的占比 └──visualize_data.bin # MindStudio Insight呈现文件更多msProf工具使用方式请参考MindStudio工具算子调优msProf中的内容。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考