[BUG][Fuzzer][ice-on-valid-code] `T.cumsum`/`T.cummax` leaks a raw NVCC `static_assert` instead of compiling or cleanly rejecting the block size

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.)

What version of TileLang are you using?

0.1.11 (source pinned to tag v0.1.11, commit cd37ed5fc35ae7a60a1277c8eb49028174ac51e6).

System information

• TileLang: 0.1.11 (source pinned to tag v0.1.11, commit cd37ed5fc35ae7a60a1277c8eb49028174ac51e6)
• Python: 3.12.3
• PyTorch: 2.12.1+cu130 (CUDA 13.0)
• GPU: NVIDIA H100 80GB HBM3 (sm_90)

The defect is in the architecture-independent scan lowering + device template (static_assert on the block thread count), not Hopper-specific. Reproduced on sm_90 (the only arch run this session); the arch is incidental — the assertion fires at NVCC compile time on the generated .cu, independent of target GPU.

Problem description

T.cumsum (and T.cummax) over a shared buffer aborts the NVCC compile with a static_assert when the kernel's block thread count is not one of {32, 64, 128, 256, 512, 1024} — e.g. T.Kernel(1, threads=96). The same block size compiles and runs everywhere else in the kernel (a T.copy at threads=96 compiles and runs fine), so 96 is a block size TileLang admits, yet a scan over that block leaks an internal device-template assertion instead of either compiling or rejecting the kernel with a front-end error:

