tile-ai · LeiWang1999 · May 20, 2026 · Jan 28, 2026 · Jan 28, 2026 · Jan 28, 2026
diff --git a/.agents/skills/tilelang-tvm-ir/SKILL.md b/.agents/skills/tilelang-tvm-ir/SKILL.md
@@ -1,19 +1,19 @@
 ---
 name: tilelang-tvm-ir
-description: Use when editing TileLang C++ passes or TVM TIR code that handles ObjectRef/NodeRef types such as For, Buffer, Var, Block, Stmt, PrimExpr, or their *Node raw node counterparts; especially when choosing function parameters, optional values, identity maps/sets, or equality checks.
+description: Use when editing TileLang C++ passes or TVM TIRX code that handles ObjectRef/NodeRef types such as For, Buffer, Var, SBlock, Stmt, PrimExpr, or their *Node raw node counterparts; especially when choosing function parameters, optional values, identity maps/sets, or equality checks.
 ---
 
 # TileLang TVM IR Handle Conventions
 
 ## Core Rule
 
-In TVM C++, `For`, `Buffer`, `Var`, `Block`, `Stmt`, `PrimExpr`, `SeqStmt`, etc. are `ObjectRef` smart handles. `ForNode`, `BufferNode`, `VarNode`, `BlockNode`, etc. are raw node structs reached through visitor callbacks, `as<TNode>()`, `operator->`, or `.get()`.
+In TVM C++, `For`, `Buffer`, `Var`, `SBlock`, `Stmt`, `PrimExpr`, `SeqStmt`, etc. are `ObjectRef` smart handles. `ForNode`, `BufferNode`, `VarNode`, `SBlockNode`, etc. are raw node structs reached through visitor callbacks, `as<TNode>()`, `operator->`, or `.get()`.
 
 When a value needs to cross a function boundary, be stored, be optional, be used as an identity key, or survive beyond a local inspection branch, prefer the handle type over `const *Node`.
 
 ## Preferred Patterns
 
-- Function parameters and return values: use handles such as `For`, `Buffer`, `Var`, `Block`, `Stmt`, or `SeqStmt`.
+- Function parameters and return values: use handles such as `For`, `Buffer`, `Var`, `SBlock`, `Stmt`, or `SeqStmt`.
 - Nullable AST values: use `Optional<For>`, `Optional<SeqStmt>`, etc., not `const ForNode* = nullptr`.
 - Identity maps and sets: use handle keys with TVM identity hashing:
 
@@ -45,7 +45,7 @@ Keep these raw pointers local to the immediate inspection or mutation site.
 
 ## Avoid
 
-- Passing `const ForNode*`, `const BufferNode*`, `const VarNode*`, or `const Block*` between helper functions when a handle exists.
+- Passing `const ForNode*`, `const BufferNode*`, `const VarNode*`, or `const SBlockNode*` between helper functions when a handle exists.
 - Storing raw node pointers in `std::unordered_map` or `std::unordered_set` for identity tracking.
 - Using `.get()` as a key unless a callee requires a raw TVM node API and the pointer is not retained.
 - Comparing handles through `.get() == other.get()`; prefer `.same_as()`.
