CANN PTO自动模式

PTO AUTO模式【免费下载链接】pto-isaParallel Tile Operation (PTO) is a virtual instruction set architecture designed by Ascend CANN, focusing on tile-level operations. This repository offers high-performance, cross-platform tile operations across Ascend platforms.项目地址: https://gitcode.com/cann/pto-isaauto模式是什么PTO AUTO是一种新的编程模式主要提供以下两点优势降低kernel的开发难度的同时能使开发者实现必要的优化确保跨代兼容不同的昇腾硬件架构更具体来说在AUTO模式下kernel开发者不用手动为tile分配内存也不用亲自手写不同pipe间的同步。作为替代PTO AUTO编译器会帮助kernel开发者在不同buffer上分配内存。而且编译器也会在PTO指令之间自动插入同步最大化pipe之间的流水线并行。最后kernel开发者也不用关心不同昇腾硬件架构之间的区别尤其是关于CUBE和VECTOR交流和同步的机制。注意auto模式目前仅支持编译器-O2选项。简单示例一个简单的示例逐元素的乘法。这展示了最关键的auto模式与manual模式的区别TMUL manual模式template typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_ __global__ AICORE void runTMul(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1) { using DynShapeDim5 Shape1, 1, 1, kGRows_, kGCols_; using DynStridDim5 Stride1, 1, 1, kGCols_, 1; using GlobalData GlobalTensorT, DynShapeDim5, DynStridDim5; using TileData TileTileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1; TileData src0Tile(kGRows_, kGCols_); TileData src1Tile(kGRows_, kGCols_); TileData dstTile(kGRows_, kGCols_); TASSIGN(src0Tile, 0x0 0x400 * block_idx); TASSIGN(src1Tile, 0x4000 0x400 * block_idx); TASSIGN(dstTile, 0x8000 0x400 * block_idx); int offset (block_idx / 4) * (64 * 16) (block_idx % 4) * 16; GlobalData src0Global(src0 offset); GlobalData src1Global(src1 offset); GlobalData dstGlobal(out offset); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TMUL(dstTile, src0Tile, src1Tile); set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); out dstGlobal.data(); }TMUL AUTO模式template typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_ __global__ AICORE void runTMul(__gm__ T __out__ *out, __gm__ T __in__ *src0, __gm__ T __in__ *src1) { using DynShapeDim5 Shape1, 1, 1, kGRows_, kGCols_; using DynStridDim5 Stride1, 1, 1, kGCols_, 1; using GlobalData GlobalTensorT, DynShapeDim5, DynStridDim5; using TileData TileTileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1; TileData src0Tile(kGRows_, kGCols_); TileData src1Tile(kGRows_, kGCols_); TileData dstTile(kGRows_, kGCols_); int offset (block_idx / 4) * (64 * 16) (block_idx % 4) * 16; GlobalData src0Global(src0 offset); GlobalData src1Global(src1 offset); GlobalData dstGlobal(out offset); TLOAD(src0Tile, src0Global); TLOAD(src1Tile, src1Global); TMUL(dstTile, src0Tile, src1Tile); TSTORE(dstGlobal, dstTile); out dstGlobal.data(); }PTO AUTO编译器特性不同架构之间的兼容性PTO AUTO编译器确保同一个PTO源码程序能在不同昇腾架构上编译和运行同时确保性能。自动同步在manual模式下程序员需要使用PTO的Event机制精准地在必要的源码位置上手动调用同步指令来确保正确的结果和性能。这非常繁琐且容易出错。Auto模式编译器让程序员避免了这个麻烦。编译器会自动在需要的位置插入同步确保正确的结果和较好的性能。Tile内存分配在manual模式下当用户声定义一个Tile变量后需要显式调用TASSIGN来为这个Tile分配在指定内存空间里的内存地址。 在auto模式下用户不需要手动分配内存只需要定义Tile变量即可编译器会自动为所有Tile在正确的buffer上分配内存地址。PTO AUTO文档更多PTO AUTO的详细文档如下PTO_AUTO_kernel_developer_rules_and_limitationsPTO_AUTO_library_developer_rules_and_limitationsPTO AUTO Code Examples【免费下载链接】pto-isaParallel Tile Operation (PTO) is a virtual instruction set architecture designed by Ascend CANN, focusing on tile-level operations. This repository offers high-performance, cross-platform tile operations across Ascend platforms.项目地址: https://gitcode.com/cann/pto-isa创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考