Merge upstream/main - fix backend import path for metal

Author: oraluben

benchmark/matmul/benchmark_matmul_intrinsic.py

@@ -62,7 +62,7 @@ Below is a simplified code snippet for a 1024 x 1024 x 1024 matrix multiplicatio
 ```python
 import tilelang
 import tilelang.language as T
-from tilelang.intrinsics import make_mma_swizzle_layout
+from tilelang.cuda.intrinsics import make_mma_swizzle_layout
 
 def matmul(M, N, K, block_M, block_N, block_K, dtype="float16", accum_dtype="float"):
     @T.prim_func

docs/deeplearning_operators/matmul.md

@@ -27,7 +27,8 @@ ROCm:
 
 ## Implementation
 
-The stubs in `src/target/stubs/` implement a lazy-loading mechanism:
+The CUDA stubs in `src/backend/cuda/stubs/` and ROCm stubs in
+`src/backend/rocm/stubs/` implement a lazy-loading mechanism:
 
 - **Lazy Loading**: Libraries are loaded via `dlopen` only upon the first API call.
 - **Global Symbol Reuse**: For `cudart` and `nvrtc`, the stubs first check the global namespace (`RTLD_DEFAULT`) to use any already loaded symbols (e.g., from PyTorch).

docs/runtime_internals/stubs.md

@@ -7,12 +7,12 @@
 import tilelang.language as T
 from tilelang import tvm as tvm
 from tvm import DataType
-from tilelang.intrinsics.mma_layout import (
+from tilelang.cuda.intrinsics.layout.mma_layout import (
     make_mma_swizzle_layout as make_swizzle_layout,
 )
 import numpy as np
 
-from tilelang.intrinsics.mma_macro_generator import (
+from tilelang.cuda.intrinsics.macro.mma_macro_generator import (
     INT4TensorCoreIntrinEmitter,
 )
 from tilelang.transform import simplify_prim_func

examples/bitnet-1.58b/kernel_benchmark/tilelang_bitnet_158_int8xint2_prefill.py

@@ -33,6 +33,8 @@
     "gemm_sp/test_example_gemm_sp.py::test_example_gemm_sp",
     # Flaky — passes when run in isolation, fails under parallel execution
     "minference/test_vs_sparse_attn.py::test_vs_sparse_attn",
+    # CuTeDSL does not yet lower DeepSeek V4 FP4 act quant conversions.
+    "deepseek_v4/test_tilelang_example_deepseek_v4.py::test_example_act_quant",
 }
 
 

examples/conftest.py

@@ -141,8 +141,8 @@ def tl_matmul_with_ladder_weight_only_transform_block_reduce_int4(
     accum_dtype,
     transform_b,
 ):
-    from tilelang.intrinsics.mma_layout import make_mma_swizzle_layout as make_swizzle_layout
-    from tilelang.intrinsics.mma_macro_generator import (
+    from tilelang.cuda.intrinsics.layout.mma_layout import make_mma_swizzle_layout as make_swizzle_layout
+    from tilelang.cuda.intrinsics.macro.mma_macro_generator import (
         TensorCoreIntrinEmitterWithLadderTransform,
     )
 

examples/dequantize_gemm/example_dequant_gemm_fine_grained.py

@@ -174,7 +174,7 @@ Below is a more advanced snippet that showcases how to apply memory layouts, ena
 import tilelang.language as T
 # `make_mma_swizzle_layout` is a python-defined layout function
 # that helps align data for MMA (Matrix Multiply-Accumulate) operations.
-from tilelang.intrinsics import make_mma_swizzle_layout as make_swizzle_layout
+from tilelang.cuda.intrinsics import make_mma_swizzle_layout as make_swizzle_layout
 
 def matmul(M, N, K, block_M, block_N, block_K, dtype=T.float16, accum_dtype=T.float):
     @T.prim_func