@@ -99,7 +99,7 @@ bool uses = UsesVar(expr, [&](const VarNode* vn) {
 When reviewing TileLang TIR passes, search for:
 
 ```bash
-rg -n "std::unordered_(set|map)<const .*Node \\*|const (For|SeqStmt|Block).*Node \\*|\\.get\\(\\) ==|\\.find\\([^\\n]*\\.get\\(\\)|\\.count\\([^\\n]*\\.get\\(\\)|\\.insert\\([^\\n]*\\.get\\(\\)" src/transform
+rg -n "std::unordered_(set|map)<const .*Node \\*|const (For|SeqStmt|SBlock).*Node \\*|\\.get\\(\\) ==|\\.find\\([^\\n]*\\.get\\(\\)|\\.count\\([^\\n]*\\.get\\(\\)|\\.insert\\([^\\n]*\\.get\\(\\)" src/transform
 ```
 
 Do not mechanically remove every raw node pointer. Keep visitor signatures, `as<TNode>()` pattern checks, and `CopyOnWrite()` mutation pointers. Refactor only the places that store, pass, compare, or key identities through raw pointers.
diff --git a/3rdparty/tvm b/3rdparty/tvm
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -349,6 +349,11 @@ endforeach()
 # tvm tries to detect gtest by default, but may fail if its header is not installed.
 set(USE_GTEST OFF)
 
+# TileLang is built as a separate shared library on POSIX and subclasses / calls
+# TVM C++ internals. TVM 0.25 defaults to hiding private symbols, which makes
+# those out-of-DSO references unresolved when libtilelang.so is dlopened.
+set(HIDE_PRIVATE_SYMBOLS OFF CACHE BOOL "Compile TVM with -fvisibility=hidden." FORCE)
+
 # Include directories for TileLang
 set(TILE_LANG_INCLUDES ${TVM_INCLUDES})
 # Add TileLang's own src/ to include path so cross-directory includes
@@ -380,8 +385,6 @@ list(APPEND TILE_LANG_SRCS
   src/runtime/error_helpers.cc
 )
 
-set(TILELANG_OUTPUT_TARGETS tilelang tvm)
-
 # Track if the user explicitly selected a backend via cache options.
 set(TILELANG_BACKEND_USER_SELECTED OFF)
 foreach(BACKEND IN LISTS TILELANG_BACKENDS)
@@ -490,9 +493,9 @@ endif()
 
 # Provide the custom LogMessageImpl / LogFatalImpl implementation to TVM,
 # since TVM_LOG_CUSTOMIZE=1 requires them to be supplied by the user.
-target_sources(tvm_objs PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src/runtime/logging.cc")
+target_sources(tvm_runtime_objs PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src/runtime/logging.cc")
 if(TILELANG_RELEASE_BUILD)
-  target_compile_definitions(tvm_objs PRIVATE TILELANG_RELEASE_BUILD=1)
+  target_compile_definitions(tvm_runtime_objs PRIVATE TILELANG_RELEASE_BUILD=1)
 endif()
 
 # Resolve compile warnings in tvm
@@ -524,8 +527,8 @@ if(TILELANG_RELEASE_BUILD)
   target_compile_definitions(tilelang_objs PRIVATE TILELANG_RELEASE_BUILD=1)
 endif()
 
-set(TILELANG_OUTPUT_TARGETS tvm)
 if(WIN32)
+  set(TILELANG_OUTPUT_TARGETS tvm)
   # On Windows, TileLang subclasses a large number of non-exported TVM C++
   # classes. Building those objects into tvm.dll avoids an otherwise fragile
   # cross-DLL ABI boundary while keeping Python-side registration semantics.
@@ -549,8 +552,9 @@ if(WIN32)
   # sure `ninja` (and `cmake --build`) actually builds it by default.
   set_property(TARGET tvm PROPERTY EXCLUDE_FROM_ALL FALSE)
 else()
+  set(TILELANG_OUTPUT_TARGETS tvm_runtime tvm_compiler)
   add_library(tilelang SHARED $<TARGET_OBJECTS:tilelang_objs>)
-  target_link_libraries(tilelang PUBLIC tvm)
+  target_link_libraries(tilelang PUBLIC tvm_compiler)
 
   # Place dev build outputs under build/lib for consistency
   set_target_properties(tilelang PROPERTIES
@@ -610,15 +614,19 @@ install(TARGETS tilelang_cython_wrapper
         RUNTIME DESTINATION tilelang/lib
         ARCHIVE DESTINATION tilelang/lib)
 
-# Copy libz3.so to build folder to workaround isolated build env issue
-if(USE_Z3 AND USE_PYPI_Z3)
+# Copy libz3 into the build-tree library folder on platforms where the
+# PyPI Z3 finder does not already normalize a build-tree runtime copy.
+if(USE_Z3 AND USE_PYPI_Z3 AND NOT APPLE)
   get_target_property(Z3_LIBRARY_PATH z3::libz3 IMPORTED_LOCATION)
-  install(FILES "${Z3_LIBRARY_PATH}" DESTINATION "${CMAKE_BINARY_DIR}/lib")
-  if(APPLE)
-    set_target_properties(tvm PROPERTIES BUILD_RPATH "@loader_path")
-  else()
-    set_target_properties(tvm PROPERTIES BUILD_RPATH "\$ORIGIN")
+  get_filename_component(Z3_LIBRARY_DIR "${Z3_LIBRARY_PATH}" DIRECTORY)
+  file(GLOB Z3_RUNTIME_LIBS
+    "${Z3_LIBRARY_DIR}/libz3${CMAKE_SHARED_LIBRARY_SUFFIX}*"
+    "${Z3_LIBRARY_DIR}/z3${CMAKE_SHARED_LIBRARY_SUFFIX}*"
+  )
+  if(NOT Z3_RUNTIME_LIBS)
+    set(Z3_RUNTIME_LIBS "${Z3_LIBRARY_PATH}")
   endif()
+  file(COPY ${Z3_RUNTIME_LIBS} DESTINATION "${CMAKE_BINARY_DIR}/lib")
 endif()
 
 if(DEFINED TILELANG_ACTIVE_BACKEND_STUB_LINK)
@@ -634,13 +642,16 @@ endif()
 
 unset(PATCHELF_EXECUTABLE CACHE)
 
+set(TILELANG_BUILD_RPATH "")
 if(APPLE)
+  set(TILELANG_BUILD_RPATH "@loader_path")
   set(TILELANG_INSTALL_RPATH "@loader_path;@loader_path/../../tvm_ffi/lib")
   if(USE_Z3 AND USE_PYPI_Z3)
     # Some z3 is placed in lib/ and some in bin/, we add both in rpath
     string(APPEND TILELANG_INSTALL_RPATH ";@loader_path/../../z3/lib;@loader_path/../../z3/bin")
   endif()
 elseif(UNIX)
+  set(TILELANG_BUILD_RPATH "\$ORIGIN")
   set(TILELANG_INSTALL_RPATH "\$ORIGIN:\$ORIGIN/../../tvm_ffi/lib")
   if(USE_Z3 AND USE_PYPI_Z3)
     string(APPEND TILELANG_INSTALL_RPATH ":\$ORIGIN/../../z3/lib")
@@ -654,8 +665,9 @@ elseif(UNIX)
   endif()
 endif()
 
-# Let libtilelang search for tvm in the same directory
+# Let TileLang and TVM split libraries find each other in the same directory
 foreach(target IN LISTS TILELANG_OUTPUT_TARGETS)
+  set_target_properties(${target} PROPERTIES BUILD_RPATH "${TILELANG_BUILD_RPATH}")
   set_target_properties(${target} PROPERTIES INSTALL_RPATH "${TILELANG_INSTALL_RPATH}")
   set_target_properties(${target} PROPERTIES
     LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib"

diff --git a/cmake/load_tvm.cmake b/cmake/load_tvm.cmake
@@ -16,7 +16,6 @@ set(TVM_INCLUDES
   ${TVM_SOURCE}/include
   ${TVM_SOURCE}/src
   ${TVM_SOURCE}/3rdparty/dlpack/include
-  ${TVM_SOURCE}/3rdparty/dmlc-core/include
 )
 
 if(EXISTS ${TVM_SOURCE}/ffi/include)

diff --git a/cmake/pypi-z3/FindZ3.cmake b/cmake/pypi-z3/FindZ3.cmake
@@ -14,9 +14,10 @@ endif()
 message("-- Find Z3 in path: ${Z3_PATH}")
 find_path(Z3_INCLUDE_DIR NO_DEFAULT_PATH NAMES z3++.h PATHS ${Z3_PATH}/include)
 if(WIN32)
-    find_file(Z3_LIBRARY NO_DEFAULT_PATH NAMES libz3.dll z3.dll PATHS ${Z3_PATH}/bin ${Z3_PATH}/lib ${Z3_PATH}/lib64)
-    set(Z3_DLL "${Z3_LIBRARY}")
-    if(Z3_DLL)
+    find_file(Z3_RUNTIME_LIBRARY NO_DEFAULT_PATH NAMES libz3.dll z3.dll PATHS ${Z3_PATH}/bin ${Z3_PATH}/lib ${Z3_PATH}/lib64)
+    find_library(Z3_IMPORT_LIBRARY NO_DEFAULT_PATH NAMES libz3 z3 PATHS ${Z3_PATH}/bin ${Z3_PATH}/lib ${Z3_PATH}/lib64)
+    set(Z3_DLL "${Z3_RUNTIME_LIBRARY}")
+    if(Z3_DLL AND NOT Z3_IMPORT_LIBRARY)
         set(_z3_import_dir "${CMAKE_BINARY_DIR}/z3-import")
         file(MAKE_DIRECTORY "${_z3_import_dir}")
         set(Z3_IMPORT_LIBRARY "${_z3_import_dir}/libz3.lib")
@@ -28,27 +29,54 @@ if(WIN32)
             endif()
         endif()
     endif()
+    set(Z3_LIBRARY "${Z3_IMPORT_LIBRARY}")
 else()
     find_library(Z3_LIBRARY NO_DEFAULT_PATH NAMES z3 libz3 PATHS ${Z3_PATH}/bin ${Z3_PATH}/lib ${Z3_PATH}/lib64)
+    if(APPLE AND Z3_LIBRARY)
+        set(Z3_ORIGINAL_LIBRARY "${Z3_LIBRARY}")
+        set(Z3_LIBRARY "${CMAKE_BINARY_DIR}/lib/libz3${CMAKE_SHARED_LIBRARY_SUFFIX}")
+        file(MAKE_DIRECTORY "${CMAKE_BINARY_DIR}/lib")
+        configure_file("${Z3_ORIGINAL_LIBRARY}" "${Z3_LIBRARY}" COPYONLY)
+        execute_process(
+            COMMAND /usr/bin/install_name_tool -id "@rpath/libz3${CMAKE_SHARED_LIBRARY_SUFFIX}" "${Z3_LIBRARY}"
+            RESULT_VARIABLE Z3_INSTALL_NAME_RESULT
+        )
+        if(NOT Z3_INSTALL_NAME_RESULT EQUAL 0)
+            message(FATAL_ERROR "Failed to normalize Z3 install name for ${Z3_LIBRARY}")
+        endif()
+    endif()
 endif()
 message("-- Found Z3 include dir: ${Z3_INCLUDE_DIR}")
 message("-- Found Z3 library: ${Z3_LIBRARY}")
+if(APPLE AND Z3_ORIGINAL_LIBRARY)
+    message("-- Found Z3 original library: ${Z3_ORIGINAL_LIBRARY}")
+endif()
+if(WIN32)
+    message("-- Found Z3 runtime library: ${Z3_RUNTIME_LIBRARY}")
+endif()
+if(NOT Z3_INCLUDE_DIR OR NOT Z3_LIBRARY)
+    message(FATAL_ERROR "Could not find Z3 library or include directory")
+endif()
+if(WIN32 AND NOT Z3_RUNTIME_LIBRARY)
+    message(FATAL_ERROR "Could not find Z3 runtime library")
+endif()
 if(WIN32)
     add_library(z3::libz3 SHARED IMPORTED GLOBAL)
 else()
     add_library(z3::libz3 SHARED IMPORTED GLOBAL)
 endif()
 set_target_properties(z3::libz3
     PROPERTIES
-    IMPORTED_LOCATION ${Z3_LIBRARY}
     INTERFACE_INCLUDE_DIRECTORIES ${Z3_INCLUDE_DIR}
 )
+if(WIN32)
+    set_target_properties(z3::libz3 PROPERTIES IMPORTED_LOCATION "${Z3_RUNTIME_LIBRARY}")
+else()
+    set_target_properties(z3::libz3 PROPERTIES IMPORTED_LOCATION "${Z3_LIBRARY}")
+endif()
 if(WIN32 AND DEFINED Z3_IMPORT_LIBRARY AND EXISTS "${Z3_IMPORT_LIBRARY}")
     set_target_properties(z3::libz3 PROPERTIES IMPORTED_IMPLIB "${Z3_IMPORT_LIBRARY}")
 endif()
-if(NOT Z3_INCLUDE_DIR OR NOT Z3_LIBRARY)
-    message(FATAL_ERROR "Could not find Z3 library or include directory")
-endif()
 set(Z3_CXX_INCLUDE_DIRS ${Z3_INCLUDE_DIR})
 set(Z3_C_INCLUDE_DIRS ${Z3_INCLUDE_DIR})
 set(Z3_FOUND TRUE)
diff --git a/docs/get_started/targets.md b/docs/get_started/targets.md
@@ -5,45 +5,43 @@ The target determines which code generator is used (CUDA, HIP, Metal, LLVM, …)
 device-specific options such as GPU architecture flags. This page summarises how to pick and customise a target
 when compiling TileLang programs.
 
-## Common target strings
+## Common targets
 
-TileLang ships with a small set of common targets; each accepts the full range of TVM options so you can fine-tune
-the generated code. The most frequent choices are listed below:
+TileLang ships with a small set of common target kinds. Use a bare string for the kind, or a TVM target config
+dictionary when you need options such as GPU architecture or CPU model. The most frequent choices are listed below:
 
 | Base name | Description |
 | --------- | ----------- |
 | `auto` | Detects CUDA → HIP → Metal in that order. Useful when running the same script across machines. |
-| `cuda` | NVIDIA GPUs. Supports options such as `-arch=sm_80`, `-max_num_threads=1024`, etc. |
-| `cutedsl` | NVIDIA CUTLASS/CuTe DSL backend. Requires `nvidia-cutlass-dsl`. `cuda` options can also be applied to this target. |
-| `hip` | AMD GPUs via ROCm. Options like `-mcpu=gfx90a` can be appended. |
+| `cuda` | NVIDIA GPUs. Use a config dict for options such as `{"kind": "cuda", "arch": "sm_80"}`. |
+| `cutedsl` | NVIDIA CUTLASS/CuTe DSL backend. Requires `nvidia-cutlass-dsl`. |
+| `hip` | AMD GPUs via ROCm. Use a config dict for options such as `{"kind": "hip", "mcpu": "gfx90a"}`. |
 | `metal` | Apple Silicon GPUs (arm64 Macs). |
-| `llvm` | CPU execution; accepts the standard TVM LLVM switches. |
+| `llvm` | CPU execution. Use a config dict for options such as `{"kind": "llvm", "mtriple": "x86_64-linux-gnu"}`. |
 | `webgpu` | Browser / WebGPU runtimes. |
 | `c` | Emit plain C source for inspection or custom toolchains. |
 
-To add options, append them after the base name, separated by spaces. For example:
+To add options, pass a target config dictionary. For example:
 
 ```python
-target = "cuda -arch=sm_90"
+target = {"kind": "cuda", "arch": "sm_90"}
 kernel = tilelang.compile(func, execution_backend="cython")
 # or
 @tilelang.jit(target=target)
 def compiled_kernel(*args):
     return func(*args)
 ```
 
-The same convention works for HIP or LLVM (e.g. `hip -mcpu=gfx940`, `llvm -mtriple=x86_64-linux-gnu`).
-
 ### Advanced: Specify Exact Hardware
 
-When you already know the precise GPU model, you can encode it in the target string—either via `-arch=sm_XX` or by
-using one of TVM’s pre-defined target tags such as `nvidia/nvidia-h100`.  Supplying this detail is optional for
+When you already know the precise GPU model, you can encode it in the target config via `arch="sm_XX"` or by
+using one of TVM’s pre-defined target tags such as `nvidia/nvidia-h100`. Supplying this detail is optional for
 TileLang in general use, but it becomes valuable when the TVM cost model is enabled (e.g. during autotuning).  The
 cost model uses the extra attributes to make better scheduling predictions.  If you skip this step (or do not use the
 cost model), generic targets like `cuda` or `auto` are perfectly fine.
 
 All CUDA compute capabilities recognised by TVM’s target registry are listed below.  Pick the one that matches your
-GPU and append it to the target string or use the corresponding target tag—for example `nvidia/nvidia-a100`.
+GPU and set it in the target config or use the corresponding target tag—for example `nvidia/nvidia-a100`.
 
 | Architecture | GPUs (examples) |
 | ------------ | ---------------- |
@@ -73,14 +71,13 @@ Refer to NVIDIA’s [CUDA GPUs](https://developer.nvidia.com/cuda-gpus) page or
 
 ## Creating targets programmatically
 
-If you prefer working with TVM’s `Target` objects, TileLang exposes the helper
-`tilelang.utils.target.determine_target` (returns a canonical target string by default, or the `Target`
-object when `return_object=True`):
+TileLang exposes the helper `tilelang.utils.target.determine_target` (returns a canonical target string or config
+by default, or the `Target` object when `return_object=True`):
 
 ```python
 from tilelang.utils.target import determine_target
 
-tvm_target = determine_target("cuda -arch=sm_80", return_object=True)
+tvm_target = determine_target({"kind": "cuda", "arch": "sm_80"}, return_object=True)
 kernel = tilelang.compile(func,>
 ```
 
@@ -90,13 +87,11 @@ You can also build targets directly through TVM:
 from tvm.target import Target
 
 target = Target("cuda", host="llvm")
-target = target.with_host(Target("llvm -mcpu=skylake"))
+target = target.with_host(Target({"kind": "llvm", "mcpu": "skylake"}))
 ```
 
-TileLang accepts either `str` or `Target` inputs; internally they are normalised and cached using the canonical
-string representation.  **In user code we strongly recommend passing target strings rather than
-`tvm.target.Target` instances—strings keep cache keys compact and deterministic across runs, whereas constructing
-fresh `Target` objects may lead to slightly higher hashing overhead or inconsistent identity semantics.**
+TileLang accepts bare target strings, target config dictionaries, and `Target` inputs. For targets with options,
+prefer config dictionaries over CLI-style strings.
 
 ## Discovering supported targets in code
 
@@ -113,9 +108,9 @@ This helper mirrors the table above and is safe to call at runtime (for example
 
 ## Troubleshooting tips
 
-- If you see `Target cuda -arch=sm_80 is not supported`, double-check the spellings and that the option is valid for
-  TVM. Any invalid switch will surface as a target-construction error.
-- Runtime errors such as “no kernel image is available” usually mean the `-arch` flag does not match the GPU you are
+- If you see `Target {'kind': 'cuda', 'arch': 'sm_80'} is not supported`, double-check the spellings and that the option is valid for
+  TVM. Any invalid option will surface as a target-construction error.
+- Runtime errors such as “no kernel image is available” usually mean the `arch` value does not match the GPU you are
   running on. Try dropping the flag or switching to the correct compute capability.
-- When targeting multiple environments, use `auto` for convenience and override with an explicit string only when
+- When targeting multiple environments, use `auto` for convenience and override with an explicit config only when
   you need architecture-specific tuning.