feat: support ptodsl

.github/workflows/build_wheel.yml

@@ -159,22 +159,8 @@ jobs:
       - name: Build PTOAS
         run: |
           export PATH="${PY_PATH}/bin:$PATH"
-          cd $PTO_SOURCE_DIR
-          cmake -C "$PTO_SOURCE_DIR/cmake/LinuxHardeningCache.cmake" -G Ninja \
-            -S . \
-            -B $PTO_BUILD_DIR \
-            -DLLVM_DIR=$LLVM_BUILD_DIR/lib/cmake/llvm \
-            -DMLIR_DIR=$LLVM_BUILD_DIR/lib/cmake/mlir \
-            -DPython3_ROOT_DIR=${PY_PATH} \
-            -DPython3_EXECUTABLE=${PY_PATH}/bin/python \
-            -DPython3_FIND_STRATEGY=LOCATION \
-            -Dpybind11_DIR=$(${PY_PATH}/bin/python -m pybind11 --cmakedir) \
-            -DMLIR_PYTHON_PACKAGE_DIR=${LLVM_BUILD_DIR}/tools/mlir/python_packages/mlir_core \
-            -DPTOAS_RELEASE_VERSION_OVERRIDE=${PTOAS_VERSION} \
-            -DCMAKE_INSTALL_PREFIX=${PTO_INSTALL_DIR} \
-            -DCMAKE_BUILD_TYPE=Release
-          ninja -C $PTO_BUILD_DIR
-          ninja -C $PTO_BUILD_DIR install
+          PTOAS_RELEASE_VERSION_OVERRIDE="${PTOAS_VERSION}" \
+          pip install . --no-build-isolation
 
       - name: Create Python wheel
         if: false
@@ -187,17 +173,17 @@ jobs:
         if: false
         run: |
           export PATH="${PY_PATH}/bin:$PATH"
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
+          export PTO_WHEEL_DIST_DIR=$PTO_SOURCE_DIR/build/wheel-dist
+          export PTO_WHEELHOUSE=$GITHUB_WORKSPACE/wheelhouse
           export LD_LIBRARY_PATH=$LLVM_BUILD_DIR/lib:$PTO_INSTALL_DIR/lib:$LD_LIBRARY_PATH
-          cd $PY_PACKAGE_DIR
-          auditwheel repair --plat manylinux_2_34_${{ matrix.arch }} dist/ptoas*.whl -w wheelhouse
+          mkdir -p "$PTO_WHEELHOUSE"
+          auditwheel repair --plat manylinux_2_34_${{ matrix.arch }} "$PTO_WHEEL_DIST_DIR"/ptoas*.whl -w "$PTO_WHEELHOUSE"
 
       - name: Test wheel installation
         if: false
         run: |
           export PATH="${PY_PATH}/bin:$PATH"
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
-          pip install $PY_PACKAGE_DIR/wheelhouse/ptoas*.whl
+          pip install $GITHUB_WORKSPACE/wheelhouse/ptoas*.whl
           bash $PTO_SOURCE_DIR/docker/test_wheel_imports.sh
 
       - name: Test ptoas CLI
@@ -208,9 +194,10 @@ jobs:
       - name: Copy wheel to workspace
         if: false
         run: |
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
           mkdir -p $GITHUB_WORKSPACE/wheelhouse
-          cp $PY_PACKAGE_DIR/wheelhouse/ptoas*.whl $GITHUB_WORKSPACE/wheelhouse/
+          if ! compgen -G "$GITHUB_WORKSPACE/wheelhouse/ptoas*.whl" >/dev/null; then
+            cp $PTO_SOURCE_DIR/build/wheel-dist/ptoas*.whl $GITHUB_WORKSPACE/wheelhouse/
+          fi
 
       - name: Upload wheel artifact
         if: false

.github/workflows/build_wheel_mac.yml

@@ -157,27 +157,13 @@ jobs:
 
       - name: Build PTOAS
         run: |
