[Question] The number of threads in the kernel inconsistent with the parameters I passed in.

[printffox]

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

Just as the title indicates, I created a Kernel by

with T.Kernel(T.ceildiv(seq_len, block_m), heads, batch, threads=128) as (bx, by, bz):

However, when I used NsightCompute to profile this Kernel, I found that the blockDim of this Kernel is (256,1, 1). Why is that?

[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

Related Issues

• [Question] Restriction on Tile Size for T.GEMM? #358
• [Question] examples/flash_attention/example_mha_bwd_bhsd.py precision error？ #1213
• [Question] tvm.error.InternalError: Check failed: (M % kMPerWarp == 0) is false: M must be divisible by 16, but got 8 #1191
• [Question] Incorrect CUDA Generation: Static Indexing of Local Fragment Array inside T.Pipelined Loop #1477

👤 Suggested Assignees

• shadowpa0327
• printffox
• lemyx
• drewjin

---

🧪 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!

[Rachmanino]

I guess you did not disable TMA in the config passed and used TMA copy in the kernel. As a result, tilelang compiler will launch an extra producer wg for TMA

[LeiWang1999]

@printffox yes, tilelang will launch an extra warp group to issue tma if that was analyzed to be enabled.