AttributeError: module 'tilelang.transform._ffi_api' has no attribute 'WarpSpecializedPipeline'

[JustinJin04]

After building BitBLAS from source according to the documentation, I tried to run the example code from quick start and encountered the problem mentioned in title.

Environment:

Docker Image: nvcr.io/nvidia/pytorch:23.01-py3

GPU: NVIDIA H100 (sm_90)

CUDA Version: 12.0.140

Driver Version: 570.86.15

GLIBC Version: 2.31

Python Version: 3.8.10

BitBLAS Commit: 0c51e34

Error log:

2025-06-04 06:28:11  [TileLang:tilelang.env:WARNING]: Composable Kernel is not installed or found in the expected path
2025-06-04 06:28:14 [BitBLAS:WARNING]: TVM target not found. Please set the TVM target environment variable using `export TVM_TARGET=<target>`, where <target> is one of the available targets can be found in the output of `tools/get_available_targets.py`.
2025-06-04 06:28:14 [BitBLAS:INFO]: Auto detected target: cuda
2025-06-04 06:28:14 [BitBLAS:WARNING]: Hopper architecture is not supported for dequantize
[06:28:20] /zhan/BitBLAS/3rdparty/tilelang/src/transform/warp_specialized_rewriter.cc:1114: Warning: WarpSpecialize will be disabled because the program uses thread tags other than threadIdx.x
If you want to use warp specialization, please refactor your program to use threadIdx.x only
Lowering result:
Module(llvm, f6823c8)
Traceback (most recent call last):
  File "/zhan/flashinfer_profiler/examples/bitblas/test.py", line 24, in <module>
    matmul = bitblas.Matmul(config=matmul_config)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/ops/general_matmul/__init__.py", line 391, in __init__
    self.dispatch(target, from_database, source_format, enable_tuning)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/ops/general_matmul/__init__.py", line 410, in dispatch
    self._build_default_module(target)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/ops/operator.py", line 286, in _build_default_module
    self._build_runtime_module(target)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/ops/operator.py", line 230, in _build_runtime_module
    wrapped_source = self.wrapper.wrap(
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/builder/wrapper/tl.py", line 420, in wrap
    wrapper = wrapper_class(self.scheduled_ir_module, c_source, self.arch)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/builder/wrapper/tl.py", line 46, in __init__
    self.parse_source_information()
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/builder/wrapper/tl.py", line 52, in parse_source_information
    device_mod = get_annotated_device_mod(self.mod, self.arch.target, backend=self.backend)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/utils/rtmod_analysis.py", line 120, in get_annotated_device_mod
    return get_annotated_device_mod_from_tl(mod, target)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/utils/rtmod_analysis.py", line 26, in get_annotated_device_mod_from_tl
    mod = tilelang.transform.WarpSpecializedPipeline()(mod)
  File "/zhan/bitblas_env/lib/python3.8/site-packages/bitblas/3rdparty/tilelang/tilelang/transform/__init__.py", line 96, in WarpSpecializedPipeline
    return _ffi_api.WarpSpecializedPipeline()  # type: ignore
AttributeError: module 'tilelang.transform._ffi_api' has no attribute 'WarpSpecializedPipeline'

[Tridu33]

[tilelang-ascend]# python "/home/shared/tilelang-ascend/examples/gemm/example_gemm.py"
Traceback (most recent call last):
  File "/home/shared/tilelang-ascend/examples/gemm/example_gemm.py", line 56, in <module>
    func = matmul(M, N, K, 128, 256, 64)
  File "/home/shared/tilelang-ascend/tilelang/jit/__init__.py", line 184, in wrapper
    kernel_result = compile(
  File "/home/shared/tilelang-ascend/tilelang/jit/__init__.py", line 70, in compile
    return cached(
  File "/home/shared/tilelang-ascend/tilelang/cache/__init__.py", line 30, in cached
    return _kernel_cache_instance.cached(
  File "/home/shared/tilelang-ascend/tilelang/cache/kernel_cache.py", line 167, in cached
    kernel = JITKernel(
  File "/home/shared/tilelang-ascend/tilelang/jit/kernel.py", line 109, in __init__
    adapter = self._compile_and_create_adapter(func, out_idx)
  File "/home/shared/tilelang-ascend/tilelang/jit/kernel.py", line 209, in _compile_and_create_adapter
    artifact = tilelang.lower(
  File "/home/shared/tilelang-ascend/tilelang/engine/lower.py", line 219, in lower
    mod = OptimizeForTarget(mod, target)
  File "/home/shared/tilelang-ascend/tilelang/engine/phase.py", line 90, in OptimizeForTarget
    mod = tilelang.transform.AscendSyncInsert()(mod)
  File "/home/shared/tilelang-ascend/tilelang/transform/__init__.py", line 366, in AscendSyncInsert
    return _ffi_api.AscendSyncInsert()  # type: ignore
AttributeError: module 'tilelang.transform._ffi_api' has no attribute 'AscendSyncInsert'

same problem here, Ascend 910B, Python 3.9, torch 2.7.1, torch_npu 2.7.1
commit 52f7ead7e52c6fd3398a888b0e95d035a128e5fc (HEAD -> ascendc_pto, origin/ascendc_pto, origin/HEAD)
Author: HanXu 46563594+ArmandAlbert@users.noreply.github.com