-          cd $PTO_SOURCE_DIR
-          cmake -G Ninja \
-            -S . \
-            -B $PTO_BUILD_DIR \
-            -DLLVM_DIR=$LLVM_BUILD_DIR/lib/cmake/llvm \
-            -DMLIR_DIR=$LLVM_BUILD_DIR/lib/cmake/mlir \
-            -DPython3_ROOT_DIR=${PY_PATH} \
-            -DPython3_EXECUTABLE=$(which python) \
-            -DPython3_FIND_STRATEGY=LOCATION \
-            -Dpybind11_DIR=$(python -m pybind11 --cmakedir) \
-            -DMLIR_PYTHON_PACKAGE_DIR=${LLVM_BUILD_DIR}/tools/mlir/python_packages/mlir_core \
-            -DPTOAS_RELEASE_VERSION_OVERRIDE=${PTOAS_VERSION} \
-            -DCMAKE_INSTALL_PREFIX=${PTO_INSTALL_DIR} \
-            -DCMAKE_BUILD_TYPE=Release
-          ninja -C $PTO_BUILD_DIR
-          ninja -C $PTO_BUILD_DIR install
+          PTOAS_RELEASE_VERSION_OVERRIDE="${PTOAS_VERSION}" \
+          pip install . --no-build-isolation
 
       - name: Create Python wheel
         if: false
         run: |
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
+          export PTO_WHEEL_DIST_DIR=$PTO_SOURCE_DIR/build/wheel-dist
           export PTOAS_PYTHON_PACKAGE_VERSION="${PTOAS_VERSION}"
           if [ "${{ matrix.arch }}" = "x86_64" ]; then
             TARGET_ARCH="x86_64"
@@ -189,7 +175,7 @@ jobs:
           export WHEEL_PLAT_NAME=$(python "$PTO_SOURCE_DIR/docker/get_macos_wheel_plat_name.py" "$TARGET_ARCH")
           bash $PTO_SOURCE_DIR/docker/create_wheel.sh
           shopt -s nullglob
-          built_wheels=("$PY_PACKAGE_DIR"/dist/ptoas*.whl)
+          built_wheels=("$PTO_WHEEL_DIST_DIR"/ptoas*.whl)
           if [ "${#built_wheels[@]}" -gt 0 ]; then
             printf 'Built wheel file: %s\n' "$(basename "${built_wheels[0]}")"
           fi
@@ -198,17 +184,18 @@ jobs:
         if: false
         run: |
           set -euo pipefail
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
-          cd $PY_PACKAGE_DIR
+          export PTO_WHEEL_DIST_DIR=$PTO_SOURCE_DIR/build/wheel-dist
+          export PTO_WHEELHOUSE=$GITHUB_WORKSPACE/wheelhouse
+          mkdir -p "$PTO_WHEELHOUSE"
           if [ "${{ matrix.arch }}" = "x86_64" ]; then
             REQUIRED_ARCH="x86_64"
           else
             REQUIRED_ARCH="arm64"
           fi
           shopt -s nullglob
-          wheels=(dist/ptoas*.whl)
+          wheels=("$PTO_WHEEL_DIST_DIR"/ptoas*.whl)
           if [ "${#wheels[@]}" -ne 1 ]; then
-            echo "Expected exactly one wheel in dist/, found ${#wheels[@]}:"
+            echo "Expected exactly one wheel in $PTO_WHEEL_DIST_DIR, found ${#wheels[@]}:"
             printf '  %s\n' "${wheels[@]}"
             exit 1
           fi
@@ -217,13 +204,13 @@ jobs:
           echo "Wheel size:"
           du -sh "${wheels[0]}"
           echo "dist contents:"
-          ls -lh dist
+          ls -lh "$PTO_WHEEL_DIST_DIR"
           echo "delocate dependency summary:"
           delocate-listdeps "${wheels[0]}" || true
 
           delocate-wheel \
             --require-archs "${REQUIRED_ARCH}" \
