量逐元素最小值運(yùn)算的數(shù)學(xué)語(yǔ)義、約束與實(shí)戰(zhàn)用法)
PTO TMINS 指令詳解Ascend Tile 與標(biāo)量逐元素最小值運(yùn)算的數(shù)學(xué)語(yǔ)義、約束與實(shí)戰(zhàn)用法【免費(fèi)下載鏈接】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.項(xiàng)目地址: https://gitcode.com/cann/pto-isa導(dǎo)讀TMINS 是 CANN pto-isaParallel Tile Operation面向昇騰平臺(tái)的 Tile 級(jí)虛擬指令集中用于執(zhí)行Tile 與標(biāo)量逐元素最小值運(yùn)算的核心向量指令它以整個(gè) Tile 的每個(gè)元素為操作對(duì)象與一個(gè)標(biāo)量值逐一比較并取較小者寫入目標(biāo) Tile。本文以 TMINS_zh.md 為骨架結(jié)合倉(cāng)庫(kù)內(nèi) NPUA2/A3、Ascend 950PR/950DT與 CPU 模擬器上的源碼實(shí)現(xiàn)及完整測(cè)試用例系統(tǒng)講解 TMINS 的數(shù)學(xué)語(yǔ)義、三級(jí)匯編形式、C 內(nèi)建接口、平臺(tái)約束、自動(dòng)/手動(dòng)兩種編程模式并給出可復(fù)制運(yùn)行的實(shí)戰(zhàn)示例。讀完本文你將能夠在 PTO 編程框架中正確聲明與調(diào)用 TMINS理解其有效區(qū)域valid region迭代規(guī)則與底層vmins指令的映射關(guān)系并能參照測(cè)試用例編寫自己的標(biāo)量歸約類算子。指令一覽TMINS 屬于 PTO 指令集中 Tile 與標(biāo)量的二元運(yùn)算族S 后綴表示 scalar即標(biāo)量操作數(shù)與 TMINTile 與 Tile 逐元素最小值互為補(bǔ)充也與 TADDS、TSUBS、TMULS、TDIVS、TMAXS 等標(biāo)量廣播類指令共享同一套編程模型與硬件流水線。數(shù)學(xué)語(yǔ)義TMINS 對(duì) Tile 的**有效區(qū)域valid region**內(nèi)每個(gè)元素執(zhí)行一次min運(yùn)算對(duì)有效區(qū)域中的每個(gè)元素(i, j)其目標(biāo)值為源 Tile 對(duì)應(yīng)元素與標(biāo)量值兩者中的較小者$$ \mathrm{dst}{i,j} \min(\mathrm{src}{i,j}, \mathrm{scalar}) $$這里scalar是標(biāo)量操作數(shù)會(huì)被隱式廣播到每個(gè)元素位置參與比較。需要注意的是TMINS 與 TMIN 的區(qū)別在于第二操作數(shù)TMIN 比較的是兩個(gè) Tile 的逐元素值dst[i][j] min(src0[i][j], src1[i][j])見(jiàn) TMIN_zh.md而 TMINS 中所有元素共享同一個(gè)標(biāo)量閾值。典型應(yīng)用場(chǎng)景包括數(shù)據(jù)裁剪clamp 上界、激活函數(shù)實(shí)現(xiàn)中的上界約束如min(x, alpha)、量化/歸一化中的閾值截?cái)嗟刃枰獙⒄麄€(gè) Tile 與一個(gè)固定常數(shù)比較的運(yùn)算。匯編語(yǔ)法三級(jí)抽象層次TMINS 在 PTO 指令集中同時(shí)存在三種匯編形式分別對(duì)應(yīng)不同的編譯/調(diào)度階段。同步形式PTO 匯編最直接的指令形式%scalar直接給出標(biāo)量值類型以f32等具體 dtype 標(biāo)注%dst tmins %src, %scalar : !pto.tile..., f32AS Level 1SSA 形式SSA靜態(tài)單賦值中間表示操作數(shù)與結(jié)果均為!pto.tile...類型的值value由編譯器后續(xù)進(jìn)行資源分配%dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...AS Level 2DPS 形式DPSDestructive/顯式 place-and-schedule形式此時(shí) Tile 已被綁定到具體緩沖區(qū)!pto.tile_buf...輸入與輸出分離聲明pto.tmins ins(%src, %scalar : !pto.tile_buf..., dtype) outs(%dst : !pto.tile_buf...)從源碼結(jié)構(gòu)看pto.tmins指令名稱中的 s 后綴貫穿三個(gè)抽象層次保持一致便于后端在編譯流水線中逐級(jí)降級(jí)映射。C 內(nèi)建接口TMINS 的 C 內(nèi)建接口模板聲明于 include/pto/common/pto_instr.hpp對(duì)開發(fā)者開放的公共包含頭為pto/pto-inst.hpptemplate typename TileDataDst, typename TileDataSrc, typename... WaitEvents PTO_INST RecordEvent TMINS(TileDataDst dst, TileDataSrc src, typename TileDataSrc::DType scalar, WaitEvents ... events);接口要點(diǎn)參數(shù)順序目標(biāo) Tiledst、源 Tilesrc、標(biāo)量scalar其類型必須等于TileDataSrc::DType、可選的事件列表events用于跨流水線同步等待返回類型PTO_INST RecordEvent即返回一個(gè)記錄事件可與后續(xù)指令構(gòu)成事件依賴鏈宏分發(fā)TMINS通過(guò)MAP_INSTR_IMPL(TMINS, dst, src, scalar)宏映射到平臺(tái)相關(guān)的TMINS_IMPL實(shí)現(xiàn)見(jiàn) include/pto/common/pto_instr.hpp同一份用戶代碼可無(wú)縫編譯到不同昇騰平臺(tái)或 CPU 模擬器。底層實(shí)現(xiàn)調(diào)用鏈在 NPU 側(cè)TMINS_IMPL最終落到硬件向量指令vminsAtlas A2/A3 平臺(tái)實(shí)現(xiàn)在 include/pto/npu/a2a3/TMins.hppMinSOp::BinSInstr直接調(diào)用vmins(dst, src0, src1, repeats, 1, 1, 8, 8)并支持通過(guò)dstRepeatStride/srcRepeatStride自定義 repeat 間步長(zhǎng)Ascend 950 平臺(tái)實(shí)現(xiàn)在 include/pto/npu/a5/TMins.hppMinSOp::BinSInstr調(diào)用vmins(reg_dst, reg_src0, src1, preg, MODE_ZEROING)即帶掩碼寄存器與清零模式zeroing mode的向量比較形式對(duì)于int64_t/uint64_t這類寬類型A5 實(shí)現(xiàn)走Int64ScalarInt64Op::Min, ...專用路徑將 64 位標(biāo)量比較拆分為內(nèi)部多次 32 位運(yùn)算組合完成見(jiàn) include/pto/npu/a5/TMins.hpp。在 CPU 模擬器側(cè)TMINS 與 TADDS/TSUBS/TMULS/TMAXS 等標(biāo)量指令共用同一套UnaryTileScalarOpImpl逐元素循環(huán)框架僅以ElementOp::OP_MINS區(qū)分運(yùn)算類型實(shí)現(xiàn)在 include/pto/cpu/TBinSOps.hpp。這意味著在無(wú)昇騰硬件的開發(fā)環(huán)境中CPU 模擬可以給出與硬件一致的數(shù)值行為便于算子邏輯先行驗(yàn)證。約束與有效區(qū)域使用 TMINS 時(shí)必須滿足以下平臺(tái)與通用約束違反約束會(huì)在編譯期static_assert或運(yùn)行期PTO_ASSERT報(bào)錯(cuò)。Atlas A2/A3 訓(xùn)練/推理系列產(chǎn)品實(shí)現(xiàn)檢查TileData::DType必須是以下類型之一int32_t、int、int16_t、half、float16_t、float、float32_t運(yùn)行時(shí)src.GetValidRow() dst.GetValidRow()且src.GetValidCol() dst.GetValidCol()。對(duì)應(yīng)實(shí)現(xiàn)中A2/A3 的TMINS_IMPL通過(guò)static_assert校驗(yàn)數(shù)據(jù)類型集合并通過(guò)PTO_ASSERT(src.GetValidCol() dst.GetValidCol(), ...)與PTO_ASSERT(src.GetValidRow() dst.GetValidRow(), ...)在運(yùn)行期強(qiáng)制行列有效邊界一致見(jiàn) include/pto/npu/a2a3/TMins.hpp。Ascend 950PR / Ascend 950DT實(shí)現(xiàn)檢查TileData::DType必須是以下類型之一uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t、int64_t、uint64_t、half、float、bfloat16_t——相比 A2/A3 大幅擴(kuò)展了整數(shù)與低精度類型覆蓋并額外支持bfloat16_t運(yùn)行時(shí)src.GetValidCol() dst.GetValidCol()。對(duì)應(yīng)實(shí)現(xiàn)中A5 的TMINS_IMPL同樣以static_assert校驗(yàn)類型集合與TileType::Vec位置以PTO_ASSERT(src0.GetValidCol() dst.GetValidCol(), ...)校驗(yàn)列數(shù)一致見(jiàn) include/pto/npu/a5/TMins.hpp。通用約束所有平臺(tái)dst與src必須使用相同的元素類型A2/A3 實(shí)現(xiàn)在類型層通過(guò)static_assert(std::is_same_vT, typename TileDataDst::DType, ...)直接強(qiáng)制標(biāo)量類型必須與 Tile 數(shù)據(jù)類型一致接口簽名中scalar類型即TileDataSrc::DTypeTile 位置必須是向量TileData::Loc TileType::Vec。有效區(qū)域TMINS 以dst.GetValidRow()/dst.GetValidCol()作為迭代域運(yùn)算只覆蓋目標(biāo) Tile 聲明的有效行列范圍內(nèi)超出部分不參與計(jì)算。因此調(diào)用方應(yīng)保證src的有效區(qū)域至少覆蓋dst的有效區(qū)域避免讀取越界數(shù)據(jù)。實(shí)戰(zhàn)示例自動(dòng)模式與手動(dòng)模式PTO 提供兩種 Tile 資源管理方式。自動(dòng)模式由編譯器/運(yùn)行時(shí)負(fù)責(zé) Tile 的內(nèi)存放置與調(diào)度手動(dòng)模式通過(guò)TASSIGN顯式將 Tile 綁定到指定地址再由指令發(fā)射。自動(dòng)Auto模式#include pto/pto-inst.hpp using namespace pto; void example_auto() { using TileT TileTileType::Vec, float, 16, 16; TileT src, dst; TMINS(dst, src, 0.0f); }聲明一個(gè) 16×16 的float向量 Tile調(diào)用TMINS將src中每個(gè)元素與0.0f比較取小寫入dst。0.0f的類型與TileT::DTypefloat一致滿足通用約束。手動(dòng)Manual模式#include pto/pto-inst.hpp using namespace pto; void example_manual() { using TileT TileTileType::Vec, float, 16, 16; TileT src, dst; TASSIGN(src, 0x1000); TASSIGN(dst, 0x2000); TMINS(dst, src, 0.0f); }先通過(guò)TASSIGN將src綁定到地址0x1000、dst綁定到0x2000再發(fā)射TMINS。這種模式適合開發(fā)者需要精確控制緩沖區(qū)布局如復(fù)用固定的片上 buffer 池的場(chǎng)景。匯編示例ASM自動(dòng)模式自動(dòng)模式下資源放置與調(diào)度由編譯器/運(yùn)行時(shí)完成指令以 SSA 值形式出現(xiàn)# 自動(dòng)模式由編譯器/運(yùn)行時(shí)負(fù)責(zé)資源放置與調(diào)度。 %dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...手動(dòng)模式手動(dòng)模式先顯式綁定資源pto.tassign將 tile 操作數(shù)綁定到物理地址再發(fā)射指令# 手動(dòng)模式先顯式綁定資源再發(fā)射指令。 # 可選當(dāng)該指令包含 tile 操作數(shù)時(shí) # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...PTO 匯編形式匯總%dst tmins %src, %scalar : !pto.tile..., f32 # AS Level 2 (DPS) pto.tmins ins(%src, %scalar : !pto.tile_buf..., dtype) outs(%dst : !pto.tile_buf...)端到端算子示例TLOAD → TMINS → TSTORE倉(cāng)庫(kù)中 tests/cpu/st/testcase/tmins/tmins_kernel.cpp 給出了一個(gè)完整的端到端 kernel將全局內(nèi)存數(shù)據(jù)加載到 Tile執(zhí)行 TMINS 標(biāo)量裁剪再存回全局內(nèi)存并用set_flag/wait_flag在 MTE2加載、V向量運(yùn)算、MTE3存儲(chǔ)三條流水線之間建立事件同步template typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_ __global__ AICORE void runTMins(__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(kTRows_, kTCols_); TileData dstTile(kTRows_, kTCols_); TASSIGN(src0Tile, 0x0 0x400); TASSIGN(dstTile, 0x8000 0x400); GlobalData src0Global(src0); GlobalData dstGlobal(out); TLOAD(src0Tile, src0Global); set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TMINS(dstTile, src0Tile, src1[0]); // 標(biāo)量來(lái)自全局內(nèi)存首元素 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); out dstGlobal.data(); }該示例同時(shí)展示了一個(gè)實(shí)用技巧標(biāo)量操作數(shù)不一定是編譯期常量也可以是src1[0]這類從全局內(nèi)存讀取的運(yùn)行時(shí)值使 TMINS 可以用于動(dòng)態(tài)閾值的裁剪邏輯。測(cè)試驗(yàn)證與 golden 機(jī)制TMINS 在倉(cāng)庫(kù)中擁有完整的多平臺(tái)測(cè)試覆蓋包括 CPU 模擬tests/cpu/st/testcase/tmins/、A2/A3tests/npu/a2a3/src/st/testcase/tmins/、A5tests/npu/a5/src/st/testcase/tmins/以及 costmodel 相關(guān)測(cè)試tests/costmodel/st/testcase/tmins/main.cpp。以 CPU 側(cè)測(cè)試 tests/cpu/st/testcase/tmins/main.cpp 為例其驗(yàn)證流程為通過(guò)gen_data.py生成隨機(jī)輸入input1.bin、標(biāo)量輸入input_scalar.bin與參考 golden 文件golden.bin使用 ACL 運(yùn)行時(shí) APIaclrtMalloc/aclrtMemcpy分配主機(jī)與設(shè)備內(nèi)存將輸入拷貝到設(shè)備啟動(dòng)LaunchTMinskernel同步流后將結(jié)果拷回主機(jī)寫為output.bin將output.bin與golden.bin逐元素比較容差 0.0001f通過(guò)EXPECT_TRUE(ret)判定測(cè)試通過(guò)。覆蓋的數(shù)據(jù)類型與形狀包括float/int32_t/int64_t/uint64_t/int16_t/half以及開啟CPU_SIM_BFLOAT_ENABLED時(shí)的bfloat16_t形狀覆蓋 64×64 與 16×256 兩種 Tile 配置見(jiàn) tests/cpu/st/testcase/tmins/main.cpp可直接作為自行擴(kuò)展 TMINS 用例的模板。TMINS 與 TMIN 的對(duì)比維度TMINSTile × 標(biāo)量TMINTile × Tile數(shù)學(xué)語(yǔ)義dst[i][j] min(src[i][j], scalar)dst[i][j] min(src0[i][j], src1[i][j])第二操作數(shù)標(biāo)量值TileDataSrc::DType另一個(gè) TileA2/A3 數(shù)據(jù)類型int32_t/int/int16_t/half/float16_t/float/float32_tint32_t/int16_t/half/float并額外要求行主序布局與靜態(tài)有效邊界檢查運(yùn)行時(shí)校驗(yàn)行列有效邊界一致src0/src1/dst的validRow/validCol相同典型用途閾值截?cái)?、上界裁剪逐元素融合比較如 ReLU6、逐點(diǎn) clamp 下限兩者都要求TileType::Vec向量位置、以dst有效區(qū)域?yàn)榈蚯揖暶饔?include/pto/common/pto_instr.hpp接口風(fēng)格完全一致可在同一 kernel 中混合使用?？偨Y(jié)TMINS 是 PTO 指令集中一個(gè)簡(jiǎn)潔但高頻使用的 Tile-標(biāo)量二元指令數(shù)學(xué)上等價(jià)于對(duì)有效區(qū)域內(nèi)每個(gè)元素執(zhí)行min(x, scalar)編程上通過(guò)TMINS(dst, src, scalar)一行即可完成整個(gè) Tile 的標(biāo)量裁剪。使用時(shí)需要特別關(guān)注兩點(diǎn)一是平臺(tái)相關(guān)的數(shù)據(jù)類型支持范圍A2/A3 與 950 系列差異明顯950 額外支持 8/16/64 位整數(shù)與bfloat16_t二是有效區(qū)域一致性約束src與dst的 valid 行列必須匹配。其底層實(shí)現(xiàn)無(wú)論是 A2/A3 的vminsrepeat 循環(huán)、A5 的帶掩碼vmins與 64 位專用路徑還是 CPU 模擬器的UnaryTileScalarOpImpl逐元素循環(huán)都保證了跨平臺(tái)一致的語(yǔ)義配合倉(cāng)庫(kù)內(nèi)覆蓋多平臺(tái)、多數(shù)據(jù)類型、多 Tile 形狀的完整測(cè)試用例開發(fā)者可以放心在昇騰 NPU 上以自動(dòng)或手動(dòng)模式使用 TMINS 實(shí)現(xiàn)高效的逐元素最小值運(yùn)算?！久赓M(fèi)下載鏈接】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.項(xiàng)目地址: https://gitcode.com/cann/pto-isa創(chuàng)作聲明：本文部分內(nèi)容由AI輔助生成（AIGC），僅供參考 PTO TMINS 指令詳解：Ascend Tile 與標(biāo)量逐元素最小值運(yùn)算的數(shù)學(xué)語(yǔ)義、約束與實(shí)戰(zhàn)用法
PTO TMINS 指令詳解Ascend Tile 與標(biāo)量逐元素最小值運(yùn)算的數(shù)學(xué)語(yǔ)義、約束與實(shí)戰(zhàn)用法【免費(fèi)下載鏈接】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.項(xiàng)目地址: https://gitcode.com/cann/pto-isa導(dǎo)讀TMINS 是 CANN pto-isaParallel Tile Operation面向昇騰平臺(tái)的 Tile 級(jí)虛擬指令集中用于執(zhí)行Tile 與標(biāo)量逐元素最小值運(yùn)算的核心向量指令它以整個(gè) Tile 的每個(gè)元素為操作對(duì)象與一個(gè)標(biāo)量值逐一比較并取較小者寫入目標(biāo) Tile。本文以 TMINS_zh.md 為骨架結(jié)合倉(cāng)庫(kù)內(nèi) NPUA2/A3、Ascend 950PR/950DT與 CPU 模擬器上的源碼實(shí)現(xiàn)及完整測(cè)試用例系統(tǒng)講解 TMINS 的數(shù)學(xué)語(yǔ)義、三級(jí)匯編形式、C 內(nèi)建接口、平臺(tái)約束、自動(dòng)/手動(dòng)兩種編程模式并給出可復(fù)制運(yùn)行的實(shí)戰(zhàn)示例。讀完本文你將能夠在 PTO 編程框架中正確聲明與調(diào)用 TMINS理解其有效區(qū)域valid region迭代規(guī)則與底層vmins指令的映射關(guān)系并能參照測(cè)試用例編寫自己的標(biāo)量歸約類算子。指令一覽TMINS 屬于 PTO 指令集中 Tile 與標(biāo)量的二元運(yùn)算族S 后綴表示 scalar即標(biāo)量操作數(shù)與 TMINTile 與 Tile 逐元素最小值互為補(bǔ)充也與 TADDS、TSUBS、TMULS、TDIVS、TMAXS 等標(biāo)量廣播類指令共享同一套編程模型與硬件流水線。數(shù)學(xué)語(yǔ)義TMINS 對(duì) Tile 的**有效區(qū)域valid region**內(nèi)每個(gè)元素執(zhí)行一次min運(yùn)算對(duì)有效區(qū)域中的每個(gè)元素(i, j)其目標(biāo)值為源 Tile 對(duì)應(yīng)元素與標(biāo)量值兩者中的較小者$$ \mathrm{dst}{i,j} \min(\mathrm{src}{i,j}, \mathrm{scalar}) $$這里scalar是標(biāo)量操作數(shù)會(huì)被隱式廣播到每個(gè)元素位置參與比較。需要注意的是TMINS 與 TMIN 的區(qū)別在于第二操作數(shù)TMIN 比較的是兩個(gè) Tile 的逐元素值dst[i][j] min(src0[i][j], src1[i][j])見(jiàn) TMIN_zh.md而 TMINS 中所有元素共享同一個(gè)標(biāo)量閾值。典型應(yīng)用場(chǎng)景包括數(shù)據(jù)裁剪clamp 上界、激活函數(shù)實(shí)現(xiàn)中的上界約束如min(x, alpha)、量化/歸一化中的閾值截?cái)嗟刃枰獙⒄麄€(gè) Tile 與一個(gè)固定常數(shù)比較的運(yùn)算。匯編語(yǔ)法三級(jí)抽象層次TMINS 在 PTO 指令集中同時(shí)存在三種匯編形式分別對(duì)應(yīng)不同的編譯/調(diào)度階段。同步形式PTO 匯編最直接的指令形式%scalar直接給出標(biāo)量值類型以f32等具體 dtype 標(biāo)注%dst tmins %src, %scalar : !pto.tile..., f32AS Level 1SSA 形式SSA靜態(tài)單賦值中間表示操作數(shù)與結(jié)果均為!pto.tile...類型的值value由編譯器后續(xù)進(jìn)行資源分配%dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...AS Level 2DPS 形式DPSDestructive/顯式 place-and-schedule形式此時(shí) Tile 已被綁定到具體緩沖區(qū)!pto.tile_buf...輸入與輸出分離聲明pto.tmins ins(%src, %scalar : !pto.tile_buf..., dtype) outs(%dst : !pto.tile_buf...)從源碼結(jié)構(gòu)看pto.tmins指令名稱中的 s 后綴貫穿三個(gè)抽象層次保持一致便于后端在編譯流水線中逐級(jí)降級(jí)映射。C 內(nèi)建接口TMINS 的 C 內(nèi)建接口模板聲明于 include/pto/common/pto_instr.hpp對(duì)開發(fā)者開放的公共包含頭為pto/pto-inst.hpptemplate typename TileDataDst, typename TileDataSrc, typename... WaitEvents PTO_INST RecordEvent TMINS(TileDataDst dst, TileDataSrc src, typename TileDataSrc::DType scalar, WaitEvents ... events);接口要點(diǎn)參數(shù)順序目標(biāo) Tiledst、源 Tilesrc、標(biāo)量scalar其類型必須等于TileDataSrc::DType、可選的事件列表events用于跨流水線同步等待返回類型PTO_INST RecordEvent即返回一個(gè)記錄事件可與后續(xù)指令構(gòu)成事件依賴鏈宏分發(fā)TMINS通過(guò)MAP_INSTR_IMPL(TMINS, dst, src, scalar)宏映射到平臺(tái)相關(guān)的TMINS_IMPL實(shí)現(xiàn)見(jiàn) include/pto/common/pto_instr.hpp同一份用戶代碼可無(wú)縫編譯到不同昇騰平臺(tái)或 CPU 模擬器。底層實(shí)現(xiàn)調(diào)用鏈在 NPU 側(cè)TMINS_IMPL最終落到硬件向量指令vminsAtlas A2/A3 平臺(tái)實(shí)現(xiàn)在 include/pto/npu/a2a3/TMins.hppMinSOp::BinSInstr直接調(diào)用vmins(dst, src0, src1, repeats, 1, 1, 8, 8)并支持通過(guò)dstRepeatStride/srcRepeatStride自定義 repeat 間步長(zhǎng)Ascend 950 平臺(tái)實(shí)現(xiàn)在 include/pto/npu/a5/TMins.hppMinSOp::BinSInstr調(diào)用vmins(reg_dst, reg_src0, src1, preg, MODE_ZEROING)即帶掩碼寄存器與清零模式zeroing mode的向量比較形式對(duì)于int64_t/uint64_t這類寬類型A5 實(shí)現(xiàn)走Int64ScalarInt64Op::Min, ...專用路徑將 64 位標(biāo)量比較拆分為內(nèi)部多次 32 位運(yùn)算組合完成見(jiàn) include/pto/npu/a5/TMins.hpp。在 CPU 模擬器側(cè)TMINS 與 TADDS/TSUBS/TMULS/TMAXS 等標(biāo)量指令共用同一套UnaryTileScalarOpImpl逐元素循環(huán)框架僅以ElementOp::OP_MINS區(qū)分運(yùn)算類型實(shí)現(xiàn)在 include/pto/cpu/TBinSOps.hpp。這意味著在無(wú)昇騰硬件的開發(fā)環(huán)境中CPU 模擬可以給出與硬件一致的數(shù)值行為便于算子邏輯先行驗(yàn)證。約束與有效區(qū)域使用 TMINS 時(shí)必須滿足以下平臺(tái)與通用約束違反約束會(huì)在編譯期static_assert或運(yùn)行期PTO_ASSERT報(bào)錯(cuò)。Atlas A2/A3 訓(xùn)練/推理系列產(chǎn)品實(shí)現(xiàn)檢查TileData::DType必須是以下類型之一int32_t、int、int16_t、half、float16_t、float、float32_t運(yùn)行時(shí)src.GetValidRow() dst.GetValidRow()且src.GetValidCol() dst.GetValidCol()。對(duì)應(yīng)實(shí)現(xiàn)中A2/A3 的TMINS_IMPL通過(guò)static_assert校驗(yàn)數(shù)據(jù)類型集合并通過(guò)PTO_ASSERT(src.GetValidCol() dst.GetValidCol(), ...)與PTO_ASSERT(src.GetValidRow() dst.GetValidRow(), ...)在運(yùn)行期強(qiáng)制行列有效邊界一致見(jiàn) include/pto/npu/a2a3/TMins.hpp。Ascend 950PR / Ascend 950DT實(shí)現(xiàn)檢查TileData::DType必須是以下類型之一uint8_t、int8_t、uint16_t、int16_t、uint32_t、int32_t、int64_t、uint64_t、half、float、bfloat16_t——相比 A2/A3 大幅擴(kuò)展了整數(shù)與低精度類型覆蓋并額外支持bfloat16_t運(yùn)行時(shí)src.GetValidCol() dst.GetValidCol()。對(duì)應(yīng)實(shí)現(xiàn)中A5 的TMINS_IMPL同樣以static_assert校驗(yàn)類型集合與TileType::Vec位置以PTO_ASSERT(src0.GetValidCol() dst.GetValidCol(), ...)校驗(yàn)列數(shù)一致見(jiàn) include/pto/npu/a5/TMins.hpp。通用約束所有平臺(tái)dst與src必須使用相同的元素類型A2/A3 實(shí)現(xiàn)在類型層通過(guò)static_assert(std::is_same_vT, typename TileDataDst::DType, ...)直接強(qiáng)制標(biāo)量類型必須與 Tile 數(shù)據(jù)類型一致接口簽名中scalar類型即TileDataSrc::DTypeTile 位置必須是向量TileData::Loc TileType::Vec。有效區(qū)域TMINS 以dst.GetValidRow()/dst.GetValidCol()作為迭代域運(yùn)算只覆蓋目標(biāo) Tile 聲明的有效行列范圍內(nèi)超出部分不參與計(jì)算。因此調(diào)用方應(yīng)保證src的有效區(qū)域至少覆蓋dst的有效區(qū)域避免讀取越界數(shù)據(jù)。實(shí)戰(zhàn)示例自動(dòng)模式與手動(dòng)模式PTO 提供兩種 Tile 資源管理方式。自動(dòng)模式由編譯器/運(yùn)行時(shí)負(fù)責(zé) Tile 的內(nèi)存放置與調(diào)度手動(dòng)模式通過(guò)TASSIGN顯式將 Tile 綁定到指定地址再由指令發(fā)射。自動(dòng)Auto模式#include pto/pto-inst.hpp using namespace pto; void example_auto() { using TileT TileTileType::Vec, float, 16, 16; TileT src, dst; TMINS(dst, src, 0.0f); }聲明一個(gè) 16×16 的float向量 Tile調(diào)用TMINS將src中每個(gè)元素與0.0f比較取小寫入dst。0.0f的類型與TileT::DTypefloat一致滿足通用約束。手動(dòng)Manual模式#include pto/pto-inst.hpp using namespace pto; void example_manual() { using TileT TileTileType::Vec, float, 16, 16; TileT src, dst; TASSIGN(src, 0x1000); TASSIGN(dst, 0x2000); TMINS(dst, src, 0.0f); }先通過(guò)TASSIGN將src綁定到地址0x1000、dst綁定到0x2000再發(fā)射TMINS。這種模式適合開發(fā)者需要精確控制緩沖區(qū)布局如復(fù)用固定的片上 buffer 池的場(chǎng)景。匯編示例ASM自動(dòng)模式自動(dòng)模式下資源放置與調(diào)度由編譯器/運(yùn)行時(shí)完成指令以 SSA 值形式出現(xiàn)# 自動(dòng)模式由編譯器/運(yùn)行時(shí)負(fù)責(zé)資源放置與調(diào)度。 %dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...手動(dòng)模式手動(dòng)模式先顯式綁定資源pto.tassign將 tile 操作數(shù)綁定到物理地址再發(fā)射指令# 手動(dòng)模式先顯式綁定資源再發(fā)射指令。 # 可選當(dāng)該指令包含 tile 操作數(shù)時(shí) # pto.tassign %arg0, tile(0x1000) # pto.tassign %arg1, tile(0x2000) %dst pto.tmins %src, %scalar : (!pto.tile..., dtype) - !pto.tile...PTO 匯編形式匯總%dst tmins %src, %scalar : !pto.tile..., f32 # AS Level 2 (DPS) pto.tmins ins(%src, %scalar : !pto.tile_buf..., dtype) outs(%dst : !pto.tile_buf...)端到端算子示例TLOAD → TMINS → TSTORE倉(cāng)庫(kù)中 tests/cpu/st/testcase/tmins/tmins_kernel.cpp 給出了一個(gè)完整的端到端 kernel將全局內(nèi)存數(shù)據(jù)加載到 Tile執(zhí)行 TMINS 標(biāo)量裁剪再存回全局內(nèi)存并用set_flag/wait_flag在 MTE2加載、V向量運(yùn)算、MTE3存儲(chǔ)三條流水線之間建立事件同步template typename T, int kGRows_, int kGCols_, int kTRows_, int kTCols_ __global__ AICORE void runTMins(__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(kTRows_, kTCols_); TileData dstTile(kTRows_, kTCols_); TASSIGN(src0Tile, 0x0 0x400); TASSIGN(dstTile, 0x8000 0x400); GlobalData src0Global(src0); GlobalData dstGlobal(out); TLOAD(src0Tile, src0Global); set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0); TMINS(dstTile, src0Tile, src1[0]); // 標(biāo)量來(lái)自全局內(nèi)存首元素 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0); TSTORE(dstGlobal, dstTile); out dstGlobal.data(); }該示例同時(shí)展示了一個(gè)實(shí)用技巧標(biāo)量操作數(shù)不一定是編譯期常量也可以是src1[0]這類從全局內(nèi)存讀取的運(yùn)行時(shí)值使 TMINS 可以用于動(dòng)態(tài)閾值的裁剪邏輯。測(cè)試驗(yàn)證與 golden 機(jī)制TMINS 在倉(cāng)庫(kù)中擁有完整的多平臺(tái)測(cè)試覆蓋包括 CPU 模擬tests/cpu/st/testcase/tmins/、A2/A3tests/npu/a2a3/src/st/testcase/tmins/、A5tests/npu/a5/src/st/testcase/tmins/以及 costmodel 相關(guān)測(cè)試tests/costmodel/st/testcase/tmins/main.cpp。以 CPU 側(cè)測(cè)試 tests/cpu/st/testcase/tmins/main.cpp 為例其驗(yàn)證流程為通過(guò)gen_data.py生成隨機(jī)輸入input1.bin、標(biāo)量輸入input_scalar.bin與參考 golden 文件golden.bin使用 ACL 運(yùn)行時(shí) APIaclrtMalloc/aclrtMemcpy分配主機(jī)與設(shè)備內(nèi)存將輸入拷貝到設(shè)備啟動(dòng)LaunchTMinskernel同步流后將結(jié)果拷回主機(jī)寫為output.bin將output.bin與golden.bin逐元素比較容差 0.0001f通過(guò)EXPECT_TRUE(ret)判定測(cè)試通過(guò)。覆蓋的數(shù)據(jù)類型與形狀包括float/int32_t/int64_t/uint64_t/int16_t/half以及開啟CPU_SIM_BFLOAT_ENABLED時(shí)的bfloat16_t形狀覆蓋 64×64 與 16×256 兩種 Tile 配置見(jiàn) tests/cpu/st/testcase/tmins/main.cpp可直接作為自行擴(kuò)展 TMINS 用例的模板。TMINS 與 TMIN 的對(duì)比維度TMINSTile × 標(biāo)量TMINTile × Tile數(shù)學(xué)語(yǔ)義dst[i][j] min(src[i][j], scalar)dst[i][j] min(src0[i][j], src1[i][j])第二操作數(shù)標(biāo)量值TileDataSrc::DType另一個(gè) TileA2/A3 數(shù)據(jù)類型int32_t/int/int16_t/half/float16_t/float/float32_tint32_t/int16_t/half/float并額外要求行主序布局與靜態(tài)有效邊界檢查運(yùn)行時(shí)校驗(yàn)行列有效邊界一致src0/src1/dst的validRow/validCol相同典型用途閾值截?cái)?、上界裁剪逐元素融合比較如 ReLU6、逐點(diǎn) clamp 下限兩者都要求TileType::Vec向量位置、以dst有效區(qū)域?yàn)榈蚯揖暶饔?include/pto/common/pto_instr.hpp接口風(fēng)格完全一致可在同一 kernel 中混合使用?？偨Y(jié)TMINS 是 PTO 指令集中一個(gè)簡(jiǎn)潔但高頻使用的 Tile-標(biāo)量二元指令數(shù)學(xué)上等價(jià)于對(duì)有效區(qū)域內(nèi)每個(gè)元素執(zhí)行min(x, scalar)編程上通過(guò)TMINS(dst, src, scalar)一行即可完成整個(gè) Tile 的標(biāo)量裁剪。使用時(shí)需要特別關(guān)注兩點(diǎn)一是平臺(tái)相關(guān)的數(shù)據(jù)類型支持范圍A2/A3 與 950 系列差異明顯950 額外支持 8/16/64 位整數(shù)與bfloat16_t二是有效區(qū)域一致性約束src與dst的 valid 行列必須匹配。其底層實(shí)現(xiàn)無(wú)論是 A2/A3 的vminsrepeat 循環(huán)、A5 的帶掩碼vmins與 64 位專用路徑還是 CPU 模擬器的UnaryTileScalarOpImpl逐元素循環(huán)都保證了跨平臺(tái)一致的語(yǔ)義配合倉(cāng)庫(kù)內(nèi)覆蓋多平臺(tái)、多數(shù)據(jù)類型、多 Tile 形狀的完整測(cè)試用例開發(fā)者可以放心在昇騰 NPU 上以自動(dòng)或手動(dòng)模式使用 TMINS 實(shí)現(xiàn)高效的逐元素最小值運(yùn)算?！久赓M(fèi)下載鏈接】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.項(xiàng)目地址: https://gitcode.com/cann/pto-isa創(chuàng)作聲明：本文部分內(nèi)容由AI輔助生成（AIGC），僅供參考