[BUG][Fuzzer][ice-on-valid-code] A `bfloat16` scalar kernel parameter aborts with internal `Check failed: (dtype.is_float())` instead of compiling like a `float16` scalar
- Dominant language
- Python
- Stars
- 7.4k
- Forks
- 742
- Avg merge
- 1d 2h
- Merged PRs (30d)
- 108
Description
### Required prerequisites
- [x] I have read the documentation .
- [x] I have searched the [Issue Tracker](https://github.com/tile-ai/tilelang/issues) and found no similar report.
### What version of TileLang are you using?
0.1.13 (latest release)
### System information
TileLang 0.1.13 · PyTorch 2.13.0+cu130 · CUDA 13.0 · NVIDIA L40S (sm_89). The failing pass runs on the host before any device codegen, so it is target-independent.
### Problem description
Declaring a scalar kernel parameter of type `bfloat16` aborts compilation with an internal check failure:
```
tvm.error.InternalError: Check failed: (dtype.is_float()) is false:
```
The identical kernel with the scalar typed `float16` or `float32` compiles, and `bfloat16` used as a *tensor element* type compiles — only a `bfloat16` **scalar parameter** trips it. The failure is an `ICHECK` (an internal assertion), not a user-facing diagnostic, so a legitimate typed input surfaces as an opaque compiler crash rather than a computed kernel or a clean error message. Not a regression — see Provenance.
### Reproducible example code
```python
import tilelang, tilelang.language as T
# CONTROL: a float16 scalar parameter compiles (float16 has DataType code kFloat)
@T.prim_func
def add_f16(A: T.Tensor((128,), "float16"), s: T.float16, B: T.Tensor((128,), "float16")):
with T.Kernel(1, threads=128):
i = T.get_thread_binding()
B[i] = A[i] + s
# CONTROL: bfloat16 as a *tensor element* type compiles
@T.prim_func
def add_bf16_tensor(A: T.Tensor((128,), "bfloat16"), B: T.Tensor((128,), "bfloat16")):
with T.Kernel(1, threads=128):
i = T.get_thread_binding()
B[i] = A[i] + T.cast(3.0, "bfloat16")
# REPRO: a bfloat16 *scalar parameter* — same kernel, only the scalar dtype differs
@T.prim_func
def add_bf16(A: T.Tensor((128,), "bfloat16"), s: T.bfloat16, B: T.Tensor((128,), "bfloat16")):
with T.Kernel(1, threads=128):
i = T.get_thread_binding()
B[i] = A[i] + s
for name, fn in [("CONTROL float16 scalar", add_f16),
("CONTROL bfloat16 tensor-only", add_bf16_tensor),
("REPRO bfloat16 scalar", add_bf16)]:
try:
tilelang.compile(fn, execution_backend="cython")
print(f"{name}: COMPILED OK")
except Exception as e:
print(f"{name}: {type(e).__name__}: {str(e).strip()[:120]}")
```
Output:
```
CONTROL float16 scalar: COMPILED OK
CONTROL bfloat16 tensor-only: COMPILED OK
REPRO bfloat16 scalar: InternalError: Check failed: (dtype.is_float()) is false:
```
### Traceback
```
[..] : Fatal: InternalError: Check failed: (dtype.is_float()) is false:
Traceback (most recent call last):
...
File ".../tilelang/engine/phase.py", line 293, in OptimizeForTarget
mod = tilelang.transform.MakePackedAPI()(mod)
...
File "", line 0, in tvm::tl::MakePackedAPI(tvm::tir::PrimFunc)
File "", line 0, in tvm::runtime::detail::LogFatalImpl(...)
tvm.error.InternalError: Check failed: (dtype.is_float()) is false:
```
### Expected behavior
A `bfloat16` scalar parameter compiles and computes, exactly as `float16`/`float32` scalar parameters already do (`bfloat16` is an exported scalar dtype — `T.bfloat16` — and it already works as a tensor element type). If the packed-API binder genuinely cannot marshal a given scalar dtype, a clear error naming the dtype would be preferable to an internal `Check failed`.
### Additional context
**Root cause.** `MakePackedAPI`'s scalar-argument binder classifies each scalar parameter as bool / int / uint / float, and its float branch guards on `DataType::is_float()`, which is true only for the `kFloat` type code — `bfloat16` has the distinct `kBFloat` code, so it falls through to the final `else` and fails `ICHECK(dtype.is_float())` at [`src/transform/make_packed_api.cc#L563`](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/transform/make_packed_api.cc#L563) (v0.1.13). The [preceding `is_bool()` and `is_int()/is_uint()` branches](https://github.com/tile-ai/tilelang/blob/8001cc4ccf6149382d2019654a19f59c1d4d0482/src/transform/make_packed_api.cc#L540-L562) enumerate the standard scalar dtypes but omit the narrow-float families (`bfloat16`, and by inspection the `float8_*` codes), so any scalar parameter of those dtypes reaches the assertion.
Sibling face: a float16 scalar parameter passes this pass but breaks downstream
`float16` has the `kFloat` code, so it clears the `is_float()` guard above — but a `float16` scalar parameter is still not usable end-to-end: with `execution_backend="tvm_ffi"` the generated host C references an undeclared `half` (`error: 'half' was not declared in this scope`), and with `execution_backend="cython"` the kernel launches but reads the scalar wrong (`cudaErrorIllegalAddress` at runtime). That is a distinct root cause (host-side scalar marshaling / host type mapping), separate from the `bfloat16` `is_float()` gate above; noting it here only because the underlying gap — narrow-float scalar parameters — is shared. The minimal repro above isolates the `bfloat16` assertion, which is the clean, target-independent failure.
**Suggested fix.** Widen the float branch's guard at `make_packed_api.cc:563` from `is_float()` to also admit `is_bfloat16()` (and the `float8_*` codes) so a `bfloat16` scalar marshals through the same `kTVMFFIFloat` path as `float16`/`float32`. If some dtype genuinely cannot be marshaled, replace the bare `ICHECK` with an error that names the dtype instead of an internal crash. This is a C++ change requiring a rebuild (not verified end-to-end).
**Generalization (tested this session, 0.1.13 / L40S).** This is a *missing-case* class: the binder enumerates bool / int / uint / `kFloat`-float and asserts on everything else, so every narrow-float scalar dtype falls through. Tested — `bfloat16` scalar, `float8_e4m3` scalar, and `float8_e5m2` scalar all abort with the identical `Check failed: (dtype.is_float()) is false`; `int32` and `uint32` scalars compile OK (controls), as do `float16`/`float32` scalars (`kFloat` code). The `float8_*` faces are no longer "by inspection" — they were run and share this one root. The title stays on `bfloat16` as the clean, target-independent worked example; the fp8 scalars are recorded here as same-root neighbors.
**Provenance.** The scalar float branch with `ICHECK(dtype.is_float())` was introduced when tvm-ffi became the default execution backend in [#1259](https://github.com/tile-ai/tilelang/pull/1259) (merged 2025-11-18); `git log -S` on the assertion string attributes it to that commit (`74da3696`). Not a regression — scalar `bfloat16` parameters have not compiled since this binder shipped.
**Dedup.** I searched the open and closed tracker and found no existing report of this defect.
**Reach.** The trigger is a scalar kernel parameter typed `bfloat16` — `T.bfloat16` is an exported scalar dtype, in the same family as `T.float16`/`T.float32` which the binder accepts. `grep` of `examples/` and `testing/` finds no kernel that passes any float-typed scalar parameter (float scalars are passed as tensors or Python constants), so no shipped test exercises the scalar-parameter binder for narrow floats — which is why CI is green.
**Impact.** The trigger is narrow — only a scalar (non-tensor) kernel parameter typed `bfloat16` (or, by inspection, `float8_*`); the accepted `float16`/`float32`/int scalar dtypes are unaffected. When it fires it is a host-side `ICHECK` abort during `MakePackedAPI`, before any codegen: the crash is loud and deterministic, blocks that one kernel from building, and cannot silently reach a running workload or corrupt any output. Fixing it closes the narrow-float scalar-parameter boundary and replaces an opaque internal assertion with either a working marshal path or a dtype-named error.
Contributor guide
Research direction
Start with src/transform/make_packed_api.cc around lines 540-563 and run the bfloat16 scalar reproducer with the listed controls. Done means the bfloat16 scalar parameter compiles and computes, or produces a clear dtype-specific error instead of the internal assertion; add coverage for the reported narrow-float cases if the test location is identified.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp, python
- Domain
- backend-api-design, compilers
- Issue type
- Bug
- Difficulty
- 3/5
- Estimated time
- 1-2 days
- Activity status
- Quiet
- Clarity
- Clearly specified
- Newbie friendliness
- 68/100