[RFC] SignednessCastOp

MLIR
#1

We have been discussing lowering the mhlo dialect, which uses signed and unsigned integer types to model signedness, to linalg and the standard dialect, which uses signless integer types and models signdness semantics via different operations.

As a first step in this direction, I propose to add a cast operation that allows to cast unsigned and signed integer types to the signless integer type.

%signless = signedness_cast(%signed) : si8 to i8
%signed = signedness_cast(%signless) : i8 to si8
%unsigned = signedness_cast(%signless) : i8 to ui8

It should also be allowed to cast between tensors of corresponding types.

%signless = signedness_cast(%signed) : tensor<32xsi8> to tensor<32xi8>
%signless = signedness_cast(%signed) : memref<32xsi8> to memref<32xi8>

A cast is hence valid if it does not modify the shape and casts from a signed/unsigned type to its signless counterpart. I would not allow direct casts between signed and unsigned but I am also not strongly opposed to allowing this.

I would place this cast in the standard dialect, as it also works on scalars. If there is a strong desire, I can also add three separate ops for scalars, tensors and memref.

With this cast operation added, we can gradually lower mhlo based IR to linalg by inserting signedness_cast operations to bridge between dialects. As a pseudo-code example, we could get

%res = mhlo.add %0, %1 : tensor<42xui32>

could be lowered to something like

