Files
tilelang/pyproject.toml
994b44eca1 [Ascend] Add PTO backend for Ascend950 (#3310)
* feat(pto): add Ascend PTO backend

Add PTO code generation, target registration, PTODSL helpers and JIT
compilation support. Share Python source generation with CuTeDSL.

Co-authored-by: cj <erhsh_165@126.com>
Co-authored-by: liggest <43201720+liggest@users.noreply.github.com>

* feat(pto): pto support c:v is 1:1 (#482)

Co-authored-by: caojian5 <caojian5@huawei.com>

* refactor(pto): clean up and simplify PTO codegen (#483)

- remove unused PTO codegen helpers and fragment state
- share PTO local variable initialization
- simplify PTO cast dispatch

Co-authored-by: caojian5 <caojian5@huawei.com>

* fix(pto): consume per-op SFU precision annotations in PTO codegen (#484)

Resolve per-op precision annotations in PTO code generation and port
the subnormal-preserving exponential, logarithm and square-root helpers.
Reject unsupported legalized unary SIMD merging forms.

Co-authored-by: caojian5 <caojian5@huawei.com>

* feat(pto): support SIMD MODE_MERGING via vsel composition (#480)

Consume the explicit destination produced by LegalizeSimdMerging and
lower MODE_MERGING to a zeroing operation composed with
pto.vsel(result, old_dst, mask), covering elementwise binary/unary,
vector-scalar, SFU, vdup/vdupv, read-modify-write, and vcvt ops.
Reductions merge only their result lanes (PAT_VL1/PAT_VL2 for
vcadd/vcmax/vcmin, PAT_VL8 for vcg*, and PAT_VL<N/2> for vcpadd's
packed low half); vcvt reinterprets the predicate at the destination
element width; FP8/FP4 results are selected through their byte
representation.

* fix(pto): emit ui8 ptrs for no-padding GM->UB, aligning with ascend (#486)

* style(pto): format PR 479 changes

* fix(pto): avoid doing identity pto.castptr for no-padding GM->UB (#490)

* feat(pto): adapt to ptodsl scalar & pto unification (#491)

* feat(pto): unify signed/unsigned scalar shift right & optimize div for signed int

Emit shift_left/shift_right as authored runtime '<<'/'>>' operators so
PTODSL selects ShRSI/ShRUI from the operand's authored int/uint type,
replacing the tl.ushr path and the shift==bits-1 sign-mask special case.

Lower scalar signed DivNode to pto.div (truncating division) instead of
Python '//' floordiv, matching C semantics.

Adapt PhiloxRNG to the unified pto surface ('>>' split, pto.cast,
pto.select, pto.cast for the uint-to-float draw conversion).

* refactor(pto): adapt to ptodsl scalar & pto unification

PTODSL unified the scalar.* and pto.* surfaces into pto.*: scalar.load/
store -> pto.load/store, scalar.cast/tl.scalar_cast -> pto.cast, scalar.
select -> pto.select, scalar.min/max -> pto.min/max, scalar.exp/log/sqrt
-> pto.exp/log/sqrt (single unary-math table, no simt/scalar split),
tl.scalar_bitcast -> pto.bitcast, scalar.muli -> pto.mul, pto.convert ->
pto.cast, and integer min/max now use the unified signed pto.min/max.

Restructure CastNode: int immediates convert through pto.const, cross-
category numeric conversions go through shape-preserving pto.cast, and
the SIMT-only pto.convert integer-to-float workaround (signless payload
normalization + widening) is subsumed by pto.cast, which infers
signedness from the authored value type.

Only integer atomic_min/atomic_max pass the signedness attribute now.

SIMT local int32 load/store keep the signless alloca with pto.cast
boundary adaptation (LLVM stack slots stay signless i32).

Python side: drop the ushr/scalar_bitcast/scalar_cast and coerce_* re-
exports, switch dcache_bypass to pto.ld_dev/st_dev, and sync's event id
to pto.cast(flag_id, pto.index) (new ptodsl exposes no pto.index_cast
function, so the deprecated-index_cast switch lands here directly).
gemm uses pto.mul instead of scalar.muli. Generated kernels import only
'from ptodsl import pto'.

* refactor(pto): switch deprecated index_cast to cast

Route every cross-core flag mode through the tl.ascend_cross_core_
set/wait_flag helper instead of emitting the named pto.set/wait_
cross_block / pto.set/wait_intra_block operations directly for modes
0/4. The helper already owns the mode/pipe/event mapping, so the
codegen-side direct path duplicated it.

Cast the flag_id to the type its consumer requires inside each mode
branch: the named block operations take an i64 surface event id (the
unified ptodsl rejects pto.index there), while the low-level
pto.sync.set/wait surface keeps the raw index operand for FFTS modes
1 and 2.

* refactor(pto): use pto.maximum / minimum for float T.max / T.min

Float min/max now emit the NaN-propagating pto.minimum/maximum surface
instead of the maxNum/minNum pto.min/max; the vectorized f32x2 fallback
applies the same operation lane-by-lane.

* style(pto): format mixed_kernel helper

* feat(pto): initialize PTO core state in generated kernels (#481)

* feat(pto): initialize PTO core state in generated kernels

* feat(pto): initialize mixed kernel core state

---------

Co-authored-by: caojian5 <caojian5@huawei.com>

* fix(pto): lower FP8 SIMT arithmetic through packed casts (#506)

* fix(pto): lower FP8 SIMT arithmetic through packed casts

* refactor(pto): use typed buffers for scalar FP8 arithmetic

Allocate scalar FP8 scratch buffers with their actual element dtype and use PTODSL contiguous load/store APIs for packed two-lane conversions.

Keep scalar values represented as bytes at the helper boundaries and remove the direct LLVM vector type-punning dependency.

* fix(pto): reject unsupported scalar FP8 operations

---------

Co-authored-by: caojian5 <caojian5@huawei.com>

* style(pto): fix pre-commit issues

* fix(pto): write_gm_bypass_dcache support const float bitcast (#513)

* fix(pto): align packed argument ABI

* fix(pto): preserve integer signedness in SIMT reductions (#522)

Adapt redux argument handling, scalar select and allreduce dispatch to
PTOAS signless integer carriers while preserving authored signedness.

* fix(pto): support standalone MX scale-factor loads (#532)

* fix(pto): lower remaining SIMD special intrinsics

* fix(pto): avoid duplicate helper symbols in multi-kernel JIT (#530)

Co-authored-by: caojian <caojian5@huawei.com>

* feat(pto): simt enhancement: support ShuffleNode codegen (#541)

* fix(pto): reset per-function RNG state in PTO codegen (#535)

* fix(jit): resolve scalar params and skip empty tensors in dynamic stride check (#542)

* fix(pto): normalize scalar int<->bool casts on the device (#546)

The raw-passthrough branch kept unnormalized values: int32(bool(x))
stayed 2 for x == 2 and non-0/1 bool backing bytes survived round-trips,
diverging from the Ascend C++ backend's (bool) normalization.

Lower int/uint -> bool to tl.as_logical_bool(value) and bool -> int/uint
to pto.select(cond, const(1), const(0)) with immediates folded. Vector
bool casts are same-width (bool is an 8-bit type) and keep the widening
branch's representation-preserving passthrough.

* fix(pto): map f16 unary math extern names to pto intrinsics (#548)

Map AscendMath f16 extern names hexp, hfabs, hlog, hsqrt and hrsqrt
to PTO unary math operations. Preserve f32 absolute-value mappings
and implement reciprocal square root using the reciprocal form.

* refactor(ascend): isolate JIT backend adapters

* refactor(pto): preserve existing Ascend adapter paths

* refactor(pto): remove unsupported TVM-FFI adapter

* refactor(pto): inline Cython adapter selection

* refactor(pto): use dedicated execution backend

* refactor: move changes in param.py to cython/adapter.py

* refactor(pto): use standalone kernel adapter

* refactor(pto): isolate source and launch wrappers

* docs(pto): address marker placement and backend overview feedback

* test(pto): include PTO in builtin backend expectations

---------

Co-authored-by: cj <erhsh_165@126.com>
Co-authored-by: liggest <43201720+liggest@users.noreply.github.com>
Co-authored-by: caojian5 <caojian5@huawei.com>
Co-authored-by: lwwangcgz <wlw9393@gmail.com>
Co-authored-by: liggest <liggest@sina.cn>
2026-09-30 12:08:45 +08:00

335 lines
10 KiB
TOML

[project]
name = "tilelang"
description = "A tile level programming language to generate high performance code."
readme = "README.md"
requires-python = ">=3.10"
authors = [{ name = "TileLang Contributors" }, { name = "Tile-AI" }]
maintainers = [{ name = "Lei Wang", email = "leiwang1999@outlook.com" }]
license = "MIT"
keywords = ["BLAS", "CUDA", "HIP", "Code Generation", "TVM"]
classifiers = [
"Development Status :: 4 - Beta",
"Environment :: GPU",
"Operating System :: Microsoft :: Windows",
"Operating System :: POSIX :: Linux",
"Operating System :: MacOS",
"Programming Language :: C++",
"Programming Language :: Python :: 3",
"Programming Language :: Python :: 3.10",
"Programming Language :: Python :: 3.11",
"Programming Language :: Python :: 3.12",
"Programming Language :: Python :: 3.13",
"Programming Language :: Python :: 3.14",
"Programming Language :: Python :: Implementation :: CPython",
"Intended Audience :: Developers",
"Intended Audience :: Science/Research",
"Topic :: Scientific/Engineering :: Artificial Intelligence",
]
dynamic = ["version"]
dependencies = [
"apache-tvm-ffi>=0.1.11,<0.1.13",
# torch-c-dlpack-ext provides prebuilt torch extensions.
# Without it, TVM FFI may require JIT compilation on first import.
"torch-c-dlpack-ext; python_version < '3.14'",
"cloudpickle",
"ml-dtypes",
"numpy>=1.23.5",
"psutil",
"torch",
"torch>=2.4; platform_system == 'Darwin'",
# jit-compiling a torch extension needs setuptools
"setuptools; platform_system == 'Darwin'",
"tqdm>=4.62.3",
"typing-extensions>=4.10.0",
"z3-solver>=4.13.0,<4.15.5",
]
[project.optional-dependencies]
# mldtypes should be greater than 0.5.1
# if you want to enable fp4
fp4 = ["ml-dtypes>=0.5.1"]
# if you want to enable layout inference visualization
vis = ["matplotlib"]
# if you want to build with CUDA NVCC support
nvcc = [
"torch",
"nvidia-cuda-nvcc>=13.0.48",
"nvidia-cuda-cccl>=13.0.50",
# On Windows, ship nvrtc through the runtime deps so `pip install tilelang`
# works without a host CUDA install (TileLang JITs CUDA kernels via nvrtc).
"nvidia-cuda-nvrtc>=13; platform_system == 'Windows'",
# cuRAND is required when kernels use tilelang.cuda.language.random
"nvidia-curand>=10.4; platform_system == 'Windows'",
]
[[tool.uv.index]]
name = "pytorch-cu130"
url = "https://download.pytorch.org/whl/cu130"
explicit = true
[tool.uv.sources]
torch = [
{ index = "pytorch-cu130", marker = "sys_platform == 'win32'", extra = "nvcc" },
]
[build-system]
requires = [
"cython>=3.1.0",
"scikit-build-core",
"z3-solver>=4.13.0,<4.15.5",
# Not for auditwheel, explicitly add patchelf for repairing libz3.so.
# See tvm's CMakeLists.txt for more information.
"patchelf>=0.17.2; platform_system == 'Linux'",
# Under PEP 517 build isolation on Windows, provide nvcc and CUDA headers so
# CMake's find_package(CUDAToolkit) succeeds without a host CUDA install.
"nvidia-cuda-nvcc>=13; platform_system == 'Windows'",
"nvidia-cuda-cccl>=13; platform_system == 'Windows'",
"nvidia-cuda-nvrtc>=13; platform_system == 'Windows'",
# CMake's MSVC helpers in FindPipCUDAToolkit.cmake only activate under the
# Ninja generator; require ninja on Windows so scikit-build-core can pick
# it up (paired with the override below that passes -G Ninja).
"ninja>=1.11; platform_system == 'Windows'",
]
build-backend = "scikit_build_core.build"
[tool.scikit-build]
wheel.py-api = "cp39"
cmake.version = ">=3.26.1"
build-dir = "build"
# Include backend and git info in version
metadata.version.provider = "version_provider"
metadata.version.provider-path = "."
experimental = true
# On Windows, force the Ninja generator so the project's MSVC environment
# activation path runs (it is guarded on `CMAKE_GENERATOR MATCHES "Ninja"`).
# Linux and macOS keep the scikit-build-core default (Makefiles).
[[tool.scikit-build.overrides]]
if.platform-system = "win32"
cmake.args = ["-G", "Ninja"]
# editable.rebuild = true
# build.verbose = true
# logging.level = "DEBUG"
[tool.scikit-build.sdist]
include = [
"./VERSION",
".git_commit.txt",
"./LICENSE",
"THIRDPARTYNOTICES.txt",
"version_provider.py",
"requirements*.txt",
"tilelang/jit/adapter/cython/cython_wrapper.pyx",
"tilelang/jit/adapter/pto/cython_wrapper.pyx",
"CMakeLists.txt",
"src/**",
"cmake/**",
# The vendored 3rdparty contents in sdist should be same as wheel.
# Need full TVM to build from source.
"3rdparty/tvm",
# CUTLASS
"3rdparty/cutlass/include",
"3rdparty/cutlass/tools",
# Vendored HIP headers (build-time only) so source builds work on hosts
# without a ROCm install. Not mapped into the wheel at runtime.
"3rdparty/hip-headers/include",
"testing/**",
"examples/**",
]
exclude = [
".git",
".github",
"**/.git",
"**/.github",
"3rdparty/**",
"build",
]
[tool.scikit-build.wheel.packages]
tilelang = "tilelang"
"tilelang/src" = "src"
# NOTE: The mapping below places the contents of '3rdparty' inside 'tilelang/3rdparty' in the wheel.
# This is necessary to find TVM shared libraries at runtime.
# The vendored 3rdparty contents in wheel should be same as sdist.
# TVM
"tilelang/3rdparty/tvm/src" = "3rdparty/tvm/src"
"tilelang/3rdparty/tvm/python" = "3rdparty/tvm/python"
"tilelang/3rdparty/tvm/include" = "3rdparty/tvm/include"
"tilelang/3rdparty/tvm/version.py" = "3rdparty/tvm/version.py"
# CUTLASS
"tilelang/3rdparty/cutlass/include" = "3rdparty/cutlass/include"
"tilelang/3rdparty/cutlass/tools" = "3rdparty/cutlass/tools"
[tool.codespell]
ignore-words = "docs/spelling_wordlist.txt"
skip = [
"build",
"3rdparty",
"dist",
".venv"
]
[tool.ruff]
target-version = "py310"
line-length = 140
output-format = "full"
exclude = [
"3rdparty",
"examples/deepseek_v32/inference",
]
[tool.ruff.format]
quote-style = "double"
indent-style = "space"
skip-magic-trailing-comma = false
line-ending = "auto"
docstring-code-format = false
docstring-code-line-length = "dynamic"
[tool.ruff.lint.per-file-ignores]
# Do not upgrade type hint in testing and examples.
# See https://github.com/tile-ai/tilelang/issues/1079 for more information.
"testing/**.py" = ["UP", "FA"]
"examples/**.py" = ["UP", "FA"]
[tool.ruff.lint]
select = [
# pycodestyle
"E", "W",
# Pyflakes
"F",
# pyupgrade
"UP", "FA",
# flake8-bugbear
"B",
# flake8-simplify
"SIM",
# isort
# "I",
]
ignore = [
# Module level import not at top of file
"E402",
# star imports
"F405", "F403",
# ambiguous name
"E741",
# line too long
"E501",
# if-else-block instead of ternary
"SIM108",
# key in dict.keys()
"SIM118",
# open file w.o. ctx manager
"SIM115",
# memory leaks
"B019",
# zip without explicit strict
"B905",
# No such file or directory
"E902",
]
[tool.pytest.ini_options]
verbosity_assertions = 3
filterwarnings = ["always"]
markers = [
"perf: benchmark and performance API tests skipped by default",
"slow: long-running correctness sweeps",
# Feature markers emitted by vendored TVM's tvm.testing.utils requires_*
# decorators (see 3rdparty/tvm/python/tvm/testing/utils.py). TVM registers
# these via its own conftest, but tilelang only imports the decorators, so
# they must be registered here to avoid PytestUnknownMarkWarning.
"pto: tests for the PTO backend",
"cpu: Mark a test as running on CPU",
"gpu: Mark a test as using a GPU",
"multi_gpu: Mark a test as using multiple GPUs",
"cuda: Mark a test as using CUDA",
"llvm: Mark a test as using LLVM",
"metal: Mark a test as using Metal",
"rocm: Mark a test as using ROCm",
]
[tool.cibuildwheel]
archs = ["auto64"]
skip = "*musllinux*"
build-frontend = "build"
environment = { PYTHONDEVMODE = "1", PYTHONUNBUFFERED = "1" }
environment-pass = [
"CUDA_VERSION",
"NO_VERSION_LABEL",
"NO_TOOLCHAIN_VERSION",
"NO_GIT_VERSION",
"COLUMNS",
"CMAKE_GENERATOR",
"CMAKE_BUILD_PARALLEL_LEVEL",
"FORCE_COLOR",
"CLICOLOR_FORCE",
]
before-build = "env -0 | sort -z | tr '\\0' '\\n'"
windows.before-build = "set"
test-command = [
'python -c "import tilelang; print(tilelang.__version__)"',
]
[tool.cibuildwheel.linux]
environment.PYTHONDEVMODE = "1"
environment.PYTHONUNBUFFERED = "1"
environment.PATH = "/usr/local/cuda/bin:$PATH"
environment.LD_LIBRARY_PATH = "/usr/local/cuda/lib64:/usr/local/cuda/lib64/stubs:$LD_LIBRARY_PATH"
# Build a wheel with CUDA, ROCm, and Ascend support. ROCm uses vendored HIP
# headers and Ascend uses local ACL declarations; both link to runtime stubs,
# so the manylinux container only needs the CUDA toolkit. Execution still
# requires the selected backend's real runtime and device.
environment.USE_CUDA = "ON"
environment.USE_ROCM = "ON"
environment.USE_ASCEND = "ON"
manylinux-x86_64-image = "manylinux_2_28" # AlmaLinux 8
manylinux-aarch64-image = "manylinux_2_34" # Z3 requires
# Install CUDA runtime and stub driver library
# manylinux_2_28 uses gcc 14, which needs CUDA >=12.8
before-all = """
set -eux
cat /etc/*-release
uname -a
case "$(uname -m)" in
"x86_64")
DEFAULT_CUDA_VERSION="13.0"
dnf config-manager --add-repo https://developer.download.nvidia.cn/compute/cuda/repos/rhel8/x86_64/cuda-rhel8.repo
;;
"aarch64")
DEFAULT_CUDA_VERSION="13.0"
dnf config-manager --add-repo https://developer.download.nvidia.com/compute/cuda/repos/rhel8/sbsa/cuda-rhel8.repo
;;
*)
exit 1
;;
esac
cudaver="$(echo "${CUDA_VERSION:-$DEFAULT_CUDA_VERSION}" | cut -d '.' -f-2)"
v="${cudaver//./-}"
dnf install -y "cuda-minimal-build-${v}" "cuda-driver-devel-${v}" "cuda-nvrtc-devel-${v}"
dnf clean all
"""
# Do not bundle libtvm_ffi.so and libz3.so* because they are shipped by the dependency packages.
repair-wheel-command = "auditwheel -v repair --exclude libtvm_ffi.so --exclude 'libz3.so*' --exclude libcuda.so.1 --exclude '/usr/local/cuda*' -w {dest_dir} {wheel}"
[tool.cibuildwheel.macos]
repair-wheel-command = "delocate-wheel --verbose --ignore-missing-dependencies --no-sanitize-rpaths --require-archs {delocate_archs} -w {dest_dir} -v {wheel}"
[[tool.cibuildwheel.overrides]]
select = "*linux*x86_64*"
# x86_64 runners in GitHub Actions have limited storage,
# pre-install torch without caching to reduce disk usage during install tilelang.
before-test = [
"pip install torch --no-cache-dir",
]