-            --wheel-dir wheelhouse \
+            --wheel-dir "$PTO_WHEELHOUSE" \
             "${wheels[0]}" &
           delocate_pid=$!
 
@@ -241,7 +228,7 @@ jobs:
           while kill -0 "$delocate_pid" 2>/dev/null; do
             echo "[$(date '+%Y-%m-%d %H:%M:%S')] delocate still running (pid=${delocate_pid})"
             ps -o pid=,ppid=,%cpu=,%mem=,etime=,command= -p "$delocate_pid" || true
-            du -sh dist wheelhouse 2>/dev/null || true
+            du -sh "$PTO_WHEEL_DIST_DIR" "$PTO_WHEELHOUSE" 2>/dev/null || true
             sleep 60
           done
 
@@ -260,20 +247,19 @@ jobs:
             echo "Filesystem usage:"
             df -h .
             echo "dist contents after failure:"
-            ls -lh dist || true
+            ls -lh "$PTO_WHEEL_DIST_DIR" || true
             echo "wheelhouse contents after failure:"
-            ls -lh wheelhouse || true
+            ls -lh "$PTO_WHEELHOUSE" || true
             exit "$rc"
           fi
 
           echo "wheelhouse contents:"
-          ls -lh wheelhouse
+          ls -lh "$PTO_WHEELHOUSE"
 
       - name: Test wheel installation
         if: false
         run: |
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
-          pip install $PY_PACKAGE_DIR/wheelhouse/ptoas*.whl
+          pip install $GITHUB_WORKSPACE/wheelhouse/ptoas*.whl
           export DYLD_LIBRARY_PATH=$LLVM_BUILD_DIR/lib:$PTO_INSTALL_DIR/lib:${DYLD_LIBRARY_PATH}
           bash $PTO_SOURCE_DIR/docker/test_wheel_imports.sh
 
@@ -285,9 +271,10 @@ jobs:
       - name: Copy wheel to workspace
         if: false
         run: |
-          export PY_PACKAGE_DIR=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
           mkdir -p $GITHUB_WORKSPACE/wheelhouse
-          cp $PY_PACKAGE_DIR/wheelhouse/ptoas*.whl $GITHUB_WORKSPACE/wheelhouse/
+          if ! compgen -G "$GITHUB_WORKSPACE/wheelhouse/ptoas*.whl" >/dev/null; then
+            cp $PTO_SOURCE_DIR/build/wheel-dist/ptoas*.whl $GITHUB_WORKSPACE/wheelhouse/
+          fi
 
       - name: Upload wheel artifact
         if: false

.github/workflows/ci.yml

@@ -204,20 +204,24 @@ jobs:
 
       - name: Build PTOAS
         run: |
