diff --git a/.github/workflows/build_wheel.yml b/.github/workflows/build_wheel.yml index 42ee5aff9e..c0d2244436 100644 --- a/.github/workflows/build_wheel.yml +++ b/.github/workflows/build_wheel.yml @@ -25,7 +25,7 @@ permissions: env: LLVM_REPOSITORY: https://github.com/LinxISA/llvm-project.git - LLVM_COMMIT: 1245c0f89aab24104ea935fa686a0ac5ad9ab2c9 + LLVM_COMMIT: b7c83f68bf84125e696a70bec4b665c70a3b584d jobs: build_wheel: @@ -82,7 +82,7 @@ jobs: - name: Install Python dependencies run: | export PATH="${PY_PATH}/bin:$PATH" - pip install --no-cache-dir numpy nanobind typing_extensions setuptools wheel auditwheel + pip install --no-cache-dir numpy 'nanobind>=2.9,<3' typing_extensions setuptools wheel auditwheel - name: Set build directories run: | diff --git a/.github/workflows/build_wheel_mac.yml b/.github/workflows/build_wheel_mac.yml index 0cda2122cc..fe94fde068 100644 --- a/.github/workflows/build_wheel_mac.yml +++ b/.github/workflows/build_wheel_mac.yml @@ -24,7 +24,7 @@ permissions: env: LLVM_REPOSITORY: https://github.com/LinxISA/llvm-project.git - LLVM_COMMIT: 1245c0f89aab24104ea935fa686a0ac5ad9ab2c9 + LLVM_COMMIT: b7c83f68bf84125e696a70bec4b665c70a3b584d jobs: build_wheel: @@ -83,7 +83,7 @@ jobs: - name: Install Python dependencies run: | - pip install --no-cache-dir numpy nanobind typing_extensions setuptools wheel delocate + pip install --no-cache-dir numpy 'nanobind>=2.9,<3' typing_extensions setuptools wheel delocate - name: Set build directories run: | diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 82c4bcd9b5..bcb720bbfd 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -94,7 +94,7 @@ jobs: runs-on: ubuntu-22.04 env: LLVM_REPOSITORY: https://github.com/LinxISA/llvm-project.git - LLVM_COMMIT: 1245c0f89aab24104ea935fa686a0ac5ad9ab2c9 + LLVM_COMMIT: b7c83f68bf84125e696a70bec4b665c70a3b584d LLVM_DIR: ${{ github.workspace }}/llvm-project/llvm/build-shared PTO_INSTALL_DIR: ${{ github.workspace }}/install MLIR_PYTHONPATH: ${{ github.workspace }}/llvm-project/llvm/build-shared/tools/mlir/python_packages/mlir_core @@ -119,7 +119,7 @@ jobs: clang lld \ libedit-dev zlib1g-dev libxml2-dev libzstd-dev python3 -m pip install --upgrade pip - python3 -m pip install nanobind typing_extensions numpy + python3 -m pip install 'nanobind>=2.9,<3' typing_extensions numpy - name: Define payload paths shell: bash @@ -232,7 +232,7 @@ jobs: include/PTO/IR/PTOOps.td \ tools/ptobc/generated/ptobc_opcodes_v0.h ctest --test-dir build --output-on-failure \ - -R 'pto_isa_v0581_contract_check|ptobc_v0580_ops_v0_encode|ptobc_v0581_contract_encode|ptobc_tpartarg_v0_encode' + -R 'pto_isa_v0583_contract_check|ptobc_v0580_ops_v0_encode|ptobc_v0583_contract_encode|ptobc_tpartarg_v0_encode' - name: Run sample tests (py -> pto -> cpp) shell: bash diff --git a/.github/workflows/isa_contract.yml b/.github/workflows/isa_contract.yml index 9a02601574..916df1d4e8 100644 --- a/.github/workflows/isa_contract.yml +++ b/.github/workflows/isa_contract.yml @@ -1,14 +1,22 @@ -name: PTO ISA v0.58.1 contract +name: PTO ISA v0.58.3 contract on: push: tags: - - linxisa-v0.58.1 + - linxisa-v0.58.3 pull_request: paths: - docker/** - .dockerignore - .github/workflows/isa_contract.yml + - include/PTO/IR/** + - lib/PTO/IR/** + - lib/PTO/Transforms/PTOToEmitC.cpp + - test/lit/pto/v0583_** + - tools/check_v058_pto_manifest.py + - tools/test_check_v058_pto_manifest.py + - tools/pto_isa_v0_58_3_** + - tools/ptobc/** - tools/check_release_delivery_contract.py - tools/test_release_delivery_contract.py workflow_dispatch: @@ -22,7 +30,9 @@ jobs: steps: - uses: actions/checkout@v4 - name: Check exact PTO ISA contract - run: python3 tools/check_v058_pto_manifest.py --ptoas-root . + run: | + python3 tools/check_v058_pto_manifest.py --ptoas-root . + python3 tools/test_check_v058_pto_manifest.py docker-builder: runs-on: ubuntu-latest diff --git a/CMakeLists.txt b/CMakeLists.txt index e7bf6e3f52..11122715ed 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -135,12 +135,12 @@ if(BUILD_TESTING) enable_testing() find_package(Python3 COMPONENTS Interpreter REQUIRED) add_test( - NAME pto_isa_v0581_contract_check + NAME pto_isa_v0583_contract_check COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/tools/check_v058_pto_manifest.py --ptoas-root ${CMAKE_CURRENT_SOURCE_DIR}) add_test( - NAME pto_isa_v0581_contract_check_negative + NAME pto_isa_v0583_contract_check_negative COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/tools/test_check_v058_pto_manifest.py) add_test( @@ -173,6 +173,48 @@ if(BUILD_TESTING) NAME emitc_lvalue_contract_check COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/tools/check_emitc_lvalue_contract.py) + option(PTOAS_ENABLE_LINX_TILEOP_INTEGRATION + "Compile generated Linx MX C++ with the exact target toolchain and TileOP" OFF) + set(PTOAS_LINX_LLVM_BUILD "" CACHE PATH "Exact merged Linx LLVM build") + set(PTOAS_LINX_SYSROOT "" CACHE PATH "Phase-C Linx musl+libc++ sysroot") + set(PTOAS_LINX_TILEOP_ROOT "" CACHE PATH "Exact merged Linx TileOP checkout") + set(PTOAS_EXPECTED_LLVM_COMMIT + "b7c83f68bf84125e696a70bec4b665c70a3b584d" CACHE STRING + "Reviewed merged Linx LLVM commit") + set(PTOAS_EXPECTED_LLVM_TREE + "c11bd80c7dd34ed4438de1da6bde0a01eae8c76d" CACHE STRING + "Reviewed merged Linx LLVM tree") + set(PTOAS_EXPECTED_TILEOP_COMMIT + "bd1ecca97ca47da0edc462c1ce19749c6940780e" CACHE STRING + "Reviewed merged Linx TileOP commit") + set(PTOAS_EXPECTED_TILEOP_TREE + "854c463c21f9b758186b718105f4b035019cdd30" CACHE STRING + "Reviewed merged Linx TileOP tree") + if(PTOAS_ENABLE_LINX_TILEOP_INTEGRATION) + foreach(required_var IN ITEMS + PTOAS_LINX_LLVM_BUILD PTOAS_LINX_SYSROOT PTOAS_LINX_TILEOP_ROOT + PTOAS_EXPECTED_LLVM_COMMIT PTOAS_EXPECTED_LLVM_TREE + PTOAS_EXPECTED_TILEOP_COMMIT PTOAS_EXPECTED_TILEOP_TREE) + if(NOT DEFINED ${required_var} OR "${${required_var}}" STREQUAL "") + message(FATAL_ERROR + "${required_var} is required for the Linx TileOP integration gate") + endif() + endforeach() + add_test( + NAME linx_mx_tileop_compile_check + COMMAND ${CMAKE_COMMAND} -E env + PTOAS_BIN=$ + LINX_LLVM_BUILD=${PTOAS_LINX_LLVM_BUILD} + LINX_SYSROOT=${PTOAS_LINX_SYSROOT} + TILEOP_ROOT=${PTOAS_LINX_TILEOP_ROOT} + EXPECTED_LLVM_COMMIT=${PTOAS_EXPECTED_LLVM_COMMIT} + EXPECTED_LLVM_TREE=${PTOAS_EXPECTED_LLVM_TREE} + EXPECTED_TILEOP_COMMIT=${PTOAS_EXPECTED_TILEOP_COMMIT} + EXPECTED_TILEOP_TREE=${PTOAS_EXPECTED_TILEOP_TREE} + bash ${CMAKE_CURRENT_SOURCE_DIR}/tools/check_linx_mx_tileop_compile.sh) + set_tests_properties(linx_mx_tileop_compile_check PROPERTIES + LABELS "linx-target-integration") + endif() add_subdirectory(tools/ptobc/tests) endif() diff --git a/README.md b/README.md index c1e82ad558..c69cfd5afe 100644 --- a/README.md +++ b/README.md @@ -3,8 +3,8 @@ ## 1. 项目简介 (Introduction) **ptoas** (`ptoas`) 是一个基于 LinxISA 受控 **LLVM/MLIR** -（ISA release `linxisa-v0.58.1`，LLVM commit -`1245c0f89aab24104ea935fa686a0ac5ad9ab2c9`）构建的专用编译器工具链， +（ISA release `linxisa-v0.58.3`，LLVM commit +`b7c83f68bf84125e696a70bec4b665c70a3b584d`）构建的专用编译器工具链， 专为 **PTO Bytecode** (Programming Tiling Operator Bytecode) 设计。 作为连接上层 AI 框架与底层各类NPU/GPGPU/CPU硬件，`ptoas` 采用 **Out-of-Tree** 架构构建，提供了完整的 C++ 与 Python 接口，主要职责包括： @@ -41,7 +41,7 @@ PTOAS/ ## 3. 构建指南 (Build Instructions) ⚠️ **重要提示**：本项目严格依赖 LinxISA LLVM -`1245c0f89aab24104ea935fa686a0ac5ad9ab2c9`。`linxisa-v0.58.1` +`b7c83f68bf84125e696a70bec4b665c70a3b584d`。`linxisa-v0.58.3` 是 ISA release tag，不是 PTOAS 产品版本；不要替换为同名上游 LLVM tag。 @@ -92,7 +92,7 @@ python3 -m pip install nanobind numpy cd $WORKSPACE_DIR git clone https://github.com/LinxISA/llvm-project.git cd $LLVM_SOURCE_DIR -git checkout --detach 1245c0f89aab24104ea935fa686a0ac5ad9ab2c9 +git checkout --detach b7c83f68bf84125e696a70bec4b665c70a3b584d # 2. 配置 CMake (构建动态库并启用 Python 绑定) cmake -G Ninja -S llvm -B $LLVM_BUILD_DIR \ @@ -221,7 +221,7 @@ ptoas --version ``` `ptoas --version` 中的 `0.41` 是 PTOAS 产品版本；括号中的 `PTO ISA -0.58.1` 是独立的 ISA contract 版本。`linx` target 对齐受管的 +0.58.3` 是独立的 ISA contract 版本。`linx` target 对齐受管的 `Linx-TileOP-API`：生成代码包含 `jcore/template_asm.hpp`，只接受 v0.58 公共 PTO 操作目录，并在 lowering 前拒绝仅属于 A3/A5 方言面的操作。 @@ -253,6 +253,36 @@ python3 ./tmatmulk.py > ./tmatmulk.pto $PTO_SOURCE_DIR/build/tools/ptoas/ptoas ./tmatmulk.pto -o ./tmatmulk.cpp ``` +Linx MX 的真实 C++ 集成门禁默认不启用，因为它必须同时绑定精确的 Linx +LLVM 构建、musl+libc++ sysroot 和 `Linx-TileOP-API` checkout。该门禁不会使用 +host shim 或替换 `jcore/template_asm.hpp` 的 stub。依赖准备好后，使用已审核并 +合并的 commit/tree 配置独立构建目录： + +```bash +cmake -G Ninja -S . -B build-linx-integration \ + -DLLVM_DIR=$LLVM_BUILD_DIR/lib/cmake/llvm \ + -DMLIR_DIR=$LLVM_BUILD_DIR/lib/cmake/mlir \ + -DPTOAS_ENABLE_LINX_TILEOP_INTEGRATION=ON \ + -DPTOAS_LINX_LLVM_BUILD=$LINX_LLVM_BUILD \ + -DPTOAS_LINX_SYSROOT=$LINX_SYSROOT \ + -DPTOAS_LINX_TILEOP_ROOT=$TILEOP_ROOT \ + -DPTOAS_EXPECTED_LLVM_COMMIT=b7c83f68bf84125e696a70bec4b665c70a3b584d \ + -DPTOAS_EXPECTED_LLVM_TREE=c11bd80c7dd34ed4438de1da6bde0a01eae8c76d \ + -DPTOAS_EXPECTED_TILEOP_COMMIT=bd1ecca97ca47da0edc462c1ce19749c6940780e \ + -DPTOAS_EXPECTED_TILEOP_TREE=854c463c21f9b758186b718105f4b035019cdd30 +ninja -C build-linx-integration ptoas +ctest --test-dir build-linx-integration \ + -L linx-target-integration --output-on-failure +``` + +门禁先核对两个 checkout 的 commit/tree，再由真实 Linx `clang++` 和指定 +sysroot 对 PTOAS 生成的 TMATMUL/TGEMV base、ACC、BIAS 的 zero/A-only/ +B-only/both 四种形式分别执行 `-fsyntax-only` 和目标对象编译。缺少路径、 +identity 不匹配或真实 TileOP 签名/约束不接受生成代码时，门禁都会失败。 +门禁最后还运行 TileOP 的 phase-C link smoke，并用 +`verify_pto_identity.py` 核对每个 Linx ELF 中唯一且精确的 PTO 0.58.3 +identity note。 + ### 5.4 上板验证 该流程用于将 `test/samples` 下生成的 `.cpp`（ptoas 输出）自动生成 NPU 验证用例，并在 NPU 上运行。下面示例直接复用 5.3 里生成的 `MatMul/tmatmulk.cpp`。 diff --git a/ReleaseNotes.md b/ReleaseNotes.md index f94982946c..99517847a8 100644 --- a/ReleaseNotes.md +++ b/ReleaseNotes.md @@ -1,5 +1,14 @@ # PTOAS (PTO Assembler & Optimizer) +## Unreleased: PTO ISA 0.58.3 alignment + +- Locks the exact PTO-SPEC v0.58.3 release identity and all four catalogs. +- Adds Linx CUBE CELL layouts (`cube_m16`, `cube_m32`, `cube_n8`) and + fail-closed CUBE operand verification. +- Updates TGEMV/TGEMVMX operand order and distinct ACC/BIAS TileOP calls. +- Records TLOAD/TSTORE byte row-stride roles and verifies the 0.58.3 + B.IOT/B.IOS/B.FPATR/DTYPE_NONE hard-break fields. + ## 版本 - 版本号：v0.1.0 - 发布日期：2026-02-14 diff --git a/docker/Dockerfile b/docker/Dockerfile index f3919570ee..06e1e16ce3 100644 --- a/docker/Dockerfile +++ b/docker/Dockerfile @@ -11,7 +11,7 @@ ARG ARCH # NOTE: change $PY_VER for different Python versions (3.8 - 3.14 available) ARG PY_VER=cp311-cp311 ARG LLVM_REPOSITORY=https://github.com/LinxISA/llvm-project.git -ARG LLVM_COMMIT=1245c0f89aab24104ea935fa686a0ac5ad9ab2c9 +ARG LLVM_COMMIT=b7c83f68bf84125e696a70bec4b665c70a3b584d ## -- usually no need to change below -- @@ -20,7 +20,10 @@ ENV PATH="${PY_PATH}/bin:${PATH}" # dependency RUN dnf install -y ninja-build cmake git ccache gcc-c++ lld zip binutils patchelf chrpath && dnf clean all -RUN pip install --no-cache-dir numpy nanobind setuptools wheel auditwheel +# LLVM 23's MLIR Python build requires the nanobind 2.9 CMake contract. +# PTOAS itself also supports nanobind 3 in the hosted prebuilt-LLVM lanes, but +# this source-build stage must stay below the nanobind 3 hard break. +RUN pip install --no-cache-dir numpy 'nanobind>=2.9,<3' setuptools wheel auditwheel COPY cmake/LinuxHardeningCache.cmake /tmp/LinuxHardeningCache.cmake diff --git a/docker/check_ptoas_cli_identity.sh b/docker/check_ptoas_cli_identity.sh index ba7f7250b2..c4b6bca29b 100644 --- a/docker/check_ptoas_cli_identity.sh +++ b/docker/check_ptoas_cli_identity.sh @@ -24,13 +24,13 @@ if [[ "$VERSION_OUTPUT" == *$'\n'* || "$VERSION_OUTPUT" == *$'\r'* ]]; then fi if [[ -n "$PTOAS_VERSION" ]]; then - EXPECTED_VERSION_OUTPUT="ptoas ${PTOAS_VERSION} (PTO ISA 0.58.1)" + EXPECTED_VERSION_OUTPUT="ptoas ${PTOAS_VERSION} (PTO ISA 0.58.3)" if [[ "$VERSION_OUTPUT" != "$EXPECTED_VERSION_OUTPUT" ]]; then echo "Error: expected '${EXPECTED_VERSION_OUTPUT}', got '${VERSION_OUTPUT}'" >&2 exit 1 fi else - EXPECTED_VERSION_OUTPUT="ptoas ${PTOAS_RELEASE_VERSION} (PTO ISA 0.58.1)" + EXPECTED_VERSION_OUTPUT="ptoas ${PTOAS_RELEASE_VERSION} (PTO ISA 0.58.3)" if [[ "$VERSION_OUTPUT" != "$EXPECTED_VERSION_OUTPUT" ]]; then echo "Error: invalid packaged ptoas identity '${VERSION_OUTPUT}'" >&2 exit 1 diff --git a/docs/PTO_IR_manual.md b/docs/PTO_IR_manual.md index c0ca602d54..301827541a 100644 --- a/docs/PTO_IR_manual.md +++ b/docs/PTO_IR_manual.md @@ -79,6 +79,7 @@ PTO IR currently recognizes the following low-precision element types: - `f8E5M2` (corresponding C++ type name: `float8_e5m2_t`) - `!pto.hif8` +- `!pto.f8E8M0` (Linx MX scale type; emitted as `__fp8_e8m0`) - `!pto.f4E1M2x2` - `!pto.f4E2M1x2` @@ -88,6 +89,7 @@ basic storage-size plumbing. Their storage size is currently modeled as: - `f8E4M3FN`: 1 byte per element - `f8E5M2`: 1 byte per element - `!pto.hif8`: 1 byte per element +- `!pto.f8E8M0`: 1 byte per element - `!pto.f4E1M2x2`: 1 byte per packed pair of FP4 values - `!pto.f4E2M1x2`: 1 byte per packed pair of FP4 values @@ -144,12 +146,12 @@ A logical partition (slice) of a `tensor_view`. Holds shape and stride informati | Parameter | Type | Description | |-----------|------|-------------| | `loc` | keyword (`vec/mat/left/right/acc/bias`) | Local memory domain (`vec` maps to UB; use `vec` in textual IR) | -| `dtype` | `element-type(i1/i8/i16/i32/f16/f32/bf16/!pto.hif8/!pto.f4E1M2x2/!pto.f4E2M1x2...)` | Element data type | +| `dtype` | `element-type(i1/i8/i16/i32/f16/f32/bf16/!pto.hif8/!pto.f8E8M0/!pto.f4E1M2x2/!pto.f4E2M1x2...)` | Element data type | | `rows` | `int64` | Physical row count | | `cols` | `int64` | Physical column count | | `v_row` | `int64` or `?` | Valid row count | | `v_col` | `int64` or `?` | Valid column count | -| `blayout` | `BLayout` mnemonic | Base layout (`row_major` / `col_major`) | +| `blayout` | `BLayout` mnemonic | Base layout (`row_major`, `col_major`, or Linx CUBE `cube_m16` / `cube_m32` / `cube_n8`) | | `slayout` | `SLayout` mnemonic | Secondary layout (`none_box` / `row_major` / `col_major`) | | `fractal` | `int32` | Fractal size | | `pad` | `PadValue` mnemonic or integer literal | Padding policy/value selector (tests commonly use `pad=0`) | @@ -266,7 +268,7 @@ Composite attribute and component enums for tile buffer configuration. | Parameter | Type | Description | |-----------|------|-------------| -| `bLayout` | `BLayoutAttr` | Base layout (RowMajor / ColMajor) | +| `bLayout` | `BLayoutAttr` | Base layout (RowMajor / ColMajor / CubeM16 / CubeM32 / CubeN8) | | `sLayout` | `SLayoutAttr` | Secondary layout (NoneBox / RowMajor / ColMajor) | | `sFractalSize` | `IntegerAttr (i32)` | Secondary fractal size | | `pad` | `PadValueAttr` | Pad value policy | @@ -279,6 +281,15 @@ Composite attribute and component enums for tile buffer configuration. |-------|-----|----------| | `RowMajor` | 0 | `row_major` | | `ColMajor` | 1 | `col_major` | +| `CubeM16` | 2 | `cube_m16` | +| `CubeM32` | 3 | `cube_m32` | +| `CubeN8` | 4 | `cube_n8` | + +The three CUBE layouts are the PTO ISA 0.58.3 CELL layouts. Linx CUBE A and +accumulator tiles use `cube_m16` or `cube_m32`; B tiles use `cube_n8`. They +must use `none_box` secondary layout. Generated Linx TileOP types preserve +these layouts and therefore select `B.DATR` CUBE conversion codes with +`DTYPE_NONE` rather than treating encoded datatype zero as inheritance. **SLayout** (Secondary layout): @@ -754,6 +765,10 @@ For each element (i, j) in the tile valid region: `partition_tensor_view` and `tile_buf` are both 2-D in this IR profile. `pto.tload` moves data from the global logical view into the local physical tile buffer. +On Linx PTO ISA 0.58.3, the row stride carried by the lowered TLOAD scalar +operand is measured in bytes. PTOAS derives that byte stride from the global +tensor descriptor; it must not pass an element count to the TileOP API. + **Arguments:** | Name | Type | Description | @@ -779,6 +794,13 @@ For each element (i, j) in the tile valid region: - **Implementation checks (A5)** - The destination tile element size must be `1`, `2`, `4`, or `8` bytes, and must match the source partition element size. - For `i64`, the destination tile `pad` must be `null` or `zero`. +- **Implementation checks (Linx PTO ISA 0.58.3)** + - The source uses `loc=gm`; the destination may use + `vec`, `left`, `right`, `acc`, `bias`, or `scaling`. + - Source and destination element sizes match. Compact FP8/FP4 and E8M0 + tiles are valid load destinations. + - `left`, `right`, and `acc` CUBE CELL destinations lower to + `TLOAD_CUBE`; other local destinations lower to `TLOAD`. **Hardware Mapping:** @@ -797,7 +819,7 @@ pto.tload ins(%pv : !pto.partition_tensor_view<16x16xf16>) **Summary:** Probes and prefetches a rectangular global tile region with TLOAD-equivalent translation, permission, fault, restart, coherence, and ordering behavior. PTO -ISA 0.58.1 makes this operation destination-free: it changes +ISA 0.58.3 makes this operation destination-free: it changes target cache or data-movement state without publishing a tile result. **Semantics:** @@ -810,7 +832,7 @@ The detailed cache placement is target-defined, but access faults are architectural rather than optional hint behavior. Unlike most generated PTO intrinsic wrappers, `TPREFETCH` does not add implicit wait-event synchronization. -For Linx ISA 0.58.1, statically known `valid_cols`, `valid_rows`, and +For Linx ISA 0.58.3, statically known `valid_cols`, `valid_rows`, and `physical_cols` values must be in `1..65535`. `physical_cols` must also be a power of two and at least `valid_cols`. Dynamic values are accepted by the IR; the producer must guarantee the same constraints at runtime before execution. @@ -865,6 +887,10 @@ pto.tprefetch ins(%address, %stride, %cols, %rows, %physical_cols : i64, index, ##### `pto.tstore` - Store Tile to Partition View +On Linx PTO ISA 0.58.3, the row stride carried by the lowered TSTORE scalar +operand is measured in bytes, matching TLOAD. CUBE CELL sources retain their +explicit `cube_m16`, `cube_m32`, or `cube_n8` layout during lowering. + **Summary:** Stores a 2-D tile buffer back to a 2-D partition view. Supports phase/atomic/relu/pre-quant controls that lower to the corresponding `TSTORE` template overload family. **Semantics:** @@ -1216,6 +1242,16 @@ For each (i, j): - Shape constraints: `lhs.rows == dst.rows`, `lhs.cols == rhs.rows`, and `rhs.cols == dst.cols`. - Tile locations: `lhs.loc=left`, `rhs.loc=right`, `dst.loc=acc`. - Runtime: `m/k/n` (taken from `lhs valid row`, `lhs valid column`, `rhs valid column`) must be in `[1, 4095]`. +- **Implementation checks (Linx PTO ISA 0.58.3)** + - `lhs.loc=left` with `blayout=cube_m16` or `cube_m32`. + - `rhs.loc=right` with `blayout=cube_n8`. + - `dst.loc=acc` with the same M16/M32 layout as `lhs`. + - All three operands use `slayout=none_box`; CUBE element widths are + limited to 4, 8, 16, or 32 bits. + - Mixed A/B input types are permitted when the selected TileOP overload + supports them. The generated call preserves architectural A/B order. + - `pto.tgemv*` keeps A/vector then B/matrix in PTO IR, but emits the public + TileOP order `TGEMV(dst, matrix_b, vector_a)`. - **Implementation checks (A5)** - The destination element type must be `i32` or `f32`. - If the destination element type is `i32`, the lhs and rhs element types must both be `i8`. @@ -1349,9 +1385,9 @@ For each (i, j): | Name | Type | Description | |------|------|-------------| | `lhs` | `pto.tile_buf` | Left matrix | -| `lhs_scale` | `pto.tile_buf` | Left scaling tile | +| `lhs_scale` | optional `pto.tile_buf` | Left scaling tile | | `rhs` | `pto.tile_buf` | Right matrix | -| `rhs_scale` | `pto.tile_buf` | Right scaling tile | +| `rhs_scale` | optional `pto.tile_buf` | Right scaling tile | | `dst` | `pto.tile_buf` | Destination | **Results:** None. Writes into `dst` via DPS pattern. @@ -1360,6 +1396,19 @@ For each (i, j): - **Implementation checks (A5)** - `m/k/n` are taken from `lhs valid row`, `lhs valid column`, and `rhs valid column`. +- **Implementation checks (Linx PTO ISA 0.58.3)** + - Each A/B input independently selects its scale schema: FP16/BF16 requires + no scale operand; a compact FP8/FP4 input requires its corresponding scale. + - Linx EmitC spells FP16/BF16 as the target frontend types `__half` and + `__bf16`; compact inputs use the corresponding `__fp8_*`/`__fp4_*` types. + - Supplying a scale for FP16/BF16, or omitting one for a compact input, is an + error. Thus zero-scale, A-only, B-only, and two-scale forms are distinct. + - Each present scale must use `!pto.f8E8M0`, emitted as `__fp8_e8m0`, in the + scaling address space with ordinary row-major layout. + - `lhs_scale.valid_shape = [M, ceil(K/32)]`. + - `rhs_scale.valid_shape = [ceil(K/32), N]`. + - Base, accumulation, and bias variants lower one-to-one to `TMATMUL_MX`, + `TMATMUL_MX_ACC`, and `TMATMUL_MX_BIAS` with no operand permutation. **Hardware Mapping:** @@ -1368,8 +1417,9 @@ For each (i, j): **Basic Example:** ```mlir -pto.tmatmul.mx ins(%a, %a_scale, %b, %b_scale : !pto.tile_buf<...>, !pto.tile_buf<...>, - !pto.tile_buf<...>, !pto.tile_buf<...>) +pto.tmatmul.mx ins(%a, %b : !pto.tile_buf<...>, !pto.tile_buf<...>) + a_scale(%a_scale : !pto.tile_buf<...>) + b_scale(%b_scale : !pto.tile_buf<...>) outs(%c : !pto.tile_buf<...>) ``` @@ -1391,9 +1441,9 @@ dst = acc_in + (lhs * rhs) // scaling tiles configure target-defined behavior |------|------|-------------| | `acc_in` | `pto.tile_buf` | Accumulator input | | `lhs` | `pto.tile_buf` | Left matrix | -| `lhs_scale` | `pto.tile_buf` | Left scaling tile | +| `lhs_scale` | optional `pto.tile_buf` | Left scaling tile | | `rhs` | `pto.tile_buf` | Right matrix | -| `rhs_scale` | `pto.tile_buf` | Right scaling tile | +| `rhs_scale` | optional `pto.tile_buf` | Right scaling tile | | `dst` | `pto.tile_buf` | Destination | **Results:** None. Writes into `dst` via DPS pattern. @@ -1410,8 +1460,9 @@ dst = acc_in + (lhs * rhs) // scaling tiles configure target-defined behavior **Basic Example:** ```mlir -pto.tmatmul.mx.acc ins(%c_in, %a, %a_scale, %b, %b_scale : !pto.tile_buf<...>, !pto.tile_buf<...>, - !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) +pto.tmatmul.mx.acc ins(%c_in, %a, %b : !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) + a_scale(%a_scale : !pto.tile_buf<...>) + b_scale(%b_scale : !pto.tile_buf<...>) outs(%c_out : !pto.tile_buf<...>) ``` @@ -1432,9 +1483,9 @@ dst = (lhs * rhs) + bias // scaling tiles configure target-defined behavior | Name | Type | Description | |------|------|-------------| | `lhs` | `pto.tile_buf` | Left matrix | -| `lhs_scale` | `pto.tile_buf` | Left scaling tile | +| `lhs_scale` | optional `pto.tile_buf` | Left scaling tile | | `rhs` | `pto.tile_buf` | Right matrix | -| `rhs_scale` | `pto.tile_buf` | Right scaling tile | +| `rhs_scale` | optional `pto.tile_buf` | Right scaling tile | | `bias` | `pto.tile_buf` | Bias tile | | `dst` | `pto.tile_buf` | Destination | @@ -1454,8 +1505,9 @@ dst = (lhs * rhs) + bias // scaling tiles configure target-defined behavior **Basic Example:** ```mlir -pto.tmatmul.mx.bias ins(%a, %a_scale, %b, %b_scale, %bias : !pto.tile_buf<...>, !pto.tile_buf<...>, - !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) +pto.tmatmul.mx.bias ins(%a, %b, %bias : !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) + a_scale(%a_scale : !pto.tile_buf<...>) + b_scale(%b_scale : !pto.tile_buf<...>) outs(%c : !pto.tile_buf<...>) ``` @@ -1655,10 +1707,10 @@ dst = gemv(a, b) // quantization/mixed-precision behavior is target-defined | Name | Type | Description | |------|------|-------------| -| `a` | `pto.tile_buf` | Matrix tile (`loc=left`) | -| `a_scale` | `pto.tile_buf` | Scale tile associated with `a` | -| `b` | `pto.tile_buf` | Vector tile (`loc=right`) | -| `b_scale` | `pto.tile_buf` | Scale tile associated with `b` | +| `a` | `pto.tile_buf` | Left vector tile (`1xK`) | +| `a_scale` | optional `pto.tile_buf` | Scale tile associated with `a` | +| `b` | `pto.tile_buf` | Right matrix tile (`KxN`) | +| `b_scale` | optional `pto.tile_buf` | Scale tile associated with `b` | | `dst` | `pto.tile_buf` | Destination accumulator tile (`loc=acc`) | **Results:** None. Writes into `dst` via DPS pattern. @@ -1666,7 +1718,10 @@ dst = gemv(a, b) // quantization/mixed-precision behavior is target-defined **Constraints & Verification:** - `a/b/dst` reuse the same GEMV shape/location checks as `pto.tgemv`. -- `a_scale` and `b_scale` must be valid tile buffers. +- The Linx zero/A-only/B-only/two-scale rules are identical to + `pto.tmatmul.mx`; each present scale uses `!pto.f8E8M0`. +- EmitC preserves TileOP's matrix-before-vector order while keeping each scale + adjacent to the input it scales. **Hardware Mapping:** @@ -1675,8 +1730,9 @@ dst = gemv(a, b) // quantization/mixed-precision behavior is target-defined **Basic Example:** ```mlir -pto.tgemv.mx ins(%a, %a_scale, %b, %b_scale : !pto.tile_buf<...>, !pto.tile_buf<...>, - !pto.tile_buf<...>, !pto.tile_buf<...>) +pto.tgemv.mx ins(%a, %b : !pto.tile_buf<...>, !pto.tile_buf<...>) + a_scale(%a_scale : !pto.tile_buf<...>) + b_scale(%b_scale : !pto.tile_buf<...>) outs(%c : !pto.tile_buf<...>) ``` @@ -1692,15 +1748,16 @@ pto.tgemv.mx ins(%a, %a_scale, %b, %b_scale : !pto.tile_buf<...>, !pto.tile_buf< dst = c_in + gemv(a, b) ``` -**Arguments:** `c_in, a, a_scale, b, b_scale, dst` +**Arguments:** `c_in, a, optional a_scale, b, optional b_scale, dst` **Hardware Mapping:** Matrix pipeline (`PIPE_M`) **Basic Example:** ```mlir -pto.tgemv.mx.acc ins(%c_in, %a, %a_scale, %b, %b_scale : !pto.tile_buf<...>, !pto.tile_buf<...>, - !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) +pto.tgemv.mx.acc ins(%c_in, %a, %b : !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) + a_scale(%a_scale : !pto.tile_buf<...>) + b_scale(%b_scale : !pto.tile_buf<...>) outs(%c_out : !pto.tile_buf<...>) ``` @@ -1716,15 +1773,16 @@ pto.tgemv.mx.acc ins(%c_in, %a, %a_scale, %b, %b_scale : !pto.tile_buf<...>, !pt dst = gemv(a, b) + bias ``` -**Arguments:** `a, a_scale, b, b_scale, bias, dst` +**Arguments:** `a, optional a_scale, b, optional b_scale, bias, dst` **Hardware Mapping:** Matrix pipeline (`PIPE_M`) **Basic Example:** ```mlir -pto.tgemv.mx.bias ins(%a, %a_scale, %b, %b_scale, %bias : !pto.tile_buf<...>, !pto.tile_buf<...>, - !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) +pto.tgemv.mx.bias ins(%a, %b, %bias : !pto.tile_buf<...>, !pto.tile_buf<...>, !pto.tile_buf<...>) + a_scale(%a_scale : !pto.tile_buf<...>) + b_scale(%b_scale : !pto.tile_buf<...>) outs(%c : !pto.tile_buf<...>) ``` diff --git a/include/PTO/IR/PTOAttrs.td b/include/PTO/IR/PTOAttrs.td index 84dbd7cc51..ee27e22617 100644 --- a/include/PTO/IR/PTOAttrs.td +++ b/include/PTO/IR/PTOAttrs.td @@ -547,7 +547,10 @@ def PTO_MaskPatternAttr : PTO_Attr<"MaskPattern", "mask_pattern"> { // ---------- BLayout / SLayout / PadValue (same style as AddressSpace via PTO_I32Enum) ---------- def PTO_BLayout_Enum : PTO_I32Enum<"BLayout", "Base layout", [ I32EnumAttrCase<"RowMajor", 0, "row_major">, - I32EnumAttrCase<"ColMajor", 1, "col_major"> + I32EnumAttrCase<"ColMajor", 1, "col_major">, + I32EnumAttrCase<"CubeM16", 2, "cube_m16">, + I32EnumAttrCase<"CubeM32", 3, "cube_m32">, + I32EnumAttrCase<"CubeN8", 4, "cube_n8"> ]>; def PTO_BLayoutAttr : PTO_Attr<"BLayout", "blayout"> { diff --git a/include/PTO/IR/PTOOps.td b/include/PTO/IR/PTOOps.td index 38cdcb74be..f3d0488994 100644 --- a/include/PTO/IR/PTOOps.td +++ b/include/PTO/IR/PTOOps.td @@ -524,7 +524,7 @@ def TPrefetchOp : PTO_TOp<"tprefetch", [ ]> { let summary = "Prefetch a rectangular tile region from a global address."; let description = [{ - PTO ISA 0.58.1 models TPREFETCH as a destination-free TLSU operation over + PTO ISA 0.58.3 models TPREFETCH as a destination-free TLSU operation over a base address, element row stride, valid shape, and physical row width. }]; @@ -706,8 +706,9 @@ def TMatmulBiasOp : PTO_TOp<"tmatmul.bias", [ }]; } -// pto.tmatmul.mx ins(%a,%a_scale,%b,%b_scale) outs(%dst) -> tensor? +// Linx MX scale operands are independently optional and governed by A/B dtype. def TMatmulMxOp : PTO_TOp<"tmatmul.mx", [ + AttrSizedOperandSegments, PTO_DpsInitOpInterface, OpPipeInterface, DeclareOpInterfaceMethods @@ -716,17 +717,18 @@ def TMatmulMxOp : PTO_TOp<"tmatmul.mx", [ let arguments = (ins PTODpsType:$a, - PTODpsType:$a_scale, + Optional:$a_scale, PTODpsType:$b, - PTODpsType:$b_scale, + Optional:$b_scale, PTODpsType:$dst); let results = (outs Optional:$result); let hasVerifier = 1; let assemblyFormat = [{ - `ins` `(` $a `,` $a_scale `,` $b `,` $b_scale - `:` type($a) `,` type($a_scale) `,` type($b) `,` type($b_scale) `)` + `ins` `(` $a `,` $b `:` type($a) `,` type($b) `)` + (`a_scale` `(` $a_scale^ `:` type($a_scale) `)`)? + (`b_scale` `(` $b_scale^ `:` type($b_scale) `)`)? `outs` `(` $dst `:` qualified(type($dst) ) `)` attr-dict (`->` qualified(type($result))^)? @@ -741,8 +743,9 @@ def TMatmulMxOp : PTO_TOp<"tmatmul.mx", [ }]; } -// pto.tmatmul.mx.acc ins(%c_in,%a,%a_scale,%b,%b_scale) outs(%dst) -> tensor? +// Accumulating Linx MX form with independent A/B scale operands. def TMatmulMxAccOp : PTO_TOp<"tmatmul.mx.acc", [ + AttrSizedOperandSegments, PTO_DpsInitOpInterface, OpPipeInterface, DeclareOpInterfaceMethods @@ -752,17 +755,18 @@ def TMatmulMxAccOp : PTO_TOp<"tmatmul.mx.acc", [ let arguments = (ins PTODpsType:$c_in, PTODpsType:$a, - PTODpsType:$a_scale, + Optional:$a_scale, PTODpsType:$b, - PTODpsType:$b_scale, + Optional:$b_scale, PTODpsType:$dst); let results = (outs Optional:$result); let hasVerifier = 1; let assemblyFormat = [{ - `ins` `(` $c_in `,` $a `,` $a_scale `,` $b `,` $b_scale - `:` type($c_in) `,` type($a) `,` type($a_scale) `,` type($b) `,` type($b_scale) `)` + `ins` `(` $c_in `,` $a `,` $b `:` type($c_in) `,` type($a) `,` type($b) `)` + (`a_scale` `(` $a_scale^ `:` type($a_scale) `)`)? + (`b_scale` `(` $b_scale^ `:` type($b_scale) `)`)? `outs` `(` $dst `:` qualified(type($dst) ) `)` attr-dict (`->` qualified(type($result))^)? @@ -777,8 +781,9 @@ def TMatmulMxAccOp : PTO_TOp<"tmatmul.mx.acc", [ }]; } -// pto.tmatmul.mx.bias ins(%a,%a_scale,%b,%b_scale,%bias) outs(%dst) -> tensor? +// Bias Linx MX form with independent A/B scale operands. def TMatmulMxBiasOp : PTO_TOp<"tmatmul.mx.bias",[ + AttrSizedOperandSegments, PTO_DpsInitOpInterface, OpPipeInterface, DeclareOpInterfaceMethods @@ -787,9 +792,9 @@ def TMatmulMxBiasOp : PTO_TOp<"tmatmul.mx.bias",[ let arguments = (ins PTODpsType:$a, - PTODpsType:$a_scale, + Optional:$a_scale, PTODpsType:$b, - PTODpsType:$b_scale, + Optional:$b_scale, PTODpsType:$bias, PTODpsType:$dst); @@ -798,8 +803,10 @@ def TMatmulMxBiasOp : PTO_TOp<"tmatmul.mx.bias",[ let assemblyFormat = [{ - `ins` `(` $a `,` $a_scale `,` $b `,` $b_scale `,` $bias - `:` type($a) `,` type($a_scale) `,` type($b) `,` type($b_scale) `,` qualified(type($bias)) `)` + `ins` `(` $a `,` $b `,` $bias + `:` type($a) `,` type($b) `,` qualified(type($bias)) `)` + (`a_scale` `(` $a_scale^ `:` type($a_scale) `)`)? + (`b_scale` `(` $b_scale^ `:` type($b_scale) `)`)? `outs` `(` $dst `:` qualified(type($dst) ) `)` attr-dict (`->` qualified(type($result))^)? @@ -989,17 +996,18 @@ def TGemvBiasOp : PTO_TOp<"tgemv.bias", [ } def TGemvMxOp : PTO_TOp<"tgemv.mx", [ + AttrSizedOperandSegments, PTO_DpsInitOpInterface, OpPipeInterface, DeclareOpInterfaceMethods ]> { - let summary = "Mixed-precision GEMV with scale tiles (tile world, ins/outs)."; + let summary = "Mixed-precision GEMV with independent optional scale tiles."; let arguments = (ins PTODpsType:$a, - PTODpsType:$a_scale, + Optional:$a_scale, PTODpsType:$b, - PTODpsType:$b_scale, + Optional:$b_scale, PTODpsType:$dst ); @@ -1007,8 +1015,9 @@ def TGemvMxOp : PTO_TOp<"tgemv.mx", [ let hasVerifier = 1; let assemblyFormat = [{ - `ins` `(` $a `,` $a_scale `,` $b `,` $b_scale - `:` type($a) `,` type($a_scale) `,` type($b) `,` type($b_scale) `)` + `ins` `(` $a `,` $b `:` type($a) `,` type($b) `)` + (`a_scale` `(` $a_scale^ `:` type($a_scale) `)`)? + (`b_scale` `(` $b_scale^ `:` type($b_scale) `)`)? `outs` `(` $dst `:` qualified(type($dst) ) `)` attr-dict (`->` qualified(type($result))^)? @@ -1022,18 +1031,19 @@ def TGemvMxOp : PTO_TOp<"tgemv.mx", [ } def TGemvMxAccOp : PTO_TOp<"tgemv.mx.acc", [ + AttrSizedOperandSegments, PTO_DpsInitOpInterface, OpPipeInterface, DeclareOpInterfaceMethods ]> { - let summary = "Mixed-precision GEMV accumulate with scale tiles (tile world, ins/outs)."; + let summary = "Mixed-precision GEMV accumulate with optional scale tiles."; let arguments = (ins PTODpsType:$c_in, PTODpsType:$a, - PTODpsType:$a_scale, + Optional:$a_scale, PTODpsType:$b, - PTODpsType:$b_scale, + Optional:$b_scale, PTODpsType:$dst ); @@ -1041,8 +1051,9 @@ def TGemvMxAccOp : PTO_TOp<"tgemv.mx.acc", [ let hasVerifier = 1; let assemblyFormat = [{ - `ins` `(` $c_in `,` $a `,` $a_scale `,` $b `,` $b_scale - `:` type($c_in) `,` type($a) `,` type($a_scale) `,` type($b) `,` type($b_scale) `)` + `ins` `(` $c_in `,` $a `,` $b `:` type($c_in) `,` type($a) `,` type($b) `)` + (`a_scale` `(` $a_scale^ `:` type($a_scale) `)`)? + (`b_scale` `(` $b_scale^ `:` type($b_scale) `)`)? `outs` `(` $dst `:` qualified(type($dst) ) `)` attr-dict (`->` qualified(type($result))^)? @@ -1056,17 +1067,18 @@ def TGemvMxAccOp : PTO_TOp<"tgemv.mx.acc", [ } def TGemvMxBiasOp : PTO_TOp<"tgemv.mx.bias", [ + AttrSizedOperandSegments, PTO_DpsInitOpInterface, OpPipeInterface, DeclareOpInterfaceMethods ]> { - let summary = "Mixed-precision GEMV with bias and scale tiles (tile world, ins/outs)."; + let summary = "Mixed-precision GEMV with bias and optional scale tiles."; let arguments = (ins PTODpsType:$a, - PTODpsType:$a_scale, + Optional:$a_scale, PTODpsType:$b, - PTODpsType:$b_scale, + Optional:$b_scale, PTODpsType:$bias, PTODpsType:$dst ); @@ -1075,8 +1087,10 @@ def TGemvMxBiasOp : PTO_TOp<"tgemv.mx.bias", [ let hasVerifier = 1; let assemblyFormat = [{ - `ins` `(` $a `,` $a_scale `,` $b `,` $b_scale `,` $bias - `:` type($a) `,` type($a_scale) `,` type($b) `,` type($b_scale) `,` qualified(type($bias)) `)` + `ins` `(` $a `,` $b `,` $bias + `:` type($a) `,` type($b) `,` qualified(type($bias)) `)` + (`a_scale` `(` $a_scale^ `:` type($a_scale) `)`)? + (`b_scale` `(` $b_scale^ `:` type($b_scale) `)`)? `outs` `(` $dst `:` qualified(type($dst) ) `)` attr-dict (`->` qualified(type($result))^)? diff --git a/include/PTO/IR/PTOTypeDefs.td b/include/PTO/IR/PTOTypeDefs.td index 6b8c0ee5c6..36ad62e6b3 100644 --- a/include/PTO/IR/PTOTypeDefs.td +++ b/include/PTO/IR/PTOTypeDefs.td @@ -214,7 +214,7 @@ def TileBufType : TypeDef { mlir::Attribute getCompactModeAttr() const; // 如果你仍然想要“数值枚举”，就提供 int getter（不会依赖 enum 类型） - int32_t getBLayoutValueI32() const; // 0 row_major, 1 col_major + int32_t getBLayoutValueI32() const; // 0/1 row/col, 2/3/4 CUBE M16/M32/N8 int32_t getSLayoutValueI32() const; // 0 none_box, 1 row_major, 2 col_major int32_t getPadValueI32() const; // 0 null, 1 zero, 2 max, 3 min int32_t getCompactModeI32() const; // 0 null, 1 normal, 2 row_plus_one @@ -281,6 +281,17 @@ def HiF8Type : TypeDef +]> { + let mnemonic = "f8E8M0"; + let summary = + "PTO ISA E8M0 scale element used independently by A and B MX inputs."; +} + // ---- !pto.f4E1M2x2 ---- Corresponding type name: float4_e1m2x2_t (1 byte = 2 FP4 packed) def F4E1M2x2Type : TypeDef(BLayout::CubeM16) || + blayout == static_cast(BLayout::CubeM32) || + blayout == static_cast(BLayout::CubeN8); + if (isCubeLayout) { + if (!isVerifierTargetLinx(op)) + return op->emitOpError() << "expects " << name + << " CUBE CELL layout only for Linx"; + if (slayout != static_cast(SLayout::NoneBox)) + return op->emitOpError() << "expects " << name + << " CUBE CELL layout to use none_box"; + return success(); + } + auto checkByteAlignment = [&](int64_t dim, StringRef layoutName, StringRef byteExpr) -> LogicalResult { if (dim == ShapedType::kDynamic) @@ -2212,9 +2241,22 @@ LogicalResult AllocTileOp::verify() { auto ty = getResult().getType(); // TileBufType Type elemTy = ty.getElementType(); - if (isPTOLowPrecisionType(elemTy)) - return emitOpError() << "result dtype " << elemTy - << " is not supported by pto.alloc_tile yet"; + if (isPTOLowPrecisionType(elemTy)) { + int32_t layout = ty.getBLayoutValueI32(); + bool isLinxCubeLayout = + isVerifierTargetLinx(getOperation()) && + (layout == static_cast(pto::BLayout::CubeM16) || + layout == static_cast(pto::BLayout::CubeM32) || + layout == static_cast(pto::BLayout::CubeN8)); + auto addressSpace = getPTOMemorySpaceEnum(ty); + bool isLinxE8M0Scale = + isVerifierTargetLinx(getOperation()) && + isa(elemTy) && addressSpace && + *addressSpace == pto::AddressSpace::SCALING; + if (!isLinxCubeLayout && !isLinxE8M0Scale) + return emitOpError() << "result dtype " << elemTy + << " is not supported by pto.alloc_tile yet"; + } if (failed(verifyTileBufLayoutConstraints(*this, ty, "result"))) return failure(); @@ -2382,15 +2424,23 @@ LogicalResult TLoadOp::verify() { return emitOpError( "expects Linx tload src to be a memref/partition tensor view and " "dst to be a memref/tile buffer"); - if (failed(verifyTileBufCommon(*this, dstTy, "dst"))) + if (failed(verifyTileBufCommon(*this, dstTy, "dst", + /*allowLowPrecision=*/true))) return failure(); auto srcSpace = getPTOMemorySpaceEnum(srcTy); auto dstSpace = getPTOMemorySpaceEnum(dstTy); if (!srcSpace || *srcSpace != pto::AddressSpace::GM) return emitOpError("expects Linx tload src to use loc=gm"); - if (!dstSpace || *dstSpace != pto::AddressSpace::VEC) - return emitOpError("expects Linx tload dst to use loc=vec"); + if (!dstSpace || + (*dstSpace != pto::AddressSpace::VEC && + *dstSpace != pto::AddressSpace::LEFT && + *dstSpace != pto::AddressSpace::RIGHT && + *dstSpace != pto::AddressSpace::ACC && + *dstSpace != pto::AddressSpace::BIAS && + *dstSpace != pto::AddressSpace::SCALING)) + return emitOpError( + "expects Linx tload dst to use loc=vec/left/right/acc/bias/scaling"); if (getElemByteSize(getElemTy(srcTy)) != getElemByteSize(getElemTy(dstTy))) return emitOpError( "expects Linx tload src and dst element sizes to match"); @@ -3243,12 +3293,43 @@ static LogicalResult verifyScaleTileMatchesOperand(Operation *op, Type scaleTy, Type operandTy, StringRef scaleName, StringRef operandName) { - if (failed(verifyTileBufCommon(op, scaleTy, scaleName))) + if (failed(verifyTileBufCommon(op, scaleTy, scaleName, + /*allowLowPrecision=*/ + isVerifierTargetLinx(op)))) return failure(); auto scaleSpace = getPTOMemorySpaceEnum(scaleTy); if (!scaleSpace || *scaleSpace != pto::AddressSpace::SCALING) return op->emitOpError() << "expects " << scaleName << " to be in the scaling address space"; + if (isVerifierTargetLinx(op)) { + if (!isa(getElemTy(scaleTy))) + return op->emitOpError() << "expects Linx MX " << scaleName + << " to use !pto.f8E8M0 elements"; + auto scaleValid = getValidShapeVec(scaleTy); + auto operandValid = getValidShapeVec(operandTy); + if (scaleValid.size() != 2 || operandValid.size() != 2) + return op->emitOpError() << "expects Linx MX " << scaleName + << " and " << operandName + << " to have rank-2 valid_shape"; + auto ceilDiv32 = [](int64_t value) { + return value == ShapedType::kDynamic ? value : (value + 31) / 32; + }; + SmallVector expected = + scaleName == "a_scale" + ? SmallVector{operandValid[0], + ceilDiv32(operandValid[1])} + : SmallVector{ceilDiv32(operandValid[0]), + operandValid[1]}; + for (unsigned i = 0; i < 2; ++i) { + if (scaleValid[i] != ShapedType::kDynamic && + expected[i] != ShapedType::kDynamic && scaleValid[i] != expected[i]) + return op->emitOpError() + << "expects Linx MX " << scaleName << " valid_shape to be " + << (scaleName == "a_scale" ? "M x ceil(K/32)" + : "ceil(K/32) x N"); + } + return success(); + } auto scaleShape = getShapeVec(scaleTy); auto operandShape = getShapeVec(operandTy); @@ -3278,6 +3359,84 @@ static LogicalResult verifyScaleTileMatchesOperand(Operation *op, Type scaleTy, return success(); } +static LogicalResult verifyScaleTileValueMatchesOperand( + Operation *op, Value scale, Value operand, StringRef scaleName, + StringRef operandName) { + if (!isVerifierTargetLinx(op)) + return verifyScaleTileMatchesOperand(op, scale.getType(), operand.getType(), + scaleName, operandName); + + Type scaleTy = scale.getType(); + if (failed(verifyTileBufCommon(op, scaleTy, scaleName, + /*allowLowPrecision=*/true))) + return failure(); + auto scaleSpace = getPTOMemorySpaceEnum(scaleTy); + if (!scaleSpace || *scaleSpace != pto::AddressSpace::SCALING) + return op->emitOpError() << "expects " << scaleName + << " to be in the scaling address space"; + if (!isa(getElemTy(scaleTy))) + return op->emitOpError() << "expects Linx MX " << scaleName + << " to use !pto.f8E8M0 elements"; + + auto scaleValid = getValidShapeVec(scale); + auto operandValid = getValidShapeVec(operand); + if (scaleValid.size() != 2 || operandValid.size() != 2) + return op->emitOpError() << "expects Linx MX " << scaleName << " and " + << operandName << " to have rank-2 valid_shape"; + auto ceilDiv32 = [](int64_t value) { + return value == ShapedType::kDynamic ? value : (value + 31) / 32; + }; + SmallVector expected = + scaleName == "a_scale" + ? SmallVector{operandValid[0], + ceilDiv32(operandValid[1])} + : SmallVector{ceilDiv32(operandValid[0]), + operandValid[1]}; + for (unsigned i = 0; i < 2; ++i) { + if (scaleValid[i] != ShapedType::kDynamic && + expected[i] != ShapedType::kDynamic && scaleValid[i] != expected[i]) + return op->emitOpError() + << "expects Linx MX " << scaleName << " valid_shape to be " + << (scaleName == "a_scale" ? "M x ceil(K/32)" + : "ceil(K/32) x N"); + } + return success(); +} + +static bool linxMxInputNeedsScale(Type type) { + return pto::isPTOFloat8Type(type) || pto::isPTOFloat4PackedType(type); +} + +static LogicalResult verifyMxScaleSide(Operation *op, Value operand, + Value scale, StringRef scaleName, + StringRef operandName) { + if (!isVerifierTargetLinx(op)) { + if (!scale) + return op->emitOpError() << "expects " << scaleName + << " for the A5 MX form"; + return verifyScaleTileValueMatchesOperand(op, scale, operand, scaleName, + operandName); + } + + Type elementType = getElemTy(operand.getType()); + bool isUnscaled = elementType.isF16() || elementType.isBF16(); + bool needsScale = linxMxInputNeedsScale(elementType); + if (!isUnscaled && !needsScale) + return op->emitOpError() << "expects Linx MX " << operandName + << " dtype to be FP16, BF16, or a compact type"; + if (needsScale && !scale) + return op->emitOpError() << "expects Linx MX " << scaleName + << " because " << operandName + << " uses a compact dtype"; + if (isUnscaled && scale) + return op->emitOpError() << "expects Linx MX " << scaleName + << " to be absent for FP16/BF16 " << operandName; + if (!scale) + return success(); + return verifyScaleTileValueMatchesOperand(op, scale, operand, scaleName, + operandName); +} + static LogicalResult verifyPartialValidPattern(Operation *op, Type src0Ty, Type src1Ty, Type dstTy) { auto src0Valid = getValidShapeVec(src0Ty); @@ -3649,8 +3808,22 @@ static LogicalResult verifyAccTileCommon(Operation *op, Type ty, StringRef name) return verifyAccTileCommonA2A3(op, ty, name); case VerifierTargetArch::A5: return verifyAccTileCommonA5(op, ty, name); - case VerifierTargetArch::Linx: - return op->emitOpError("Linx ACC tile legality is not implemented"); + case VerifierTargetArch::Linx: { + if (failed(verifyTileBufCommon(op, ty, name, + /*allowLowPrecision=*/true))) + return failure(); + auto tile = dyn_cast(ty); + auto addressSpace = getPTOMemorySpaceEnum(ty); + if (!tile || !addressSpace || *addressSpace != pto::AddressSpace::ACC) + return op->emitOpError() << "expects Linx " << name + << " to be an explicit acc tile_buf"; + int32_t layout = tile.getBLayoutValueI32(); + if (layout != static_cast(pto::BLayout::CubeM16) && + layout != static_cast(pto::BLayout::CubeM32)) + return op->emitOpError() << "expects Linx " << name + << " to use cube_m16 or cube_m32 blayout"; + return success(); + } } return failure(); } @@ -3720,9 +3893,126 @@ static LogicalResult verifyMatTileOperandsA5(Operation *op, Type lhsTy, return success(); } -static LogicalResult verifyMatTileOperands(Operation *op, Type lhsTy, Type rhsTy, - Type dstTy, +static bool isLinxCubeElementType(Type type) { + if (isa(type)) + return true; + if (auto integer = dyn_cast(type)) { + unsigned width = integer.getWidth(); + return width == 4 || width == 8 || width == 16 || width == 32; + } + if (type.isBF16()) + return true; + if (auto floating = dyn_cast(type)) { + unsigned width = floating.getWidth(); + return width == 8 || width == 16 || width == 32; + } + return false; +} + +static pto::TileBufConfigAttr getEffectiveTileConfig(Value value) { + if (auto tile = dyn_cast(value.getType())) + return tile.getConfigAttr(); + if (auto bind = value.getDefiningOp()) + return bind.getConfigAttr(); + return {}; +} + +static LogicalResult verifyLinxCubeTile(Operation *op, Value value, + StringRef name, + pto::AddressSpace addressSpace, + ArrayRef layouts, + int32_t &acceptedLayout) { + Type type = value.getType(); + if (failed(verifyTileBufCommon(op, type, name, + /*allowLowPrecision=*/true))) + return failure(); + auto config = getEffectiveTileConfig(value); + if (!config) + return op->emitOpError() << "expects Linx CUBE " << name + << " to be an explicit tile or governed bind_tile"; + auto actualSpace = getPTOMemorySpaceEnum(type); + if (!actualSpace || *actualSpace != addressSpace) + return op->emitOpError() << "expects Linx CUBE " << name + << " to use its architectural address space"; + auto layoutAttr = dyn_cast_or_null(config.getBLayout()); + if (!layoutAttr) + return op->emitOpError() << "expects Linx CUBE " << name + << " to carry a governed BLayout attribute"; + auto layout = layoutAttr.getValue(); + if (!llvm::is_contained(layouts, layout)) + return op->emitOpError() << "expects Linx CUBE " << name + << " to use its architectural CUBE CELL blayout"; + auto sLayoutAttr = dyn_cast_or_null(config.getSLayout()); + if (!sLayoutAttr || sLayoutAttr.getValue() != pto::SLayout::NoneBox) + return op->emitOpError() << "expects Linx CUBE " << name + << " to use none_box slayout"; + if (!isLinxCubeElementType(getElemTy(type))) + return op->emitOpError() << "expects Linx CUBE " << name + << " dtype width to be 4, 8, 16, or 32 bits"; + auto shape = getShapeVec(type); + if (layout == pto::BLayout::CubeM16 && shape[0] != ShapedType::kDynamic && + shape[0] > 16) + return op->emitOpError() << "expects cube_m16 " << name + << " to have at most 16 rows"; + if (layout == pto::BLayout::CubeM32 && shape[0] != ShapedType::kDynamic && + shape[0] > 32) + return op->emitOpError() << "expects cube_m32 " << name + << " to have at most 32 rows"; + acceptedLayout = static_cast(layout); + return success(); +} + +static LogicalResult verifyLinxAccumulatorPair(Operation *op, Value acc, + Value dst) { + int32_t accLayout = 0; + int32_t dstLayout = 0; + if (failed(verifyLinxCubeTile( + op, acc, "acc_in", pto::AddressSpace::ACC, + {pto::BLayout::CubeM16, pto::BLayout::CubeM32}, accLayout)) || + failed(verifyLinxCubeTile( + op, dst, "dst", pto::AddressSpace::ACC, + {pto::BLayout::CubeM16, pto::BLayout::CubeM32}, dstLayout))) + return failure(); + if (accLayout != dstLayout) + return op->emitOpError( + "expects Linx CUBE acc_in and dst to use the same M16/M32 layout"); + return success(); +} + +static LogicalResult verifyMatTileOperandsLinx(Operation *op, Value lhs, + Value rhs, Value dst) { + int32_t lhsLayout = 0; + int32_t rhsLayout = 0; + int32_t dstLayout = 0; + if (failed(verifyLinxCubeTile( + op, lhs, "lhs", pto::AddressSpace::LEFT, + {pto::BLayout::CubeM16, pto::BLayout::CubeM32}, lhsLayout)) || + failed(verifyLinxCubeTile(op, rhs, "rhs", pto::AddressSpace::RIGHT, + {pto::BLayout::CubeN8}, rhsLayout)) || + failed(verifyLinxCubeTile( + op, dst, "dst", pto::AddressSpace::ACC, + {pto::BLayout::CubeM16, pto::BLayout::CubeM32}, dstLayout))) + return failure(); + if (lhsLayout != dstLayout) + return op->emitOpError( + "expects Linx CUBE lhs and dst to use the same M16/M32 layout"); + auto lhsShape = getValidShapeVec(lhs); + auto rhsShape = getValidShapeVec(rhs); + auto dstShape = getValidShapeVec(dst); + if (!hasCompatibleKnownExtent(lhsShape[0], dstShape[0]) || + !hasCompatibleKnownExtent(lhsShape[1], rhsShape[0]) || + !hasCompatibleKnownExtent(rhsShape[1], dstShape[1])) + return op->emitOpError( + "expects Linx CUBE shapes lhs[M,K], rhs[K,N], and dst[M,N]"); + return success(); +} + +static LogicalResult verifyMatTileOperands(Operation *op, Value lhs, Value rhs, + Value dst, bool allowLowPrecisionInputs) { + Type lhsTy = lhs.getType(); + Type rhsTy = rhs.getType(); + Type dstTy = dst.getType(); switch (getVerifierTargetArch(op)) { case VerifierTargetArch::A2A3: return verifyMatTileOperandsA2A3(op, lhsTy, rhsTy, dstTy, @@ -3731,7 +4021,7 @@ static LogicalResult verifyMatTileOperands(Operation *op, Type lhsTy, Type rhsTy return verifyMatTileOperandsA5(op, lhsTy, rhsTy, dstTy, allowLowPrecisionInputs); case VerifierTargetArch::Linx: - return op->emitOpError("Linx CUBE operand legality is not implemented"); + return verifyMatTileOperandsLinx(op, lhs, rhs, dst); } return failure(); } @@ -3783,9 +4073,12 @@ static LogicalResult verifyGemvTileOperandsA5(Operation *op, Type lhsTy, allowLowPrecisionInputs); } -static LogicalResult verifyGemvTileOperands(Operation *op, Type lhsTy, Type rhsTy, - Type dstTy, +static LogicalResult verifyGemvTileOperands(Operation *op, Value lhs, Value rhs, + Value dst, bool allowLowPrecisionInputs) { + Type lhsTy = lhs.getType(); + Type rhsTy = rhs.getType(); + Type dstTy = dst.getType(); switch (getVerifierTargetArch(op)) { case VerifierTargetArch::A2A3: return verifyGemvTileOperandsA2A3(op, lhsTy, rhsTy, dstTy, @@ -3794,7 +4087,7 @@ static LogicalResult verifyGemvTileOperands(Operation *op, Type lhsTy, Type rhsT return verifyGemvTileOperandsA5(op, lhsTy, rhsTy, dstTy, allowLowPrecisionInputs); case VerifierTargetArch::Linx: - return op->emitOpError("Linx CUBE GEMV legality is not implemented"); + return verifyMatTileOperandsLinx(op, lhs, rhs, dst); } return failure(); } @@ -3829,22 +4122,47 @@ static LogicalResult verifyMatBiasTileA5(Operation *op, Type biasTy, Type dstTy, return success(); } -static LogicalResult verifyMatBiasTile(Operation *op, Type biasTy, Type dstTy, +static LogicalResult verifyMatBiasTile(Operation *op, Value bias, Value dst, bool requireFloatBias) { + Type biasTy = bias.getType(); + Type dstTy = dst.getType(); switch (getVerifierTargetArch(op)) { case VerifierTargetArch::A2A3: return verifyMatBiasTileA2A3(op, biasTy, dstTy, requireFloatBias); case VerifierTargetArch::A5: return verifyMatBiasTileA5(op, biasTy, dstTy, requireFloatBias); - case VerifierTargetArch::Linx: - return op->emitOpError("Linx CUBE bias legality is not implemented"); + case VerifierTargetArch::Linx: { + if (failed(verifyTileBufCommon(op, biasTy, "bias", + /*allowLowPrecision=*/true))) + return failure(); + if (!getEffectiveTileConfig(bias)) + return op->emitOpError( + "expects Linx CUBE bias to be an explicit tile or governed bind_tile"); + if (requireFloatBias && !getElemTy(biasTy).isF32()) + return op->emitOpError("expects Linx CUBE bias to have element type f32"); + auto biasShape = getShapeVec(biasTy); + auto dstShape = getShapeVec(dstTy); + if (!hasCompatibleKnownExtent(biasShape[1], dstShape[1])) + return op->emitOpError( + "expects Linx CUBE bias and dst to have the same column shape"); + return success(); + } } return failure(); } static LogicalResult verifyMatmulTypeTriple(Operation *op, Type lhsElemTy, Type rhsElemTy, Type dstElemTy) { - bool isA5 = getVerifierTargetArch(op) == VerifierTargetArch::A5; + VerifierTargetArch arch = getVerifierTargetArch(op); + bool isA5 = arch == VerifierTargetArch::A5; + if (arch == VerifierTargetArch::Linx) { + if (isLinxCubeElementType(lhsElemTy) && + isLinxCubeElementType(rhsElemTy) && + isLinxCubeElementType(dstElemTy)) + return success(); + return op->emitOpError( + "expects Linx CUBE element widths to be 4, 8, 16, or 32 bits"); + } auto isInt8 = [](Type ty) { return ty.isInteger(8); }; @@ -5334,10 +5652,19 @@ static LogicalResult verifyA5MxTypeTriple(Operation *op, Type lhsTy, Type rhsTy, Type rhsElem = getElemTy(rhsTy); Type dstElem = getElemTy(dstTy); - if (!isA5MxInputType(lhsElem) || !isA5MxInputType(rhsElem)) + if (isVerifierTargetLinx(op)) { + auto isLinxMxInputType = [](Type type) { + return type.isF16() || type.isBF16() || linxMxInputNeedsScale(type); + }; + if (!isLinxMxInputType(lhsElem) || !isLinxMxInputType(rhsElem)) + return op->emitOpError() + << "expects Linx MX operands " << lhsName << " and " << rhsName + << " to use FP16, BF16, or compact element types"; + } else if (!isA5MxInputType(lhsElem) || !isA5MxInputType(rhsElem)) { return op->emitOpError() << "expects A5 mx operands " << lhsName << " and " << rhsName << " to use fp8 element types"; + } if (!dstElem.isF32()) return op->emitOpError() @@ -6249,9 +6576,8 @@ LogicalResult RlsBufOp::verify() { // ---- TOp ---- LogicalResult TGemvBiasOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyGemvTileOperands(*this, getA().getType(), getB().getType(), - getDst().getType())) || - failed(verifyMatBiasTile(*this, getBias().getType(), getDst().getType()))) + if (failed(verifyGemvTileOperands(*this, getA(), getB(), getDst())) || + failed(verifyMatBiasTile(*this, getBias(), getDst()))) return failure(); if (failed(verifyMatmulTypeTriple(*this, getElemTy(getA().getType()), getElemTy(getB().getType()), @@ -6261,7 +6587,8 @@ LogicalResult TGemvBiasOp::verify() { getDst().getType()); }; auto verifyA5 = [&]() -> LogicalResult { return verifyA2A3(); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA2A3); } LogicalResult TGemvMxOp::verify() { @@ -6269,12 +6596,9 @@ LogicalResult TGemvMxOp::verify() { return emitOpError("tgemv.mx is only supported on A5 targets"); }; auto verifyA5 = [&]() -> LogicalResult { - if (failed(verifyScaleTileMatchesOperand(*this, getAScale().getType(), - getA().getType(), "a_scale", "a")) || - failed(verifyScaleTileMatchesOperand(*this, getBScale().getType(), - getB().getType(), "b_scale", "b")) || - failed(verifyGemvTileOperands(*this, getA().getType(), getB().getType(), - getDst().getType(), + if (failed(verifyMxScaleSide(*this, getA(), getAScale(), "a_scale", "a")) || + failed(verifyMxScaleSide(*this, getB(), getBScale(), "b_scale", "b")) || + failed(verifyGemvTileOperands(*this, getA(), getB(), getDst(), /*allowLowPrecisionInputs=*/true))) return failure(); if (failed(verifyA5MxTypeTriple(*this, getA().getType(), getB().getType(), @@ -6283,7 +6607,8 @@ LogicalResult TGemvMxOp::verify() { return verifyMatmulLike(*this, getA().getType(), getB().getType(), getDst().getType()); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA5); } LogicalResult TGemvMxAccOp::verify() { @@ -6291,13 +6616,15 @@ LogicalResult TGemvMxAccOp::verify() { return emitOpError("tgemv.mx.acc is only supported on A5 targets"); }; auto verifyA5 = [&]() -> LogicalResult { - if (failed(verifyAccTileCommon(*this, getCIn().getType(), "c_in")) || - failed(verifyScaleTileMatchesOperand(*this, getAScale().getType(), - getA().getType(), "a_scale", "a")) || - failed(verifyScaleTileMatchesOperand(*this, getBScale().getType(), - getB().getType(), "b_scale", "b")) || - failed(verifyGemvTileOperands(*this, getA().getType(), getB().getType(), - getDst().getType(), + LogicalResult accCheck = isVerifierTargetLinx(getOperation()) + ? verifyLinxAccumulatorPair(*this, getCIn(), + getDst()) + : verifyAccTileCommon( + *this, getCIn().getType(), "c_in"); + if (failed(accCheck) || + failed(verifyMxScaleSide(*this, getA(), getAScale(), "a_scale", "a")) || + failed(verifyMxScaleSide(*this, getB(), getBScale(), "b_scale", "b")) || + failed(verifyGemvTileOperands(*this, getA(), getB(), getDst(), /*allowLowPrecisionInputs=*/true))) return failure(); if (failed(verifyA5MxTypeTriple(*this, getA().getType(), getB().getType(), @@ -6311,7 +6638,8 @@ LogicalResult TGemvMxAccOp::verify() { return verifyMatmulLike(*this, getA().getType(), getB().getType(), getDst().getType()); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA5); } LogicalResult TGemvMxBiasOp::verify() { @@ -6319,14 +6647,11 @@ LogicalResult TGemvMxBiasOp::verify() { return emitOpError("tgemv.mx.bias is only supported on A5 targets"); }; auto verifyA5 = [&]() -> LogicalResult { - if (failed(verifyScaleTileMatchesOperand(*this, getAScale().getType(), - getA().getType(), "a_scale", "a")) || - failed(verifyScaleTileMatchesOperand(*this, getBScale().getType(), - getB().getType(), "b_scale", "b")) || - failed(verifyGemvTileOperands(*this, getA().getType(), getB().getType(), - getDst().getType(), + if (failed(verifyMxScaleSide(*this, getA(), getAScale(), "a_scale", "a")) || + failed(verifyMxScaleSide(*this, getB(), getBScale(), "b_scale", "b")) || + failed(verifyGemvTileOperands(*this, getA(), getB(), getDst(), /*allowLowPrecisionInputs=*/true)) || - failed(verifyMatBiasTile(*this, getBias().getType(), getDst().getType(), + failed(verifyMatBiasTile(*this, getBias(), getDst(), /*requireFloatBias=*/true))) return failure(); if (failed(verifyA5MxTypeTriple(*this, getA().getType(), getB().getType(), @@ -6339,20 +6664,25 @@ LogicalResult TGemvMxBiasOp::verify() { if (biasShape[1] != ShapedType::kDynamic && dstShape[1] != ShapedType::kDynamic && biasShape[1] != dstShape[1]) return emitOpError("expects bias and dst to have the same column shape"); - if (failed(verifyTileBufSameValidShape(*this, getBias().getType(), + bool linxDecodedMemrefs = + isVerifierTargetLinx(getOperation()) && + isa(getBias().getType()) && + isa(getDst().getType()); + if (!linxDecodedMemrefs && + failed(verifyTileBufSameValidShape(*this, getBias().getType(), getDst().getType(), "bias", "dst"))) return failure(); return verifyMatmulLike(*this, getA().getType(), getB().getType(), getDst().getType()); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA5); } LogicalResult TMatmulBiasOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyMatTileOperands(*this, getA().getType(), getB().getType(), - getDst().getType())) || - failed(verifyMatBiasTile(*this, getBias().getType(), getDst().getType()))) + if (failed(verifyMatTileOperands(*this, getA(), getB(), getDst())) || + failed(verifyMatBiasTile(*this, getBias(), getDst()))) return failure(); if (failed(verifyMatmulTypeTriple(*this, getElemTy(getA().getType()), getElemTy(getB().getType()), @@ -6362,13 +6692,16 @@ LogicalResult TMatmulBiasOp::verify() { getDst().getType()); }; auto verifyA5 = [&]() -> LogicalResult { return verifyA2A3(); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA2A3); } LogicalResult TMatmulMxOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyTileBufCommon(*this, getAScale().getType(), "a_scale")) || - failed(verifyTileBufCommon(*this, getBScale().getType(), "b_scale"))) + if (failed(verifyMxScaleSide(*this, getA(), getAScale(), "a_scale", "a")) || + failed(verifyMxScaleSide(*this, getB(), getBScale(), "b_scale", "b")) || + failed(verifyMatTileOperands(*this, getA(), getB(), getDst(), + /*allowLowPrecisionInputs=*/true))) return failure(); return verifyMatmulLike(*this, getA().getType(), getB().getType(), getDst().getType()); @@ -6379,14 +6712,22 @@ LogicalResult TMatmulMxOp::verify() { return verifyA5MxTypeTriple(*this, getA().getType(), getB().getType(), getDst().getType(), "a", "b", "dst"); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA5); } LogicalResult TMatmulMxAccOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyAccTileCommon(*this, getCIn().getType(), "c_in")) || - failed(verifyTileBufCommon(*this, getAScale().getType(), "a_scale")) || - failed(verifyTileBufCommon(*this, getBScale().getType(), "b_scale"))) + LogicalResult accCheck = isVerifierTargetLinx(getOperation()) + ? verifyLinxAccumulatorPair(*this, getCIn(), + getDst()) + : verifyAccTileCommon( + *this, getCIn().getType(), "c_in"); + if (failed(accCheck) || + failed(verifyMxScaleSide(*this, getA(), getAScale(), "a_scale", "a")) || + failed(verifyMxScaleSide(*this, getB(), getBScale(), "b_scale", "b")) || + failed(verifyMatTileOperands(*this, getA(), getB(), getDst(), + /*allowLowPrecisionInputs=*/true))) return failure(); return success(); }; @@ -6403,15 +6744,16 @@ LogicalResult TMatmulMxAccOp::verify() { return failure(); return success(); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA5); } LogicalResult TMatmulMxBiasOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyTileBufCommon(*this, getAScale().getType(), "a_scale")) || - failed(verifyTileBufCommon(*this, getBScale().getType(), "b_scale")) || - failed(verifyMatTileOperands(*this, getA().getType(), getB().getType(), - getDst().getType())) || - failed(verifyMatBiasTile(*this, getBias().getType(), getDst().getType(), + if (failed(verifyMxScaleSide(*this, getA(), getAScale(), "a_scale", "a")) || + failed(verifyMxScaleSide(*this, getB(), getBScale(), "b_scale", "b")) || + failed(verifyMatTileOperands(*this, getA(), getB(), getDst(), + /*allowLowPrecisionInputs=*/true)) || + failed(verifyMatBiasTile(*this, getBias(), getDst(), /*requireFloatBias=*/true))) return failure(); return verifyMatmulLike(*this, getA().getType(), getB().getType(), @@ -6423,7 +6765,8 @@ LogicalResult TMatmulMxBiasOp::verify() { return verifyA5MxTypeTriple(*this, getA().getType(), getB().getType(), getDst().getType(), "a", "b", "dst"); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA5); } // ---- TSetValOp ---- LogicalResult TSetValOp::verify() { @@ -9353,8 +9696,7 @@ mlir::LogicalResult mlir::pto::TPrintOp::verify() { LogicalResult mlir::pto::TMatmulOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyMatTileOperands(*this, getLhs().getType(), getRhs().getType(), - getDst().getType()))) + if (failed(verifyMatTileOperands(*this, getLhs(), getRhs(), getDst()))) return failure(); if (failed(verifyMatmulTypeTriple(*this, getElemTy(getLhs().getType()), getElemTy(getRhs().getType()), @@ -9364,13 +9706,13 @@ LogicalResult mlir::pto::TMatmulOp::verify() { getDst().getType()); }; auto verifyA5 = [&]() -> LogicalResult { return verifyA2A3(); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA2A3); } LogicalResult mlir::pto::TGemvOp::verify() { auto verifyA2A3 = [&]() -> LogicalResult { - if (failed(verifyGemvTileOperands(*this, getLhs().getType(), getRhs().getType(), - getDst().getType()))) + if (failed(verifyGemvTileOperands(*this, getLhs(), getRhs(), getDst()))) return failure(); if (failed(verifyMatmulTypeTriple(*this, getElemTy(getLhs().getType()), getElemTy(getRhs().getType()), @@ -9380,15 +9722,24 @@ LogicalResult mlir::pto::TGemvOp::verify() { getDst().getType()); }; auto verifyA5 = [&]() -> LogicalResult { return verifyA2A3(); }; - return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5); + return dispatchVerifierByArch(getOperation(), verifyA2A3, verifyA5, + verifyA2A3); } LogicalResult mlir::pto::TMatmulAccOp::verify() { if (shouldBypassDecodedMemrefVerifier(getOperation())) return success(); - if (failed(verifyAccTileCommon(*this, getAccIn().getType(), "acc_in")) || - failed(verifyMatTileOperands(*this, getLhs().getType(), getRhs().getType(), - getDst().getType()))) + LogicalResult accCheck = isVerifierTargetLinx(getOperation()) + ? verifyLinxAccumulatorPair(*this, getAccIn(), + getDst()) + : verifyAccTileCommon( + *this, getAccIn().getType(), "acc_in"); + if (failed(accCheck) || + failed(verifyMatTileOperands(*this, getLhs(), getRhs(), getDst())) || + failed(verifyTileBufSameElemType(*this, getAccIn().getType(), + getDst().getType(), "acc_in", "dst")) || + failed(verifyTileBufSameValidShape(*this, getAccIn().getType(), + getDst().getType(), "acc_in", "dst"))) return failure(); return success(); } @@ -9396,9 +9747,17 @@ LogicalResult mlir::pto::TMatmulAccOp::verify() { LogicalResult mlir::pto::TGemvAccOp::verify() { if (shouldBypassDecodedMemrefVerifier(getOperation())) return success(); - if (failed(verifyAccTileCommon(*this, getAccIn().getType(), "acc_in")) || - failed(verifyGemvTileOperands(*this, getLhs().getType(), getRhs().getType(), - getDst().getType()))) + LogicalResult accCheck = isVerifierTargetLinx(getOperation()) + ? verifyLinxAccumulatorPair(*this, getAccIn(), + getDst()) + : verifyAccTileCommon( + *this, getAccIn().getType(), "acc_in"); + if (failed(accCheck) || + failed(verifyGemvTileOperands(*this, getLhs(), getRhs(), getDst())) || + failed(verifyTileBufSameElemType(*this, getAccIn().getType(), + getDst().getType(), "acc_in", "dst")) || + failed(verifyTileBufSameValidShape(*this, getAccIn().getType(), + getDst().getType(), "acc_in", "dst"))) return failure(); return success(); } @@ -10227,6 +10586,13 @@ static void addEffect( if (operand) effects.emplace_back(effect, operand, SideEffects::DefaultResource::get()); } + +static void addEffect( + SmallVectorImpl> &effects, + MutableOperandRange operands, MemoryEffects::Effect *effect) { + for (OpOperand &operand : operands) + addEffect(effects, &operand, effect); +} // 针对结果 (Result) 的重载 static void addEffect( @@ -10921,9 +11287,9 @@ void TGemvBiasOp::getEffects(SmallVectorImpl> &effects) { addEffect(effects, &getAMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getAScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getAScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getBScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getBScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getDstMutable(), MemoryEffects::Write::get()); } @@ -10932,9 +11298,9 @@ void TGemvMxOp::getEffects(SmallVectorImpl> &effects) { addEffect(effects, &getCInMutable(), MemoryEffects::Read::get()); addEffect(effects, &getAMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getAScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getAScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getBScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getBScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getDstMutable(), MemoryEffects::Write::get()); } @@ -10942,9 +11308,9 @@ void TGemvMxAccOp::getEffects(SmallVectorImpl> &effects) { addEffect(effects, &getAMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getAScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getAScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getBScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getBScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBiasMutable(), MemoryEffects::Read::get()); addEffect(effects, &getDstMutable(), MemoryEffects::Write::get()); } @@ -10952,9 +11318,9 @@ void TGemvMxBiasOp::getEffects(SmallVectorImpl> &effects) { addEffect(effects, &getAMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getAScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getAScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getBScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getBScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getDstMutable(), MemoryEffects::Write::get()); } @@ -10963,9 +11329,9 @@ void TMatmulMxOp::getEffects(SmallVectorImpl> &effects) { addEffect(effects, &getCInMutable(), MemoryEffects::Read::get()); addEffect(effects, &getAMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getAScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getAScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getBScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getBScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getDstMutable(), MemoryEffects::Write::get()); } @@ -10973,9 +11339,9 @@ void TMatmulMxAccOp::getEffects(SmallVectorImpl> &effects) { addEffect(effects, &getAMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getAScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getAScaleMutable(), MemoryEffects::Read::get()); addEffect(effects, &getBMutable(), MemoryEffects::Read::get()); - addEffect(effects, &getBScaleMutable(), MemoryEffects::Read::get()); + addEffect(effects, getBScaleMutable(), MemoryEffects::Read::get()); // 这里的 bias 是必选的 AnyType:$bias，所以是 Singleton addEffect(effects, &getBiasMutable(), MemoryEffects::Read::get()); addEffect(effects, &getDstMutable(), MemoryEffects::Write::get()); diff --git a/lib/PTO/IR/PTOTypeUtils.cpp b/lib/PTO/IR/PTOTypeUtils.cpp index 32128f12f5..391f7df066 100644 --- a/lib/PTO/IR/PTOTypeUtils.cpp +++ b/lib/PTO/IR/PTOTypeUtils.cpp @@ -32,12 +32,15 @@ bool mlir::pto::isPTOFloat8Type(Type t) { bool mlir::pto::isPTOHiFloat8Type(Type t) { return isa(t); } +bool mlir::pto::isPTOFloat8E8M0Type(Type t) { return isa(t); } + bool mlir::pto::isPTOFloat4PackedType(Type t) { return isa(t); } bool mlir::pto::isPTOLowPrecisionType(Type t) { - return isPTOFloat8Type(t) || isPTOHiFloat8Type(t) || isPTOFloat4PackedType(t); + return isPTOFloat8Type(t) || isPTOHiFloat8Type(t) || + isPTOFloat8E8M0Type(t) || isPTOFloat4PackedType(t); } unsigned mlir::pto::getPTOStorageElemByteSize(Type t) { diff --git a/lib/PTO/Transforms/GraphSyncSolver/MemInfo.cpp b/lib/PTO/Transforms/GraphSyncSolver/MemInfo.cpp index 50d0024cb6..757230207e 100644 --- a/lib/PTO/Transforms/GraphSyncSolver/MemInfo.cpp +++ b/lib/PTO/Transforms/GraphSyncSolver/MemInfo.cpp @@ -11,6 +11,7 @@ #include "PTO/Transforms/GraphSyncSolver/MemInfo.h" #include "PTO/IR/PTO.h" +#include "PTO/IR/PTOTypeUtils.h" #include "../Utils.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/IR/BuiltinTypeInterfaces.h" @@ -29,7 +30,7 @@ static std::optional getBufferBitSize(Value value) { return ShapedType::kDynamic; } Type elementType = shaped.getElementType(); - auto bitWidth = elementType.getIntOrFloatBitWidth(); + auto bitWidth = pto::getPTOStorageElemByteSize(elementType) * pto::kBitsToByte; if (bitWidth == 0) { return ShapedType::kDynamic; } diff --git a/lib/PTO/Transforms/InferPTOLayout.cpp b/lib/PTO/Transforms/InferPTOLayout.cpp index 5d9f4d6b76..77dae36e0c 100644 --- a/lib/PTO/Transforms/InferPTOLayout.cpp +++ b/lib/PTO/Transforms/InferPTOLayout.cpp @@ -207,6 +207,10 @@ static std::optional tileBLayoutToGlobalLayout(Type tileLikeTy) { return Layout::ND; case BLayout::ColMajor: return Layout::DN; + case BLayout::CubeM16: + case BLayout::CubeM32: + case BLayout::CubeN8: + return std::nullopt; } return std::nullopt; } diff --git a/lib/PTO/Transforms/InsertSync/PTOIRTranslator.cpp b/lib/PTO/Transforms/InsertSync/PTOIRTranslator.cpp index 523a38b903..e5edb41855 100644 --- a/lib/PTO/Transforms/InsertSync/PTOIRTranslator.cpp +++ b/lib/PTO/Transforms/InsertSync/PTOIRTranslator.cpp @@ -192,7 +192,7 @@ static std::pair getStaticOffsetAndSize(Operation *op, Value s auto srcType = dyn_cast(src.getType()); if (!srcType) return {0, 0}; - int64_t elemSize = srcType.getElementType().getIntOrFloatBitWidth() / 8; + int64_t elemSize = pto::getPTOStorageElemByteSize(srcType.getElementType()); if (elemSize == 0) elemSize = 1; // === Case 1: memref.subview === @@ -395,7 +395,7 @@ LogicalResult PTOIRTranslator::UpdateAllocTileOpMemInfo(pto::AllocTileOp op) { } if (isStatic) { - int64_t elemSize = tileType.getElementType().getIntOrFloatBitWidth() / 8; + int64_t elemSize = pto::getPTOStorageElemByteSize(tileType.getElementType()); int64_t numElements = 1; for (auto dim : shape) numElements *= dim; sizeInBytes = numElements * elemSize; @@ -440,7 +440,7 @@ LogicalResult PTOIRTranslator::UpdatePointerCastOpMemInfo(pto::PointerCastOp op) uint64_t sizeInBytes = 0; if (memRefType.hasStaticShape()) { - int64_t elemSize = memRefType.getElementType().getIntOrFloatBitWidth() / 8; + int64_t elemSize = pto::getPTOStorageElemByteSize(memRefType.getElementType()); int64_t numElements = 1; for (auto dim : memRefType.getShape()) numElements *= dim; sizeInBytes = numElements * elemSize; @@ -474,7 +474,7 @@ PTOIRTranslator::UpdateDeclareTileMemRefOpMemInfo(pto::DeclareTileMemRefOp op) { uint64_t sizeInBytes = 0; if (memRefType.hasStaticShape()) { - int64_t elemSize = memRefType.getElementType().getIntOrFloatBitWidth() / 8; + int64_t elemSize = pto::getPTOStorageElemByteSize(memRefType.getElementType()); if (elemSize == 0) elemSize = 1; @@ -885,7 +885,7 @@ LogicalResult PTOIRTranslator::UpdateMemrefAllocOpMemInfo(memref::AllocOp op) { // 1. 计算大小 (Bytes) uint64_t sizeInBytes = 0; if (memRefType.hasStaticShape()) { - int64_t elemSize = memRefType.getElementType().getIntOrFloatBitWidth() / 8; + int64_t elemSize = pto::getPTOStorageElemByteSize(memRefType.getElementType()); if (elemSize == 0) elemSize = 1; // bool case int64_t numElements = 1; diff --git a/lib/PTO/Transforms/PTOPlanMemory.cpp b/lib/PTO/Transforms/PTOPlanMemory.cpp index dc8c2a5022..2d371f95ef 100644 --- a/lib/PTO/Transforms/PTOPlanMemory.cpp +++ b/lib/PTO/Transforms/PTOPlanMemory.cpp @@ -10,6 +10,7 @@ //===----------------------------------------------------------------------===// #include "PTOPlanMemory.h" +#include "PTO/IR/PTOTypeUtils.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" @@ -836,7 +837,9 @@ BufferInfo MemLivenessAnalysis::GetBufferInfo(Operation *op, Value operand, llvm::report_fatal_error("failed to obtain buffer static shape size"); bufferInfo.constBits = totalStaticSize.value() * - static_cast(memRefType.getElementTypeBitWidth()); + static_cast( + pto::getPTOStorageElemByteSize(memRefType.getElementType()) * + kBitsPerByte); return bufferInfo; } diff --git a/lib/PTO/Transforms/PTOToEmitC.cpp b/lib/PTO/Transforms/PTOToEmitC.cpp index 57465e4e55..315def4d3c 100644 --- a/lib/PTO/Transforms/PTOToEmitC.cpp +++ b/lib/PTO/Transforms/PTOToEmitC.cpp @@ -287,20 +287,24 @@ static bool isEmitCTileLikeType(Type ty) { } static std::string getEmitCScalarTypeToken(Type elemTy) { + bool isLinx = + getPTOParserTargetArch(elemTy.getContext()) == PTOParserTargetArch::Linx; if (pto::isPTOFloat8E4M3FamilyType(elemTy)) - return "float8_e4m3_t"; + return isLinx ? "__fp8_e4m3" : "float8_e4m3_t"; if (pto::isPTOFloat8E5M2FamilyType(elemTy)) - return "float8_e5m2_t"; + return isLinx ? "__fp8_e5m2" : "float8_e5m2_t"; if (isa(elemTy)) - return "hifloat8_t"; + return isLinx ? "__hif8" : "hifloat8_t"; + if (isa(elemTy)) + return "__fp8_e8m0"; if (isa(elemTy)) - return "float4_e1m2x2_t"; + return isLinx ? "__fp4_e1m2x2" : "float4_e1m2x2_t"; if (isa(elemTy)) - return "float4_e2m1x2_t"; + return isLinx ? "__fp4_e2m1x2" : "float4_e2m1x2_t"; if (elemTy.isF16()) - return "half"; + return isLinx ? "__half" : "half"; if (elemTy.isBF16()) - return "bfloat16_t"; + return isLinx ? "__bf16" : "bfloat16_t"; if (elemTy.isF32()) return "float"; if (elemTy.isF64()) @@ -434,27 +438,43 @@ class PTOToEmitCTypeConverter : public TypeConverter { // --------------------------------------------------------- // 1. 基本类型 (f32, i32, index) // --------------------------------------------------------- - addConversion([Ctx](FloatType type) -> Type { + addConversion([Ctx, targetArch](FloatType type) -> Type { if (pto::isPTOFloat8E4M3FamilyType(type)) - return emitc::OpaqueType::get(Ctx, "float8_e4m3_t"); + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__fp8_e4m3" + : "float8_e4m3_t"); if (pto::isPTOFloat8E5M2FamilyType(type)) - return emitc::OpaqueType::get(Ctx, "float8_e5m2_t"); + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__fp8_e5m2" + : "float8_e5m2_t"); if (type.isF32()) return emitc::OpaqueType::get(Ctx, "float"); - if (type.isF16()) return emitc::OpaqueType::get(Ctx, "half"); - if (type.isBF16()) return emitc::OpaqueType::get(Ctx, "bfloat16_t"); + if (type.isF16()) + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__half" : "half"); + if (type.isBF16()) + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__bf16" : "bfloat16_t"); if (type.isF64()) return emitc::OpaqueType::get(Ctx, "double"); llvm::errs() << "[Debug] Unsupported FloatType: " << type << "\n"; return Type{}; }); - addConversion([Ctx](pto::HiF8Type) -> Type { - return emitc::OpaqueType::get(Ctx, "hifloat8_t"); + addConversion([Ctx, targetArch](pto::HiF8Type) -> Type { + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__hif8" : "hifloat8_t"); }); - addConversion([Ctx](pto::F4E1M2x2Type) -> Type { - return emitc::OpaqueType::get(Ctx, "float4_e1m2x2_t"); + addConversion([Ctx](pto::F8E8M0Type) -> Type { + return emitc::OpaqueType::get(Ctx, "__fp8_e8m0"); }); - addConversion([Ctx](pto::F4E2M1x2Type) -> Type { - return emitc::OpaqueType::get(Ctx, "float4_e2m1x2_t"); + addConversion([Ctx, targetArch](pto::F4E1M2x2Type) -> Type { + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__fp4_e1m2x2" + : "float4_e1m2x2_t"); + }); + addConversion([Ctx, targetArch](pto::F4E2M1x2Type) -> Type { + return emitc::OpaqueType::get( + Ctx, targetArch == PTOArch::Linx ? "__fp4_e2m1x2" + : "float4_e2m1x2_t"); }); addConversion([Ctx](IntegerType type) -> Type { @@ -576,7 +596,8 @@ class PTOToEmitCTypeConverter : public TypeConverter { // --------------------------------------------------------- // 3. MemRef 转换 (Debug 重点) // --------------------------------------------------------- - addConversion([this, Ctx](MemRefType type) -> std::optional { + addConversion([this, Ctx, + targetArch](MemRefType type) -> std::optional { LLVM_DEBUG(llvm::dbgs() << "Converting MemRef: " << type << "\n"); // A. 转换元素类型 @@ -609,7 +630,11 @@ class PTOToEmitCTypeConverter : public TypeConverter { qualifier = "__gm__"; // Fallback } - std::string finalTypeStr = qualifier + " " + elemTypeStr; + if (targetArch == PTOArch::Linx) + qualifier.clear(); + + std::string finalTypeStr = + qualifier.empty() ? elemTypeStr : qualifier + " " + elemTypeStr; LLVM_DEBUG(llvm::dbgs() << " [Success] -> " << finalTypeStr << "*\n"); return emitc::PointerType::get(emitc::OpaqueType::get(Ctx, finalTypeStr)); @@ -3452,29 +3477,7 @@ struct SubviewToEmitCPattern : public OpConversionPattern { //===----------------------------------------------------------------------===// static std::string getElemTypeStringForGT(Type elemTy) { - if (elemTy.isF16()) return "half"; - if (elemTy.isBF16()) return "bfloat16_t"; - if (elemTy.isF32()) return "float"; - if (elemTy.isF64()) return "double"; - if (elemTy.isInteger(8)) { - if (elemTy.isSignlessInteger(8) || elemTy.isSignedInteger(8)) - return "int8_t"; - return "uint8_t"; - } - if (elemTy.isInteger(16)) { - if (elemTy.isSignlessInteger(16) || elemTy.isSignedInteger(16)) - return "int16_t"; - return "uint16_t"; - } - if (elemTy.isInteger(32)) { - if (elemTy.isSignlessInteger(32) || elemTy.isSignedInteger(32)) - return "int32_t"; - return "uint32_t"; - } - if (elemTy.isInteger(64)) { - return cast(elemTy).isUnsigned() ? "uint64_t" : "int64_t"; - } - return "float"; + return getEmitCScalarTypeToken(elemTy); } static bool hasStaticShape(MemRefType mrTy) { @@ -3766,12 +3769,22 @@ static Value materializeTensorViewDataPointer( } static std::string tileBufBLayoutToken(pto::TileBufConfigAttr configAttr) { - std::string blTok = "BLayout::RowMajor"; - if (auto blAttr = dyn_cast(configAttr.getBLayout())) { - if (static_cast(blAttr.getValue()) == 1) - blTok = "BLayout::ColMajor"; - } - return blTok; + auto blAttr = dyn_cast(configAttr.getBLayout()); + pto::BLayout layout = + blAttr ? blAttr.getValue() : pto::BLayout::RowMajor; + switch (layout) { + case pto::BLayout::RowMajor: + return "BLayout::RowMajor"; + case pto::BLayout::ColMajor: + return "BLayout::ColMajor"; + case pto::BLayout::CubeM16: + return "BLayout::CubeM16"; + case pto::BLayout::CubeM32: + return "BLayout::CubeM32"; + case pto::BLayout::CubeN8: + return "BLayout::CubeN8"; + } + llvm_unreachable("unknown PTO BLayout"); } static std::string tileBufSLayoutToken(pto::TileBufConfigAttr configAttr) { @@ -4022,8 +4035,28 @@ struct PointerCastConversion : public OpConversionPattern { if (auto attr = dyn_cast(config.getBLayout())) blVal = static_cast(attr.getValue()); - if (blVal == 1) layoutParams = "BLayout::ColMajor"; - blayout = blVal == 1 ? pto::BLayout::ColMajor : pto::BLayout::RowMajor; + switch (blVal) { + case 1: + layoutParams = "BLayout::ColMajor"; + blayout = pto::BLayout::ColMajor; + break; + case 2: + layoutParams = "BLayout::CubeM16"; + blayout = pto::BLayout::CubeM16; + break; + case 3: + layoutParams = "BLayout::CubeM32"; + blayout = pto::BLayout::CubeM32; + break; + case 4: + layoutParams = "BLayout::CubeN8"; + blayout = pto::BLayout::CubeN8; + break; + default: + layoutParams = "BLayout::RowMajor"; + blayout = pto::BLayout::RowMajor; + break; + } int32_t slVal = 0; if (auto attr = dyn_cast(config.getSLayout())) @@ -4236,8 +4269,22 @@ struct PTOTLoadToTLOAD : public OpConversionPattern { } } + StringRef callee = "TLOAD"; + if (getPTOParserTargetArch(rewriter.getContext()) == + PTOParserTargetArch::Linx) { + if (auto tileTy = dyn_cast(op.getDst().getType())) { + if (auto space = dyn_cast_or_null( + tileTy.getMemorySpace())) { + auto value = space.getAddressSpace(); + if (value == pto::AddressSpace::LEFT || + value == pto::AddressSpace::RIGHT || + value == pto::AddressSpace::ACC) + callee = "TLOAD_CUBE"; + } + } + } rewriter.create( - op.getLoc(), TypeRange{}, "TLOAD", + op.getLoc(), TypeRange{}, callee, ArrayAttr{}, ArrayAttr{}, ValueRange{dst, srcArg}); @@ -4485,15 +4532,16 @@ struct PTOTGemvToTGEMV : public OpConversionPattern { LogicalResult matchAndRewrite(pto::TGemvOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { // 1. 获取操作数 (剥离 Cast) - Value lhs = peelUnrealized(adaptor.getLhs()); // A (Matrix) - Value rhs = peelUnrealized(adaptor.getRhs()); // B (Vector) + Value lhs = peelUnrealized(adaptor.getLhs()); // A (Vector) + Value rhs = peelUnrealized(adaptor.getRhs()); // B (Matrix) Value dst = peelUnrealized(adaptor.getDst()); // C (Result) - // 2. 直接生成函数调用 TGEMV(dst, lhs, rhs) + // TileOP API order is destination, matrix (architectural B), vector + // (architectural A). PTO IR keeps lhs=A/vector and rhs=B/matrix. rewriter.create( op.getLoc(), TypeRange{}, "TGEMV", ArrayAttr{}, ArrayAttr{}, - ValueRange{dst, lhs, rhs}); + ValueRange{dst, rhs, lhs}); // 3. 处理 Op 替换/删除 if (op->getNumResults() == 1) { @@ -4518,15 +4566,15 @@ struct PTOTGemvAccToTGEMVACC : public OpConversionPattern { // 1. 获取操作数 Value accIn = peelUnrealized(adaptor.getAccIn()); // AccOld - Value lhs = peelUnrealized(adaptor.getLhs()); // A (Matrix) - Value rhs = peelUnrealized(adaptor.getRhs()); // B (Vector) + Value lhs = peelUnrealized(adaptor.getLhs()); // A (Vector) + Value rhs = peelUnrealized(adaptor.getRhs()); // B (Matrix) Value dst = peelUnrealized(adaptor.getDst()); // AccNew - // 2. 直接生成函数调用 TGEMV_ACC(dst, accIn, lhs, rhs) + // Preserve the TileOP matrix-before-vector public API order. rewriter.create( op.getLoc(), TypeRange{}, "TGEMV_ACC", ArrayAttr{}, ArrayAttr{}, - ValueRange{dst, accIn, lhs, rhs}); + ValueRange{dst, accIn, rhs, lhs}); // 3. 处理 Op 替换/删除 if (op->getNumResults() == 1) { @@ -9267,7 +9315,7 @@ struct PTOTGemvBiasToTGEMV_BIAS Value dst = peelUnrealized(adaptor.getDst()); replaceOrEraseWithOpaqueCall(op.getOperation(), "TGEMV_BIAS", - {dst, a, b, bias}, rewriter); + {dst, b, a, bias}, rewriter); return success(); } }; @@ -9279,18 +9327,24 @@ struct PTOTGemvMXToTGEMV_MX LogicalResult matchAndRewrite(pto::TGemvMxOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Value a = peelUnrealized(adaptor.getA()); - Value aScale = peelUnrealized(adaptor.getAScale()); + Value aScale = adaptor.getAScale() ? peelUnrealized(adaptor.getAScale()) : Value(); Value b = peelUnrealized(adaptor.getB()); - Value bScale = peelUnrealized(adaptor.getBScale()); + Value bScale = adaptor.getBScale() ? peelUnrealized(adaptor.getBScale()) : Value(); Value dst = peelUnrealized(adaptor.getDst()); + SmallVector operands{dst, b}; + if (bScale) + operands.push_back(bScale); + operands.push_back(a); + if (aScale) + operands.push_back(aScale); replaceOrEraseWithOpaqueCallAndReturnDst(op.getOperation(), dst, "TGEMV_MX", - {dst, a, aScale, b, bScale}, rewriter); + operands, rewriter); return success(); } }; -struct PTOTGemvMXAccToTGEMV_MX +struct PTOTGemvMXAccToTGEMV_MX_ACC : public OpConversionPattern { using OpConversionPattern::OpConversionPattern; @@ -9298,32 +9352,45 @@ struct PTOTGemvMXAccToTGEMV_MX ConversionPatternRewriter &rewriter) const override { Value cIn = peelUnrealized(adaptor.getCIn()); Value a = peelUnrealized(adaptor.getA()); - Value aScale = peelUnrealized(adaptor.getAScale()); + Value aScale = adaptor.getAScale() ? peelUnrealized(adaptor.getAScale()) : Value(); Value b = peelUnrealized(adaptor.getB()); - Value bScale = peelUnrealized(adaptor.getBScale()); + Value bScale = adaptor.getBScale() ? peelUnrealized(adaptor.getBScale()) : Value(); Value dst = peelUnrealized(adaptor.getDst()); - replaceOrEraseWithOpaqueCallAndReturnDst(op.getOperation(), dst, "TGEMV_MX", - {dst, cIn, a, aScale, b, bScale}, rewriter); + SmallVector operands{dst, cIn, b}; + if (bScale) + operands.push_back(bScale); + operands.push_back(a); + if (aScale) + operands.push_back(aScale); + replaceOrEraseWithOpaqueCallAndReturnDst( + op.getOperation(), dst, "TGEMV_MX_ACC", operands, rewriter); return success(); } }; -struct PTOTGemvMXBiasToTGEMV_MX +struct PTOTGemvMXBiasToTGEMV_MX_BIAS : public OpConversionPattern { using OpConversionPattern::OpConversionPattern; LogicalResult matchAndRewrite(pto::TGemvMxBiasOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Value a = peelUnrealized(adaptor.getA()); - Value aScale = peelUnrealized(adaptor.getAScale()); + Value aScale = adaptor.getAScale() ? peelUnrealized(adaptor.getAScale()) : Value(); Value b = peelUnrealized(adaptor.getB()); - Value bScale = peelUnrealized(adaptor.getBScale()); + Value bScale = adaptor.getBScale() ? peelUnrealized(adaptor.getBScale()) : Value(); Value bias = peelUnrealized(adaptor.getBias()); Value dst = peelUnrealized(adaptor.getDst()); - replaceOrEraseWithOpaqueCallAndReturnDst(op.getOperation(), dst, "TGEMV_MX", - {dst, a, aScale, b, bScale, bias}, rewriter); + SmallVector operands{dst, b}; + if (bScale) + operands.push_back(bScale); + operands.push_back(a); + if (aScale) + operands.push_back(aScale); + operands.push_back(bias); + replaceOrEraseWithOpaqueCallAndReturnDst( + op.getOperation(), dst, "TGEMV_MX_BIAS", operands, rewriter); return success(); } }; @@ -9352,13 +9419,19 @@ struct PTOTMatmulMXToTMATMUL_MX LogicalResult matchAndRewrite(pto::TMatmulMxOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Value a = peelUnrealized(adaptor.getA()); - Value aScale = peelUnrealized(adaptor.getAScale()); + Value aScale = adaptor.getAScale() ? peelUnrealized(adaptor.getAScale()) : Value(); Value b = peelUnrealized(adaptor.getB()); - Value bScale = peelUnrealized(adaptor.getBScale()); + Value bScale = adaptor.getBScale() ? peelUnrealized(adaptor.getBScale()) : Value(); Value dst = peelUnrealized(adaptor.getDst()); - replaceOrEraseWithOpaqueCall(op.getOperation(), "TMATMUL_MX", - {dst, a, aScale, b, bScale}, rewriter); + SmallVector operands{dst, a}; + if (aScale) + operands.push_back(aScale); + operands.push_back(b); + if (bScale) + operands.push_back(bScale); + replaceOrEraseWithOpaqueCall(op.getOperation(), "TMATMUL_MX", operands, + rewriter); return success(); } }; @@ -9371,13 +9444,19 @@ struct PTOTMatmulMXAccToTMATMUL_MX_ACC ConversionPatternRewriter &rewriter) const override { Value cIn = peelUnrealized(adaptor.getCIn()); Value a = peelUnrealized(adaptor.getA()); - Value aScale = peelUnrealized(adaptor.getAScale()); + Value aScale = adaptor.getAScale() ? peelUnrealized(adaptor.getAScale()) : Value(); Value b = peelUnrealized(adaptor.getB()); - Value bScale = peelUnrealized(adaptor.getBScale()); + Value bScale = adaptor.getBScale() ? peelUnrealized(adaptor.getBScale()) : Value(); Value dst = peelUnrealized(adaptor.getDst()); - replaceOrEraseWithOpaqueCall(op.getOperation(), "TMATMUL_MX", - {dst, cIn, a, aScale, b, bScale}, rewriter); + SmallVector operands{dst, cIn, a}; + if (aScale) + operands.push_back(aScale); + operands.push_back(b); + if (bScale) + operands.push_back(bScale); + replaceOrEraseWithOpaqueCall(op.getOperation(), "TMATMUL_MX_ACC", operands, + rewriter); return success(); } }; @@ -9389,14 +9468,21 @@ struct PTOTMatmulMXBiasToTMATMUL_MX_BIAS LogicalResult matchAndRewrite(pto::TMatmulMxBiasOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override { Value a = peelUnrealized(adaptor.getA()); - Value aScale = peelUnrealized(adaptor.getAScale()); + Value aScale = adaptor.getAScale() ? peelUnrealized(adaptor.getAScale()) : Value(); Value b = peelUnrealized(adaptor.getB()); - Value bScale = peelUnrealized(adaptor.getBScale()); + Value bScale = adaptor.getBScale() ? peelUnrealized(adaptor.getBScale()) : Value(); Value bias = peelUnrealized(adaptor.getBias()); Value dst = peelUnrealized(adaptor.getDst()); - replaceOrEraseWithOpaqueCall(op.getOperation(), "TMATMUL_MX", - {dst, a, aScale, b, bScale, bias}, rewriter); + SmallVector operands{dst, a}; + if (aScale) + operands.push_back(aScale); + operands.push_back(b); + if (bScale) + operands.push_back(bScale); + operands.push_back(bias); + replaceOrEraseWithOpaqueCall(op.getOperation(), "TMATMUL_MX_BIAS", operands, + rewriter); return success(); } }; @@ -10345,11 +10431,7 @@ struct PTOBindTileToEmitC : public OpConversionPattern { if (rows == ShapedType::kDynamic || cols == ShapedType::kDynamic) return failure(); - std::string blTok = "BLayout::RowMajor"; - if (auto blAttr = dyn_cast(configAttr.getBLayout())) { - if (static_cast(blAttr.getValue()) == 1) - blTok = "BLayout::ColMajor"; - } + std::string blTok = tileBufBLayoutToken(configAttr); pto::BLayout blayout = getTileBufBLayoutValue(configAttr); if (isSubView) { @@ -10767,7 +10849,8 @@ struct PTOAllocTileToEmitC } Value addr = adaptor.getAddr(); - if (addr) { + bool isLinx = getPTOParserTargetArch(ctx) == PTOParserTargetArch::Linx; + if (addr && !isLinx) { addr = peelUnrealized(addr); auto u64Ty = emitc::OpaqueType::get(ctx, "uint64_t"); if (isa(addr.getType()) || @@ -11987,8 +12070,8 @@ static void populatePTOToEmitCPatterns(RewritePatternSet &patterns, PTOTMatmulMXBiasToTMATMUL_MX_BIAS, PTOTGemvBiasToTGEMV_BIAS, PTOTGemvMXToTGEMV_MX, - PTOTGemvMXAccToTGEMV_MX, - PTOTGemvMXBiasToTGEMV_MX, + PTOTGemvMXAccToTGEMV_MX_ACC, + PTOTGemvMXBiasToTGEMV_MX_BIAS, PTOBarrierToEmitC >(typeConverter, ctx); @@ -12025,6 +12108,12 @@ struct EmitPTOManualPass LLVM_DEBUG(llvm::dbgs() << "DEBUG: Start PTOToEmitC Pass\n"); MLIRContext *ctx = &getContext(); ModuleOp mop = getOperation(); + PTOParserTargetArch parserArch = + targetArch == PTOArch::Linx + ? PTOParserTargetArch::Linx + : targetArch == PTOArch::A5 ? PTOParserTargetArch::A5 + : PTOParserTargetArch::A3; + ScopedPTOParserTargetArch parserArchScope(ctx, parserArch); if (failed(pto::validatePTOEntryFunctions(mop))) return signalPassFailure(); diff --git a/lib/PTO/Transforms/PTOViewToMemref.cpp b/lib/PTO/Transforms/PTOViewToMemref.cpp index 4db2aa6246..a5de57bfb4 100644 --- a/lib/PTO/Transforms/PTOViewToMemref.cpp +++ b/lib/PTO/Transforms/PTOViewToMemref.cpp @@ -1777,7 +1777,7 @@ struct PTOViewToMemrefPass rewriter.setInsertionPoint(op); rewriter.replaceOpWithNewOp( op, TypeRange{}, - op->getOperand(0), op->getOperand(1), op->getOperand(2), op->getOperand(3), op->getOperand(4)); + op.getA(), op.getAScale(), op.getB(), op.getBScale(), op.getDst()); } // --- TMatmulMxAccOp --- @@ -1788,7 +1788,8 @@ struct PTOViewToMemrefPass rewriter.setInsertionPoint(op); rewriter.replaceOpWithNewOp( op, TypeRange{}, - op->getOperand(0), op->getOperand(1), op->getOperand(2), op->getOperand(3), op->getOperand(4), op->getOperand(5)); + op.getCIn(), op.getA(), op.getAScale(), op.getB(), op.getBScale(), + op.getDst()); } // --- TMatmulMxBiasOp --- @@ -1799,7 +1800,8 @@ struct PTOViewToMemrefPass rewriter.setInsertionPoint(op); rewriter.replaceOpWithNewOp( op, TypeRange{}, - op->getOperand(0), op->getOperand(1), op->getOperand(2), op->getOperand(3), op->getOperand(4), op->getOperand(5)); + op.getA(), op.getAScale(), op.getB(), op.getBScale(), op.getBias(), + op.getDst()); } // --- TGemvOp [Lhs, Rhs, Dst] --- @@ -1847,7 +1849,7 @@ struct PTOViewToMemrefPass rewriter.setInsertionPoint(op); rewriter.replaceOpWithNewOp( op, TypeRange{}, - op->getOperand(0), op->getOperand(1), op->getOperand(2), op->getOperand(3), op->getOperand(4)); + op.getA(), op.getAScale(), op.getB(), op.getBScale(), op.getDst()); } // --- TGemvMxAccOp [CIn, A, AScale, B, BScale, Dst] --- @@ -1858,8 +1860,8 @@ struct PTOViewToMemrefPass rewriter.setInsertionPoint(op); rewriter.replaceOpWithNewOp( op, TypeRange{}, - op->getOperand(0), op->getOperand(1), op->getOperand(2), - op->getOperand(3), op->getOperand(4), op->getOperand(5)); + op.getCIn(), op.getA(), op.getAScale(), op.getB(), op.getBScale(), + op.getDst()); } // --- TGemvMxBiasOp [A, AScale, B, BScale, Bias, Dst] --- @@ -1870,8 +1872,8 @@ struct PTOViewToMemrefPass rewriter.setInsertionPoint(op); rewriter.replaceOpWithNewOp( op, TypeRange{}, - op->getOperand(0), op->getOperand(1), op->getOperand(2), - op->getOperand(3), op->getOperand(4), op->getOperand(5)); + op.getA(), op.getAScale(), op.getB(), op.getBScale(), op.getBias(), + op.getDst()); } // --- TMovOp [Src, Dst] --- diff --git a/test/lit/pto/tgemv_mx_emitc.pto b/test/lit/pto/tgemv_mx_emitc.pto index 96293f6d6c..1a52b1c157 100644 --- a/test/lit/pto/tgemv_mx_emitc.pto +++ b/test/lit/pto/tgemv_mx_emitc.pto @@ -7,7 +7,7 @@ module { %b_scale = pto.alloc_tile : !pto.tile_buf %dst = pto.alloc_tile : !pto.tile_buf - pto.tgemv.mx ins(%a, %a_scale, %b, %b_scale : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst : !pto.tile_buf) + pto.tgemv.mx ins(%a, %b : !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst : !pto.tile_buf) return } diff --git a/test/lit/pto/tgemv_mx_variants_emitc.pto b/test/lit/pto/tgemv_mx_variants_emitc.pto index e5b3e88295..85f9dc0075 100644 --- a/test/lit/pto/tgemv_mx_variants_emitc.pto +++ b/test/lit/pto/tgemv_mx_variants_emitc.pto @@ -10,8 +10,8 @@ module { %dst0 = pto.alloc_tile : !pto.tile_buf %dst1 = pto.alloc_tile : !pto.tile_buf - pto.tgemv.mx.acc ins(%c_in, %a, %a_scale, %b, %b_scale : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst0 : !pto.tile_buf) - pto.tgemv.mx.bias ins(%a, %a_scale, %b, %b_scale, %bias : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst1 : !pto.tile_buf) + pto.tgemv.mx.acc ins(%c_in, %a, %b : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst0 : !pto.tile_buf) + pto.tgemv.mx.bias ins(%a, %b, %bias : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst1 : !pto.tile_buf) return } diff --git a/test/lit/pto/tmatmul_mx_emitc.pto b/test/lit/pto/tmatmul_mx_emitc.pto index 1e6eb05430..c3a2b593dc 100644 --- a/test/lit/pto/tmatmul_mx_emitc.pto +++ b/test/lit/pto/tmatmul_mx_emitc.pto @@ -14,7 +14,7 @@ module { %a_scale = pto.alloc_tile : !pto.tile_buf %b_scale = pto.alloc_tile : !pto.tile_buf %dst = pto.alloc_tile : !pto.tile_buf - pto.tmatmul.mx ins(%a, %a_scale, %b, %b_scale : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst : !pto.tile_buf) + pto.tmatmul.mx ins(%a, %b : !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst : !pto.tile_buf) return } @@ -24,7 +24,7 @@ module { %a_scale = pto.alloc_tile : !pto.tile_buf %b_scale = pto.alloc_tile : !pto.tile_buf %dst = pto.alloc_tile : !pto.tile_buf - pto.tmatmul.mx.acc ins(%c_in, %a, %a_scale, %b, %b_scale : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst : !pto.tile_buf) + pto.tmatmul.mx.acc ins(%c_in, %a, %b : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst : !pto.tile_buf) return } @@ -34,9 +34,9 @@ module { %b_scale = pto.alloc_tile : !pto.tile_buf %bias = pto.alloc_tile : !pto.tile_buf %dst = pto.alloc_tile : !pto.tile_buf - pto.tmatmul.mx.bias ins(%a, %a_scale, %b, %b_scale, %bias : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst : !pto.tile_buf) + pto.tmatmul.mx.bias ins(%a, %b, %bias : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst : !pto.tile_buf) return } } -// A5: error: 'pto.tmatmul.mx.bias' op lhs: dtype 'f8E4M3FN' is not supported by this op yet +// A5: error: 'pto.alloc_tile' op result dtype 'f8E4M3FN' is not supported by pto.alloc_tile yet diff --git a/test/lit/pto/v0581_linx_tload_decoded_contract.pto b/test/lit/pto/v0581_linx_tload_decoded_contract.pto index 1f9b56d816..6bd9372107 100644 --- a/test/lit/pto/v0581_linx_tload_decoded_contract.pto +++ b/test/lit/pto/v0581_linx_tload_decoded_contract.pto @@ -9,4 +9,4 @@ module attributes {pto.target_arch = "linx"} { } } -// CHECK: error: 'pto.tload' op expects Linx tload dst to use loc=vec +// CHECK: error: 'pto.tload' op expects Linx tload dst to use loc=vec/left/right/acc/bias/scaling diff --git a/test/lit/pto/v0581_linx_contract.pto b/test/lit/pto/v0583_linx_contract.pto similarity index 97% rename from test/lit/pto/v0581_linx_contract.pto rename to test/lit/pto/v0583_linx_contract.pto index 3e450af80f..8fc2e13be4 100644 --- a/test/lit/pto/v0581_linx_contract.pto +++ b/test/lit/pto/v0583_linx_contract.pto @@ -9,7 +9,7 @@ // RUN: ptoas --pto-arch=linx %s | FileCheck %s module attributes {pto.target_arch = "linx"} { - func.func @v0581_contract() attributes {pto.kernel_kind = #pto.kernel_kind} { + func.func @v0583_contract() attributes {pto.kernel_kind = #pto.kernel_kind} { %address = arith.constant 4096 : i64 %row_stride = arith.constant 64 : index %valid_cols = arith.constant 32 : index diff --git a/test/lit/pto/v0583_linx_cube_contract.pto b/test/lit/pto/v0583_linx_cube_contract.pto new file mode 100644 index 0000000000..3347715401 --- /dev/null +++ b/test/lit/pto/v0583_linx_cube_contract.pto @@ -0,0 +1,21 @@ +// RUN: ptoas --pto-arch=linx %s | FileCheck %s + +module attributes {pto.target_arch = "linx"} { + func.func @matmul_and_gemv() attributes {pto.kernel_kind = #pto.kernel_kind} { + %a = pto.alloc_tile : !pto.tile_buf + %v = pto.alloc_tile : !pto.tile_buf + %b = pto.alloc_tile : !pto.tile_buf + %d = pto.alloc_tile : !pto.tile_buf + %g = pto.alloc_tile : !pto.tile_buf + + pto.tmatmul ins(%a, %b : !pto.tile_buf, !pto.tile_buf) outs(%d : !pto.tile_buf) + pto.tgemv ins(%v, %b : !pto.tile_buf, !pto.tile_buf) outs(%g : !pto.tile_buf) + return + } +} + +// CHECK: Tile&1 | FileCheck %s + +module attributes {pto.target_arch = "linx"} { + func.func @reject_row_major_a() attributes {pto.kernel_kind = #pto.kernel_kind} { + %a = pto.alloc_tile : !pto.tile_buf + %b = pto.alloc_tile : !pto.tile_buf + %d = pto.alloc_tile : !pto.tile_buf + pto.tmatmul ins(%a, %b : !pto.tile_buf, !pto.tile_buf) outs(%d : !pto.tile_buf) + return + } +} + +// CHECK: error: 'pto.tmatmul' op expects Linx CUBE lhs to use its architectural CUBE CELL blayout diff --git a/test/lit/pto/v0583_linx_mx_scale_schema_negative.pto b/test/lit/pto/v0583_linx_mx_scale_schema_negative.pto new file mode 100644 index 0000000000..f4101fa628 --- /dev/null +++ b/test/lit/pto/v0583_linx_mx_scale_schema_negative.pto @@ -0,0 +1,60 @@ +// RUN: awk '/^\/\/--- missing-a/{f=1;next}/^\/\/---/{f=0}f' %s > %t.missing-a.pto +// RUN: awk '/^\/\/--- missing-b/{f=1;next}/^\/\/---/{f=0}f' %s > %t.missing-b.pto +// RUN: awk '/^\/\/--- wrong-side/{f=1;next}/^\/\/---/{f=0}f' %s > %t.wrong-side.pto +// RUN: awk '/^\/\/--- wrong-type/{f=1;next}/^\/\/---/{f=0}f' %s > %t.wrong-type.pto +// RUN: not ptoas --pto-arch=linx %t.missing-a.pto 2>&1 | FileCheck %s --check-prefix=MISSING-A +// RUN: not ptoas --pto-arch=linx %t.missing-b.pto 2>&1 | FileCheck %s --check-prefix=MISSING-B +// RUN: not ptoas --pto-arch=linx %t.wrong-side.pto 2>&1 | FileCheck %s --check-prefix=WRONG-SIDE +// RUN: not ptoas --pto-arch=linx %t.wrong-type.pto 2>&1 | FileCheck %s --check-prefix=WRONG-TYPE + +//--- missing-a.pto +module attributes {pto.target_arch = "linx"} { + func.func @missing_a() attributes {pto.kernel_kind = #pto.kernel_kind} { + %a = pto.alloc_tile : !pto.tile_buf + %b = pto.alloc_tile : !pto.tile_buf + %d = pto.alloc_tile : !pto.tile_buf + pto.tmatmul.mx ins(%a, %b : !pto.tile_buf, !pto.tile_buf) outs(%d : !pto.tile_buf) + return + } +} +// MISSING-A: error: 'pto.tmatmul.mx' op expects Linx MX a_scale because a uses a compact dtype + +//--- missing-b.pto +module attributes {pto.target_arch = "linx"} { + func.func @missing_b() attributes {pto.kernel_kind = #pto.kernel_kind} { + %c = pto.alloc_tile : !pto.tile_buf + %a = pto.alloc_tile : !pto.tile_buf + %b = pto.alloc_tile : !pto.tile_buf + %d = pto.alloc_tile : !pto.tile_buf + pto.tmatmul.mx.acc ins(%c, %a, %b : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%d : !pto.tile_buf) + return + } +} +// MISSING-B: error: 'pto.tmatmul.mx.acc' op expects Linx MX b_scale because b uses a compact dtype + +//--- wrong-side.pto +module attributes {pto.target_arch = "linx"} { + func.func @wrong_side() attributes {pto.kernel_kind = #pto.kernel_kind} { + %a = pto.alloc_tile : !pto.tile_buf + %sa = pto.alloc_tile : !pto.tile_buf + %b = pto.alloc_tile : !pto.tile_buf + %bias = pto.alloc_tile : !pto.tile_buf + %d = pto.alloc_tile : !pto.tile_buf + pto.tgemv.mx.bias ins(%a, %b, %bias : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%sa : !pto.tile_buf) outs(%d : !pto.tile_buf) + return + } +} +// WRONG-SIDE: error: 'pto.tgemv.mx.bias' op expects Linx MX a_scale to be absent for FP16/BF16 a + +//--- wrong-type.pto +module attributes {pto.target_arch = "linx"} { + func.func @wrong_type() attributes {pto.kernel_kind = #pto.kernel_kind} { + %a = pto.alloc_tile : !pto.tile_buf + %sa = pto.alloc_tile : !pto.tile_buf + %b = pto.alloc_tile : !pto.tile_buf + %d = pto.alloc_tile : !pto.tile_buf + pto.tmatmul.mx ins(%a, %b : !pto.tile_buf, !pto.tile_buf) a_scale(%sa : !pto.tile_buf) outs(%d : !pto.tile_buf) + return + } +} +// WRONG-TYPE: error: 'pto.tmatmul.mx' op expects Linx MX a_scale to use !pto.f8E8M0 elements diff --git a/test/lit/pto/v0583_linx_tgemv_mx_optional_scales.pto b/test/lit/pto/v0583_linx_tgemv_mx_optional_scales.pto new file mode 100644 index 0000000000..9e0059f02d --- /dev/null +++ b/test/lit/pto/v0583_linx_tgemv_mx_optional_scales.pto @@ -0,0 +1,76 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// RUN: ptoas --pto-arch=linx %s | FileCheck %s + +!a0 = !pto.tile_buf +!a1 = !pto.tile_buf +!sa = !pto.tile_buf +!b0 = !pto.tile_buf +!b1 = !pto.tile_buf +!sb = !pto.tile_buf +!d = !pto.tile_buf +!bias = !pto.tile_buf + +module attributes {pto.target_arch = "linx"} { + func.func @linx_gemv_mx_optional_scales( + %gm_a0: memref<1x32xf16, #pto.address_space>, + %gm_a1: memref<1x32xf8E4M3FN, #pto.address_space>, + %gm_sa: memref<4x32x!pto.f8E8M0, #pto.address_space>, + %gm_b0: memref<32x16xbf16, #pto.address_space>, + %gm_b1: memref<32x16xf8E4M3FN, #pto.address_space>, + %gm_sb: memref<4x32x!pto.f8E8M0, #pto.address_space>, + %gm_c: memref<1x16xf32, #pto.address_space>, + %gm_bias: memref<8x16xf32, #pto.address_space>) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %a0 = pto.alloc_tile : !a0 + %a1 = pto.alloc_tile : !a1 + %sa = pto.alloc_tile : !sa + %b0 = pto.alloc_tile : !b0 + %b1 = pto.alloc_tile : !b1 + %sb = pto.alloc_tile : !sb + %c = pto.alloc_tile : !d + %bias = pto.alloc_tile : !bias + %d = pto.alloc_tile : !d + pto.tload ins(%gm_a0 : memref<1x32xf16, #pto.address_space>) outs(%a0 : !a0) + pto.tload ins(%gm_a1 : memref<1x32xf8E4M3FN, #pto.address_space>) outs(%a1 : !a1) + pto.tload ins(%gm_sa : memref<4x32x!pto.f8E8M0, #pto.address_space>) outs(%sa : !sa) + pto.tload ins(%gm_b0 : memref<32x16xbf16, #pto.address_space>) outs(%b0 : !b0) + pto.tload ins(%gm_b1 : memref<32x16xf8E4M3FN, #pto.address_space>) outs(%b1 : !b1) + pto.tload ins(%gm_sb : memref<4x32x!pto.f8E8M0, #pto.address_space>) outs(%sb : !sb) + pto.tload ins(%gm_c : memref<1x16xf32, #pto.address_space>) outs(%c : !d) + pto.tload ins(%gm_bias : memref<8x16xf32, #pto.address_space>) outs(%bias : !bias) + pto.tgemv.mx ins(%a0, %b0 : !a0, !b0) outs(%d : !d) + pto.tgemv.mx ins(%a1, %b0 : !a1, !b0) a_scale(%sa : !sa) outs(%d : !d) + pto.tgemv.mx ins(%a0, %b1 : !a0, !b1) b_scale(%sb : !sb) outs(%d : !d) + pto.tgemv.mx ins(%a1, %b1 : !a1, !b1) a_scale(%sa : !sa) b_scale(%sb : !sb) outs(%d : !d) + pto.tgemv.mx.acc ins(%c, %a0, %b0 : !d, !a0, !b0) outs(%d : !d) + pto.tgemv.mx.acc ins(%c, %a1, %b0 : !d, !a1, !b0) a_scale(%sa : !sa) outs(%d : !d) + pto.tgemv.mx.acc ins(%c, %a0, %b1 : !d, !a0, !b1) b_scale(%sb : !sb) outs(%d : !d) + pto.tgemv.mx.acc ins(%c, %a1, %b1 : !d, !a1, !b1) a_scale(%sa : !sa) b_scale(%sb : !sb) outs(%d : !d) + pto.tgemv.mx.bias ins(%a0, %b0, %bias : !a0, !b0, !bias) outs(%d : !d) + pto.tgemv.mx.bias ins(%a1, %b0, %bias : !a1, !b0, !bias) a_scale(%sa : !sa) outs(%d : !d) + pto.tgemv.mx.bias ins(%a0, %b1, %bias : !a0, !b1, !bias) b_scale(%sb : !sb) outs(%d : !d) + pto.tgemv.mx.bias ins(%a1, %b1, %bias : !a1, !b1, !bias) a_scale(%sa : !sa) b_scale(%sb : !sb) outs(%d : !d) + return + } +} + +// CHECK-NOT: TASSIGN( +// CHECK: TGEMV_MX({{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_ACC({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_ACC({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_ACC({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_ACC({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_BIAS({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_BIAS({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_BIAS({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); +// CHECK: TGEMV_MX_BIAS({{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,]+}}, {{[^,\)]+}}); diff --git a/test/lit/pto/v0583_linx_tmatmul_mx_variants.pto b/test/lit/pto/v0583_linx_tmatmul_mx_variants.pto new file mode 100644 index 0000000000..a50f38f08e --- /dev/null +++ b/test/lit/pto/v0583_linx_tmatmul_mx_variants.pto @@ -0,0 +1,76 @@ +// Copyright (c) 2026 Huawei Technologies Co., Ltd. +// This program is free software, you can redistribute it and/or modify it under the terms and conditions of +// CANN Open Software License Agreement Version 2.0 (the "License"). +// Please refer to the License for details. You may not use this file except in compliance with the License. +// THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +// INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +// See LICENSE in the root of the software repository for the full text of the License. + +// RUN: ptoas --pto-arch=linx %s | FileCheck %s + +!a0 = !pto.tile_buf +!a1 = !pto.tile_buf +!sa = !pto.tile_buf +!b0 = !pto.tile_buf +!b1 = !pto.tile_buf +!sb = !pto.tile_buf +!d = !pto.tile_buf +!bias = !pto.tile_buf +module attributes {pto.target_arch = "linx"} { + func.func @linx_mx_variants( + %gm_a0: memref<16x32xf16, #pto.address_space>, + %gm_a1: memref<16x32xf8E4M3FN, #pto.address_space>, + %gm_sa: memref<16x32x!pto.f8E8M0, #pto.address_space>, + %gm_b0: memref<32x16xbf16, #pto.address_space>, + %gm_b1: memref<32x16xf8E4M3FN, #pto.address_space>, + %gm_sb: memref<4x32x!pto.f8E8M0, #pto.address_space>, + %gm_c: memref<16x16xf32, #pto.address_space>, + %gm_bias: memref<8x16xf32, #pto.address_space>) + attributes {pto.kernel_kind = #pto.kernel_kind} { + %a0 = pto.alloc_tile : !a0 + %a1 = pto.alloc_tile : !a1 + %sa = pto.alloc_tile : !sa + %b0 = pto.alloc_tile : !b0 + %b1 = pto.alloc_tile : !b1 + %sb = pto.alloc_tile : !sb + %c = pto.alloc_tile : !d + %bias = pto.alloc_tile : !bias + %d = pto.alloc_tile : !d + pto.tload ins(%gm_a0 : memref<16x32xf16, #pto.address_space>) outs(%a0 : !a0) + pto.tload ins(%gm_a1 : memref<16x32xf8E4M3FN, #pto.address_space>) outs(%a1 : !a1) + pto.tload ins(%gm_sa : memref<16x32x!pto.f8E8M0, #pto.address_space>) outs(%sa : !sa) + pto.tload ins(%gm_b0 : memref<32x16xbf16, #pto.address_space>) outs(%b0 : !b0) + pto.tload ins(%gm_b1 : memref<32x16xf8E4M3FN, #pto.address_space>) outs(%b1 : !b1) + pto.tload ins(%gm_sb : memref<4x32x!pto.f8E8M0, #pto.address_space>) outs(%sb : !sb) + pto.tload ins(%gm_c : memref<16x16xf32, #pto.address_space>) outs(%c : !d) + pto.tload ins(%gm_bias : memref<8x16xf32, #pto.address_space>) outs(%bias : !bias) + pto.tmatmul.mx ins(%a0, %b0 : !a0, !b0) outs(%d : !d) + pto.tmatmul.mx ins(%a1, %b0 : !a1, !b0) a_scale(%sa : !sa) outs(%d : !d) + pto.tmatmul.mx ins(%a0, %b1 : !a0, !b1) b_scale(%sb : !sb) outs(%d : !d) + pto.tmatmul.mx ins(%a1, %b1 : !a1, !b1) a_scale(%sa : !sa) b_scale(%sb : !sb) outs(%d : !d) + pto.tmatmul.mx.acc ins(%c, %a0, %b0 : !d, !a0, !b0) outs(%d : !d) + pto.tmatmul.mx.acc ins(%c, %a1, %b0 : !d, !a1, !b0) a_scale(%sa : !sa) outs(%d : !d) + pto.tmatmul.mx.acc ins(%c, %a0, %b1 : !d, !a0, !b1) b_scale(%sb : !sb) outs(%d : !d) + pto.tmatmul.mx.acc ins(%c, %a1, %b1 : !d, !a1, !b1) a_scale(%sa : !sa) b_scale(%sb : !sb) outs(%d : !d) + pto.tmatmul.mx.bias ins(%a0, %b0, %bias : !a0, !b0, !bias) outs(%d : !d) + pto.tmatmul.mx.bias ins(%a1, %b0, %bias : !a1, !b0, !bias) a_scale(%sa : !sa) outs(%d : !d) + pto.tmatmul.mx.bias ins(%a0, %b1, %bias : !a0, !b1, !bias) b_scale(%sb : !sb) outs(%d : !d) + pto.tmatmul.mx.bias ins(%a1, %b1, %bias : !a1, !b1, !bias) a_scale(%sa : !sa) b_scale(%sb : !sb) outs(%d : !d) + return + } +} + +// CHECK-COUNT-2: Tile&2 + exit 1 + fi +done +if [[ ! -x "${PTOAS_BIN}" ]]; then + echo "PTOAS_BIN must be executable" >&2 + exit 1 +fi + +LLVM_SOURCE_DIR=$(sed -n 's/^LLVM_SOURCE_DIR:STATIC=//p' "${LLVM_CACHE}") +if [[ -z "${LLVM_SOURCE_DIR}" ]] || + ! LLVM_REPO=$(git -C "${LLVM_SOURCE_DIR}" rev-parse --show-toplevel 2>/dev/null); then + echo "LINX_LLVM_BUILD does not identify its LLVM source checkout" >&2 + exit 1 +fi +ACTUAL_LLVM_COMMIT=$(git -C "${LLVM_REPO}" rev-parse HEAD) +ACTUAL_LLVM_TREE=$(git -C "${LLVM_REPO}" rev-parse 'HEAD^{tree}') +if [[ "${ACTUAL_LLVM_COMMIT}" != "${EXPECTED_LLVM_COMMIT}" || + "${ACTUAL_LLVM_TREE}" != "${EXPECTED_LLVM_TREE}" ]]; then + echo "Linx LLVM identity mismatch: expected ${EXPECTED_LLVM_COMMIT}/${EXPECTED_LLVM_TREE}, got ${ACTUAL_LLVM_COMMIT}/${ACTUAL_LLVM_TREE}" >&2 + exit 1 +fi + +cmake --build "${LINX_LLVM_BUILD}" --target clang \ + --parallel "${PTOAS_LINX_BUILD_JOBS:-2}" +if [[ ! -x "${LINX_CXX}" ]]; then + echo "exact Linx clang++ was not built at ${LINX_CXX}" >&2 + exit 1 +fi + +ACTUAL_TILEOP_COMMIT=$(git -C "${TILEOP_ROOT}" rev-parse HEAD) +ACTUAL_TILEOP_TREE=$(git -C "${TILEOP_ROOT}" rev-parse 'HEAD^{tree}') +if [[ "${ACTUAL_TILEOP_COMMIT}" != "${EXPECTED_TILEOP_COMMIT}" || + "${ACTUAL_TILEOP_TREE}" != "${EXPECTED_TILEOP_TREE}" ]]; then + echo "TileOP identity mismatch: expected ${EXPECTED_TILEOP_COMMIT}/${EXPECTED_TILEOP_TREE}, got ${ACTUAL_TILEOP_COMMIT}/${ACTUAL_TILEOP_TREE}" >&2 + exit 1 +fi + +TMP_DIR=$(mktemp -d "${TMPDIR:-/tmp}/ptoas-linx-mx-tileop.XXXXXX") +trap 'rm -rf "${TMP_DIR}"' EXIT + +"${PTOAS_BIN}" --pto-arch=linx \ + "${SOURCE_ROOT}/test/lit/pto/v0583_linx_tmatmul_mx_variants.pto" \ + >"${TMP_DIR}/tmatmul.cpp" +"${PTOAS_BIN}" --pto-arch=linx \ + "${SOURCE_ROOT}/test/lit/pto/v0583_linx_tgemv_mx_optional_scales.pto" \ + >"${TMP_DIR}/tgemv.cpp" + +for callee in TMATMUL_MX TMATMUL_MX_ACC TMATMUL_MX_BIAS; do + [[ $(grep -c " ${callee}(" "${TMP_DIR}/tmatmul.cpp") -eq 4 ]] +done +for callee in TGEMV_MX TGEMV_MX_ACC TGEMV_MX_BIAS; do + [[ $(grep -c " ${callee}(" "${TMP_DIR}/tgemv.cpp") -eq 4 ]] +done + +FLAGS=( + --target=linx64-unknown-linux-musl + --sysroot="${LINX_SYSROOT}" + -nostdinc++ + -isystem "${LINX_SYSROOT}/include/c++/v1" + -fenable-matrix + -O2 + -std=c++20 + -D__linx + -DENABLE_TENSOR_INSTR + -Werror + -I "${TILEOP_ROOT}/include" +) + +for generated in "${TMP_DIR}/tmatmul.cpp" "${TMP_DIR}/tgemv.cpp"; do + "${LINX_CXX}" "${FLAGS[@]}" -fsyntax-only "${generated}" + object=${generated%.cpp}.o + "${LINX_CXX}" "${FLAGS[@]}" -c "${generated}" -o "${object}" + if [[ ! -s "${object}" ]]; then + echo "Linx object gate produced no object: ${object}" >&2 + exit 1 + fi +done + +TC_DIR=${LINX_LLVM_BUILD}/bin LINX_SYSROOT=${LINX_SYSROOT} \ + bash "${TILEOP_ROOT}/test/tileop_api/verify_target_cxx_frontend.sh" + +( + cd "${TILEOP_ROOT}/test/tileop_api" + COMPILER_DIR=${LINX_LLVM_BUILD}/bin \ + LINX_SYSROOT=${LINX_SYSROOT} \ + LINX_TARGET=linx64-unknown-linux-musl \ + bash ./compile.all link-smoke +) diff --git a/tools/check_python_binding_link_contract.py b/tools/check_python_binding_link_contract.py index e77e10f415..d2e7e3d095 100644 --- a/tools/check_python_binding_link_contract.py +++ b/tools/check_python_binding_link_contract.py @@ -70,26 +70,29 @@ def main() -> int: contract = re.compile( r"if\s*\(UNIX\s+AND\s+NOT\s+APPLE\).*?" - r"target_link_options\s*\(nanobind-mlir\s+PRIVATE\s+" + r"set\s*\(PTOAS_NANOBIND_RUNTIME_TARGET\s+\"\"\s*\).*?" + r"nanobind-mlir.*?nanobind.*?" + r"target_link_options\s*\(\$\{PTOAS_NANOBIND_RUNTIME_TARGET\}\s+PRIVATE\s+" r'"LINKER:-z,undefs"\s*\).*?endif\s*\(\)', re.DOTALL, ) if not contract.search(source): raise SystemExit( - "error: the ELF nanobind-mlir target must override -z defs with " - "LINKER:-z,undefs so CPython resolves its ABI symbols at module load" + "error: the ELF nanobind shared runtime resolver must support " + "nanobind 2.x/3.x target names and override -z defs with " + "LINKER:-z,undefs" ) install_contract = re.compile( - r"install\s*\(\s*TARGETS\s+nanobind-mlir\s+" + r"install\s*\(\s*TARGETS\s+\$\{PTOAS_NANOBIND_RUNTIME_TARGET\}\s+" r"LIBRARY\s+DESTINATION\s+lib\s+" r"(?:COMPONENT\s+PTOASPythonRuntime\s+)?\)", re.DOTALL, ) if not install_contract.search(source): raise SystemExit( - "error: install nanobind-mlir into lib so auditwheel can locate " - "the NB_SHARED runtime" + "error: install the resolved nanobind shared runtime into lib so " + "auditwheel can locate it" ) workflow_path = args.ptoas_root / ".github" / "workflows" / "build_wheel.yml" diff --git a/tools/check_release_delivery_contract.py b/tools/check_release_delivery_contract.py index 2b503a15fa..b08192927f 100644 --- a/tools/check_release_delivery_contract.py +++ b/tools/check_release_delivery_contract.py @@ -18,6 +18,12 @@ from pathlib import Path +EXPECTED_LLVM_COMMIT = "b7c83f68bf84125e696a70bec4b665c70a3b584d" +EXPECTED_LLVM_TREE = "c11bd80c7dd34ed4438de1da6bde0a01eae8c76d" +EXPECTED_TILEOP_COMMIT = "bd1ecca97ca47da0edc462c1ce19749c6940780e" +EXPECTED_TILEOP_TREE = "854c463c21f9b758186b718105f4b035019cdd30" + + def local_copy_sources(dockerfile: str) -> list[tuple[int, str]]: """Return local Docker COPY sources and reject unvalidated syntax.""" sources: list[tuple[int, str]] = [] @@ -90,8 +96,8 @@ def validate_packaging_identity_mode(script: Path) -> None: """Exercise a packaging script with valid and adversarial identity outputs.""" for output, product_version in ( - ("ptoas 0.41 (PTO ISA 0.58.1)", "0.41"), - ("ptoas 0.41 (PTO ISA 0.58.1)", ""), + ("ptoas 0.41 (PTO ISA 0.58.3)", "0.41"), + ("ptoas 0.41 (PTO ISA 0.58.3)", ""), ): result = run_packaging_identity_mode(script, output, product_version) if result.returncode != 0: @@ -101,10 +107,10 @@ def validate_packaging_identity_mode(script: Path) -> None: invalid_outputs = ( "ptoas 0.41", - "ptoas 0.40 (PTO ISA 0.58.1)", + "ptoas 0.40 (PTO ISA 0.58.3)", "ptoas 0.41 (PTO ISA 0.58.0)", - "warning\nptoas 0.41 (PTO ISA 0.58.1)", - "ptoas 0.41 (PTO ISA 0.58.1)\nptoas 0.40", + "warning\nptoas 0.41 (PTO ISA 0.58.3)", + "ptoas 0.41 (PTO ISA 0.58.3)\nptoas 0.40", ) for product_version in ("0.41", ""): for output in invalid_outputs: @@ -132,6 +138,53 @@ def main() -> int: raise SystemExit("Docker build must COPY the reviewed PTOAS checkout") if "git clone https://github.com/zhangstevenunity/PTOAS.git" in dockerfile: raise SystemExit("Docker build must not clone an unpinned PTOAS fork") + if "'nanobind>=2.9,<3'" not in dockerfile: + raise SystemExit( + "Docker LLVM source build must pin nanobind to the MLIR 23 " + "compatible >=2.9,<3 range" + ) + + llvm_pin_paths = ( + root / "README.md", + root / "docker/Dockerfile", + root / ".github/workflows/ci.yml", + root / ".github/workflows/build_wheel.yml", + root / ".github/workflows/build_wheel_mac.yml", + ) + for path in llvm_pin_paths: + text = path.read_text() + if EXPECTED_LLVM_COMMIT not in text: + raise SystemExit( + f"{path.relative_to(root)} does not pin reviewed merged LLVM " + f"commit {EXPECTED_LLVM_COMMIT}" + ) + + integration_pin_text = (root / "CMakeLists.txt").read_text() + ( + root / "README.md" + ).read_text() + for identity in ( + EXPECTED_LLVM_COMMIT, + EXPECTED_LLVM_TREE, + EXPECTED_TILEOP_COMMIT, + EXPECTED_TILEOP_TREE, + ): + if identity not in integration_pin_text: + raise SystemExit( + f"Linx target integration gate does not pin reviewed identity {identity}" + ) + + nanobind_pin_paths = ( + root / "docker/Dockerfile", + root / ".github/workflows/ci.yml", + root / ".github/workflows/build_wheel.yml", + root / ".github/workflows/build_wheel_mac.yml", + ) + for path in nanobind_pin_paths: + if "nanobind>=2.9,<3" not in path.read_text(): + raise SystemExit( + f"{path.relative_to(root)} must pin nanobind >=2.9,<3 for " + "the LLVM 23 MLIR Python source build" + ) if ( 'test "$(git -C pto-isa rev-parse HEAD)" = "${PTO_ISA_COMMIT}"' not in dockerfile @@ -143,8 +196,8 @@ def main() -> int: raise SystemExit("Docker README must document repository-root context") workflow = (root / ".github/workflows/isa_contract.yml").read_text() - if "linxisa-v0.58.1" not in workflow: - raise SystemExit("raw linxisa-v0.58.1 tag is not an ISA checker trigger") + if "linxisa-v0.58.3" not in workflow: + raise SystemExit("raw linxisa-v0.58.3 tag is not an ISA checker trigger") exact_command = "python3 tools/check_v058_pto_manifest.py --ptoas-root ." if exact_command not in workflow: raise SystemExit("ISA tag workflow does not run the exact contract checker") @@ -161,7 +214,7 @@ def main() -> int: for name in ("build_wheel.yml", "build_wheel_mac.yml"): text = (root / ".github/workflows" / name).read_text() - if 'GITHUB_REF_NAME}" = "linxisa-v0.58.1"' in text: + if 'GITHUB_REF_NAME}" = "linxisa-v0.58.3"' in text: raise SystemExit( f"{name} conflates ISA identity with PTOAS product version" ) diff --git a/tools/check_v058_pto_manifest.py b/tools/check_v058_pto_manifest.py index e25139de9a..23443c6509 100644 --- a/tools/check_v058_pto_manifest.py +++ b/tools/check_v058_pto_manifest.py @@ -7,7 +7,7 @@ # INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. # See LICENSE in the root of the software repository for the full text of the License. -"""Check PTOAS PTO-op contracts against the LinxISA v0.58.1 PTO manifest. +"""Check PTOAS PTO-op contracts against the LinxISA v0.58.3 PTO manifest. PTOAS is an MLIR PTO dialect-to-EmitC compiler, not a Linx scalar assembler. This check validates all 109 public operations' exact Linx operand roles/arity, @@ -26,30 +26,30 @@ EXPECTED_LOCK = { - "release": "0.58.1", - "encoding_abi": "pto-isa-0.58.1-mode-function-v1", - "encoding_projection_sha256": "89b872d6eaf0252200bc9349d49b9346e2a69d894cdcc2dcd0fd71911c1e0b8c", - "content_sha256": "693e8c0734b48598ac35ffe7fe6f2a01037788fba30ebe026895808d23139f2c", - "release_manifest_sha256": "acea87af67301173e6d1c6e04014a8dc6e2f658cd2992ed658ab2589cddc7841", - "source_commit": "c381465b2b8e457e162a4246ee58bb9a2c5b49fd", - "source_tree": "463a19db3d6ba70022f18bdbca0d4b2c6ed586e4", + "release": "0.58.3", + "encoding_abi": "pto-isa-0.58.3-mode-function-v1", + "encoding_projection_sha256": "8a48b80e04484c70870f155bf9efc79d2a805cf99e809f4e4e8a7e6a7eb34172", + "content_sha256": "f299fe3d256c5d071e57bb4aaa2be2de2e4a386ae090048df1f73ae92d392678", + "release_manifest_sha256": "ebe8c75e2f8159634e49a9bbf81aba7141e15b45e63b3a5606e758ed7bc22023", + "source_commit": "e599a3d36ebfad43362ff591ea5e128816c684c7", + "source_tree": "abb6899d2e664e378ac9c1b77062670daa4d31b4", "hardware_profile_path": "spec/hardware-conformance-profile.json", - "hardware_profile_id": "pto-hardware-numeric-0.58.1-ieee-v1", - "hardware_profile_sha256": "170deadacb174c933c287231fb67da1046d7989f84b6852bf353d68a495d1755", - "numeric_vectors_path": "spec/evidence/pto-isa-0581-hardware-numeric-vectors.json", - "numeric_vectors_sha256": "59c96cc2f45f8e8f3eebb8230338b21ec3a77a99e8fb5e1c7c7b391819a6aa81", + "hardware_profile_id": "pto-hardware-numeric-0.58.3-ieee-v1", + "hardware_profile_sha256": "117cd8e61f1e82001755cdb1ee98ef224fbd624ddf2adb8577c32d6acf833575", + "numeric_vectors_path": "spec/evidence/pto-isa-0583-hardware-numeric-vectors.json", + "numeric_vectors_sha256": "09d863e39e5fcd932353f4dc3bc5a7d2eed91ec9818c83886747aa69d28c3890", "command_forms_path": "spec/catalog/command-forms.json", - "command_forms_sha256": "300a3a57a8728e6c4770da6fff0202b372ec2830edb8dc978dc141d1c26424d0", + "command_forms_sha256": "fa3c8a6ca86d0fc273052b77d4f977ca69f3b8da6fe94bf6dd0dad44e0dd01e4", "command_forms_count": 74, "scalar_forms_path": "spec/catalog/scalar-forms.json", - "scalar_forms_sha256": "9f3841d568ffa73fcb43bf4fd365d3c4dba42d27acffa7e273e0f403c0f0c602", - "scalar_forms_count": 474, + "scalar_forms_sha256": "bdfcb4df19da4329c5ff0184b34daebf258d832992fa06fb3e0c34ca891c5923", + "scalar_forms_count": 466, "tile_operations_path": "spec/catalog/tile-operations.json", - "tile_operations_sha256": "f163dea8be281fd67173713d373b60f95a9c3c4e558adcdf8034cc213507a1a3", + "tile_operations_sha256": "07c5cf6f6e59916f3cbbecb0b83fe364704e1a79e1a43fa83f2997b6f7242207", "tile_operation_count": 109, "extension_encoding_reservations_path": "spec/catalog/extension-encoding-reservations.json", - "extension_encoding_reservations_sha256": "bdb82b839b98984779d9a1394f6b308f141052ef0b520e5bedb8e87dadd883d4", - "extension_encoding_reservations_count": 32, + "extension_encoding_reservations_sha256": "f1b424060d3aae9432934724ba66908eab0476dcc4880447eaaca44fd016be8b", + "extension_encoding_reservations_count": 40, "release_manifest_path": "spec/release-manifest.json", "source_repository": "https://github.com/PTO-ISA/pto-spec.git", } @@ -79,6 +79,17 @@ "TSUBSC", } +RETIRED_SCALAR_FORMS = { + "B.EQ", + "B.GE", + "B.GEU", + "B.LT", + "B.LTU", + "B.NE", + "B.NZ", + "B.Z", +} + def normalize(name: str) -> str: return name.replace(".", "").replace("_", "").upper() @@ -88,7 +99,7 @@ def sha256(path: Path) -> str: def load_lock(ptoas_root: Path) -> dict: - lock_path = ptoas_root / "tools/pto_isa_v0_58_1_lock.json" + lock_path = ptoas_root / "tools/pto_isa_v0_58_3_lock.json" lock = json.loads(lock_path.read_text()) errors = [] for key in ("release", "encoding_abi", "encoding_projection_sha256", "content_sha256"): @@ -152,7 +163,7 @@ def load_lock(ptoas_root: Path) -> dict: if entry.get(key) != EXPECTED_LOCK[expected_key]: errors.append(f"catalogs.{catalog}.{key} mismatch") if errors: - raise SystemExit("unexpected PTO ISA 0.58.1 lock:\n " + "\n ".join(errors)) + raise SystemExit("unexpected PTO ISA 0.58.3 lock:\n " + "\n ".join(errors)) return lock @@ -235,15 +246,135 @@ def validate_source_tree(source_root: Path, lock: dict) -> None: ): raise SystemExit("PTO ISA release manifest command form count mismatch") + validate_release_catalog_semantics(source_root) + + +def _field(form: dict, name: str) -> dict: + for field in form.get("fields", []): + if field.get("name") == name: + return field + raise SystemExit( + f"PTO ISA 0.58.3 {form.get('mnemonic', '')} missing {name} field" + ) + + +def _constraint_values(form: dict, name: str) -> list[int]: + for constraint in form.get("constraints", []): + if constraint.get("field") == name and constraint.get("operator") == "one-of": + return constraint.get("values", []) + raise SystemExit( + f"PTO ISA 0.58.3 {form.get('mnemonic', '')} missing {name} constraint" + ) + + +def _require_contiguous_field(form: dict, name: str, lsb: int, width: int) -> None: + field = _field(form, name) + pieces = field.get("pieces", []) + expected = [{"instruction_lsb": lsb, "value_lsb": 0, "width": width}] + if field.get("width") != width or pieces != expected: + raise SystemExit( + f"PTO ISA 0.58.3 {form.get('mnemonic')} {name} encoding mismatch: " + f"expected lsb={lsb}/width={width}, got {field}" + ) + + +def validate_release_catalog_semantics(source_root: Path) -> None: + """Lock the 0.58.3 hard-break details PTOAS downstreams rely on.""" + + catalog_root = source_root / "spec/catalog" + command_forms = json.loads((catalog_root / "command-forms.json").read_text())["forms"] + scalar_forms = json.loads((catalog_root / "scalar-forms.json").read_text())["forms"] + reservations = json.loads( + (catalog_root / "extension-encoding-reservations.json").read_text() + )["reservations"] + tile_operations = json.loads( + (catalog_root / "tile-operations.json").read_text() + )["operations"] + + active_scalar_names = {form["mnemonic"] for form in scalar_forms} + retired_active = sorted(RETIRED_SCALAR_FORMS & active_scalar_names) + if retired_active: + raise SystemExit( + "PTO ISA 0.58.3 retired scalar forms are still active: " + + ", ".join(retired_active) + ) + reservation_names = {entry["mnemonic"] for entry in reservations} + missing_reservations = sorted(RETIRED_SCALAR_FORMS - reservation_names) + if missing_reservations: + raise SystemExit( + "PTO ISA 0.58.3 retired scalar reservations are missing: " + + ", ".join(missing_reservations) + ) + + by_mnemonic: dict[str, list[dict]] = {} + for form in command_forms: + by_mnemonic.setdefault(form["mnemonic"], []).append(form) + + iot_forms = by_mnemonic.get("B.IOT", []) + if len(iot_forms) != 5: + raise SystemExit(f"PTO ISA 0.58.3 expected five B.IOT forms, got {len(iot_forms)}") + for form in iot_forms: + _require_contiguous_field(form, "PEMode", 9, 3) + if _constraint_values(form, "PEMode") != list(range(8)): + raise SystemExit("PTO ISA 0.58.3 B.IOT PEMode must accept codes 0..7") + if "SizeCode" in {field["name"] for field in form["fields"]}: + _require_contiguous_field(form, "SizeCode", 15, 4) + _require_contiguous_field(form, "DstTile", 7, 2) + if _constraint_values(form, "SizeCode") != list(range(1, 11)): + raise SystemExit("PTO ISA 0.58.3 B.IOT SizeCode must accept codes 1..10") + + ios_forms = by_mnemonic.get("B.IOS", []) + if len(ios_forms) != 1: + raise SystemExit(f"PTO ISA 0.58.3 expected one B.IOS form, got {len(ios_forms)}") + ios = ios_forms[0] + _require_contiguous_field(ios, "SizeCode", 15, 4) + _require_contiguous_field(ios, "PEMode", 9, 3) + if _constraint_values(ios, "SizeCode") != list(range(13)): + raise SystemExit("PTO ISA 0.58.3 B.IOS SizeCode must accept codes 0..12") + if _constraint_values(ios, "PEMode") != list(range(8)): + raise SystemExit("PTO ISA 0.58.3 B.IOS PEMode must accept codes 0..7") + + fpatr_forms = by_mnemonic.get("B.FPATR", []) + if len(fpatr_forms) != 1: + raise SystemExit( + f"PTO ISA 0.58.3 expected one B.FPATR form, got {len(fpatr_forms)}" + ) + fpatr = fpatr_forms[0] + _require_contiguous_field(fpatr, "TransA", 7, 1) + _require_contiguous_field(fpatr, "TransB", 8, 1) + for name in ("TransA", "TransB"): + if _constraint_values(fpatr, name) != [0, 1]: + raise SystemExit(f"PTO ISA 0.58.3 B.FPATR {name} must be one bit") + + datr_forms = by_mnemonic.get("B.DATR", []) + if len(datr_forms) != 1: + raise SystemExit(f"PTO ISA 0.58.3 expected one B.DATR form, got {len(datr_forms)}") + datr = datr_forms[0] + _require_contiguous_field(datr, "DataType", 20, 5) + if 31 not in _constraint_values(datr, "DataType"): + raise SystemExit("PTO ISA 0.58.3 B.DATR must admit DTYPE_NONE code 31") + + tile_by_name = {operation["name"]: operation for operation in tile_operations} + for name in ("TLOAD", "TSTORE"): + operation = tile_by_name[name] + roles = {operand["field"]: operand["role"] for operand in operation["operands"]} + if roles.get("scalar0") != "row-stride-bytes": + raise SystemExit(f"PTO ISA 0.58.3 {name} scalar0 must be row-stride-bytes") + datr_contract = operation.get("datr_contract", {}) + if datr_contract.get("allowed_nonzero_fields") != ["PadValueOrByteId", "Layout"]: + raise SystemExit(f"PTO ISA 0.58.3 {name} CUBE DATR contract mismatch") + if datr_contract.get("pad_union") != "pad-value": + raise SystemExit(f"PTO ISA 0.58.3 {name} CUBE padding contract mismatch") + def load_manifest(linx_root: Path) -> dict[str, dict]: manifest_path = linx_root / "isa/v0.58/state/pto_ops.json" manifest = json.loads(manifest_path.read_text()) if manifest["profile"] != "v0.58" or manifest["operation_count"] != 109: - raise SystemExit(f"unexpected v0.58.1 manifest header in {manifest_path}") + raise SystemExit(f"unexpected v0.58.3 manifest header in {manifest_path}") source_lock = manifest.get("source_lock") if source_lock != "isa/v0.58/pto-spec.lock.json": - raise SystemExit(f"unexpected v0.58.1 source_lock in {manifest_path}: {source_lock}") + raise SystemExit(f"unexpected v0.58.3 source_lock in {manifest_path}: {source_lock}") operations = manifest["operations"] names = [entry["name"] for entry in operations] if len(names) != len(set(names)): @@ -285,7 +416,7 @@ def validate_linx_identity(linx_root: Path, lock: dict) -> None: def load_expected_contracts(ptoas_root: Path) -> tuple[dict[str, dict], dict[str, dict]]: - contract_path = ptoas_root / "tools/pto_isa_v0_58_1_operation_contracts.json" + contract_path = ptoas_root / "tools/pto_isa_v0_58_3_operation_contracts.json" contract_file = json.loads(contract_path.read_text()) for field in ( "release", @@ -339,6 +470,16 @@ def load_ptoas_ops(ptoas_root: Path) -> dict[str, dict]: body = text[start:end] match = re.search(r"let\s+arguments\s*=\s*\(ins(.*?)\);", body, re.DOTALL) arguments = tuple(re.findall(r"\$([A-Za-z0-9_]+)", match.group(1))) if match else () + optional_arguments = ( + tuple( + re.findall( + r"Optional<[^>]+>\s*:\s*\$([A-Za-z0-9_]+)", + match.group(1), + ) + ) + if match + else () + ) key = normalize(mnemonic) if key in operations: raise SystemExit(f"duplicate normalized PTOAS mnemonic in {ods_path}: {mnemonic}") @@ -349,6 +490,7 @@ def load_ptoas_ops(ptoas_root: Path) -> dict[str, dict]: operations[key] = { "mnemonic": mnemonic, "arguments": arguments, + "optional_arguments": optional_arguments, "linx_engine": linx_engine.group(1) if linx_engine else None, } return operations @@ -367,7 +509,7 @@ def validate_linx_target_surface(ptoas_root: Path) -> None: if token not in cli_text and token not in lowering_text: errors.append(f"missing Linx target implementation token: {token}") contracts = json.loads( - (ptoas_root / "tools/pto_isa_v0_58_1_operation_contracts.json").read_text() + (ptoas_root / "tools/pto_isa_v0_58_3_operation_contracts.json").read_text() ) for mnemonic_name in contracts.get("linx_rejected_mnemonics", []): mnemonic = f'"pto.{mnemonic_name}"' @@ -377,8 +519,54 @@ def validate_linx_target_surface(ptoas_root: Path) -> None: errors.append("GMOV lowering must match Linx-TileOP-API order (dst, peer_tid, src)") if 'ValueRange{dst, src0, src1, src2}' not in lowering_text: errors.append("TFMA lowering must match Linx-TileOP-API order (dst, src0, src1, src2)") + if 'ValueRange{dst, rhs, lhs}' not in lowering_text: + errors.append("TGEMV lowering must match TileOP order (dst, matrix-B, vector-A)") + if 'ValueRange{dst, accIn, rhs, lhs}' not in lowering_text: + errors.append( + "TGEMV_ACC lowering must match TileOP order (dst, acc, matrix-B, vector-A)" + ) + cube_variant_mappings = { + "PTOTGemvMXToTGEMV_MX": ( + "TGEMV_MX", + "SmallVectoroperands{dst,b};if(bScale)operands.push_back(bScale);operands.push_back(a);if(aScale)operands.push_back(aScale);", + ), + "PTOTGemvMXAccToTGEMV_MX_ACC": ( + "TGEMV_MX_ACC", + "SmallVectoroperands{dst,cIn,b};if(bScale)operands.push_back(bScale);operands.push_back(a);if(aScale)operands.push_back(aScale);", + ), + "PTOTGemvMXBiasToTGEMV_MX_BIAS": ( + "TGEMV_MX_BIAS", + "SmallVectoroperands{dst,b};if(bScale)operands.push_back(bScale);operands.push_back(a);if(aScale)operands.push_back(aScale);operands.push_back(bias);", + ), + "PTOTMatmulMXToTMATMUL_MX": ( + "TMATMUL_MX", + "SmallVectoroperands{dst,a};if(aScale)operands.push_back(aScale);operands.push_back(b);if(bScale)operands.push_back(bScale);", + ), + "PTOTMatmulMXAccToTMATMUL_MX_ACC": ( + "TMATMUL_MX_ACC", + "SmallVectoroperands{dst,cIn,a};if(aScale)operands.push_back(aScale);operands.push_back(b);if(bScale)operands.push_back(bScale);", + ), + "PTOTMatmulMXBiasToTMATMUL_MX_BIAS": ( + "TMATMUL_MX_BIAS", + "SmallVectoroperands{dst,a};if(aScale)operands.push_back(aScale);operands.push_back(b);if(bScale)operands.push_back(bScale);operands.push_back(bias);", + ), + } + for class_name, (callee, operands) in cube_variant_mappings.items(): + match = re.search( + rf"struct\s+{class_name}\b(?P.*?)\n\}};", + lowering_text, + re.DOTALL, + ) + if match is None: + errors.append(f"missing Linx CUBE lowering class {class_name}") + continue + compact_body = re.sub(r"\s+", "", match.group("body")) + if f'"{callee}"' not in compact_body or operands not in compact_body: + errors.append( + f"{class_name} must preserve the exact optional-scale order for {callee}" + ) if errors: - raise SystemExit("invalid PTOAS Linx v0.58.1 target surface:\n " + "\n ".join(errors)) + raise SystemExit("invalid PTOAS Linx v0.58.3 target surface:\n " + "\n ".join(errors)) def main() -> int: @@ -407,7 +595,7 @@ def main() -> int: ) boundary_error = bool(deleted_present) if deleted_present: - print("PTOAS has deleted PTO ISA 0.58.1 names active in the dialect:") + print("PTOAS has deleted PTO ISA 0.58.3 names active in the dialect:") for name in deleted_present: print(f" - {name}") @@ -433,6 +621,16 @@ def main() -> int: expected = tuple(contract["ptoas_arguments"]) if actual != expected: contract_errors.append(f"{name}: expected PTOAS arguments {expected}, got {actual}") + if "ptoas_optional_arguments" in contract: + actual_optional = ptoas_ops[normalize(contract["ptoas_mnemonic"])][ + "optional_arguments" + ] + expected_optional = tuple(contract["ptoas_optional_arguments"]) + if actual_optional != expected_optional: + contract_errors.append( + f"{name}: expected optional PTOAS arguments {expected_optional}, " + f"got {actual_optional}" + ) for key, contract in expected_dialect_only.items(): actual = ptoas_ops[key]["arguments"] expected = tuple(contract["ptoas_arguments"]) @@ -442,7 +640,7 @@ def main() -> int: f"expected PTOAS arguments {expected}, got {actual}" ) if contract_errors: - print("PTOAS has incorrect PTO ISA 0.58.1 operation roles/arity:") + print("PTOAS has incorrect PTO ISA 0.58.3 operation roles/arity:") for error in contract_errors: print(f" - {error}") return 1 @@ -463,7 +661,7 @@ def main() -> int: if args.linx_root is None: print( - "PTOAS v0.58.1 PTO lock/dialect check OK: all 109 public " + "PTOAS v0.58.3 PTO lock/dialect check OK: all 109 public " f"operation argument contracts and {len(expected_dialect_only)} explicit " "dialect-only contracts match; hardware numeric profile/vectors are " "identity metadata only (execution conformance not evaluated)" @@ -476,7 +674,7 @@ def main() -> int: expected_names = set(expected_public) actual_names = set(manifest) if actual_names != expected_names: - print("LinxISA v0.58.1 public operation boundary mismatch:") + print("LinxISA v0.58.3 public operation boundary mismatch:") for name in sorted(expected_names - actual_names): print(f" - missing manifest operation: {name}") for name in sorted(actual_names - expected_names): @@ -503,13 +701,13 @@ def main() -> int: f"{name}: expected ISA operands {expected}, got {actual}" ) if role_errors: - print("LinxISA v0.58.1 manifest role/arity mismatch:") + print("LinxISA v0.58.3 manifest role/arity mismatch:") for error in role_errors: print(f" - {error}") return 1 print( - "PTOAS v0.58.1 PTO manifest contract check OK: " + "PTOAS v0.58.3 PTO manifest contract check OK: " f"all {len(manifest)} public operations match exact Linx roles/arity; " f"{len(expected_dialect_only)} dialect-only operations are explicitly bounded; " "hardware numeric profile/vectors are identity metadata only " diff --git a/tools/pto_isa_v0_58_3_lock.json b/tools/pto_isa_v0_58_3_lock.json new file mode 100644 index 0000000000..29902c04be --- /dev/null +++ b/tools/pto_isa_v0_58_3_lock.json @@ -0,0 +1,47 @@ +{ + "$schema": "https://docs.openclaw.ai/schemas/linxisa/pto_spec_lock.v1.json", + "catalogs": { + "command_forms": { + "path": "spec/catalog/command-forms.json", + "sha256": "fa3c8a6ca86d0fc273052b77d4f977ca69f3b8da6fe94bf6dd0dad44e0dd01e4", + "count": 74 + }, + "scalar_forms": { + "path": "spec/catalog/scalar-forms.json", + "sha256": "bdfcb4df19da4329c5ff0184b34daebf258d832992fa06fb3e0c34ca891c5923", + "count": 466 + }, + "tile_operations": { + "path": "spec/catalog/tile-operations.json", + "sha256": "07c5cf6f6e59916f3cbbecb0b83fe364704e1a79e1a43fa83f2997b6f7242207", + "count": 109 + }, + "extension_encoding_reservations": { + "path": "spec/catalog/extension-encoding-reservations.json", + "sha256": "f1b424060d3aae9432934724ba66908eab0476dcc4880447eaaca44fd016be8b", + "count": 40 + } + }, + "content_sha256": "f299fe3d256c5d071e57bb4aaa2be2de2e4a386ae090048df1f73ae92d392678", + "encoding_abi": "pto-isa-0.58.3-mode-function-v1", + "encoding_projection_sha256": "8a48b80e04484c70870f155bf9efc79d2a805cf99e809f4e4e8a7e6a7eb34172", + "hardware_conformance_profile": { + "path": "spec/hardware-conformance-profile.json", + "profile_id": "pto-hardware-numeric-0.58.3-ieee-v1", + "sha256": "117cd8e61f1e82001755cdb1ee98ef224fbd624ddf2adb8577c32d6acf833575" + }, + "numeric_conformance_vectors": { + "path": "spec/evidence/pto-isa-0583-hardware-numeric-vectors.json", + "sha256": "09d863e39e5fcd932353f4dc3bc5a7d2eed91ec9818c83886747aa69d28c3890" + }, + "release": "0.58.3", + "release_manifest": { + "path": "spec/release-manifest.json", + "sha256": "ebe8c75e2f8159634e49a9bbf81aba7141e15b45e63b3a5606e758ed7bc22023" + }, + "source": { + "commit": "e599a3d36ebfad43362ff591ea5e128816c684c7", + "tree": "abb6899d2e664e378ac9c1b77062670daa4d31b4", + "repository": "https://github.com/PTO-ISA/pto-spec.git" + } +} diff --git a/tools/pto_isa_v0_58_3_operation_contracts.json b/tools/pto_isa_v0_58_3_operation_contracts.json new file mode 100644 index 0000000000..556feaea1f --- /dev/null +++ b/tools/pto_isa_v0_58_3_operation_contracts.json @@ -0,0 +1,2837 @@ +{ + "release": "0.58.3", + "encoding_abi": "pto-isa-0.58.3-mode-function-v1", + "encoding_projection_sha256": "8a48b80e04484c70870f155bf9efc79d2a805cf99e809f4e4e8a7e6a7eb34172", + "source_commit": "e599a3d36ebfad43362ff591ea5e128816c684c7", + "source_tree": "abb6899d2e664e378ac9c1b77062670daa4d31b4", + "public_operation_count": 109, + "operations": { + "GMOV": { + "ptoas_mnemonic": "gmov", + "ptoas_arguments": [ + "src", + "peer_tid", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "resolved-peer-source" + ], + [ + "scalar0", + "peer-tid" + ] + ], + "engine": "TLSU" + }, + "MGATHER": { + "ptoas_mnemonic": "mgather", + "ptoas_arguments": [ + "mem", + "idx", + "dst", + "gatherOob" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "address", + "base-address" + ], + [ + "source0", + "indices" + ] + ], + "engine": "TLSU" + }, + "MGATHER_CAS": { + "ptoas_mnemonic": "mgather_cas", + "ptoas_arguments": [ + "mem", + "idx", + "expected", + "replacement", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "address", + "base-address" + ], + [ + "source0", + "indices" + ], + [ + "source1", + "expected" + ], + [ + "source2", + "replacement" + ] + ], + "engine": "TLSU" + }, + "MGATHER_MASK": { + "ptoas_mnemonic": "mgather_mask", + "ptoas_arguments": [ + "mem", + "idx", + "mask", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "address", + "base-address" + ], + [ + "source0", + "indices" + ], + [ + "source1", + "mask" + ] + ], + "engine": "TLSU" + }, + "MSCATTER": { + "ptoas_mnemonic": "mscatter", + "ptoas_arguments": [ + "src", + "idx", + "mem", + "scatterAtomicOp", + "scatterOob" + ], + "isa_operands": [ + [ + "address", + "base-address" + ], + [ + "source0", + "source" + ], + [ + "source1", + "byte-displacement-indices" + ] + ], + "engine": "TLSU" + }, + "MSCATTER_MASK": { + "ptoas_mnemonic": "mscatter_mask", + "ptoas_arguments": [ + "src", + "idx", + "mask", + "mem" + ], + "isa_operands": [ + [ + "address", + "base-address" + ], + [ + "source0", + "source" + ], + [ + "source1", + "byte-displacement-indices" + ], + [ + "source2", + "exact-predicate-mask" + ] + ], + "engine": "TLSU" + }, + "TABS": { + "ptoas_mnemonic": "tabs", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ] + ], + "engine": "VEC" + }, + "TADD": { + "ptoas_mnemonic": "tadd", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TADDS": { + "ptoas_mnemonic": "tadds", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TAND": { + "ptoas_mnemonic": "tand", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TANDS": { + "ptoas_mnemonic": "tands", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TCI": { + "ptoas_mnemonic": "tci", + "ptoas_arguments": [ + "S", + "dst", + "descending" + ], + "isa_operands": [ + [ + "destination0", + "new Local S32, S16, U32, or U16 destination" + ], + [ + "scalar0", + "typed sequence start" + ], + [ + "flag0", + "ascending or descending direction" + ] + ], + "engine": "SFU" + }, + "TCMP": { + "ptoas_mnemonic": "tcmp", + "ptoas_arguments": [ + "src0", + "src1", + "dst", + "cmpMode" + ], + "isa_operands": [ + [ + "destination0", + "predicate destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ], + [ + "comparison", + "comparison" + ] + ], + "engine": "VEC" + }, + "TCMPS": { + "ptoas_mnemonic": "tcmps", + "ptoas_arguments": [ + "src", + "scalar", + "cmpMode", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new packed Local predicate destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ], + [ + "comparison", + "six-mode comparison" + ] + ], + "engine": "VEC" + }, + "TCOLARGMAX": { + "ptoas_mnemonic": "tcolargmax", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local S32 or U32 index destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TCOLARGMIN": { + "ptoas_mnemonic": "tcolargmin", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local S32 or U32 index destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPAND": { + "ptoas_mnemonic": "tcolexpand", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDADD": { + "ptoas_mnemonic": "tcolexpandadd", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDDIV": { + "ptoas_mnemonic": "tcolexpanddiv", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDEXPDIF": { + "ptoas_mnemonic": "tcolexpandexpdif", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local destination with DstDataType" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDMAX": { + "ptoas_mnemonic": "tcolexpandmax", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDMIN": { + "ptoas_mnemonic": "tcolexpandmin", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDMUL": { + "ptoas_mnemonic": "tcolexpandmul", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLEXPANDSUB": { + "ptoas_mnemonic": "tcolexpandsub", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-row broadcast source" + ] + ], + "engine": "SFU" + }, + "TCOLMAX": { + "ptoas_mnemonic": "tcolmax", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TCOLMIN": { + "ptoas_mnemonic": "tcolmin", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TCOLPROD": { + "ptoas_mnemonic": "tcolprod", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TCOLSUM": { + "ptoas_mnemonic": "tcolsum", + "ptoas_arguments": [ + "src", + "tmp", + "dst", + "isBinary" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TCONCAT": { + "ptoas_mnemonic": "tconcat", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "SFU" + }, + "TCVT": { + "ptoas_mnemonic": "tcvt", + "ptoas_arguments": [ + "src", + "dst", + "rmode", + "sat_mode" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ], + [ + "numeric_control", + "rounding-and-saturation" + ] + ], + "engine": "VEC" + }, + "TDEQUANT": { + "ptoas_mnemonic": "tdequant", + "ptoas_arguments": [ + "src", + "scale", + "offset", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new FP32 destination" + ], + [ + "source0", + "persistent S8 or U8 source" + ], + [ + "scalar0", + "positive finite FP32 multiplier" + ], + [ + "scalar1", + "source-typed integer zero point" + ], + [ + "numeric_control", + "rounding" + ] + ], + "engine": "SFU" + }, + "TDIV": { + "ptoas_mnemonic": "tdiv", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "numerator" + ], + [ + "source1", + "denominator" + ] + ], + "engine": "SFU" + }, + "TDIVS": { + "ptoas_mnemonic": "tdivs", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "SFU" + }, + "TEXP": { + "ptoas_mnemonic": "texp", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local floating destination" + ], + [ + "source0", + "persistent Local floating source" + ] + ], + "engine": "SFU" + }, + "TEXPANDS": { + "ptoas_mnemonic": "texpands", + "ptoas_arguments": [ + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TEXTRACT": { + "ptoas_mnemonic": "textract", + "ptoas_arguments": [ + "src", + "indexRow", + "indexCol", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ], + [ + "natural0", + "row-offset" + ], + [ + "natural1", + "column-offset" + ] + ], + "engine": "SFU" + }, + "TFILLPAD": { + "ptoas_mnemonic": "tfillpad", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ], + [ + "scalar0", + "padding" + ] + ], + "engine": "SFU" + }, + "TFMA": { + "ptoas_mnemonic": "tfma", + "ptoas_arguments": [ + "src0", + "src1", + "src2", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "multiplicand-left" + ], + [ + "source1", + "multiplicand-right" + ], + [ + "source2", + "addend" + ] + ], + "engine": "VEC" + }, + "TGATHER": { + "ptoas_mnemonic": "tgather", + "ptoas_arguments": [ + "src", + "dst", + "cdst", + "indices", + "tmp", + "kValue", + "maskPattern", + "cmpMode", + "offset" + ], + "isa_operands": [ + [ + "destination0", + "new Local value destination" + ], + [ + "source0", + "persistent Local value source" + ], + [ + "source1", + "persistent Local S16, U16, S32, U32, S64, or U64 row-index source" + ] + ], + "engine": "SFU" + }, + "TGEMV": { + "ptoas_mnemonic": "tgemv", + "ptoas_arguments": [ + "lhs", + "rhs", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left-vector" + ], + [ + "source1", + "right-matrix" + ] + ], + "engine": "CUBE" + }, + "TGEMV_ACC": { + "ptoas_mnemonic": "tgemv.acc", + "ptoas_arguments": [ + "acc_in", + "lhs", + "rhs", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "accumulator" + ], + [ + "source1", + "left-vector" + ], + [ + "source2", + "right-matrix" + ] + ], + "engine": "CUBE" + }, + "TGEMV_BIAS": { + "ptoas_mnemonic": "tgemv.bias", + "ptoas_arguments": [ + "a", + "b", + "bias", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left-vector" + ], + [ + "source1", + "right-matrix" + ], + [ + "source2", + "bias" + ] + ], + "engine": "CUBE" + }, + "TGEMV_MX": { + "ptoas_mnemonic": "tgemv.mx", + "ptoas_arguments": [ + "a", + "a_scale", + "b", + "b_scale", + "dst" + ], + "ptoas_optional_arguments": ["a_scale", "b_scale"], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left-vector" + ], + [ + "source1", + "row-scale" + ], + [ + "source2", + "right-matrix" + ], + [ + "source3", + "column-scale" + ] + ], + "engine": "CUBE" + }, + "TGEMV_MX_ACC": { + "ptoas_mnemonic": "tgemv.mx.acc", + "ptoas_arguments": [ + "c_in", + "a", + "a_scale", + "b", + "b_scale", + "dst" + ], + "ptoas_optional_arguments": ["a_scale", "b_scale"], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "accumulator" + ], + [ + "source1", + "left-vector" + ], + [ + "source2", + "row-scale" + ], + [ + "source3", + "right-matrix" + ], + [ + "source4", + "column-scale" + ] + ], + "engine": "CUBE" + }, + "TGEMV_MX_BIAS": { + "ptoas_mnemonic": "tgemv.mx.bias", + "ptoas_arguments": [ + "a", + "a_scale", + "b", + "b_scale", + "bias", + "dst" + ], + "ptoas_optional_arguments": ["a_scale", "b_scale"], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left-vector" + ], + [ + "source1", + "row-scale" + ], + [ + "source2", + "right-matrix" + ], + [ + "source3", + "column-scale" + ], + [ + "source4", + "bias" + ] + ], + "engine": "CUBE" + }, + "THISTOGRAM": { + "ptoas_mnemonic": "thistogram", + "ptoas_arguments": [ + "src", + "idx", + "dst", + "isMSB" + ], + "isa_operands": [ + [ + "destination0", + "new Local U32 prefix-histogram destination" + ], + [ + "source0", + "persistent Local U16 or U32 source" + ], + [ + "source1", + "persistent Local U8 prefix filter" + ], + [ + "selected_byte", + "B.DATR ByteId zero through three" + ] + ], + "engine": "SFU" + }, + "TIMG2COL": { + "ptoas_mnemonic": "timg2col", + "ptoas_arguments": [ + "src", + "posM", + "posK", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local Matrix destination" + ], + [ + "source0", + "persistent Local Matrix feature-map source" + ], + [ + "natural0", + "unsigned low-sixteen-bit posM" + ], + [ + "natural1", + "unsigned low-sixteen-bit posK" + ] + ], + "engine": "SFU" + }, + "TINSERT": { + "ptoas_mnemonic": "tinsert", + "ptoas_arguments": [ + "src", + "indexRow", + "indexCol", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "persistent old destination" + ], + [ + "source1", + "persistent insertion source" + ], + [ + "natural0", + "row-offset" + ], + [ + "natural1", + "column-offset" + ] + ], + "engine": "SFU" + }, + "TLOAD": { + "ptoas_mnemonic": "tload", + "ptoas_arguments": [ + "src", + "dst", + "pad_mode", + "pad_value", + "left_padding_num", + "right_padding_num", + "init_out_buffer", + "init_condition" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "address", + "base-address" + ], + [ + "scalar0", + "row-stride-bytes" + ] + ], + "engine": "TLSU" + }, + "TLOG": { + "ptoas_mnemonic": "tlog", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local floating destination" + ], + [ + "source0", + "persistent Local floating source" + ] + ], + "engine": "SFU" + }, + "TMATMUL": { + "ptoas_mnemonic": "tmatmul", + "ptoas_arguments": [ + "lhs", + "rhs", + "dst", + "accPhase" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left" + ], + [ + "source1", + "right" + ] + ], + "engine": "CUBE" + }, + "TMATMUL_ACC": { + "ptoas_mnemonic": "tmatmul.acc", + "ptoas_arguments": [ + "acc_in", + "lhs", + "rhs", + "dst", + "accPhase" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "accumulator" + ], + [ + "source1", + "left" + ], + [ + "source2", + "right" + ] + ], + "engine": "CUBE" + }, + "TMATMUL_BIAS": { + "ptoas_mnemonic": "tmatmul.bias", + "ptoas_arguments": [ + "a", + "b", + "bias", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left" + ], + [ + "source1", + "right" + ], + [ + "source2", + "bias" + ] + ], + "engine": "CUBE" + }, + "TMATMUL_MX": { + "ptoas_mnemonic": "tmatmul.mx", + "ptoas_arguments": [ + "a", + "a_scale", + "b", + "b_scale", + "dst" + ], + "ptoas_optional_arguments": ["a_scale", "b_scale"], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left" + ], + [ + "source1", + "row-scale" + ], + [ + "source2", + "right" + ], + [ + "source3", + "column-scale" + ] + ], + "engine": "CUBE" + }, + "TMATMUL_MX_ACC": { + "ptoas_mnemonic": "tmatmul.mx.acc", + "ptoas_arguments": [ + "c_in", + "a", + "a_scale", + "b", + "b_scale", + "dst" + ], + "ptoas_optional_arguments": ["a_scale", "b_scale"], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "accumulator" + ], + [ + "source1", + "left" + ], + [ + "source2", + "row-scale" + ], + [ + "source3", + "right" + ], + [ + "source4", + "column-scale" + ] + ], + "engine": "CUBE" + }, + "TMATMUL_MX_BIAS": { + "ptoas_mnemonic": "tmatmul.mx.bias", + "ptoas_arguments": [ + "a", + "a_scale", + "b", + "b_scale", + "bias", + "dst" + ], + "ptoas_optional_arguments": ["a_scale", "b_scale"], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "left" + ], + [ + "source1", + "row-scale" + ], + [ + "source2", + "right" + ], + [ + "source3", + "column-scale" + ], + [ + "source4", + "bias" + ] + ], + "engine": "CUBE" + }, + "TMAX": { + "ptoas_mnemonic": "tmax", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TMAXS": { + "ptoas_mnemonic": "tmaxs", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TMIN": { + "ptoas_mnemonic": "tmin", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TMINS": { + "ptoas_mnemonic": "tmins", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TMOV": { + "ptoas_mnemonic": "tmov", + "ptoas_arguments": [ + "src", + "dst", + "fp", + "preQuantScalar", + "accToVecMode", + "reluPreMode" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ] + ], + "engine": "TLSU" + }, + "TMRGSORT": { + "ptoas_mnemonic": "tmrgsort", + "ptoas_arguments": [ + "srcs", + "blockLen", + "dsts", + "tmp", + "excuted", + "exhausted" + ], + "isa_operands": [ + [ + "destination0", + "new Local merged destination" + ], + [ + "source0", + "persistent sorted Local left source" + ], + [ + "source1", + "persistent sorted Local right source" + ], + [ + "flag0", + "ascending or descending selection" + ] + ], + "engine": "SFU" + }, + "TMUL": { + "ptoas_mnemonic": "tmul", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TMULS": { + "ptoas_mnemonic": "tmuls", + "ptoas_arguments": [ + "src0", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TNEG": { + "ptoas_mnemonic": "tneg", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ] + ], + "engine": "VEC" + }, + "TNOT": { + "ptoas_mnemonic": "tnot", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ] + ], + "engine": "VEC" + }, + "TOR": { + "ptoas_mnemonic": "tor", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TORS": { + "ptoas_mnemonic": "tors", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TPARTADD": { + "ptoas_mnemonic": "tpartadd", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local union destination" + ], + [ + "source0", + "persistent Local left source anchored at origin" + ], + [ + "source1", + "persistent Local right source anchored at origin" + ] + ], + "engine": "SFU" + }, + "TPARTMAX": { + "ptoas_mnemonic": "tpartmax", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local union destination" + ], + [ + "source0", + "persistent Local left source anchored at origin" + ], + [ + "source1", + "persistent Local right source anchored at origin" + ] + ], + "engine": "SFU" + }, + "TPARTMIN": { + "ptoas_mnemonic": "tpartmin", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local union destination" + ], + [ + "source0", + "persistent Local left source anchored at origin" + ], + [ + "source1", + "persistent Local right source anchored at origin" + ] + ], + "engine": "SFU" + }, + "TPARTMUL": { + "ptoas_mnemonic": "tpartmul", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local union destination" + ], + [ + "source0", + "persistent Local left source anchored at origin" + ], + [ + "source1", + "persistent Local right source anchored at origin" + ] + ], + "engine": "SFU" + }, + "TPREFETCH": { + "ptoas_mnemonic": "tprefetch", + "ptoas_arguments": [ + "address", + "row_stride", + "valid_cols", + "valid_rows", + "physical_cols" + ], + "isa_operands": [ + [ + "address", + "base-address" + ], + [ + "scalar0", + "row-stride-elements" + ], + [ + "positive0", + "valid-columns" + ], + [ + "positive1", + "valid-rows" + ], + [ + "positive2", + "physical-columns" + ] + ], + "engine": "TLSU" + }, + "TQUANT": { + "ptoas_mnemonic": "tquant", + "ptoas_arguments": [ + "src", + "fp", + "offset", + "dst", + "quant_type" + ], + "isa_operands": [ + [ + "destination0", + "new S8 or U8 destination" + ], + [ + "source0", + "persistent FP32 source" + ], + [ + "scalar0", + "positive finite FP32 multiplier" + ], + [ + "scalar1", + "destination-typed integer zero point" + ], + [ + "numeric_control", + "rounding and saturation" + ] + ], + "engine": "SFU" + }, + "TRECIP": { + "ptoas_mnemonic": "trecip", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local floating destination" + ], + [ + "source0", + "persistent Local floating source" + ] + ], + "engine": "SFU" + }, + "TRELU": { + "ptoas_mnemonic": "trelu", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ] + ], + "engine": "VEC" + }, + "TREM": { + "ptoas_mnemonic": "trem", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "dividend" + ], + [ + "source1", + "divisor" + ] + ], + "engine": "SFU" + }, + "TREMS": { + "ptoas_mnemonic": "trems", + "ptoas_arguments": [ + "src", + "scalar", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "SFU" + }, + "TROWARGMAX": { + "ptoas_mnemonic": "trowargmax", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local S32 or U32 index destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TROWARGMIN": { + "ptoas_mnemonic": "trowargmin", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local S32 or U32 index destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TROWEXPAND": { + "ptoas_mnemonic": "trowexpand", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDADD": { + "ptoas_mnemonic": "trowexpandadd", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDDIV": { + "ptoas_mnemonic": "trowexpanddiv", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDEXPDIF": { + "ptoas_mnemonic": "trowexpandexpdif", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local destination with DstDataType" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDMAX": { + "ptoas_mnemonic": "trowexpandmax", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDMIN": { + "ptoas_mnemonic": "trowexpandmin", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDMUL": { + "ptoas_mnemonic": "trowexpandmul", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWEXPANDSUB": { + "ptoas_mnemonic": "trowexpandsub", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local full-shape numeric source" + ], + [ + "source1", + "persistent Local one-column broadcast source" + ] + ], + "engine": "SFU" + }, + "TROWMAX": { + "ptoas_mnemonic": "trowmax", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TROWMIN": { + "ptoas_mnemonic": "trowmin", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TROWPROD": { + "ptoas_mnemonic": "trowprod", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TROWSUM": { + "ptoas_mnemonic": "trowsum", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local same-type numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ] + ], + "engine": "SFU" + }, + "TRSQRT": { + "ptoas_mnemonic": "trsqrt", + "ptoas_arguments": [ + "src", + "dst", + "tmp" + ], + "isa_operands": [ + [ + "destination0", + "new Local floating destination" + ], + [ + "source0", + "persistent Local floating source" + ] + ], + "engine": "SFU" + }, + "TSCATTER": { + "ptoas_mnemonic": "tscatter", + "ptoas_arguments": [ + "src", + "indexes", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new zero-initialized Local value destination" + ], + [ + "source0", + "persistent Local value source" + ], + [ + "source1", + "persistent Local S16, U16, S32, U32, S64, or U64 row-index source" + ] + ], + "engine": "SFU" + }, + "TSEL": { + "ptoas_mnemonic": "tsel", + "ptoas_arguments": [ + "mask", + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "numeric destination" + ], + [ + "source0", + "packed predicate mask" + ], + [ + "source1", + "source-true" + ], + [ + "source2", + "source-false" + ] + ], + "engine": "VEC" + }, + "TSELS": { + "ptoas_mnemonic": "tsels", + "ptoas_arguments": [ + "mask", + "src", + "tmp", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "packed Local predicate mask" + ], + [ + "source1", + "persistent Local source selected by one" + ], + [ + "scalar0", + "per-participating-PE scalar selected by zero" + ] + ], + "engine": "VEC" + }, + "TSHL": { + "ptoas_mnemonic": "tshl", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "value-source" + ], + [ + "source1", + "shift-count-source" + ] + ], + "engine": "VEC" + }, + "TSHLS": { + "ptoas_mnemonic": "tshls", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TSHR": { + "ptoas_mnemonic": "tshr", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "value-source" + ], + [ + "source1", + "shift-count-source" + ] + ], + "engine": "VEC" + }, + "TSHRS": { + "ptoas_mnemonic": "tshrs", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TSORT": { + "ptoas_mnemonic": "tsort", + "ptoas_arguments": [ + "src", + "dst", + "dst_indices", + "descending" + ], + "isa_operands": [ + [ + "destination0", + "new Local sorted-value destination" + ], + [ + "destination1", + "new Local U32 original-index destination" + ], + [ + "source0", + "persistent Local source" + ], + [ + "sort_width", + "LB0 row-group width" + ], + [ + "flag0", + "ascending or descending selection" + ] + ], + "engine": "SFU" + }, + "TSQRT": { + "ptoas_mnemonic": "tsqrt", + "ptoas_arguments": [ + "src", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local floating destination" + ], + [ + "source0", + "persistent Local floating source" + ] + ], + "engine": "SFU" + }, + "TSTORE": { + "ptoas_mnemonic": "tstore", + "ptoas_arguments": [ + "src", + "dst", + "preQuantScalar", + "stPhase", + "atomicType", + "reluPreMode" + ], + "isa_operands": [ + [ + "address", + "base-address" + ], + [ + "scalar0", + "row-stride-bytes" + ], + [ + "source0", + "source" + ] + ], + "engine": "TLSU" + }, + "TSUB": { + "ptoas_mnemonic": "tsub", + "ptoas_arguments": [ + "src0", + "src1", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TSUBS": { + "ptoas_mnemonic": "tsubs", + "ptoas_arguments": [ + "src", + "scalar", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + }, + "TTRANS": { + "ptoas_mnemonic": "ttrans", + "ptoas_arguments": [ + "src", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source" + ] + ], + "engine": "SFU" + }, + "TTRI": { + "ptoas_mnemonic": "ttri", + "ptoas_arguments": [ + "diagonal", + "dst", + "upperOrLower" + ], + "isa_operands": [ + [ + "destination0", + "new Local triangular destination" + ], + [ + "flag0", + "lower or upper orientation" + ], + [ + "diagonal", + "signed diagonal displacement" + ] + ], + "engine": "SFU" + }, + "TXOR": { + "ptoas_mnemonic": "txor", + "ptoas_arguments": [ + "src0", + "src1", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "destination" + ], + [ + "source0", + "source-left" + ], + [ + "source1", + "source-right" + ] + ], + "engine": "VEC" + }, + "TXORS": { + "ptoas_mnemonic": "txors", + "ptoas_arguments": [ + "src", + "scalar", + "tmp", + "dst" + ], + "isa_operands": [ + [ + "destination0", + "new Local numeric destination" + ], + [ + "source0", + "persistent Local numeric source" + ], + [ + "scalar0", + "per-participating-PE private-GPR scalar" + ] + ], + "engine": "VEC" + } + }, + "linx_rejected_mnemonics": [ + "acccvt", + "talloc", + "taxpy", + "tdeinterleave", + "tfree", + "tgatherb", + "tinterleave", + "tpartargmax", + "tpartargmin", + "tpop", + "tprelu", + "tpush", + "treshape" + ], + "dialect_only_operations": { + "acccvt": [ + "dst" + ], + "talloc": [ + "entry", + "pipe_handle", + "split" + ], + "talloc_to_aic": [ + "id", + "split" + ], + "talloc_to_aiv": [ + "id", + "split" + ], + "taxpy": [ + "src", + "scalar", + "dst" + ], + "tconcatidx": [ + "src0", + "src1", + "src0Idx", + "src1Idx", + "dst" + ], + "tdeinterleave": [ + "src", + "dst0", + "dst1" + ], + "textract_fp": [ + "src", + "fp", + "indexRow", + "indexCol", + "dst" + ], + "tfillpad_expand": [ + "src", + "dst" + ], + "tfillpad_inplace": [ + "src", + "dst" + ], + "tfree": [ + "entry", + "pipe_handle", + "split" + ], + "tfree_from_aic": [ + "entry", + "id", + "split" + ], + "tfree_from_aiv": [ + "entry", + "id", + "split" + ], + "tgatherb": [ + "src", + "offsets", + "dst" + ], + "tget_scale_addr": [ + "src", + "dst" + ], + "tgetval": [ + "src", + "offset" + ], + "tinsert_fp": [ + "src", + "fp", + "indexRow", + "indexCol", + "dst" + ], + "tinterleave": [ + "src0", + "src1", + "dst" + ], + "tmov.fp": [ + "src", + "fp", + "dst" + ], + "tpartargmax": [ + "src0", + "src1", + "src0Idx", + "src1Idx", + "dst", + "dstIdx" + ], + "tpartargmin": [ + "src0", + "src1", + "src0Idx", + "src1Idx", + "dst", + "dstIdx" + ], + "tpop": [ + "tile", + "pipe_handle", + "split" + ], + "tpop_from_aic": [ + "valid_row", + "valid_col", + "id", + "split" + ], + "tpop_from_aiv": [ + "valid_row", + "valid_col", + "id", + "split" + ], + "tprelu": [ + "src0", + "src1", + "tmp", + "dst" + ], + "tprint": [ + "src" + ], + "tpush": [ + "tile", + "pipe_handle", + "split" + ], + "tpush_to_aic": [ + "tile", + "id", + "split" + ], + "tpush_to_aiv": [ + "tile", + "id", + "split" + ], + "treshape": [ + "src" + ], + "tsetval": [ + "dst", + "offset", + "val" + ], + "tstore_fp": [ + "src", + "fp", + "dst" + ], + "tsync": [ + "events" + ] + } +} diff --git a/tools/ptoas/ptoas.cpp b/tools/ptoas/ptoas.cpp index 6c9d5f5e52..37618be003 100644 --- a/tools/ptoas/ptoas.cpp +++ b/tools/ptoas/ptoas.cpp @@ -56,7 +56,7 @@ using namespace pto; #endif static void printPTOASVersion(llvm::raw_ostream &os) { - os << "ptoas " << PTOAS_RELEASE_VERSION << " (PTO ISA 0.58.1)\n"; + os << "ptoas " << PTOAS_RELEASE_VERSION << " (PTO ISA 0.58.3)\n"; } static LogicalResult reorderEmitCFunctions(ModuleOp module) { @@ -263,7 +263,7 @@ static LogicalResult validateLinxOperationBoundary(ModuleOp module) { StringRef name = op->getName().getStringRef(); op->emitError() << "'" << name << "' op is dialect-only and not part of the active Linx " - "PTO ISA v0.58.1 target"; + "PTO ISA v0.58.3 target"; } return invalid.empty() ? success() : failure(); } diff --git a/tools/ptobc/MAINTENANCE.md b/tools/ptobc/MAINTENANCE.md index ac2baf1fb4..61b5c7420f 100644 --- a/tools/ptobc/MAINTENANCE.md +++ b/tools/ptobc/MAINTENANCE.md @@ -23,12 +23,20 @@ python3 tools/ptobc/update_v0581_schema.py \ The second command is the non-mutating CI audit. Historical 0.58.0 opcode assignments and their round-trip test remain release evidence. +PTO ISA 0.58.3 keeps those operation opcodes. Its PTO-BC v0 deltas are the new +`BLayout` enum values `cube_m16=2`, `cube_m32=3`, and `cube_n8=4`, plus dynamic +operand counts for the `TGEMV` opcode family and `TMATMUL_MX` opcode family. +The latter preserves the independently optional A/B MX scales; the ordinary +attribute dictionary retains `operandSegmentSizes`, so A-only and B-only forms +remain distinguishable during decode. The Linx CUBE lit tests and the full +PTO-BC round-trip gate cover all zero/A-only/B-only/two-scale forms. + ## Required gates Run (or rely on CI): - `ctest -R ptobc_stage9_e2e` - `ctest -R ptobc_to_ptoas_smoke` - `ctest -R ptobc_opcode_coverage_check` -- `ctest -R ptobc_v0581_contract_encode` +- `ctest -R ptobc_v0583_contract_encode` ## Notes - `ptobc_opcode_coverage_check` is a heuristic based on `mnemonic = "..."` occurrences. diff --git a/tools/ptobc/generated/ptobc_opcodes_v0.h b/tools/ptobc/generated/ptobc_opcodes_v0.h index 6594f784a9..f31791ff6d 100644 --- a/tools/ptobc/generated/ptobc_opcodes_v0.h +++ b/tools/ptobc/generated/ptobc_opcodes_v0.h @@ -119,7 +119,7 @@ inline constexpr OpInfo kOpTable[] = { {0x1025, "pto.tfillpad_inplace", 0, 0x00, 0x00, 2, 0, 0, 0x00}, {0x1028, "pto.tgather", 0, 0x00, 0x02, 0, 0, 0, 0x00}, {0x1029, "pto.tgatherb", 0, 0x00, 0x00, 3, 0, 0, 0x00}, - {0x102A, "pto.tgemv", 1, 0x00, 0x01, 0, 0, 0, 0x00}, + {0x102A, "pto.tgemv", 1, 0x00, 0x02, 0, 0, 0, 0x00}, {0x102B, "pto.tgetval", 0, 0x01, 0x00, 2, 1, 0, 0x00}, {0x102C, "pto.timg2col", 0, 0x00, 0x00, 4, 0, 0, 0x00}, {0x102D, "pto.tinsert", 0, 0x00, 0x00, 4, 0, 0, 0x00}, @@ -127,7 +127,7 @@ inline constexpr OpInfo kOpTable[] = { {0x102F, "pto.tload", 0, 0x00, 0x00, 2, 0, 0, 0x00}, {0x1030, "pto.tlog", 0, 0x00, 0x00, 2, 0, 0, 0x00}, {0x1032, "pto.tmatmul", 1, 0x00, 0x01, 0, 0, 0, 0x00}, - {0x1033, "pto.tmatmul.mx", 1, 0x00, 0x01, 0, 0, 0, 0x00}, + {0x1033, "pto.tmatmul.mx", 1, 0x00, 0x02, 0, 0, 0, 0x00}, {0x1034, "pto.tmax", 0, 0x00, 0x00, 3, 0, 0, 0x00}, {0x1035, "pto.tmaxs", 0, 0x00, 0x00, 3, 0, 0, 0x00}, {0x1036, "pto.tmin", 0, 0x00, 0x00, 3, 0, 0, 0x00}, diff --git a/tools/ptobc/testdata/recent_ops_v0_roundtrip.pto b/tools/ptobc/testdata/recent_ops_v0_roundtrip.pto index b1324400bd..b23fe2c8ff 100644 --- a/tools/ptobc/testdata/recent_ops_v0_roundtrip.pto +++ b/tools/ptobc/testdata/recent_ops_v0_roundtrip.pto @@ -44,9 +44,9 @@ module attributes {pto.target_arch = "a5"} { %dst_mx = pto.alloc_tile : !pto.tile_buf %dst_mx_acc = pto.alloc_tile : !pto.tile_buf %dst_mx_bias = pto.alloc_tile : !pto.tile_buf - pto.tgemv.mx ins(%a, %a_scale, %b, %b_scale : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst_mx : !pto.tile_buf) - pto.tgemv.mx.acc ins(%c_in, %a, %a_scale, %b, %b_scale : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst_mx_acc : !pto.tile_buf) - pto.tgemv.mx.bias ins(%a, %a_scale, %b, %b_scale, %bias_mx : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) outs(%dst_mx_bias : !pto.tile_buf) + pto.tgemv.mx ins(%a, %b : !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst_mx : !pto.tile_buf) + pto.tgemv.mx.acc ins(%c_in, %a, %b : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst_mx_acc : !pto.tile_buf) + pto.tgemv.mx.bias ins(%a, %b, %bias_mx : !pto.tile_buf, !pto.tile_buf, !pto.tile_buf) a_scale(%a_scale : !pto.tile_buf) b_scale(%b_scale : !pto.tile_buf) outs(%dst_mx_bias : !pto.tile_buf) return } } diff --git a/tools/ptobc/tests/CMakeLists.txt b/tools/ptobc/tests/CMakeLists.txt index eb3344cf00..366e4d2be2 100644 --- a/tools/ptobc/tests/CMakeLists.txt +++ b/tools/ptobc/tests/CMakeLists.txt @@ -44,12 +44,12 @@ add_test(NAME ptobc_v0580_ops_v0_encode bash ${CMAKE_CURRENT_LIST_DIR}/v0580_ops_v0_encode.sh ) -add_test(NAME ptobc_v0581_contract_encode +add_test(NAME ptobc_v0583_contract_encode COMMAND ${CMAKE_COMMAND} -E env PTOBC_BIN=$ TEST_INPUT_DIR=${CMAKE_SOURCE_DIR}/test/lit/pto - OUT_DIR=${CMAKE_CURRENT_BINARY_DIR}/ptobc_v0581_contract_out - bash ${CMAKE_CURRENT_LIST_DIR}/ptobc_v0581_contract_encode.sh + OUT_DIR=${CMAKE_CURRENT_BINARY_DIR}/ptobc_v0583_contract_out + bash ${CMAKE_CURRENT_LIST_DIR}/ptobc_v0583_contract_encode.sh ) add_test(NAME ptobc_trowexpandsub_v0_encode diff --git a/tools/ptobc/tests/ptobc_v0581_contract_encode.sh b/tools/ptobc/tests/ptobc_v0581_contract_encode.sh deleted file mode 100755 index 932b7c5228..0000000000 --- a/tools/ptobc/tests/ptobc_v0581_contract_encode.sh +++ /dev/null @@ -1,26 +0,0 @@ -#!/usr/bin/env bash -# Copyright (c) 2026 Huawei Technologies Co., Ltd. -# This program is free software, you can redistribute it and/or modify it under the terms and conditions of -# CANN Open Software License Agreement Version 2.0 (the "License"). -# Please refer to the License for details. You may not use this file except in compliance with the License. -# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, -# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. -# See LICENSE in the root of the software repository for the full text of the License. - -set -euo pipefail - -: "${PTOBC_BIN:?PTOBC_BIN not set}" -: "${TEST_INPUT_DIR:?TEST_INPUT_DIR not set}" - -OUT_DIR=${OUT_DIR:-"${PWD}/ptobc_v0581_contract_out"} -mkdir -p "${OUT_DIR}" - -input="${TEST_INPUT_DIR}/v0581_linx_contract.pto" -bytecode="${OUT_DIR}/v0581_linx_contract.ptobc" -decoded="${OUT_DIR}/v0581_linx_contract.roundtrip.pto" -"${PTOBC_BIN}" encode "${input}" -o "${bytecode}" -"${PTOBC_BIN}" decode "${bytecode}" -o "${decoded}" - -grep -F "pto.timg2col ins(" "${decoded}" >/dev/null -grep -F "pto.tinsert ins(" "${decoded}" >/dev/null -grep -F "pto.tprefetch ins(" "${decoded}" >/dev/null diff --git a/tools/ptobc/tests/ptobc_v0583_contract_encode.sh b/tools/ptobc/tests/ptobc_v0583_contract_encode.sh new file mode 100755 index 0000000000..983a0d9f94 --- /dev/null +++ b/tools/ptobc/tests/ptobc_v0583_contract_encode.sh @@ -0,0 +1,62 @@ +#!/usr/bin/env bash +# Copyright (c) 2026 Huawei Technologies Co., Ltd. +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of +# CANN Open Software License Agreement Version 2.0 (the "License"). +# Please refer to the License for details. You may not use this file except in compliance with the License. +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. +# See LICENSE in the root of the software repository for the full text of the License. + +set -euo pipefail + +: "${PTOBC_BIN:?PTOBC_BIN not set}" +: "${TEST_INPUT_DIR:?TEST_INPUT_DIR not set}" + +OUT_DIR=${OUT_DIR:-"${PWD}/ptobc_v0583_contract_out"} +mkdir -p "${OUT_DIR}" + +input="${TEST_INPUT_DIR}/v0583_linx_contract.pto" +bytecode="${OUT_DIR}/v0583_linx_contract.ptobc" +decoded="${OUT_DIR}/v0583_linx_contract.roundtrip.pto" +"${PTOBC_BIN}" encode "${input}" -o "${bytecode}" +"${PTOBC_BIN}" decode "${bytecode}" -o "${decoded}" + +grep -F "pto.timg2col ins(" "${decoded}" >/dev/null +grep -F "pto.tinsert ins(" "${decoded}" >/dev/null +grep -F "pto.tprefetch ins(" "${decoded}" >/dev/null + +cube_input="${TEST_INPUT_DIR}/v0583_linx_cube_contract.pto" +cube_bytecode="${OUT_DIR}/v0583_linx_cube_contract.ptobc" +cube_decoded="${OUT_DIR}/v0583_linx_cube_contract.roundtrip.pto" +"${PTOBC_BIN}" encode "${cube_input}" -o "${cube_bytecode}" +"${PTOBC_BIN}" decode "${cube_bytecode}" -o "${cube_decoded}" + +grep -F "blayout=cube_m16" "${cube_decoded}" >/dev/null +grep -F "blayout=cube_n8" "${cube_decoded}" >/dev/null +grep -F "pto.tmatmul ins(" "${cube_decoded}" >/dev/null +grep -F "pto.tgemv ins(" "${cube_decoded}" >/dev/null + +mx_input="${TEST_INPUT_DIR}/v0583_linx_tmatmul_mx_variants.pto" +mx_bytecode="${OUT_DIR}/v0583_linx_tmatmul_mx_variants.ptobc" +mx_decoded="${OUT_DIR}/v0583_linx_tmatmul_mx_variants.roundtrip.pto" +"${PTOBC_BIN}" encode "${mx_input}" -o "${mx_bytecode}" +"${PTOBC_BIN}" decode "${mx_bytecode}" -o "${mx_decoded}" + +grep -F "!pto.f8E8M0" "${mx_decoded}" >/dev/null +grep -F "pto.tmatmul.mx ins(" "${mx_decoded}" >/dev/null +grep -F "pto.tmatmul.mx.acc ins(" "${mx_decoded}" >/dev/null +grep -F "pto.tmatmul.mx.bias ins(" "${mx_decoded}" >/dev/null +grep -F "a_scale(" "${mx_decoded}" >/dev/null +grep -F "b_scale(" "${mx_decoded}" >/dev/null + +gemv_mx_input="${TEST_INPUT_DIR}/v0583_linx_tgemv_mx_optional_scales.pto" +gemv_mx_bytecode="${OUT_DIR}/v0583_linx_tgemv_mx_optional_scales.ptobc" +gemv_mx_decoded="${OUT_DIR}/v0583_linx_tgemv_mx_optional_scales.roundtrip.pto" +"${PTOBC_BIN}" encode "${gemv_mx_input}" -o "${gemv_mx_bytecode}" +"${PTOBC_BIN}" decode "${gemv_mx_bytecode}" -o "${gemv_mx_decoded}" + +grep -F "pto.tgemv.mx ins(" "${gemv_mx_decoded}" >/dev/null +grep -F "pto.tgemv.mx.acc ins(" "${gemv_mx_decoded}" >/dev/null +grep -F "pto.tgemv.mx.bias ins(" "${gemv_mx_decoded}" >/dev/null +grep -F "a_scale(" "${gemv_mx_decoded}" >/dev/null +grep -F "b_scale(" "${gemv_mx_decoded}" >/dev/null diff --git a/tools/test_check_python_binding_link_contract.py b/tools/test_check_python_binding_link_contract.py index 3b2fd59e74..156dbff549 100644 --- a/tools/test_check_python_binding_link_contract.py +++ b/tools/test_check_python_binding_link_contract.py @@ -20,8 +20,15 @@ CHECKER = Path(__file__).with_name("check_python_binding_link_contract.py") VALID_PRODUCER = """ if(UNIX AND NOT APPLE) - target_link_options(nanobind-mlir PRIVATE "LINKER:-z,undefs") - install(TARGETS nanobind-mlir + set(PTOAS_NANOBIND_RUNTIME_TARGET "") + foreach(_candidate nanobind-mlir nanobind-abi3-ft nanobind-ft nanobind-abi3 nanobind) + if(TARGET ${_candidate}) + set(PTOAS_NANOBIND_RUNTIME_TARGET ${_candidate}) + break() + endif() + endforeach() + target_link_options(${PTOAS_NANOBIND_RUNTIME_TARGET} PRIVATE "LINKER:-z,undefs") + install(TARGETS ${PTOAS_NANOBIND_RUNTIME_TARGET} LIBRARY DESTINATION lib ) endif() @@ -70,13 +77,16 @@ def test_rejects_shared_runtime_not_installed_in_auditwheel_search_path( result = self.run_checker( """ if(UNIX AND NOT APPLE) - target_link_options(nanobind-mlir PRIVATE "LINKER:-z,undefs") + set(PTOAS_NANOBIND_RUNTIME_TARGET "") + foreach(_candidate nanobind-mlir nanobind) + endforeach() + target_link_options(${PTOAS_NANOBIND_RUNTIME_TARGET} PRIVATE "LINKER:-z,undefs") endif() """, "export LD_LIBRARY_PATH=$LLVM_BUILD_DIR/lib:$PTO_INSTALL_DIR/lib:$LD_LIBRARY_PATH\n", ) self.assertNotEqual(result.returncode, 0) - self.assertIn("install nanobind-mlir into lib", result.stderr) + self.assertIn("install the resolved nanobind shared runtime", result.stderr) def test_accepts_installed_runtime_and_matching_auditwheel_search_path( self, diff --git a/tools/test_check_v058_pto_manifest.py b/tools/test_check_v058_pto_manifest.py index 7f35243351..75ee289bcd 100755 --- a/tools/test_check_v058_pto_manifest.py +++ b/tools/test_check_v058_pto_manifest.py @@ -9,6 +9,7 @@ import importlib.util import json +import re import tempfile import unittest from pathlib import Path @@ -22,9 +23,22 @@ class LinxIdentityTest(unittest.TestCase): + def test_local_0583_identity_mismatch_fails_closed(self): + local_lock = json.loads( + Path(__file__).with_name("pto_isa_v0_58_3_lock.json").read_text() + ) + local_lock["encoding_projection_sha256"] = "0" * 64 + with tempfile.TemporaryDirectory() as directory: + ptoas_root = Path(directory) + tools = ptoas_root / "tools" + tools.mkdir() + (tools / "pto_isa_v0_58_3_lock.json").write_text(json.dumps(local_lock)) + with self.assertRaisesRegex(SystemExit, "unexpected PTO ISA 0.58.3 lock"): + CHECKER.load_lock(ptoas_root) + def test_mismatched_release_fails_closed(self): local_lock = json.loads( - Path(__file__).with_name("pto_isa_v0_58_1_lock.json").read_text() + Path(__file__).with_name("pto_isa_v0_58_3_lock.json").read_text() ) mismatched = json.loads(json.dumps(local_lock)) mismatched["release"] = "0.58.0" @@ -58,6 +72,102 @@ def test_pipe_v_only_mapping_is_not_accepted_as_linx_engine(self): operations = CHECKER.load_ptoas_ops(fake_root) self.assertIsNone(operations["TDIV"]["linx_engine"]) + def test_linx_target_surface_has_0583_cube_call_contracts(self): + ptoas_root = Path(__file__).resolve().parents[1] + CHECKER.validate_linx_target_surface(ptoas_root) + + def test_mx_scales_are_independently_optional_in_ods(self): + ptoas_root = Path(__file__).resolve().parents[1] + operations = CHECKER.load_ptoas_ops(ptoas_root) + for mnemonic in ( + "TGEMVMX", + "TGEMVMXACC", + "TGEMVMXBIAS", + "TMATMULMX", + "TMATMULMXACC", + "TMATMULMXBIAS", + ): + self.assertEqual( + operations[mnemonic]["optional_arguments"], + ("a_scale", "b_scale"), + ) + + def test_linx_mx_compile_gate_uses_real_target_inputs(self): + ptoas_root = Path(__file__).resolve().parents[1] + gate = (ptoas_root / "tools/check_linx_mx_tileop_compile.sh").read_text() + self.assertNotIn("linx_host_type_shim", gate) + self.assertNotIn("linx_mx_tileop_overlay", gate) + for required in ( + 'LINX_CXX=${LINX_LLVM_BUILD}/bin/clang++', + '"${TILEOP_ROOT}/include/jcore/template_asm.hpp"', + '"${TILEOP_ROOT}/test/tileop_api/verify_pto_identity.py"', + '--target=linx64-unknown-linux-musl', + '-fsyntax-only', + '-c "${generated}"', + "EXPECTED_LLVM_COMMIT", + "EXPECTED_LLVM_TREE", + "EXPECTED_TILEOP_COMMIT", + "EXPECTED_TILEOP_TREE", + "bash ./compile.all link-smoke", + ): + self.assertIn(required, gate) + + def test_misrouted_matmul_mx_variant_fails_closed(self): + ptoas_root = Path(__file__).resolve().parents[1] + lowering = (ptoas_root / "lib/PTO/Transforms/PTOToEmitC.cpp").read_text() + lowering, count = re.subn( + r'(struct\s+PTOTMatmulMXToTMATMUL_MX\b.*?\n\s*' + r'replaceOrEraseWithOpaqueCall\([^\n]*?)"TMATMUL_MX"', + r'\1"TMATMUL_MX_ACC"', + lowering, + count=1, + flags=re.DOTALL, + ) + self.assertEqual(count, 1) + with tempfile.TemporaryDirectory() as directory: + fake_root = Path(directory) + for relative in ( + "tools/ptoas/ptoas.cpp", + "tools/pto_isa_v0_58_3_operation_contracts.json", + ): + destination = fake_root / relative + destination.parent.mkdir(parents=True, exist_ok=True) + destination.write_text((ptoas_root / relative).read_text()) + destination = fake_root / "lib/PTO/Transforms/PTOToEmitC.cpp" + destination.parent.mkdir(parents=True, exist_ok=True) + destination.write_text(lowering) + with self.assertRaisesRegex(SystemExit, "PTOTMatmulMXToTMATMUL_MX"): + CHECKER.validate_linx_target_surface(fake_root) + + def test_misordered_tgemv_mx_optional_scale_fails_closed(self): + ptoas_root = Path(__file__).resolve().parents[1] + lowering = (ptoas_root / "lib/PTO/Transforms/PTOToEmitC.cpp").read_text() + lowering, count = re.subn( + r"SmallVector operands\{dst, b\};\s*" + r"if \(bScale\)\s*operands\.push_back\(bScale\);\s*" + r"operands\.push_back\(a\);", + "SmallVector operands{dst, b};\n" + " operands.push_back(a);\n" + " if (bScale) operands.push_back(bScale);", + lowering, + count=1, + ) + self.assertEqual(count, 1) + with tempfile.TemporaryDirectory() as directory: + fake_root = Path(directory) + for relative in ( + "tools/ptoas/ptoas.cpp", + "tools/pto_isa_v0_58_3_operation_contracts.json", + ): + destination = fake_root / relative + destination.parent.mkdir(parents=True, exist_ok=True) + destination.write_text((ptoas_root / relative).read_text()) + destination = fake_root / "lib/PTO/Transforms/PTOToEmitC.cpp" + destination.parent.mkdir(parents=True, exist_ok=True) + destination.write_text(lowering) + with self.assertRaisesRegex(SystemExit, "PTOTGemvMXToTGEMV_MX"): + CHECKER.validate_linx_target_surface(fake_root) + if __name__ == "__main__": unittest.main() diff --git a/tools/test_release_delivery_contract.py b/tools/test_release_delivery_contract.py index ee2d75a546..ecf4d81b83 100644 --- a/tools/test_release_delivery_contract.py +++ b/tools/test_release_delivery_contract.py @@ -18,6 +18,10 @@ from pathlib import Path from check_release_delivery_contract import ( + EXPECTED_LLVM_COMMIT, + EXPECTED_LLVM_TREE, + EXPECTED_TILEOP_COMMIT, + EXPECTED_TILEOP_TREE, local_copy_sources, run_packaging_identity_mode, validate_local_copy_sources, @@ -57,6 +61,45 @@ def test_documented_build_uses_repository_root_context(self) -> None: readme = (ROOT / "docker/README.md").read_text() self.assertIn("docker build -f docker/Dockerfile .", readme) + def test_docker_llvm_source_build_pins_nanobind_2_9(self) -> None: + dockerfile = (ROOT / "docker/Dockerfile").read_text() + self.assertIn("'nanobind>=2.9,<3'", dockerfile) + + def test_all_llvm_source_build_lanes_pin_nanobind_2_9(self) -> None: + paths = ( + ROOT / "docker/Dockerfile", + ROOT / ".github/workflows/ci.yml", + ROOT / ".github/workflows/build_wheel.yml", + ROOT / ".github/workflows/build_wheel_mac.yml", + ) + for path in paths: + with self.subTest(path=path.relative_to(ROOT)): + self.assertIn("nanobind>=2.9,<3", path.read_text()) + + def test_all_delivery_lanes_pin_merged_reviewed_llvm(self) -> None: + paths = ( + ROOT / "README.md", + ROOT / "docker/Dockerfile", + ROOT / ".github/workflows/ci.yml", + ROOT / ".github/workflows/build_wheel.yml", + ROOT / ".github/workflows/build_wheel_mac.yml", + ) + for path in paths: + with self.subTest(path=path.relative_to(ROOT)): + self.assertIn(EXPECTED_LLVM_COMMIT, path.read_text()) + + def test_target_integration_pins_merged_llvm_and_tileop_trees(self) -> None: + text = (ROOT / "CMakeLists.txt").read_text() + ( + ROOT / "README.md" + ).read_text() + for identity in ( + EXPECTED_LLVM_COMMIT, + EXPECTED_LLVM_TREE, + EXPECTED_TILEOP_COMMIT, + EXPECTED_TILEOP_TREE, + ): + self.assertIn(identity, text) + def test_hosted_builder_stage_gate_uses_buildkit(self) -> None: workflow = (ROOT / ".github/workflows/isa_contract.yml").read_text() self.assertIn("docker/setup-buildx-action@", workflow) @@ -80,22 +123,22 @@ def run( invalid_outputs = ( "ptoas 0.41", - "ptoas 0.40 (PTO ISA 0.58.1)", + "ptoas 0.40 (PTO ISA 0.58.3)", "ptoas 0.41 (PTO ISA 0.58.0)", - "warning\nptoas 0.41 (PTO ISA 0.58.1)", - "ptoas 0.41 (PTO ISA 0.58.1)\nptoas 0.40", + "warning\nptoas 0.41 (PTO ISA 0.58.3)", + "ptoas 0.41 (PTO ISA 0.58.3)\nptoas 0.40", ) for script in scripts: with self.subTest( script=script.name, output="valid", product_version="0.41" ): self.assertEqual( - run(script, "ptoas 0.41 (PTO ISA 0.58.1)", "0.41").returncode, + run(script, "ptoas 0.41 (PTO ISA 0.58.3)", "0.41").returncode, 0, ) with self.subTest(script=script.name, output="valid", product_version=""): self.assertEqual( - run(script, "ptoas 0.41 (PTO ISA 0.58.1)").returncode, 0 + run(script, "ptoas 0.41 (PTO ISA 0.58.3)").returncode, 0 ) for product_version in ("0.41", ""): for invalid in invalid_outputs: