Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
65 changes: 20 additions & 45 deletions .github/workflows/ci_sim.yml
Original file line number Diff line number Diff line change
Expand Up @@ -61,7 +61,6 @@ jobs:
- 'test/tilelang_st/**'
- 'test/dsl-st/**'
- 'test/tilelib-st/**'
- 'tilelang-dsl/**'
- 'ptodsl/**'
- 'tools/ptoas/**'
- 'lib/TileOps/**'
Expand Down Expand Up @@ -115,8 +114,7 @@ jobs:
PTO_WHEELHOUSE: ${{ github.workspace }}/.work/ptoas-wheelhouse
PTOAS_VENV: ${{ github.workspace }}/.work/ptoas-sim-venv
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
TILELIB_ST_WORKSPACE: ${{ github.workspace }}/.work/tilelib-st-ci
PYPTO_REF: ef6ce7cd8bd33b4c93b58dc34830a20b145736ac
PYPTO_WORKSPACE: ${{ github.workspace }}/.work/pypto-ci
PYPTO_RUN_WORKSPACE: ${{ github.workspace }}/.work/pypto-run-ci
Expand Down Expand Up @@ -253,7 +251,7 @@ jobs:
rm -rf "${PTO_WHEELHOUSE}"
rm -rf "${PTOAS_VENV}"
rm -rf "${VPTO_SIM_WORKSPACE}"
rm -rf "${TILELANG_DSL_WORKSPACE}"
rm -rf "${TILELIB_ST_WORKSPACE}"
rm -rf "${PYPTO_WORKSPACE}"
rm -rf "${PYPTO_RUN_WORKSPACE}"
rm -rf "${PTO_ISA_ROOT}"
Expand Down Expand Up @@ -703,23 +701,23 @@ jobs:
echo "PTO_DSL_ST_PYTHON_TAG=${PYTHON_TAG}" >> "${GITHUB_ENV}"
echo "PTO_DSL_ST_LLVM_DIR=${PTO_DSL_ST_LLVM_DIR}" >> "${GITHUB_ENV}"

- name: Run TileLang DSL CI
- name: Run TileOp ST CI
shell: bash
run: |
set -euo pipefail
mkdir -p "${TILELANG_DSL_WORKSPACE}"
mkdir -p "${TILELIB_ST_WORKSPACE}"
export LLVM_BUILD_DIR="${LLVM_DIR}"
export MLIR_PYTHON_ROOT="${MLIR_PYTHONPATH}"
if [[ "${{ github.event_name }}" == "pull_request" ]]; then
ASCEND_HOME_PATH="${ASCEND_HOME_PATH}" \
PTOAS_BIN="${PTOAS_BIN}" \
bash test/tilelang_st/script/run_ci.sh -r sim -v a5 --jobs 64 --smoke \
2>&1 | tee "${TILELANG_DSL_WORKSPACE}/run_ci.log"
2>&1 | tee "${TILELIB_ST_WORKSPACE}/run_ci.log"
else
ASCEND_HOME_PATH="${ASCEND_HOME_PATH}" \
PTOAS_BIN="${PTOAS_BIN}" \
bash test/tilelang_st/script/run_ci.sh -r sim -v a5 --jobs 64 \
2>&1 | tee "${TILELANG_DSL_WORKSPACE}/run_ci.log"
2>&1 | tee "${TILELIB_ST_WORKSPACE}/run_ci.log"
fi

- name: Restore PTODSL LLVM build cache
Expand Down Expand Up @@ -822,7 +820,7 @@ jobs:
shell: bash
run: |
set -euo pipefail
mkdir -p "${TILELANG_DSL_WORKSPACE}"
mkdir -p "${TILELIB_ST_WORKSPACE}"
export TORCH_DEVICE_BACKEND_AUTOLOAD=0
source "${ASCEND_HOME_PATH}/bin/setenv.bash"

