Summary
When a T.symbolic dimension (lowered to int64_t) is combined with a Python int constant via T.min(...) to compute a dynamic tail length for a copy, codegen produces C++ that fails to compile with two distinct errors:
call to 'min' is ambiguous — codegen emits min(<int32 literal>, <int64 expr>). The CCE/bisheng runtime only provides min(int,int) and min(long,long), so the mixed-width call has no unique best match.
use of undeclared identifier 'T' — the dynamic length value (T.min(...)) is streamed verbatim (Python expression text) into a compile-time template argument of copy_gm_to_ub<...>.
Minimal reproducer
import tilelang
from tilelang import language as T
@tilelang.jit(out_idx=[1])
def dyncopy(TILE=512, MAX_CORES=24, dtype="float32"):
total = T.symbolic("TOTAL") # -> int64_t
@T.prim_func
def main(X: T.Tensor((total,), dtype), Y: T.Tensor((total,), dtype)):
with T.Kernel(MAX_CORES, is_npu=True) as (cid, vid):
ub = T.alloc_ub([TILE], dtype)
ntiles = T.ceildiv(total, TILE)
with T.Scope("V"):
for it in T.serial(T.ceildiv(ntiles, MAX_CORES)):
tid = cid + it * MAX_CORES
if tid < ntiles:
s = tid * TILE
L = T.min(TILE, total - s) # int32 TILE vs int64 (total - s)
T.copy(X[s:s + L], ub[0:L])
T.copy(ub[0:L], Y[s:s + L])
return main
dyncopy() # raises at bisheng compile
Generated C++ (the failing lines)
tl::ascend::copy_gm_to_ub<float, T.min(512, TOTAL - cid * 512 - it * 12288)>( // <- "T.min(...)" leaked into template arg
ub[0], X[((it * 12288) + (cid * 512))], TOTAL, 1,
min(512, ((TOTAL - (cid * 512)) - (it * 12288))), 0.000000e+00f); // <- min(int, int64_t) ambiguous
Compiler output
error: use of undeclared identifier 'T'
error: call to 'min' is ambiguous
CCE_INTRINSIC[aicore] long min(long in1, long in2)
CCE_INTRINSIC[aicore] int min(int in1, int in2)
Root cause
T.symbolic dims lower to int64_t; a Python int tile size stays int32. The scalar MinNode is printed by the base codegen's PrintBinaryExpr(op, "min", ...) as a bare min(a, b) without unifying operand integer widths, so the mixed min(int, int64_t) is ambiguous against the only available overloads.
- Separately, a runtime length used as the last buffer extent is emitted into the
copy_gm_to_ub<dtype, EXTENT> template parameter (must be a compile-time constant), where PrintExpr of the dynamic value yields the literal text T.min(...).
Proposed fix
- For scalar
MinNode/MaxNode with mixed integer operand widths, cast operands to the common (wider) type, e.g. emit min(static_cast<int64_t>(512), TOTAL - ...) — or promote the int literal during type inference so both operands are int64_t.
- For the template-arg path: a dynamic (non-const) extent should be rejected with a clear diagnostic (or supported via the runtime-extent argument only), not streamed as Python expression text into a template parameter.
Environment
- tilelang-ascend commit
e74c06b4d903f2c418a7a4d5712d5afa9b80f841 (branch ascendc_pto)
- bisheng
--npu-arch=dav-2201, -std=c++17
- A fixed compile-time tile extent (no
T.min, full-tile copies) compiles fine; only the dynamic-tail-length pattern triggers this.
Summary
When a
T.symbolicdimension (lowered toint64_t) is combined with a Pythonintconstant viaT.min(...)to compute a dynamic tail length for a copy, codegen produces C++ that fails to compile with two distinct errors:call to 'min' is ambiguous— codegen emitsmin(<int32 literal>, <int64 expr>). The CCE/bisheng runtime only providesmin(int,int)andmin(long,long), so the mixed-width call has no unique best match.use of undeclared identifier 'T'— the dynamic length value (T.min(...)) is streamed verbatim (Python expression text) into a compile-time template argument ofcopy_gm_to_ub<...>.Minimal reproducer
Generated C++ (the failing lines)
Compiler output
Root cause
T.symbolicdims lower toint64_t; a Pythoninttile size staysint32. The scalarMinNodeis printed by the base codegen'sPrintBinaryExpr(op, "min", ...)as a baremin(a, b)without unifying operand integer widths, so the mixedmin(int, int64_t)is ambiguous against the only available overloads.copy_gm_to_ub<dtype, EXTENT>template parameter (must be a compile-time constant), wherePrintExprof the dynamic value yields the literal textT.min(...).Proposed fix
MinNode/MaxNodewith mixed integer operand widths, cast operands to the common (wider) type, e.g. emitmin(static_cast<int64_t>(512), TOTAL - ...)— or promote the int literal during type inference so both operands areint64_t.Environment
e74c06b4d903f2c418a7a4d5712d5afa9b80f841(branchascendc_pto)--npu-arch=dav-2201,-std=c++17T.min, full-tile copies) compiles fine; only the dynamic-tail-length pattern triggers this.