-          export PYBIND11_CMAKE_DIR="$(python3 -m pybind11 --cmakedir)"
-          cmake -C "${GITHUB_WORKSPACE}/cmake/LinuxHardeningCache.cmake" -G Ninja -S . -B "${PTO_BUILD_DIR}" \
-            -DLLVM_DIR="${LLVM_DIR}/lib/cmake/llvm" \
-            -DMLIR_DIR="${LLVM_DIR}/lib/cmake/mlir" \
-            -DPython3_EXECUTABLE=python3 \
-            -DPython3_FIND_STRATEGY=LOCATION \
-            -Dpybind11_DIR="${PYBIND11_CMAKE_DIR}" \
-            -DMLIR_ENABLE_BINDINGS_PYTHON=ON \
-            -DMLIR_PYTHON_PACKAGE_DIR="${LLVM_DIR}/tools/mlir/python_packages/mlir_core" \
-            -DCMAKE_INSTALL_PREFIX="${PTO_INSTALL_DIR}" \
-            -DCMAKE_BUILD_TYPE=Debug
-          ninja -C "${PTO_BUILD_DIR}" ptoas
-          ninja -C "${PTO_BUILD_DIR}" ptobc
-          ninja -C "${PTO_BUILD_DIR}" install
+          rm -rf "${PTO_INSTALL_DIR}"
+          # LLVM_BUILD_DIR is the env var read by the build backend (_ptoas_build_backend.py).
+          LLVM_BUILD_DIR="${LLVM_DIR}" \
+            PTO_INSTALL_DIR="${PTO_INSTALL_DIR}" \
+            python3 -m pip install . --no-build-isolation --no-deps --ignore-installed --prefix "${PTO_INSTALL_DIR}"
+
+      - name: Run PTODSL Python tests
+        shell: bash
+        env:
+          LLVM_DIR: ${{ env.LLVM_DIR }}
+          PTO_INSTALL_DIR: ${{ env.PTO_INSTALL_DIR }}
+          MLIR_PYTHONPATH: ${{ env.MLIR_PYTHONPATH }}
+        run: |
+          PTOAS_PYTHON_SITE="$(PTO_INSTALL_DIR="${PTO_INSTALL_DIR}" python3 -c 'import os, sysconfig; prefix = os.environ["PTO_INSTALL_DIR"]; print(sysconfig.get_path("purelib", vars={"base": prefix, "platbase": prefix}))')"
+          export PATH="${LLVM_DIR}/bin:${PTO_INSTALL_DIR}/bin:${PATH}"
+          export LD_LIBRARY_PATH="${LLVM_DIR}/lib:${PTO_INSTALL_DIR}/lib:${LD_LIBRARY_PATH:-}"
+          export PYTHONPATH="${PTOAS_PYTHON_SITE}:${PTO_INSTALL_DIR}:${MLIR_PYTHONPATH}:${PYTHONPATH:-}"
+          ctest --test-dir "${PTO_BUILD_DIR}" --output-on-failure -L PTODSL
 
       - name: Run lit tests
         shell: bash
@@ -226,9 +230,10 @@ jobs:
           PTO_INSTALL_DIR: ${{ env.PTO_INSTALL_DIR }}
           MLIR_PYTHONPATH: ${{ env.MLIR_PYTHONPATH }}
         run: |
+          PTOAS_PYTHON_SITE="$(PTO_INSTALL_DIR="${PTO_INSTALL_DIR}" python3 -c 'import os, sysconfig; prefix = os.environ["PTO_INSTALL_DIR"]; print(sysconfig.get_path("purelib", vars={"base": prefix, "platbase": prefix}))')"
           export PATH="${LLVM_DIR}/bin:${PTO_INSTALL_DIR}/bin:${PATH}"
           export LD_LIBRARY_PATH="${LLVM_DIR}/lib:${PTO_INSTALL_DIR}/lib:${LD_LIBRARY_PATH:-}"
-          export PYTHONPATH="${MLIR_PYTHONPATH}:${PTO_INSTALL_DIR}:${PYTHONPATH:-}"
+          export PYTHONPATH="${PTOAS_PYTHON_SITE}:${PTO_INSTALL_DIR}:${MLIR_PYTHONPATH}:${PYTHONPATH:-}"
           ninja -C "${PTO_BUILD_DIR}" check-pto
 
       - name: Run sample tests (py -> pto -> cpp)
@@ -241,11 +246,12 @@ jobs:
           PTO_BUILD_DIR: ${{ env.PTO_BUILD_DIR }}
           PYTHON_BIN: /usr/bin/python3
           MLIR_PYTHON_ROOT: ${{ env.MLIR_PYTHONPATH }}
-          PTO_PYTHON_ROOT: ${{ env.PTO_INSTALL_DIR }}/
+          PTO_INSTALL_DIR: ${{ env.PTO_INSTALL_DIR }}
         run: |
           set -euo pipefail