src/tl_templates/cuda/scan.h(102): error: static assertion failed
    static_assert(threads == 1024 or threads == 512 or threads == 256 or

This is a diagnostic-quality defect, not a silent miscompile: the compile aborts loudly, but into an internal header instead of surfacing a TileLang-level message.

Reproducible example code

import torch, tilelang
import tilelang.language as T

N = 128

def cumsum_kernel(threads):
    @T.prim_func
    def kern(A: T.Tensor((N,), "float32"), B: T.Tensor((N,), "float32")):
        with T.Kernel(1, threads=threads):
            s = T.alloc_shared((N,), "float32")
            T.copy(A, s)
            T.cumsum(src=s, dst=s, dim=0)
            T.copy(s, B)
    return kern

def copy_only(threads):                       # control: same block size, no scan
    @T.prim_func
    def kern(A: T.Tensor((N,), "float32"), B: T.Tensor((N,), "float32")):
        with T.Kernel(1, threads=threads):
            s = T.alloc_shared((N,), "float32")
            T.copy(A, s)
            T.copy(s, B)
    return kern

def build(k):
    try:
        return tilelang.compile(k, out_idx=[1], target="cuda"), None
    except Exception as e:
        return None, str(e)

A = torch.arange(1, N + 1, dtype=torch.float32, device="cuda")

jit, _ = build(cumsum_kernel(128))            # control-defined
print("threads=128 + cumsum:", "OK, matches torch =", torch.allclose(jit(A), torch.cumsum(A, 0)))

jit, _ = build(copy_only(96))                 # control-admitted: 96 is a legal block size
print("threads=96 copy-only :", "OK, 96 admitted =", torch.allclose(jit(A), A))

jit, err = build(cumsum_kernel(96))           # trigger
print("threads=96 + cumsum  :", "OK" if jit else "COMPILE ABORT: " + err.strip().splitlines()[-1])

threads=128 + cumsum: OK, matches torch = True
threads=96 copy-only : OK, 96 admitted = True
threads=96 + cumsum  : COMPILE ABORT: 1 error detected in the compilation of "…/tvm_kernels.cu".

.../src/tl_templates/cuda/scan.h(102): error: static assertion failed
    static_assert(threads == 1024 or threads == 512 or threads == 256 or
          detected during:
            instantiation of class "tl::InclusiveScan1D<Reducer, threads, reverse>
              [with Reducer=tl::ScanSumOp, threads=96, reverse=false]" at line 152
            instantiation of "void tl::CumSum1D<threads, reverse>::run(...)
              [with threads=96, reverse=false, T=float, SEG=32]"
1 error detected in the compilation of ".../tvm_kernels.cu".

Traceback

The failure is an nvcc compile abort surfaced by TileLang as a `CompilationError`; the load-bearing line is:

src/tl_templates/cuda/scan.h(102): error: static assertion failed
    static_assert(threads == 1024 or threads == 512 or threads == 256 or ...)
      detected during instantiation of tl::CumSum1D<threads=96, ...>::run

Expected behavior

The T.cumsum/T.cummax docstring places no restriction on the block thread count — it shows threads=128 and threads=256 examples and never lists a supported set — and T.copy and other tile ops already accept T.Kernel(threads=96) (probe B compiles and runs). So 96 is an in-contract block size, and a scan over the same block should behave consistently — either:

• compile and run the scan at that block size, or
• reject at the front end with a clear message (e.g. "cumsum/cummax requires block threads in {32,64,128,256,512,1024}").

Leaking an internal device-template static_assert from a private header is neither — the user gets a raw NVCC assertion pointing into TileLang's generated code, with no indication that the block thread count is the cause.

Additional context

Root cause. The shared-scan lowering forwards the full block thread count straight into the device template's threads parameter, and the template static_asserts that the value is one of {32,64,128,256,512,1024} — with no front-end guard restricting T.Kernel(threads=...) to that set for a scan. LowerSharedScan reads the block size at src/backend/common/op/scan.h:39 (auto threads = T.thread_bounds->extent;) and emits it verbatim into the template symbol at scan.h:54 (→ tl::CumSum1D<96, false>). The device template then asserts on it at src/tl_templates/cuda/scan.h:102 (1-D) and scan.h:115 (2-D). T.cumsum and T.cummax share the identical LowerSharedScan path, so both fail identically.

Suggested fix (proposed, not built). Two options:

• Add a front-end guard in LowerSharedScan that rejects (or normalizes) a block thread count outside the supported set before it reaches the template, so the user sees a TileLang error rather than an nvcc assertion.
• Or relax the template. The 1-D scan uses a single 32-lane warp — run early-returns for threadIdx.x >= SEG (SEG = 32, scan.h:107) — so for the 1-D case the threads template parameter is effectively only a partition/tiling hint; it could be clamped rather than asserted.

Provenance. Both the thread_bounds->extent forward and the static_assert shipped together in #2262 "[TileOP] Add scan operators" (commit 146a1d3b, merged 2026-05-25) — the PR that introduced the scan operators; introduced, not relocated. Verified via git log -S.

I searched the open and closed tracker (cumsum, cummax threads, scan.h, static_assert scan) and found no report of this compile-time block-thread-count assertion. The three nearby scan issues are all distinct — they are triggered by the buffer geometry / region, not the block thread count, and none aborts at the scan.h static_assert:

• #2523 (open) — T.cumsum/T.cummax over a non-contiguous 2-D sub-region silently returns wrong numbers (silent miscompile, region-stride trigger).
• #2536 (open) — T.cumsum/T.cummax over a row-offset 2-D sub-region silently scans the wrong rows (silent miscompile, non-zero region min).
• #2284 (closed) — T.cumsum(dim=0) on a rectangular 2-D shared tile (e.g. (32, 64)) hits a runtime CUDA_ERROR_ILLEGAL_ADDRESS from the 2-D axis-extent mapping.

All three keep the block thread count in the supported set and fail on the data layout; this report is the orthogonal case (a supported geometry with an unsupported block thread count, failing at NVCC compile time). Whatever resolved the closed #2284 left this path untouched — the static_assert still fires on v0.1.11.

Start with LowerSharedScan in src/backend/common/op/scan.h, then inspect the scan template assertions in src/tl_templates/cuda/scan.h at lines 102 and 115. Reproduce the provided Python kernels with threads=96 and 128 for T.cumsum and T.cummax. Done means unsupported block sizes no longer leak a raw NVCC assertion: they either compile successfully or produce a clear TileLang-level diagnostic.