[BUG] T.While is broken for target="c" / CPU mode

### Required prerequisites

- [x] I have read the documentation <https://tilelang.com>.
- [x] I have searched the [Issue Tracker](https://github.com/tile-ai/tilelang/issues) that this hasn't already been reported. (comment there if it has.)

### What version of TileLang are you using?

0.1.9

### Problem description

The use of T.While loops within tilelang, currently causes compilation issues for CPU / C based backend. But not for GPU based backends

### Reproducible example code

The Python snippets:

```python
from __future__ import annotations

import sys
from pathlib import Path

import tilelang
import tilelang.language as T

def test_cpu_kernel_source_generation_with_while_reproduces_issue() -> None:
    @T.prim_func
    def main(flag: T.Buffer((1,), "int32"), out: T.Buffer((1,), "int32")):
        with T.Kernel(1, is_cpu=True):
            state = T.alloc_fragment((1,), "int32")
            state[0] = 0

            with T.While(state[0] == 0):
                state[0] = flag[0]

            out[0] = state[0]

    # This currently aborts or fails during CPU lowering instead of compiling cleanly.
    tilelang.compile(main, target="c")
```


### Traceback

```pytb
F                                                                                                        [100%]
=================================================== FAILURES ===================================================
________________________ test_cpu_kernel_source_generation_with_while_reproduces_issue _________________________

    def test_cpu_kernel_source_generation_with_while_reproduces_issue() -> None:
        @T.prim_func
        def main(flag: T.Buffer((1,), "int32"), out: T.Buffer((1,), "int32")):
            with T.Kernel(1, is_cpu=True):
                state = T.alloc_fragment((1,), "int32")
                state[0] = 0
    
                with T.While(state[0] == 0):
                    state[0] = flag[0]
    
                out[0] = state[0]
    
        # This currently aborts or fails during CPU lowering instead of compiling cleanly.
>       tilelang.compile(main, target="c")

test/tl/device_runtime/test_tilelang_cpu_while_repro.py:28: 
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 
.venv/lib/python3.13/site-packages/tilelang/jit/__init__.py:115: in compile
    return cached(
.venv/lib/python3.13/site-packages/tilelang/cache/__init__.py:74: in cached
    return _dispatch_map[execution_backend].cached(
.venv/lib/python3.13/site-packages/tilelang/cache/kernel_cache.py:345: in cached
    kernel = JITKernel(
.venv/lib/python3.13/site-packages/tilelang/jit/kernel.py:136: in __init__
    adapter = self._compile_and_create_adapter(func, out_idx)
              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
.venv/lib/python3.13/site-packages/tilelang/jit/kernel.py:247: in _compile_and_create_adapter
    artifact = tilelang.lower(
.venv/lib/python3.13/site-packages/tilelang/engine/lower.py:311: in lower
    mod = OptimizeForTarget(mod, target)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
.venv/lib/python3.13/site-packages/tilelang/engine/phase.py:293: in OptimizeForTarget
    mod = tilelang.transform.MakePackedAPI()(mod)
          ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
.venv/lib/python3.13/site-packages/tilelang/3rdparty/tvm/python/tvm/ir/transform.py:167: in __call__
    return _ffi_transform_api.RunPass(self, mod)
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ 

>   ???
E   tvm.error.InternalError: Check failed: undefined.size() == 0 (1 vs. 0) : In PrimFunc main variables [v_thread] are used, but are not passed in as API arguments

python/tvm_ffi/cython/function.pxi:968: InternalError
--------------------------------------------- Captured stdout call ---------------------------------------------
2026-05-14 01:41:34  [TileLang:tilelang.jit.kernel:INFO] (kernel.py:133): TileLang begins to compile kernel `main` with `out_idx=None`
--------------------------------------------- Captured stderr call ---------------------------------------------
[01:41:34] : Fatal: InternalError: Check failed: undefined.size() == 0 (1 vs. 0) : In PrimFunc main variables [v_thread] are used, but are not passed in as API arguments
=============================================== warnings summary ===============================================
.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'span' in 'ir.TupleType' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'span' in 'ir.FuncType' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'span' in 'ir.TensorMapType' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'name' in 'relax.TEPlaceholderOp' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'tag' in 'relax.TEPlaceholderOp' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'attrs' in 'relax.TEPlaceholderOp' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'shape' in 'relax.TEPlaceholderOp' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'dtype' in 'relax.TEPlaceholderOp' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tilelang/language/eager/builder.py:888
.venv/lib/python3.13/site-packages/tilelang/language/eager/builder.py:888
test/tl/device_runtime/test_tilelang_cpu_while_repro.py::test_cpu_kernel_source_generation_with_while_reproduces_issue
test/tl/device_runtime/test_tilelang_cpu_while_repro.py::test_cpu_kernel_source_generation_with_while_reproduces_issue
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tilelang/language/eager/builder.py:888: DeprecationWarning: Failing to pass a value to the 'type_params' parameter of 'typing._eval_type' is deprecated, as it leads to incorrect behaviour when calling typing._eval_type on a stringified annotation that references a PEP 695 type parameter. It will be disallowed in Python 3.15.
    hints[name] = _eval_type(value, globalns=globalns, localns=localns)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'policy_type' in 'tl.GemmSPWarpPolicy' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'm_warp' in 'tl.GemmSPWarpPolicy' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tvm_ffi/registry.py:85: UserWarning: Field 'n_warp' in 'tl.GemmSPWarpPolicy' duplicates an ancestor field. Child types should not re-register inherited fields.
    info = core._register_object_by_index(type_index, cls)

test/tl/device_runtime/test_tilelang_cpu_while_repro.py::test_cpu_kernel_source_generation_with_while_reproduces_issue
test/tl/device_runtime/test_tilelang_cpu_while_repro.py::test_cpu_kernel_source_generation_with_while_reproduces_issue
  /Users/featherless-ai/istack/tillelang/.venv/lib/python3.13/site-packages/tilelang/utils/deprecated.py:34: DeprecationWarning: T.Buffer(...) is deprecated, use T.Tensor(...) instead
    deprecated_warning(method_name, new_method_name, phaseout_version)

-- Docs: https://docs.pytest.org/en/stable/how-to/capture-warnings.html
=========================================== short test summary info ============================================
FAILED test/tl/device_runtime/test_tilelang_cpu_while_repro.py::test_cpu_kernel_source_generation_with_while_reproduces_issue - tvm.error.InternalError: Check failed: undefined.size() == 0 (1 vs. 0) : In PrimFunc main variables [v_thre...
1 failed, 17 warnings in 1.66s
zsh: bus error  pytest -q test/tl/device_runtime/test_tilelang_cpu_while_repro.py
```

### System information
```
Collecting environment information...
PyTorch version: 2.10.0
Is debug build: False
CUDA used to build PyTorch: None
ROCM used to build PyTorch: N/A

OS: macOS 26.4 (arm64)
GCC version: Could not collect
Clang version: 17.0.0 (clang-1700.6.3.2)
CMake version: version 4.3.2
Libc version: N/A

Python version: 3.13.12 (main, Feb  4 2026, 00:05:32) [Clang 21.1.4 ] (64-bit runtime)
Python platform: macOS-26.4-arm64-arm-64bit-Mach-O
Is CUDA available: False
CUDA runtime version: No CUDA
CUDA_MODULE_LOADING set to: N/A
GPU models and configuration: No CUDA
Nvidia driver version: No CUDA
cuDNN version: No CUDA
Is XPU available: False
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True
Caching allocator config: N/A
```

### Expected behavior

It compiles, without error, and is runnable

### Additional context

_No response_

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[BUG] T.While is broken for target="c" / CPU mode #2202

Required prerequisites

What version of TileLang are you using?

Problem description

Reproducible example code

Traceback

System information

Expected behavior

Additional context

Metadata

Assignees

Labels

Type

Fields

Projects

Milestone

Relationships

Development

[BUG] T.While is broken for target="c" / CPU mode #2202

Description

Required prerequisites

What version of TileLang are you using?

Problem description

Reproducible example code

Traceback

System information

Expected behavior

Additional context

Metadata

Metadata

Assignees

Labels

Type

Fields

Projects

Milestone

Relationships

Development

Issue actions