Expand Down Expand Up @@ -850,7 +848,7 @@ jobs:
PY
}

PROBE_LOG="${TILELANG_DSL_WORKSPACE}/ptodsl-python-probe.log"
PROBE_LOG="${TILELIB_ST_WORKSPACE}/ptodsl-python-probe.log"
: > "${PROBE_LOG}"
echo "Probing PTODSL DSL ST Python: ${PTO_DSL_ST_PYTHON_BIN}" | tee -a "${PROBE_LOG}"
echo "Probing PTODSL DSL ST ptoas: ${PTO_DSL_ST_PTOAS_BIN}" | tee -a "${PROBE_LOG}"
Expand Down Expand Up @@ -895,26 +893,24 @@ jobs:
shell: bash
run: |
set -euo pipefail
mkdir -p "${TILELANG_DSL_WORKSPACE}"
mkdir -p "${TILELIB_ST_WORKSPACE}"
export PYTHON_BIN="${PTO_DSL_ST_PYTHON_BIN}"
export PTO_PYTHON_BIN="${PTO_DSL_ST_PYTHON_BIN}"
export PATH="${PTO_DSL_ST_BIN_DIR}:${PATH}"
export TORCH_DEVICE_BACKEND_AUTOLOAD=0
source "${ASCEND_HOME_PATH}/bin/setenv.bash"

TILELIB_ST_OUTPUT_ROOT="${TILELANG_DSL_WORKSPACE}/tilelib-st-a5" \
TILELIB_ST_OUTPUT_ROOT="${TILELIB_ST_WORKSPACE}/tilelib-st-a5" \
ASCEND_HOME_PATH="${ASCEND_HOME_PATH}" \
PTOAS_BIN="${PTO_DSL_ST_PTOAS_BIN}" \
"${PTO_DSL_ST_PYTHON_BIN}" -m pytest \
-n "${TILELIB_ST_JOBS:-16}" --dist=worksteal -v \
test/tilelib-st/test_tilelib_st.py \
2>&1 | tee "${TILELANG_DSL_WORKSPACE}/tilelib-st-a5.log"
scripts/sim_dsl.sh test/tilelib-st/run_tilelib_st.py -- test/tilelib-st/a5 \
2>&1 | tee "${TILELIB_ST_WORKSPACE}/tilelib-st-a5.log"

- name: Run PTODSL ST CI
shell: bash
run: |
set -euo pipefail
mkdir -p "${TILELANG_DSL_WORKSPACE}"
mkdir -p "${TILELIB_ST_WORKSPACE}"
export PYTHON_BIN="${PTO_DSL_ST_PYTHON_BIN}"
export PTO_PYTHON_BIN="${PTO_DSL_ST_PYTHON_BIN}"
export PATH="${PTO_DSL_ST_BIN_DIR}:${PATH}"
Expand All @@ -924,39 +920,18 @@ jobs:
ASCEND_HOME_PATH="${ASCEND_HOME_PATH}" \
PTOAS_BIN="${PTO_DSL_ST_PTOAS_BIN}" \
scripts/sim_dsl.sh test/dsl-st \
2>&1 | tee "${TILELANG_DSL_WORKSPACE}/ptodsl-dsl-st.log"
2>&1 | tee "${TILELIB_ST_WORKSPACE}/ptodsl-dsl-st.log"

