CANN / pto-isa PTO Tile 内部函数编程模型

PTO Tile Intrinsics Programming Model【免费下载链接】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-isaPTO Tile Lib providestile-granularityintrinsics that map to the PTO ISA. The model is designed for:Portability across device generations: hardware may change (instruction details, storage layout, scheduling), but the programming model remains stable.Near-hardware performance: the Tile and GlobalTensor abstractions are low-level enough to express efficient data movement and compute.Two user profiles: a productive “compiler does the hard work” style, and an expert “I control placement and sync” style.For the abstract execution model (core/device/host), seedocs/machine/abstract-machine.md.Core conceptsTile: a fixed-capacity 2-D on-chip buffer (conceptually a tile register / SRAM block) and the unit of computation for most PTO instructions. Seedocs/coding/Tile.md.GlobalTensor: a lightweight view of global memory (GM) as a 5-D tensor with shape/stride/layout metadata, used by memory instructions such asTLOADandTSTORE. Seedocs/coding/GlobalTensor.md.Scalar: immediate values and enumerations that parameterize instructions (rounding modes, comparison modes, atomic modes, etc.). Seedocs/coding/Scalar.md.Event: explicit dependency tokens between pipeline classes, used to order operations without introducing a full barrier everywhere. Seedocs/coding/Event.md.Two development stylesPTO-AutoPTO-Auto targets developers who prefer a simple, portable programming experience:The compiler/runtime chooses memory placement and address binding.The compiler inserts required synchronization.The compiler schedules operations and applies fusions when possible.This mode is a good starting point for correctness and portability.PTO-ManualPTO-Manual targets developers who need full control for performance tuning:The developer controls memory placement and binding (for example viaTASSIGN).The developer explicitly expresses ordering (events and/orTSYNC).The developer controls the operation schedule.This mode enables expert tuning on critical kernels while still using the shared Tile/GlobalTensor abstractions.Execution models: SPMD and MPMDPTO supports bothSPMDandMPMDexecution models.These execution models describehow work is mapped onto cores. They are orthogonal to theAuto vs Manualdevelopment styles (you can write SPMD-Auto, SPMD-Manual, MPMD-Auto, or MPMD-Manual code).SPMD (Single Program, Multiple Data)In SPMD, all participating cores run the same entry function, and each core selects its own data region using its runtime identity (for exampleblock_idx).When sub-block decomposition exists, a stable “virtual id” can be constructed:auto cid get_block_idx(); auto vid get_block_idx() * get_subblockdim() get_subblockid();SPMD is a good fit for regular tensor tiling (GEMM, softmax-by-rows, elementwise ops).MPMD (Multiple Program, Multiple Data)In MPMD, different cores (or groups of cores) may executedifferent tile programsas part of the same overall tile graph. Conceptually, theDevice Machine schedulerchooses which “program” a core runs.One portable way to express this is to pass a scheduler-providedtask idinto the kernel entry function and dispatch based on it:__global__ __aicore__ void KernelMPMD(__gm__ float* out, __gm__ const float* in, uint32_t task_id) { switch (task_id) { case 0: return ProducerStage(out, in); case 1: return ConsumerStage(out, in); default: return; } }Notes:The exact mechanism that deliverstask_idis platform/runtime dependent; the abstract model only requires that the Device Machine can schedule different tile blocks onto available cores.If you prefer, MPMD can also be expressed asmultiple entry points(multiple kernels) rather than a single kernel with aswitch.【免费下载链接】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），仅供参考