Repository navigation
[BUG] cute.ceil_div throws internal error with int64 input #3734
Description
Which component has the problem?
CuTe DSL
Bug Report
Describe the bug
cute.ceil_div doesn't work with cutlass.Int64
Steps/Code to reproduce bug
from cutlass import cute, Int32, Int64
def run(name, fn, a, b):
try:
cute.compile(fn, a, b)
print(f"{name:45s} OK")
except Exception as e:
lines = [l for l in str(e).splitlines() if l.strip()]
print(f"{name:45s} FAIL: " + next((l for l in lines if "mismatch" in l or "INTERNAL" in l), lines[0])[:110])
@cute.jit
def only_cdiv_64(a: Int64, b: Int64):
cute.ceil_div(a, b) # result unused, nothing else in the function
@cute.jit
def only_cdiv_32(a: Int64, b: Int32):
cute.ceil_div(a, b)
@cute.jit
def cdiv_then_add_64(a: Int64, b: Int64):
q = cute.ceil_div(a, b)
_ = q + Int64(1) # plain arithmetic on the result
run("ceil_div(Int64, Int64) only, unused", only_cdiv_64, Int64(100), Int64(7))
run("ceil_div(Int64, Int32) only, unused", only_cdiv_32, Int64(100), Int32(7))
run("ceil_div(Int64, Int64) + arithmetic", cdiv_then_add_64, Int64(100), Int64(7))Gives:
ceil_div(Int64, Int64) only, unused FAIL: error[INTERNAL]: The compiler could not build valid IR for this code.
ceil_div(Int64, Int32) only, unused OK
ceil_div(Int64, Int64) + arithmetic FAIL: error[INTERNAL]: The compiler could not build valid IR for this code.
Expected behavior
Either support int64 or give a better error message (I'm not sure if it's intended for cdiv to not support int64)
Environment details (please complete the following information):
nvidia-cutlass-dsl==4.8.0
Additional context
Add any other context about the problem here.
Activity
yunweili3 commentedon Oct 13, 2026
Reproduced on 4.8.0 (and 4.5.3, so not a regression). Root cause is in how
_pack_tilebuilds the tile type:_cute_ir.pack_tileonly reads a leaf's bit-width from an int_tuple-typed value (?{i64}), and a plainInt64argument reaches it as a barei64value, which it types as a 32-bit leaf. The type becomes[?:1]while thecute.make_tileoperand isi64, hence the verifier mismatch. The input side ofceil_divgoes throughpack_int_tuple, which handlesi64— that is whyceil_div(Int64, Int32)works.The same failure hits every function that packs a tiler with a dynamic
Int64leaf (zipped_divide,tiled_divide,logical_divide,ceil_div((a, 64), (b, 4))).Fix (Python side, wrap 64-bit leaves as a depth-0 int_tuple before
pack_tile, same asIntValue.get_typed_value): PR #3736. With it the three cases compile andceil_div(2**40, 2**33 + 1)gives 128 on a B200.One related thing I noticed while bisecting: reading an
Int64layout shape back truncates —cute.make_layout((a, 2)).shape[0]witha = 2**33 + 5is anInt32and evaluates to 5 on device, whilecute.size(L)is correct. That comes fromunpack_x_tupleand is independent of this issue.