[Question] Does tilelang support GEMM on sm_120 now?

[GoldenStain]

Required prerequisites

• I have read the documentation https://tilelang.com.
• I have searched the Issue Tracker that this hasn't already been reported. (comment there if it has.)

Questions

I tried using examples/deepseek_v32/fp8_lighting_indexer.py, but encountered this error:

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

[coderabbitai]

📝 CodeRabbit Plan Mode

Generate an implementation plan and prompts that you can use with your favorite coding agent.

• Create Plan

Examples

• Example 1
• Example 2

---

🔗 Similar Issues

Possible Duplicates

• ISSUE-920

Related Issues

• Cannot run on Blackwell (SM100) GPUs #920
• [Feature Request] Introduce T.gemm_sp_v2 for 2:4 Sparse Tensor Core GEMM #935
• The GEMM performance cannot reach triton and pytorch #492
• [Feature Request] Upgrade cutlass to v3.5.0+ for fp8 T.gemm #143
• [Question] Restriction on Tile Size for T.GEMM? #358
• [BUG] I tried to run deepseek-v32's fp8_lighting_indexer.py on L20 with cuda 12.4, but met some bugs, is it because fp8_lighting_indexer is not supported to run on the L20? #1163
• can anyone find the function 'make_fp8_e4_16_t'? #926
• 'fp8_lighting_indexer' & 'topk_selector' seem not well supported by the latest tilelang library #923
• [Question] some mistake with 0.1.6.post1+cuda.gitb2acfc37 #1087

🔗 Related PRs

#1047 - Allow mma gemm for all cuda arch [merged]
#1246 - [Bugfix] Fix fp8 dtype for some cases [merged]
#1408 - [Typo] fix typo for SM120 [merged]

👤 Suggested Assignees

• LeiWang1999
• oncleJules
• chfwinter
• botbw
• woshildh

---

🧪 Issue enrichment is currently in open beta.

You can configure auto-planning by selecting labels in the issue_enrichment configuration.

To disable automatic issue enrichment, add the following to your .coderabbit.yaml:

issue_enrichment:
  auto_enrich:
    enabled: false

💬 Have feedback or questions? Drop into our discord or schedule a call!

[LeiWang1999]

@GoldenStain Yes, TileLang does support SM120, it's strange but I do not have sm120 avaliable currently.

https://github.com/tile-ai/tilelang/blob/main/src/op/gemm_py.cc#L127-L145

[GoldenStain]

@GoldenStain Yes, TileLang does support SM120, it's strange but I do not have sm120 avaliable currently.

https://github.com/tile-ai/tilelang/blob/main/src/op/gemm_py.cc#L127-L145

tks, I'll double check later.

[LeiWang1999]

@GoldenStain I saw that this issues was closed, has the reported problem been resolved? :)

[GoldenStain]

@GoldenStain I saw that this issues was closed, has the reported problem been resolved? :)

Not yet. We switched to Hopper temporarily, but since I’ve concluded this is likely an internal TVM issue and haven't found a workaround yet, it might be time to reopen this. I’m still looking for a solution within TVM. 😢

[PannenetsF]

@GoldenStain I did not get the error you got, but a forever hanging. cc @LeiWang1999

➜  tilelang git:(dev) ✗ python examples/deepseek_v32/fp8_lighting_indexer.py 
2025-12-24 11:50:06  [TileLang:tilelang.env:WARNING]: Loading tilelang libs from dev root: /data/tilelang/build
/usr/local/lib/python3.12/dist-packages/torch/cuda/__init__.py:63: FutureWarning: The pynvml package is deprecated. Please install nvidia-ml-py instead. If you did not install pynvml directly, please report this to the maintainers of the package that installed pynvml for you.
  import pynvml  # type: ignore[import]
2025-12-24 11:50:19  [TileLang:tilelang.language.v2.builder:WARNING]: Binding a for frame to variable may cause undefined behavior in tilelang.
Stack (most recent call last):
  File "/data/tilelang/examples/deepseek_v32/fp8_lighting_indexer.py", line 314, in <module>
    test_fp8_lighting_indexer()
  File "/data/tilelang/examples/deepseek_v32/fp8_lighting_indexer.py", line 259, in test_fp8_lighting_indexer
    logits_tl = mqa_attn_return_logits_interface(q=q_fp8, kv=kv_fp8, kv_scales=kv_scales, weights=weights, cu_seqlen_ks=ks, cu_seqlen_ke=ke)
  File "/data/tilelang/examples/deepseek_v32/fp8_lighting_indexer.py", line 208, in mqa_attn_return_logits_interface
    clean_logits_kernel = clean_logits_()
  File "/data/tilelang/tilelang/jit/__init__.py", line 423, in __call__
    kernel = self.compile(*args, **kwargs, **tune_params)
  File "/data/tilelang/tilelang/jit/__init__.py", line 354, in compile
    func = self.get_tir(*args, **kwargs)
  File "/data/tilelang/tilelang/jit/__init__.py", line 302, in get_tir
    tir = self.func(*args, **kwargs)
  File "/data/tilelang/examples/deepseek_v32/fp8_lighting_indexer.py", line 184, in clean_logits_
    @T.prim_func
  File "/data/tilelang/tilelang/language/v2/builder.py", line 875, in prim_func
    return impl(func) if func is not None else impl
  File "/data/tilelang/tilelang/language/v2/builder.py", line 864, in impl
    return prim_func_generator(**args)
  File "/data/tilelang/tilelang/language/v2/builder.py", line 651, in __call__
    self.ir_gen.gen(builder)(*args, **kwargs)
  File "/data/tilelang/examples/deepseek_v32/fp8_lighting_indexer.py", line 191, in clean_logits_kernel
    tx = T.thread_binding(0, threads, thread="threadIdx.x")
  File "/data/tilelang/tilelang/language/v2/builder.py", line 402, in bind
    res = self.bind_immutable(name, value)

[PannenetsF]

BTW, I used commit cfccd63 (12/24 17pm) and #1522 did not affect that, as the kernel has already launched (but been hanging).

[oraluben]

@PannenetsF pip install torch-c-dlpack-ext might help (if you haven't installed it yet), or you can clear tvm-ffi's cache and try again.