ITADN

[BUG] [CuTeDSL] `storage.<field>.get_tensor(...)` inside a dynamic `if` block fails with "encountered a user-defined Python object"

#3268Opencheshire 创建于 2026-05-24
bug? - Needs TriageCuTe DSL
C
cheshirecommented
### Which component has the problem? CuTe DSL ### Bug Report **Describe the bug** Calling `storage.<field>.get_tensor(...)` from inside a dynamic `if` block fails at DSL trace time with: ``` DSLRuntimeError: The 'if' statement encountered a user-defined Python object, which cannot be automatically converted into an dynamic expression. ``` This blocks the natural warp-specialization pattern, where each `if warp_idx == <role>:` branch extracts the SMEM tile it uses. **Steps/Code to reproduce bug** ```python import cutlass import cutlass.cute as cute import cutlass.utils as utils from cutlass import Int32, Float32 from cutlass.cute.runtime import make_fake_tensor @cute.struct class Smem: sA: cute.struct.MemRange[Float32, 128] @cute.kernel def k(gA: cute.Tensor): tidx, _, _ = cute.arch.thread_idx() smem_alloc = utils.SmemAllocator() storage = smem_alloc.allocate(Smem) sA_layout = cute.make_layout((128,), stride=(1,)) if tidx == Int32(0): sA = storage.sA.get_tensor(sA_layout) # fails here sA[0] = gA[0] @cute.jit def entry(gA: cute.Tensor): k(gA).launch(grid=(1, 1, 1), block=(32, 1, 1), smem=512) gA = make_fake_tensor(Float32, (128,), stride=(1,), assumed_align=4) cute.compile(entry, gA) ``` **Expected behavior** `get_tensor` should be callable from inside a dynamic `if`. The SMEM itself was already allocated by `smem_alloc.allocate(Smem)`; the `get_tensor` call only constructs a typed view over it. **Environment details (please complete the following information):** - Environment location: Docker (`nvidia/cuda:12.8.1` base image; no GPU required for the trace-time failure) - `nvidia-cutlass-dsl` versions tested: **4.4.2** and **4.5.1** (both fail on 4.5.1 with the Python-level message above; 4.4.2 produced a more cryptic MLIR-level message) **Additional context** Workaround: hoist every `storage.<field>.get_tensor(...)` call to outer kernel scope, before the warp-dispatch chain. Each role then references the resulting view variable from inside its branch. If the constraint stays, the diagnostic could name the fix ("extract storage views at outer kernel scope before the dynamic branch") instead of just saying "user-defined Python object." Related: #3243: `cute.FastDivmodDivisor.divisor` can reference an SSA value from outside isolated region.
0 条评论