using the bundled installation or the installation from prebuilt tvm will cause the SM120 error

[hantao-zhou]

tvm.error.InternalError: Check failed: (0) is false: Unsupported target for gemm: cuda -keys=cuda,gpu -arch=sm_120 -max_num_threads=1024 -thread_warp_size=32

[hantao-zhou]

(mla-bench) hans@fedora:~/workspace/mla-bench$ python external/tilelang/examples/gemm/test_example_gemm.py
================================================================================================================================ test session starts =================================================================================================================================
platform linux -- Python 3.14.0, pytest-8.4.2, pluggy-1.6.0
rootdir: /home/hans/workspace/mla-bench/external/tilelang/examples
configfile: pytest.ini
collected 1 item

external/tilelang/examples/gemm/test_example_gemm.py F [100%]

====================================================================================================================================== FAILURES ======================================================================================================================================
_________________________________________________________________________________________________________________________________ test_example_gemm __________________________________________________________________________________________________________________________________

def test_example_gemm():

  example_gemm.main()

external/tilelang/examples/gemm/test_example_gemm.py:22:

---

external/tilelang/examples/gemm/example_gemm.py:31: in main
kernel = matmul(1024, 1024, 1024, 128, 128, 32)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
external/tilelang/tilelang/jit/init.py:216: in wrapper
kernel_result = compile(
external/tilelang/tilelang/jit/init.py:81: in compile
return cached(
external/tilelang/tilelang/cache/init.py:28: in cached
return _kernel_cache_instance.cached(
external/tilelang/tilelang/cache/kernel_cache.py:184: in cached
kernel = JITKernel(
external/tilelang/tilelang/jit/kernel.py:129: in init
adapter = self._compile_and_create_adapter(func, out_idx)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
external/tilelang/tilelang/jit/kernel.py:227: in _compile_and_create_adapter
artifact = tilelang.lower(
external/tilelang/tilelang/engine/lower.py:226: in lower
mod = LowerAndLegalize(mod, target)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
external/tilelang/tilelang/engine/phase.py:106: in LowerAndLegalize
mod = tilelang.transform.LayoutInference()(mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
external/tilelang/3rdparty/tvm/python/tvm/ir/transform.py:167: in call
return _ffi_transform_api.RunPass(self, mod)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

---

???
E tvm.error.InternalError: Check failed: (0) is false: Unsupported target for gemm: cuda -keys=cuda,gpu -arch=sm_120 -max_num_threads=1024 -thread_warp_size=32

tvm/ffi/cython/function.pxi:228: InternalError
-------------------------------------------------------------------------------------------------------------------------------- Captured stdout call --------------------------------------------------------------------------------------------------------------------------------
2025-10-12 12:25:46 [TileLang:tilelang.jit.kernel:INFO]: TileLang begins to compile kernel gemm with out_idx=[-1]
============================================================================================================================== short test summary info ===============================================================================================================================
FAILED external/tilelang/examples/gemm/test_example_gemm.py::test_example_gemm - tvm.error.InternalError: Check failed: (0) is false: Unsupported target for gemm: cuda -keys=cuda,gpu -arch=sm_120 -max_num_threads=1024 -thread_warp_size=32
================================================================================================================================= 1 failed in 0.21s ==================================================================================================================================
(mla-bench) hans@fedora:~/workspace/mla-bench$

i used the bundled part, but it still won't work

[LeiWang1999]

@hantao-zhou Which TileLang version were you using?

[hantao-zhou]

the one from the github main branch , i followed the install from bundled tvm

[hantao-zhou]

@LeiWang1999
(mla-bench) hans@fedora:~/workspace/mla-bench$ python
Python 3.12.11 | packaged by conda-forge | (main, Jun 4 2025, 14:45:31) [GCC 13.3.0] on linux
Type "help", "copyright", "credits" or "license" for more information.
Ctrl click to launch VS Code Native REPL

import tilelang
tilelang.version
'0.1.6.post1+b0b5347a'

i think this presents the exact version number

[hantao-zhou]

@LeiWang1999 i have tried the latest version in the release, still have the same error

[oraluben]

Can you share your nvidia-smi output?

[hantao-zhou]

Can you share your nvidia-smi output?

(mla-bench) hans@fedora:~/workspace/mla-bench$ nvidia-smi
Tue Oct 14 13:12:17 2025
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 580.95.05 Driver Version: 580.95.05 CUDA Version: 13.0 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA GeForce RTX 5090 Off | 00000000:01:00.0 Off | N/A |
| 0% 44C P8 21W / 600W | 4MiB / 32607MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+
| 1 NVIDIA GeForce RTX 5090 Off | 00000000:03:00.0 Off | N/A |
| 0% 37C P8 23W / 600W | 4MiB / 32607MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
+-----------------------------------------------------------------------------------------+
(mla-bench) hans@fedora:~/workspace/mla-bench$

[oraluben]

@hantao-zhou Current tilelang wheel do not support CUDA 13 with driver version 580. You could build from source. cc @LeiWang1999

[hantao-zhou]

@hantao-zhou Current tilelang wheel do not support CUDA 13 with driver version 580. You could build from source. cc @LeiWang1999

Thanks for the reply, but i was indeed installing from the source, as i have mentioned in my initial post, it was following the instrucitons of building from source, using the bundled TVM.

[oraluben]

Thanks for your clarification. Currently I don't have sm120 device with driver 130, and cannot reproduce your issue. However I do found the possible reason of your issue.

Would you mind trying code from #1047?

[hantao-zhou]

Thanks for the effort and the reply, yes, this one solves the permission of compiling to blackwell. But another problem was raised

`
(opfusionbench) hans@fedora:/workspace/opfusionbench/external/tilelang$ python examples/gemm
gemm/ gemm_fp8/ gemm_sm100/ gemm_sp/ gemm_splitk/ gemm_streamk/
(opfusionbench) hans@fedora:/workspace/opfusionbench/external/tilelang$ python examples/gemm
gemm/ gemm_fp8/ gemm_sm100/ gemm_sp/ gemm_splitk/ gemm_streamk/
(opfusionbench) hans@fedora:~/workspace/opfusionbench/external/tilelang$ python examples/gemm/test_example_gemm.py
==================================== test session starts ====================================
platform linux -- Python 3.12.12, pytest-8.4.2, pluggy-1.6.0
rootdir: /home/hans/workspace/opfusionbench/external/tilelang/examples
configfile: pytest.ini
collected 4 items

examples/gemm/test_example_gemm.py F... [100%]

========================================= FAILURES ==========================================
________________________________ test_example_gemm_autotune _________________________________

def test_example_gemm_autotune():
    # enable roller for fast tuning

  example_gemm_autotune.main(M=1024, N=1024, K=1024, with_roller=True)

examples/gemm/test_example_gemm.py:10:

---

examples/gemm/example_gemm_autotune.py:246: in main
result = get_best_config(M, N, K, with_roller)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
examples/gemm/example_gemm_autotune.py:150: in get_best_config
kernel=kernel, configs=get_configs(M, N, K, with_roller)).set_compile_args(
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

---

M = 1024, N = 1024, K = 1024, with_roller = True, topk = 20

def get_configs(M, N, K, with_roller=False, topk=20):
    """
    Generate a list of kernel tuning configuration dictionaries for a tiled matrix-multiply.

    When with_roller is True this queries the MatmulTemplate roller to produce up to `topk` recommended
    configurations (device-specific TensorCore-friendly tilings). Each returned dict contains:
      - block_M, block_N, block_K: tile sizes
      - num_stages: pipeline staging (0 means no explicit staging)
      - thread_num: total threads used for the block
      - enable_rasteration: whether a rasterization/swizzle layout was recommended (note spelling)

    When with_roller is False this returns the Cartesian product of a fixed set of candidate
    parameters; the returned dicts use the backward-compatible key name "enable_rasteration" for that flag.

    Parameters:
        M, N, K (int): GEMM dimensions used to generate valid tile sizes.
        with_roller (bool): If True, use MatmulTemplate's roller to generate device-aware hints;
            otherwise use a predefined candidate grid.
        topk (int): Maximum number of roller hints to request when with_roller is True.

    Returns:
        List[dict]: A list of configuration dictionaries as described above.

    Raises:
        ValueError: if with_roller is True but the roller returns no hints.
    """
    if with_roller:
        arch = CUDA("cuda") if torch.version.hip is None else CDNA("hip")
        carve_template = MatmulTemplate(
            M=M,
            N=N,
            K=K,
            in_dtype="float16",
            out_dtype="float16",
            accum_dtype="float",
        ).with_arch(arch)

        func = carve_template.equivalent_function()
        assert func is not None, "Function is None"
        roller_hints = carve_template.recommend_hints(topk=topk)
        if roller_hints is None:
            raise ValueError("No Roller Hints Found for TensorCore Scheduling")
        configs = []
        for hint in roller_hints:
            config = {}
            block_m, block_n = hint.block

          warp_m, warp_n = hint.warp

            ^^^^^^^^^^^^^^

E ValueError: not enough values to unpack (expected 2, got 0)

examples/gemm/example_gemm_autotune.py:68: ValueError
===================================== warnings summary ======================================
gemm/test_example_gemm.py::test_example_gemm
/home/hans/workspace/opfusionbench/.pixi/envs/default/lib/python3.12/site-packages/torch/profiler/profiler.py:509: UserWarning: Profiler won't be using warmup, this can skew profiler results
warn("Profiler won't be using warmup, this can skew profiler results")

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html
================================== short test summary info ==================================
FAILED examples/gemm/test_example_gemm.py::test_example_gemm_autotune - ValueError: not enough values to unpack (expected 2, got 0)
========================== 1 failed, 3 passed, 1 warning in 19.77s ==========================
(opfusionbench) hans@fedora:~/workspace/opfusionbench/external/tilelang$
`

this is the output for the example of gemm, now there seems to be some error with the autotune

Thanks for your clarification. Currently I don't have sm120 device with driver 130, and cannot reproduce your issue. However I do found the possible reason of your issue.

Would you mind trying code from #1047?

[oraluben]

That seems to be a bug in example code