mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
* Add language dialect overlays * Refactor language dialects into static backend layers * Make language facade a static re-export and extract sparse layout leaf Replace the runtime-activated language facade with a static `from tilelang.cuda.language import *` re-export: drop the .pyi stubs + generator + py.typed, the globals()-based __all__ scraping, and the _activate_cuda_facade() setattr. Backend dialects (cuda/metal/rocm) now build on tilelang.language.common with per-backend tir overlays. Import-time dtype defaults in the mma/wgmma/mfma macro generators are pinned to the dtypes leaf so they no longer touch the half-initialized facade during bootstrap. Metal generators import tilelang.metal.language (not the cuda facade) so the eager builder's real getattr resolves the simdgroup ops; a lazy __getattr__ breaks the metal.language <-> intrinsics cycle. Extract the 2:4 sparsity layout metadata into tilelang/cuda/intrinsics/sparse_layout.py (a dtypes-only leaf) so mma_sp imports it directly and tilelang.utils.sparse drops out of the bootstrap closure, letting its @tilelang.jit compressor go back to a module-level decorator. Exports unchanged: common=558, cuda/facade=683. * Remove unwired BackendModule scaffolding The BackendModule manifest and per-backend backend.py registrations were included by mistake: nothing imports the per-backend backend.py files, the ensure_*_loaded lazy-loaders have no call sites, and the mechanism duplicates the register_lazy_execution_backends wiring already in tilelang/backend/__init__.py. Drop the dead code and its unit test. * Import tvm.arith in metal macro generator The simdgroup stride check uses arith.Analyzer but the module was missing the tvm.arith import. * Use string dtype defaults in CUDA macro generators Default dtype arguments are evaluated at import time, when the language facade is only partially initialized, so they cannot reference T.*. They previously used the tilelang.language.dtypes leaf; since every consumer treats these as strings (annotated str, wrapped via DataType before use), plain string literals are equivalent and drop the _dtypes dependency. * Drop MPSIntrinEmitter from metal language facade The emitter is only consumed directly from tilelang.metal.intrinsics, never via the language surface, and exposing it here required a lazy __getattr__ to break the metal.language <-> metal.intrinsics import cycle. Removing it from __all__ eliminates the cycle and the deferred import. * Restore dtype forwarding for ptx_ldmatrix T.ptx_ldmatrix previously accepted a dtype= keyword via _dtype_forward (from the old language/ast/ir.py path). The CUDA dialect alias dropped the wrapper while every neighboring ptx_* operator kept it, changing the public call contract. Wrap it to match. * Re-export tir Select/min/max wrappers and dedup metal __all__ Add Select/min/max to tilelang.language.tir.ir __all__ so the common facade re-exports the documented wrappers instead of the raw TVM symbols. Build metal __all__ via dict.fromkeys to match the sibling dialects and satisfy PLE0604. --------- Co-authored-by: SiriusNEO <chaofan@deepseek.com>
140 lines
4.3 KiB
Python
140 lines
4.3 KiB
Python
import importlib
|
|
|
|
import pytest
|
|
|
|
import tilelang
|
|
import tilelang.language as default_language
|
|
import tilelang.language.common as common_language
|
|
from tilelang.language.tir.exports import (
|
|
CLASSIFIED_VENDOR_TIR_EXPORTS,
|
|
CUDA_ONLY_TIR_EXPORTS,
|
|
METAL_ONLY_TIR_EXPORTS,
|
|
SHARED_LEGACY_TIR_EXPORTS,
|
|
)
|
|
|
|
|
|
CUDA_ONLY_NAMES = {
|
|
"ClusterKernel",
|
|
"CUDASourceCodeKernel",
|
|
"pdl_trigger",
|
|
"rng_init",
|
|
"tcgen05_mma",
|
|
"tma_copy",
|
|
"wgmma_mma",
|
|
"ws",
|
|
}
|
|
|
|
ROCM_ONLY_NAMES = {
|
|
"MatrixCoreIntrinEmitter",
|
|
"mfma",
|
|
"rdna_wmma",
|
|
"tvm_mfma",
|
|
}
|
|
|
|
|
|
def test_default_language_is_static_cuda_facade():
|
|
from tilelang.cuda import language as cuda_language
|
|
|
|
assert tilelang.language is default_language
|
|
assert importlib.import_module("tilelang.language") is default_language
|
|
assert default_language.__tilelang_dialect__ == "cuda"
|
|
assert cuda_language.__tilelang_dialect__ == "cuda"
|
|
assert set(default_language.__all__) == set(cuda_language.__all__)
|
|
|
|
for name in default_language.__all__:
|
|
assert getattr(default_language, name) is getattr(cuda_language, name)
|
|
|
|
|
|
def test_common_language_preserves_special_dsl_exports():
|
|
assert common_language.__tilelang_dialect__ == "common"
|
|
assert "__log" in common_language.__all__
|
|
assert hasattr(common_language, "__log")
|
|
assert CUDA_ONLY_TIR_EXPORTS.isdisjoint(common_language.__all__)
|
|
assert METAL_ONLY_TIR_EXPORTS.isdisjoint(common_language.__all__)
|
|
|
|
|
|
def test_cuda_language_composes_common_and_cuda_symbols():
|
|
from tilelang.cuda import language as T
|
|
from tilelang.cuda import debug as cuda_debug
|
|
|
|
assert T.copy is common_language.copy
|
|
assert T.tcgen05_mma is T.tcgen05_gemm
|
|
assert T.tcgen05_mma_blockscaled is T.tcgen05_gemm_blockscaled
|
|
assert T.wgmma_mma is T.wgmma_gemm
|
|
assert T.device_assert is cuda_debug.device_assert
|
|
assert set(T.__all__) >= CUDA_ONLY_NAMES
|
|
assert set(T.__all__) >= CUDA_ONLY_TIR_EXPORTS
|
|
assert hasattr(T, "TCGEN05TensorCoreIntrinEmitter")
|
|
assert hasattr(T, "WGMMATensorCoreIntrinEmitter")
|
|
|
|
|
|
def test_rocm_language_composes_common_and_rocm_symbols():
|
|
from tilelang.rocm import language as T
|
|
|
|
assert T.__tilelang_dialect__ == "rocm"
|
|
assert T.copy is common_language.copy
|
|
assert T.mfma is T.tvm_mfma
|
|
assert T.mfma_store is T.tvm_mfma_store
|
|
assert set(T.__all__) >= ROCM_ONLY_NAMES
|
|
assert hasattr(T, "MatrixCoreIntrinEmitter")
|
|
assert hasattr(T, "make_mfma_swizzle_layout")
|
|
assert set(T.__all__) >= SHARED_LEGACY_TIR_EXPORTS
|
|
|
|
|
|
def test_metal_language_owns_simdgroup_tir_exports():
|
|
from tilelang.metal import language as T
|
|
|
|
assert set(T.__all__) >= METAL_ONLY_TIR_EXPORTS
|
|
|
|
|
|
def test_upstream_vendor_tir_exports_are_classified():
|
|
import tvm.tirx.script.parser as upstream_tir_parser
|
|
|
|
vendor_exports = {name for name in upstream_tir_parser.__all__ if name.startswith("ptx_") or "simdgroup" in name}
|
|
assert vendor_exports <= CLASSIFIED_VENDOR_TIR_EXPORTS
|
|
|
|
|
|
@pytest.mark.parametrize(
|
|
"module_name",
|
|
[
|
|
"tilelang.cpu.language",
|
|
"tilelang.metal.language",
|
|
"tilelang.rocm.language",
|
|
"tilelang.webgpu.language",
|
|
],
|
|
)
|
|
def test_non_cuda_dialects_do_not_export_cuda_symbols(module_name):
|
|
module = importlib.import_module(module_name)
|
|
assert CUDA_ONLY_NAMES.isdisjoint(module.__all__)
|
|
|
|
|
|
@pytest.mark.parametrize(
|
|
"module_name",
|
|
[
|
|
"tilelang.cpu.language",
|
|
"tilelang.cuda.language",
|
|
"tilelang.metal.language",
|
|
"tilelang.webgpu.language",
|
|
],
|
|
)
|
|
def test_non_rocm_dialects_do_not_export_rocm_symbols(module_name):
|
|
module = importlib.import_module(module_name)
|
|
assert ROCM_ONLY_NAMES.isdisjoint(module.__all__)
|
|
|
|
|
|
def test_common_only_dialects_match_common_surface():
|
|
from tilelang.cpu import language as cpu_language
|
|
from tilelang.webgpu import language as webgpu_language
|
|
|
|
assert set(cpu_language.__all__) == set(common_language.__all__)
|
|
assert set(webgpu_language.__all__) == set(common_language.__all__)
|
|
|
|
|
|
def test_cuda_whole_module_implementations_live_under_cuda():
|
|
from tilelang.cuda import language as T
|
|
|
|
assert T.pdl_trigger.__module__ == "tilelang.cuda.language.pdl"
|
|
assert T.cluster_sync.__module__ == "tilelang.cuda.language.cluster"
|
|
assert T.rng_init.__module__ == "tilelang.cuda.language.random"
|
|
assert T.ws.__module__ == "tilelang.cuda.language.warpgroup"
|