mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-02 06:34:36 +08:00
* 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>
335 lines
10 KiB
TOML
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",
|
|
]
|