+          PTOAS_PYTHON_SITE="$(PTO_INSTALL_DIR="${PTO_INSTALL_DIR}" python3 -c 'import os, sysconfig; prefix = os.environ["PTO_INSTALL_DIR"]; print(sysconfig.get_path("purelib", vars={"base": prefix, "platbase": prefix}))')"
           export PATH="${PTO_INSTALL_DIR}/bin:${PATH}"
-          export PYTHONPATH="${MLIR_PYTHON_ROOT}:${PTO_PYTHON_ROOT}:${PYTHONPATH:-}"
+          export PYTHONPATH="${PTOAS_PYTHON_SITE}:${PTO_INSTALL_DIR}:${MLIR_PYTHON_ROOT}:${PYTHONPATH:-}"
           export LD_LIBRARY_PATH="${LLVM_DIR}/lib:${PTO_INSTALL_DIR}/lib:${LD_LIBRARY_PATH:-}"
           export PTOAS_OUT_DIR="${PAYLOAD_DIR}/test/samples"
           if [[ "${CI_EVENT_NAME}" == "workflow_dispatch" || "${CI_EVENT_NAME}" == "schedule" ]]; then
@@ -300,6 +306,7 @@ jobs:
     env:
       LLVM_REPO: https://github.com/vpto-dev/llvm-project.git
       LLVM_REF: feature-vpto
+      PTO_INSTALL_DIR: ${{ github.workspace }}/install
       VPTO_SIM_WORKSPACE: ${{ github.workspace }}/.work/vpto-sim-ci
       TILELANG_DSL_WORKSPACE: ${{ github.workspace }}/.work/tilelang-dsl-ci
       TILELANG_DSL_UT_WORKSPACE: ${{ github.workspace }}/.work/tilelang-dsl-ut-ci
@@ -367,6 +374,7 @@ jobs:
         run: |
           set -euo pipefail
           rm -rf "${GITHUB_WORKSPACE}/build"
+          rm -rf "${PTO_INSTALL_DIR}"
           rm -rf "${VPTO_SIM_WORKSPACE}"
           rm -rf "${TILELANG_DSL_WORKSPACE}"
 
@@ -405,17 +413,11 @@ jobs:
         shell: bash
         run: |
           set -euo pipefail
-          export PYBIND11_CMAKE_DIR="$(python3 -m pybind11 --cmakedir)"
-          cmake -G Ninja -S . -B build \
-            -DLLVM_DIR="${LLVM_DIR}/lib/cmake/llvm" \
-            -DMLIR_DIR="${LLVM_DIR}/lib/cmake/mlir" \
-            -DPython3_EXECUTABLE=python3 \
-            -DPython3_FIND_STRATEGY=LOCATION \
-            -Dpybind11_DIR="${PYBIND11_CMAKE_DIR}" \
-            -DMLIR_ENABLE_BINDINGS_PYTHON=ON \
-            -DMLIR_PYTHON_PACKAGE_DIR="${LLVM_DIR}/tools/mlir/python_packages/mlir_core" \
-            -DCMAKE_BUILD_TYPE=Release
-          ninja -C build ptoas
+          rm -rf "${PTO_INSTALL_DIR}"
+          # LLVM_BUILD_DIR is the env var read by the build backend (_ptoas_build_backend.py).
+          LLVM_BUILD_DIR="${LLVM_DIR}" \
+            PTO_INSTALL_DIR="${PTO_INSTALL_DIR}" \
+            python3 -m pip install . --no-build-isolation --no-deps --ignore-installed --prefix "${PTO_INSTALL_DIR}"
 
       - name: Resolve simulator environment
         shell: bash
@@ -475,13 +477,26 @@ jobs:
               2>&1 | tee "${TILELANG_DSL_WORKSPACE}/run_ci.log"
           fi
 
