[QST] [CuTeDSL] operand #0 does not dominate this use

[HanGuo97]

What is your question?

I'm running into this on some of the kernel code. It's not clear to me what's the cause of this, and hence not sure how to construct a minimal script, but this happens after doing CTX-wide reduction.

MLIRError: Verification failed:
error: unknown: operand #0 does not dominate this use
 note: unknown: see current operation: %622 = "arith.addf"(%737, %621) <{fastmath = #arith.fastmath<none>}> : (f32, f32) -> f32
 note: unknown: operand defined here (op in a child region)

[fengxie]

It could be a bug. Usually some places that you can check.

• a variable defined inside if-else, for-loop but referenced outside
• a variable defined in one "cute.jit" function locally get referenced in another

Or a bug inside compiler. This can be verified by CUTE_DSL_KEEP_IR=1 to dump IR. If dump IR is in verified form, then you have legal IR but compiler generate something illegal in the middle.

some example:

// Verified form
%2 = cute.slice(%0, %1) : !cute.coord<...>, !cute.memref<...>

// Generic form
%2 = "cute.slice"(%0, %1): (!cute.coord<...>, !cute.memref<...>) -> ...

The main difference is cute.slice vs. "cute.slice" with double quote. And generic form is more verbose that verified form.

[anakinxc]

Hi @HanGuo97

There is a pattern yields this kind of error that I aware of

@cute.jit
def f2(n):
    xs = [cutlass.Int32(1), cutlass.Int32(2), cutlass.Int32(3)]
    if xs[0] == 1:
        xs[0] = n + 1
        # xs = [n + 1, cutlass.Int32(2), cutlass.Int32(3)]

    cute.printf("xs: {}", xs[0])

f2(10)

So basically, assign a list/tuple element through indexing to a local variable in a dynamic control flow can let that variable leaks out of the control flow and eventually generates invalid IR. This will be fixed in next release.

As a temporary workaround, use the commented out line.

[fengxie]

Thanks for your effort to create this reproduce case. @anakinxc is this the issue on your radar?

[HanGuo97]

Can confirm what @anakinxc suggested is indeed the reason. Thanks --- it'd be nearly impossible for me to figure that out!

I tried a few variants without hard-coding list length, and here's what I found:

@cute.jit
def f(n):
    xs = [cutlass.Int32(1), cutlass.Int32(2), cutlass.Int32(3)]
    if xs[0] == 1:

        # 1. MLIRError: Verification failed
        # for i in cutlass.range_constexpr(len(xs)):
        #     xs[i] = n + 1

        # 2. DSLRuntimeError: range_constexpr should be preprocessed by preprocessor
        # xs = [
        #     n + 1
        #     for i in cutlass.range_constexpr(len(xs))
        # ]

        # 3. works
        new_xs = []
        for i in cutlass.range_constexpr(len(xs)):
            new_xs.append(n + 1)
        xs = new_xs

    cute.printf("xs: {}", xs[0])

f(10)

[HanGuo97]

Well, the above one causes other problems down the line. In some cases, xs was a list inside the if branch, but it becomes a cute.Numeric when exiting the branch.

This also seems to work.

import cutlass
import cutlass.cute as cute


@cute.jit
def f2(xs, n):
    for i in cutlass.range_constexpr(len(xs)):
        xs[i] = n + 1
    return xs
    
@cute.jit
def f(n):
    xs = [cutlass.Int32(1), cutlass.Int32(2), cutlass.Int32(3)]
    if xs[0] == 1:
        # note: we have to reassign xs even f2 is inplace
        xs = f2(xs, n)
    print(type(xs))
    cute.printf(“xs: {}“, xs[0])

f(10)

[HanGuo97]

erm, that didn't work either in my actual use case, somehow the type of xs was a list before entering the branching, but inside it, it becomes a numeric.

EDIT: ended up refactoring the code to not do structure manipulation inside branching

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.

[fengxie]

@anakinxc is this still an issue in latest codebase?

[anakinxc]

@fengxie no, fix is already in.

[github-actions]

This issue has been labeled inactive-30d due to no recent activity in the past 30 days. Please close this issue if no further response or action is needed. Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed. This issue will be labeled inactive-90d if there is no activity in the next 60 days.