- name: Upload TileLang DSL logs
- name: Upload TileLib ST logs
if: always()
uses: actions/upload-artifact@v4
with:
name: tilelang-dsl-ci-${{ github.run_id }}
name: tilelib-st-ci-${{ github.run_id }}
path: |
${{ env.TILELANG_DSL_WORKSPACE }}/run_ci.log
${{ env.TILELANG_DSL_WORKSPACE }}/ptodsl-dsl-st.log
${{ env.TILELANG_DSL_WORKSPACE }}/tilelib-st-a5.log
${{ env.TILELANG_DSL_WORKSPACE }}/tilelib-st-a5/**
${{ env.TILELANG_DSL_WORKSPACE }}/ptodsl-python-probe.log
if-no-files-found: warn

- name: Run TileLang DSL unit tests
shell: bash
run: |
set -euo pipefail
mkdir -p "${TILELANG_DSL_UT_WORKSPACE}"
export PATH="${PTOAS_VENV}/bin:${PATH}"
export LD_LIBRARY_PATH="${LLVM_DIR}/lib:${PTO_INSTALL_DIR}/lib:${GITHUB_WORKSPACE}/build/lib:${LD_LIBRARY_PATH:-}"
cd tilelang-dsl/python
python3 -m unittest discover -s ../tests -p 'test_*.py' \
2>&1 | tee "${TILELANG_DSL_UT_WORKSPACE}/unittest.log"

- name: Upload TileLang DSL unit test logs
if: always()
uses: actions/upload-artifact@v4
with:
name: tilelang-dsl-ut-ci-${{ github.run_id }}
path: |
${{ env.TILELANG_DSL_UT_WORKSPACE }}/unittest.log
${{ env.TILELIB_ST_WORKSPACE }}/run_ci.log
${{ env.TILELIB_ST_WORKSPACE }}/ptodsl-dsl-st.log
${{ env.TILELIB_ST_WORKSPACE }}/tilelib-st-a5.log
${{ env.TILELIB_ST_WORKSPACE }}/ptodsl-python-probe.log
if-no-files-found: warn

- name: Upload VPTO SIM logs
Expand Down
20 changes: 0 additions & 20 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -247,7 +247,6 @@ if(PTO_ENABLE_PYTHON_BINDING)

add_subdirectory(python)
add_subdirectory(ptodsl)
add_subdirectory(tilelang-dsl)
endif()

# =========================================================
Expand Down Expand Up @@ -283,25 +282,6 @@ if(BUILD_TESTING)
)

add_subdirectory(ptodsl/tests)
add_test(
NAME tilelang_dsl_import
COMMAND "${Python3_EXECUTABLE}"
"${CMAKE_CURRENT_SOURCE_DIR}/tilelang-dsl/tests/import_tilelang_dsl.py"
)
set_tests_properties(tilelang_dsl_import PROPERTIES
LABELS "PTODSL"
ENVIRONMENT "${_pto_python_test_env}"
)
add_test(
NAME tilelang_dsl_unittest
COMMAND "${Python3_EXECUTABLE}" -m unittest discover
-s "${CMAKE_CURRENT_SOURCE_DIR}/tilelang-dsl/tests"
-p "test_*.py"
)
set_tests_properties(tilelang_dsl_unittest PROPERTIES
LABELS "PTODSL"
ENVIRONMENT "${_pto_python_test_env}"
)
add_test(
NAME pto_memory_consistency_bindings
COMMAND "${Python3_EXECUTABLE}"
Expand Down
5 changes: 4 additions & 1 deletion _ptoas_build_backend.py
Original file line number Diff line number Diff line change
Expand Up @@ -63,6 +63,7 @@ def _find_llvm_dir():
)
_BUILD_DIR = Path(os.environ.get("PTO_BUILD_DIR", str(_REPO / "build")))
_PTODSL_SOURCE_ROOT = _REPO / "ptodsl"
_TILEOPS_SOURCE_ROOT = _REPO / "lib"
_MLIR_PY_PKG = None
if "MLIR_PYTHON_PACKAGE_DIR" in os.environ:
_MLIR_PY_PKG = Path(os.environ["MLIR_PYTHON_PACKAGE_DIR"])
Expand Down Expand Up @@ -303,9 +304,11 @@ def build_editable(wheel_directory, config_settings=None, metadata_directory=Non
# Prefer the repository PTODSL sources so editable installs pick up
# local Python edits instead of staged/install-tree copies.
str(_PTODSL_SOURCE_ROOT),
# PTODSL TileLib templates live in the top-level TileOps package.
str(_TILEOPS_SOURCE_ROOT),
# Handwritten Python sources (pto/dialects/pto.py, etc.).
str(_REPO / "python"),
# Installed PTOAS runtime overlay (TileOps/resources when present).
# Installed PTOAS runtime overlay.
str(_PTO_INSTALL_DIR),
# Generated files (_pto.so, _pto_ops_gen.py) under mlir/ namespace.
str(_BUILD_DIR / "python"),
Expand Down
16 changes: 3 additions & 13 deletions docker/collect_ptoas_dist.sh
Original file line number Diff line number Diff line change
Expand Up @@ -24,8 +24,7 @@
# ptoas_wheel_bootstrap.py - Compatibility bootstrap module
# lib/ptoas.so - Shared runtime loaded by the launcher
# lib/*.so* - Required shared library dependencies
# share/ptoas/TileOps - TileLang template library
# tilelang_dsl/ - TileLang DSL Python package
# share/ptoas/TileOps - PTODSL TileOps template package
#
# This compiler-oriented binary artifact intentionally does not bundle the
# PTODSL Python package. PTODSL is provided only by the `ptoas` wheel.
Expand Down Expand Up @@ -56,8 +55,6 @@ PTOAS_DEPS_DIR="${PTOAS_DIST_DIR}/lib"
PTOAS_SHARED_MODULE_DIST_PATH="${PTOAS_DEPS_DIR}/ptoas.so"
PTOAS_TILEOPS_SRC_DIR="${PTO_INSTALL_DIR}/share/ptoas/TileOps"
PTOAS_TILEOPS_DIST_DIR="${PTOAS_DIST_DIR}/share/ptoas/TileOps"
PTOAS_TILELANG_DSL_SRC_DIR="${PTO_INSTALL_DIR}/tilelang_dsl"
PTOAS_TILELANG_DSL_DIST_DIR="${PTOAS_DIST_DIR}/tilelang_dsl"
PTOAS_WRAPPER_PKG_SRC_DIR="${PTO_INSTALL_DIR}/ptoas"
PTOAS_WRAPPER_PKG_DIST_DIR="${PTOAS_DIST_DIR}/ptoas"
PTOAS_WHEEL_BOOTSTRAP_SRC="${PTO_INSTALL_DIR}/ptoas_wheel_bootstrap.py"
Expand Down Expand Up @@ -227,18 +224,13 @@ while read -r packaged; do
harden_elf "$packaged"
done < <(find "${PTOAS_DEPS_DIR}" -type f | sort)

echo "Copying TileLang runtime resources..."
echo "Copying TileOps runtime resources..."
if [ ! -d "${PTOAS_TILEOPS_SRC_DIR}" ]; then
echo "Error: TileOps resource directory not found at ${PTOAS_TILEOPS_SRC_DIR}" >&2
exit 1
fi
if [ ! -d "${PTOAS_TILELANG_DSL_SRC_DIR}" ]; then
echo "Error: tilelang_dsl package directory not found at ${PTOAS_TILELANG_DSL_SRC_DIR}" >&2
exit 1
fi
rm -rf "${PTOAS_TILEOPS_DIST_DIR}" "${PTOAS_TILELANG_DSL_DIST_DIR}"
rm -rf "${PTOAS_TILEOPS_DIST_DIR}"
cp -R "${PTOAS_TILEOPS_SRC_DIR}" "${PTOAS_TILEOPS_DIST_DIR}"
cp -R "${PTOAS_TILELANG_DSL_SRC_DIR}" "${PTOAS_TILELANG_DSL_DIST_DIR}"

echo "Smoke testing packaged ptoas dist..."
VERSION_OUTPUT="$(env -u PYTHONPATH -u DYLD_LIBRARY_PATH -u LD_LIBRARY_PATH \
Expand All @@ -256,7 +248,6 @@ else
fi

test -d "${PTOAS_TILEOPS_DIST_DIR}"
test -f "${PTOAS_TILELANG_DSL_DIST_DIR}/__init__.py"
test -f "${PTOAS_SHARED_MODULE_DIST_PATH}"
test -f "${PTOAS_WRAPPER_PKG_DIST_DIR}/_runtime_entry.py"
test -f "${PTOAS_WHEEL_BOOTSTRAP_DIST_PATH}"
Expand All @@ -271,7 +262,6 @@ echo "=== ptoas distribution contents ==="
ls -la "${PTOAS_DIST_DIR}/"
ls -la "${PTOAS_DIST_DIR}/bin/"
ls -la "${PTOAS_DIST_DIR}/share/ptoas/"
ls -la "${PTOAS_TILELANG_DSL_DIST_DIR}"
SO_COUNT=$(find "${PTOAS_DEPS_DIR}" -name "*.so*" 2>/dev/null | wc -l)
echo "=== Collected .so dependencies (${SO_COUNT} files) ==="
du -sh "${PTOAS_DEPS_DIR}/"
Expand Down
16 changes: 3 additions & 13 deletions docker/collect_ptoas_dist_mac.sh
Original file line number Diff line number Diff line change
Expand Up @@ -24,8 +24,7 @@
# ptoas_wheel_bootstrap.py - Compatibility bootstrap module
# lib/ptoas.so - Shared runtime loaded by the launcher
# lib/*.dylib - Required shared library dependencies
# share/ptoas/TileOps - TileLang template library
# tilelang_dsl/ - TileLang DSL Python package
# share/ptoas/TileOps - PTODSL TileOps template package
#
# This compiler-oriented binary artifact intentionally does not bundle the
# PTODSL Python package. PTODSL is provided only by the `ptoas` wheel.
Expand Down Expand Up @@ -54,8 +53,6 @@ PTOAS_DEPS_DIR="${PTOAS_DIST_DIR}/lib"
PTOAS_SHARED_MODULE_DIST_PATH="${PTOAS_DEPS_DIR}/ptoas.so"
PTOAS_TILEOPS_SRC_DIR="${PTO_INSTALL_DIR}/share/ptoas/TileOps"
PTOAS_TILEOPS_DIST_DIR="${PTOAS_DIST_DIR}/share/ptoas/TileOps"
PTOAS_TILELANG_DSL_SRC_DIR="${PTO_INSTALL_DIR}/tilelang_dsl"
PTOAS_TILELANG_DSL_DIST_DIR="${PTOAS_DIST_DIR}/tilelang_dsl"
PTOAS_WRAPPER_PKG_SRC_DIR="${PTO_INSTALL_DIR}/ptoas"
PTOAS_WRAPPER_PKG_DIST_DIR="${PTOAS_DIST_DIR}/ptoas"
PTOAS_WHEEL_BOOTSTRAP_SRC="${PTO_INSTALL_DIR}/ptoas_wheel_bootstrap.py"
Expand Down Expand Up @@ -274,18 +271,13 @@ PY
echo "Collecting dylib dependencies..."
collect_dylibs "${PTOAS_SHARED_MODULE_DIST_PATH}"

echo "Copying TileLang runtime resources..."
echo "Copying TileOps runtime resources..."
if [[ ! -d "${PTOAS_TILEOPS_SRC_DIR}" ]]; then
echo "Error: TileOps resource directory not found at ${PTOAS_TILEOPS_SRC_DIR}" >&2
exit 1
fi
if [[ ! -d "${PTOAS_TILELANG_DSL_SRC_DIR}" ]]; then
echo "Error: tilelang_dsl package directory not found at ${PTOAS_TILELANG_DSL_SRC_DIR}" >&2
exit 1
fi
rm -rf "${PTOAS_TILEOPS_DIST_DIR}" "${PTOAS_TILELANG_DSL_DIST_DIR}"
rm -rf "${PTOAS_TILEOPS_DIST_DIR}"
cp -R "${PTOAS_TILEOPS_SRC_DIR}" "${PTOAS_TILEOPS_DIST_DIR}"
cp -R "${PTOAS_TILELANG_DSL_SRC_DIR}" "${PTOAS_TILELANG_DSL_DIST_DIR}"
if [[ -e "${PTOAS_DIST_DIR}/ptodsl" ]]; then
echo "Error: compiler-oriented ptoas dist must not bundle PTODSL" >&2
exit 1
Expand Down Expand Up @@ -376,7 +368,6 @@ else
echo "$VERSION_OUTPUT" | grep -Eq '^ptoas [0-9]+\.[0-9]+$'
fi
test -d "${PTOAS_TILEOPS_DIST_DIR}"
test -f "${PTOAS_TILELANG_DSL_DIST_DIR}/__init__.py"
test -f "${PTOAS_SHARED_MODULE_DIST_PATH}"
test -f "${PTOAS_WRAPPER_PKG_DIST_DIR}/_runtime_entry.py"
test -f "${PTOAS_WHEEL_BOOTSTRAP_DIST_PATH}"
Expand All @@ -390,7 +381,6 @@ echo "=== ptoas distribution contents ==="
ls -la "${PTOAS_DIST_DIR}/"
ls -la "${PTOAS_DIST_DIR}/bin/"
ls -la "${PTOAS_DIST_DIR}/share/ptoas/"
ls -la "${PTOAS_TILELANG_DSL_DIST_DIR}"
DYLIB_COUNT=$(find "${PTOAS_DEPS_DIR}" -name "*.dylib" 2>/dev/null | wc -l)
echo "=== Collected .dylib dependencies (${DYLIB_COUNT} files) ==="
du -sh "${PTOAS_DEPS_DIR}/"
Expand Down
5 changes: 2 additions & 3 deletions docker/create_wheel.sh
Original file line number Diff line number Diff line change
Expand Up @@ -139,9 +139,8 @@ echo "Overlaying PTO native extension..."
mkdir -p "${WHEEL_STAGING_DIR}/mlir/_mlir_libs"
find "${PTO_INSTALL_DIR}/mlir/_mlir_libs" -maxdepth 1 -type f -name '_pto*' -exec cp {} "${WHEEL_STAGING_DIR}/mlir/_mlir_libs/" \;

echo "Copying TileLang resources..."
rm -rf "${WHEEL_STAGING_DIR}/tilelang_dsl" "${WHEEL_STAGING_DIR}/TileOps"
cp -R "${PTO_INSTALL_DIR}/tilelang_dsl" "${WHEEL_STAGING_DIR}/tilelang_dsl"
echo "Copying TileOps resources..."
rm -rf "${WHEEL_STAGING_DIR}/TileOps"
cp -R "${PTO_INSTALL_DIR}/share/ptoas/TileOps" "${WHEEL_STAGING_DIR}/TileOps"

echo "Copying ptodsl package..."
Expand Down
4 changes: 2 additions & 2 deletions lib/PTO/Transforms/PTOToEmitC.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -6465,7 +6465,7 @@ struct PTOSyncSetToEmitC : public OpConversionPattern<mlir::pto::SyncSetOp> {
return success();
}

[[maybe_unused]] PTOArch targetArch;
PTOArch targetArch;
};

struct PTOSyncWaitToEmitC : public OpConversionPattern<mlir::pto::SyncWaitOp> {
Expand Down Expand Up @@ -6500,7 +6500,7 @@ struct PTOSyncWaitToEmitC : public OpConversionPattern<mlir::pto::SyncWaitOp> {
return success();
}

[[maybe_unused]] PTOArch targetArch;
PTOArch targetArch;
};

// GetBlockIdxOp Lowering (pto.get_block_idx -> get_block_idx())
Expand Down
Loading
Loading