%sl0 = signedness_cast %0 : tensor<42xui32> to tensor<42xi32>
%sl1 = signedness_cast %1 : tensor<42xui32> to tensor<42xi32>
%0 = linalg.init_tensor [42] : tensor<42xi32>
%res = linalg.generic {
      indexing_maps = [#map, #map, #map], 
      iterator_types = ["parallel"] }
      ins(%si0, %si1 : tensor<42xi32>, tensor<42xi32>)
      outs(%0 : tensor<42xi32>) {
    ^bb0(%arg2: i32, %arg3: i32, %arg4: i32):  // no predecessors
      %2 = addi %arg2, %arg3 : i32
      linalg.yield %2 : i32
    } -> tensor<42xi32>
%res = signedness_cast %res : tensor<42xi32> to tensor<42xui32>
#2

Is there any reason for MHLO to continue to use signed types? If not, can we instead draw a plan to move MHLO away from signed integer to align with the rest of the ecosystem?

#3

The “rest of the ecosystem” depends a bit on the perspective. From a low-level LLVM standpoint, everything is signless (i.e. “op carries the signedness”). From a higher level standpoint, I am not aware of any ML source system which models it that way: something will need to realize this switch. As a gateway dialect, I personally feel that MHLO is more useful as a representation closer to the problem domain vs the low level compiler stack.

Put more concretely, even if MHLO flips to the signless view of the world, I suspect that it is just going to push the mismatch up a level and we likely will find that we still want the casts that herhut proposes in some fashion. As another point in the space, consider what TOSA will need. It would also be useful to look at how mlir-onnx does this, since I believe they also have to navigate the sign switch (and their reference stack is more direct).

(side note: I wish that MHLO had a more open design process. I work with it a lot and am not even sure how to propose a change like this without using back-channels)

1 Like
#4

Isn’t it a very different “switch” to go from signed->signless representation compared to going back and forth in the signed domain between signed and unsigned?

Up one level you have something like TensorFlow, which already manage it own solution for all of this. When we lower from TF → MHLO you’d need to erase the “signed-ness” aspect into signless types, but that never involve this conversion operation as far as I know.

At the moment, emailing mlir@tensorflow.org is best to discuss such design question on MHLO believe.

#5

You need to erase them into signless types and ops, and when I look at MHLO, it is pretty clear from its op coverage that it was designed in the “types carry the signedness” tradition. That can be switched, but a proposal for it would involve a re-modeling of a bunch of operations and inbound conversions much more than it would change anything about the types (i.e. need sign/signless variants of div/compare/etc, convert semantics no longer work, etc). For MHLO, the owners could conceivably make this change, but other peer dialects like TOSA are spec and CTS governed in line with their source domain. I’m not aware of any such examples at that level which designed their ops with the expressiveness to erase the sign bit from the type.

@sjarus for TOSA

I ultimately don’t have a super strong opinion about how/whether MHLO decides to align one way or another, but I don’t think that it choosing to do so eliminates the need to model this case. I haven’t fully thought through the implications of what herhut proposes, but I do like standardizing/having one way that we navigate this representation switch (presuming we can’t just factor away the need entirely – which I am taking as a given).

1 Like
#6

I’m not saying that the need does not exist, but I question how suitable the solution is :slight_smile:
In particular I believe that these types are intended as “frontend types”, and were introduced just to expose the types to the frontends but with the expectation that frontends (or other dialects) would have to model the ops they need to manipulate these types, while std (or the collection resulting from std) would stay signless. So I’m mainly concerned about starting to add support for signed integer in our core dialects in general.

#7

Yeah, agreed - that stood out to me as possibly crossing a line: “standard is defined for signless”. In previous discussions, it seems that we sometimes make accommodations for casting ops and other fidgety things needed to use the type system while lowering between levels and preserving legality or IR. Since the signed types are defined as part of the standard types, I can see the argument that ops facilitating legalization of those types across boundaries may also belong there, presuming that there is any commonality (versus every dialect defining its own casting fidgets). I’m not sure I buy that argument yet, though :slight_smile:

I wish we had better names for these classes of dialects. The “core dialects” can mean so many things, and I take your meaning to be part of the code generation stack, separate from frontend or interface dialects. From my perspective, MHLO is a frontend dialect – but that is open to debate I am sure.

#8

Thanks for the call out @stellaraccident!

This is a good perspective on what TOSA deals with top down from an ML source framework (e.g. TensorFlow or ONNX). Signless ints require something else to express sign information, and the TOSA spec defines sign information to be datatype carried, not op carried.

However, TOSA defined integer types are all signed int types, except for uint8 which is an exception because (older ?) ML networks were ui8 quantized. But as the convert_tfl_uint8 pass of our legalization flow does, the spec defines that tosa.rescale be used to convert such networks to signed int prior to further processing; all other ops implicitly deal with signed types or perform actions independent of type.

In effect, except for disambiguating fp vs int, all ops implicitly treat the int as a signed int, e.g. tosa.add is a signed add. The TOSA spec explicitly addresses the question of networks that are unsigned int quantized by defining tosa.rescale to express them as signed ints.

#9

Thanks for the comments.

In my view, the mhlo dialect is a front-end dialect and hence it is reasonable to model signdness via types rather than operations, as the dialect is still fairly close to the programming languages it models and using types for operation dispatch is the common way to handle this in languages. So I would not want to change mhlo.

I wanted to place these ops in standard, as the signed types are builtin types and we do not have a specific dialect for operations on builtin types.

I am also fine with putting these into the mhlo dialect as that is the dialect that requires them for conversion to standard. That will prevent reuse by other dialects, hence the rfc. We can always move them if the need arises.

#10

+1 on adding this. It’s needed somewhere in the stack (e.g. for materilizations when converting types), and mhlo->linalg can be a guiding example.

If you add ElementwiseMappable to it (which all the other primitive scalar ops have) then it should work for vectors/scalars/tensors nicely.

If you want, you could also lower it to

%sl0 = signedness_cast %0 : tensor<42xui32> to tensor<42xi32>
%sl1 = signedness_cast %1 : tensor<42xui32> to tensor<42xi32>
%res = std.addi %sl0, %sl1 : tensor<42xi32>
%res = signedness_cast %res : tensor<42xi32> to tensor<42xui32>

And then use convert-elementwise-to-linalg. (the elementwise stuff will all fuse away later anyway)

#11

Thanks Stephan. This also happens to be something that just started blocking me, so glad to see others taking a look at this as well :slight_smile:

When I was prototyping how to handle the lowering from MHLO → linalg, I had intended to just use a type converter as part of dialect conversion. I think that would mean that this cast op wasn’t strictly necessary, assuming the conversion happens all at once. This ties back to earlier discussions about an “anycast” op and whether it’s better to realize type casting as ops in the IR or keep it within dialect conversion. Personally, I prefer the former, and so +1 from me. I think this op is less controversial than anycast as well, since it has clearly defined semantics of adding or removing some metadata.

I agree that I’m not sure where this best belongs, but I think given that it’s operating on builtin types it might make sense in standard. Perhaps, the idea that standard is only for signless types may be less of an issue as the ops that operate on integers move into their own dialects? It is what’s currently documented in the rationale document (which should probably be updated anyway). I don’t feel strongly that this has to go to std, but I think it would be a shame to put a common thing like this only within MHLO since [un]signed is a builtin concept and the way it’s handled seems likely to be common.

I like this approach. It means we can have MHLO → std (or the dialects it becomes) and std → linalg for most cases, with the latter in core and reused across other lowerings. I guess there may still be some cases where you need to go directly from MHLO → linalg though (I haven’t done a thorough survey)

#12

In general I’m pretty -1 on adding this operation.

Given that this seems to be solely used for lowering from a system that has signed-ness types to a system with signless types, what is preventing you from just using the builtin UnrealizedConversionCastOp operation? That operation seems to cover this exact use case.

Being a “builtin” type does not mean that the “core” dialects necessarily have to support it, e.g. look at the rationale surrounding the builtin tuple type (which is the same for signed types, albeit not as explicitly spelled out). I wouldn’t expect any potential integer dialect to be non-signless, and I don’t see the splitting of standard dialect as relaxing any of the design restrictions that are already in place.

– River

#13

I think I may have missed that this op got added after the “anycast” discussion :slight_smile: This approach is more in line with what I was previously trying out in converting MHLO → linalg. One reason we might want to use it in this case, however, is that a cast between signless and [un]signed types has well defined runtime semantics (that is, none), and there may be cases where we want to leverage this property. It feels somehow more natural to use the lowering strategy Sean described with a “real” op than with explicit insertion of an UnrealizedConversionCastOp, but I don’t think I can articulate a good reason other than a vague feeling :grin:

Fair. My recollection with the tuple type, though I can no longer find the discussion, was that one of the objections with having core ops that operate on them was that the way they were used in different because of lack of clarity on what tuple storage semantics were. I don’t think that objection necessarily exists with [un]signed types. Here’s the discussion on adding signedness semantics to integers (https://groups.google.com/a/tensorflow.org/g/mlir/c/XmkV8HOPWpo/m/CYuiE3MrAgAJ), in which tuples are explicitly cited as prior art.

Lot’s of negatives there, but I think I agree :stuck_out_tongue: My point was that if various ops migrate out of standard then what remains in standard might not need to carry the restrictions forwards because it would just be “support” ops. I don’t think we should change standard or an integer dialect to operate on signed types.

#14

I would likely agree here, but I don’t think the need for such a thing has materialized yet. I am wary of adding these “interface” style operations when the need is just for lowering. This was part of the justifcation/rationale for the unrealized conversion cast, which is to avoid the explosion of 100 different cast operations on the edge.

One of the main aspects of the tuple type discussion is that MLIR natively supports multiple result types, so there isn’t a need or desire to use tuples in ways that other systems use them(the main use case for tuples at the time mostly coming from XLA). The discussion at the time was that “this is useful for downstream/interface dialects, but we don’t want to do that to permeate through the core dialects”. The discussion around signed/unsigned is the same IIRC, which is these are left for interface dialects and we don’t want them to permeate.

Apologies, this is just a knee-jerk reaction to the standard dialect (lots of dark and painful memories here). AFAIK the end goal is for standard dialect to completely die, i.e. not have anything inside it at all. I have more reservation to adding operations that don’t have a clear home or place in the post standard world. If we do end up adding these types of operations, they have to be in a place where it is clear that they aren’t blessed from a “core” dialect point of view. If we put these in standard or integer, it kind of establishes an expectation that MLIR is promoting using signed types (at least IMO).

– River

#15

Yeah these concerns makes sense to me. I started proposing a “cast” dialect in my earlier post, but then couldn’t really articulate why such a dialect would make sense and wouldn’t become a dumping ground of unrelated stuff.

Stephan, do you think we could just use the UnrealizedConversionCastOp here?

I think the only place where I’m unsure how this conversion will work is with whether we should be converting function signatures as part of this pass. Something might be expecting public function signatures to remain unchanged, but if we don’t change them then we’ve got leftover cast ops. For IREE’s purposes, I think we’d like the whole module and its functions to get converted, but I don’t think we want to necessarily impose that on all users and especially that may make less sense if someone has a mix of dialects. Maybe converting function signatures should be a separate pass.

To clarify, I meant “lots of negatives” in the “double negative” sense, so I had some difficulty parsing, not in the sense that your post was overly negative :grin:

#16

I think it makes sense to use UnrealizedConversionCast if possible. UnrealizedConversionCast is mainly useful if converting everything in the module (so that all the casts “cancel out”). I’m not sure if what is being proposed here involves doing that in-depth of a conversion from signed to signless types. It seems plausible to have an op with actual runtime semantics like signedness_cast for that scenario where signedness still exists at ABI boundaries and such (don’t know if that applies here though).

FWIW, I think it does make sense to have the conversion from signed to signless be a complete conversion of everything (function signatures, bbargs, values flowing through scf.for, etc.), in which case UnrealizedConversionCast would work. It’s just more work to do that – but seems fairly straightforward to implement after what we’ve learned/developed in the bufferization context, such as using our existing StructuralTypeConversions (e.g. for scf.for)).