+      - name: Run PTODSL DSL ST CI
+        shell: bash
+        run: |
+          set -euo pipefail
+          mkdir -p "${TILELANG_DSL_WORKSPACE}"
+          export LLVM_BUILD_DIR="${LLVM_DIR}"
+          export PYTHON_BIN="python3"
+          ASCEND_HOME_PATH="${ASCEND_HOME_PATH}" \
+          PTOAS_BIN="${PTOAS_BIN}" \
+          scripts/sim_dsl.sh test/dsl-st/predicate_pack_launch.py \
+            2>&1 | tee "${TILELANG_DSL_WORKSPACE}/ptodsl-dsl-st.log"
+
       - name: Upload TileLang DSL logs
         if: always()
         uses: actions/upload-artifact@v4
         with:
           name: tilelang-dsl-ci-${{ github.run_id }}
           path: |
             ${{ env.TILELANG_DSL_WORKSPACE }}/run_ci.log
+            ${{ env.TILELANG_DSL_WORKSPACE }}/ptodsl-dsl-st.log
           if-no-files-found: warn
 
       - name: Run TileLang DSL unit tests

.gitignore

@@ -1,6 +1,13 @@
 # Build artifacts
 build/
+build_plain/
+build_plan/
 install/
+
+# TileLang ST standalone build outputs (see temp_docs/standalone_st.md)
+test/tilelang_st/npu/a5/src/st/build/
+test/tilelang_st/npu/a5/src/st/build_plain/
+test/tilelang_st/npu/a5/src/st/build_plan/
 cmake-build-*/
 CMakeFiles/
 CMakeCache.txt
@@ -49,13 +56,21 @@ venv/
 dist/
 
 # Logs/temp
+tmp/
 *.log
 *.tmp
 *.swp
 *.swo
 .cache/
 .pytest_cache/
 
+# PTODSL JIT / msprof simulator artifacts
+.ptodsl_jit/
+.ptodsl_cache/
+msprof_res/
+ptodsl/examples/jit/.cache/
+ptodsl/examples/jit/msprof_res/
+
 # Remote/NPU validation artifacts
 /payload/
 /payload.tgz

CMakeLists.txt

@@ -133,6 +133,7 @@ add_subdirectory(tools)
 if(BUILD_TESTING)
   enable_testing()
   if(PTO_ENABLE_PYTHON_BINDING)
+    add_subdirectory(ptodsl/tests)
     add_test(
       NAME tilelang_dsl_import
       COMMAND "${Python3_EXECUTABLE}"

README.md

@@ -149,16 +149,12 @@ $PTO_SOURCE_DIR/build/python/
 │       ├── pto.py
 │       └── _pto_ops_gen.py
 
-# install 输出（Python 方言文件）
+# install 输出（Python 方言文件和原生扩展）
 $PTO_INSTALL_DIR/
 └── mlir
-    └── dialects
-        ├── pto.py
-        └── _pto_ops_gen.py
-
-# 安装到 MLIR Python 包中的原生扩展
-$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core/
-└── mlir
+    ├── dialects
+    │   ├── pto.py
+    │   └── _pto_ops_gen.py
     └── _mlir_libs
         └── _pto.cpython-*.so
 
@@ -181,7 +177,7 @@ $PTO_SOURCE_DIR/build/tools/ptobc/ptobc
 #    这样在 python 中 import mlir.dialects.pto 时能正确找到
 export MLIR_PYTHON_ROOT=$LLVM_BUILD_DIR/tools/mlir/python_packages/mlir_core
 export PTO_PYTHON_ROOT=$PTO_INSTALL_DIR/
-export PYTHONPATH=$MLIR_PYTHON_ROOT:$PTO_PYTHON_ROOT:$PYTHONPATH
+export PYTHONPATH=$PTO_PYTHON_ROOT:$MLIR_PYTHON_ROOT:$PYTHONPATH
 
 # 2. Library Path: 确保能加载 LLVM 和 PTO 的动态库
 export LD_LIBRARY_PATH=$LLVM_BUILD_DIR/lib:$PTO_INSTALL_DIR/lib:$LD_LIBRARY_PATH