Page MenuHomePhabricator

[MLIR][std] Introduce bitcast operation
ClosedPublic

Authored by GMNGeoffrey on Jul 2 2021, 2:58 PM.

Details

Summary

This patch introduces a bitcast operation to the standard dialect.
RFC: https://llvm.discourse.group/t/rfc-introduce-a-bitcast-op/3774

Diff Detail

Event Timeline

GMNGeoffrey created this revision.Jul 2 2021, 2:58 PM
GMNGeoffrey requested review of this revision.Jul 2 2021, 2:58 PM

Add folds and canonicalization

inb4 River asks me to remove trivial braces

LG in general, some comments inline

mlir/include/mlir/Dialect/StandardOps/IR/Ops.td
463

The argument is "AnyType" but the description makes me think that if should be IntOrFloat or something like that.

Similarly for the result type, there is no verifier here to check on the "equal bit width"?

mlir/lib/Dialect/StandardOps/IR/Ops.cpp
503

There is another fold here: the identity bitcast can just return its input as-is.

Edit: actually I see a test for it, but I don't know how it works? Is there a base-class / interface for cast operation that takes care of it?

538

Can't we change the operand in-place instead of creating a new operation here?

mehdi_amini added inline comments.Jul 2 2021, 5:13 PM
mlir/include/mlir/Dialect/StandardOps/IR/Ops.td
466

Actually you support vector as well according to the tests.

467

I'm not sure about the noop part: a machine can have different registers for floating point and integer and would require a copy across register files.

silvas added a subscriber: silvas.Jul 7 2021, 11:09 AM

Please add lowering to LLVM for this too (doesn't have to be this patch precisely, but a fast follow-on please).

mlir/include/mlir/Dialect/StandardOps/IR/Ops.td
467

Yeah, LLVM uses the terminology that it's like storing the value to memory with one type and reading it back in another type. Not sure how we can phrase it.

mlir/lib/Dialect/StandardOps/IR/Ops.cpp
491

you can assert/ignore this. Otherwise the op is verifier-invalid.

499

Is this folder correct? What if target endianness is different -- how will it match the target behavior?

Generally speaking, I would suggest that we follow how LLVM does this, and leave a comment saying that we are attempting to follow the LLVM folding behavior when we have similar information available (LLVM always has datalayout).

517

are all of these code paths covered in the tests?

For example, I don't think I see one that would trigger nontrivial float->float conversion. f16->bf16 might be a case to exercise that.

GMNGeoffrey added inline comments.Jul 13 2021, 11:28 AM
mlir/include/mlir/Dialect/StandardOps/IR/Ops.td
463

I followed the style of the other cast operations here where this is all done in the cast compatible interface method. I'm not sure that that's actually better, but it seems to be what's done here. I could pre-factor or follow up to express more of this in the op definition for all of these ops

466

Yeah again I was following the rest of the ops here which just talk about the scalar aspect. Maybe they should all be updated to talk about vectors as well.

mlir/lib/Dialect/StandardOps/IR/Ops.cpp
503

I believe it's encapsulated in the interface.

517

I've got a cast that uses bf16, but it's testing the cast-of-cast path. I don't think we actually have a way to trigger same-width different integer types, do we?

mehdi_amini added inline comments.Jul 13 2021, 11:53 AM
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
517

signed integer to unsigned integer maybe?

538

(I don't know if you missed this comment?)

GMNGeoffrey added inline comments.Jul 13 2021, 12:40 PM
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
517

I think standard ops aren't supposed to operated on the [un]signed integer variants?

538

No just haven't gotten to it yet and didn't have any questions about it :-)

GMNGeoffrey retitled this revision from [MLIR] Introduce bitcast operation to [MLIR][std] Introduce bitcast operation.Jul 13 2021, 1:32 PM
GMNGeoffrey added inline comments.
mlir/include/mlir/Dialect/StandardOps/IR/Ops.td
463
466

Add more tests and clean up canonicalizations

GMNGeoffrey added inline comments.
mlir/include/mlir/Dialect/StandardOps/IR/Ops.td
467

LLVM says

It is always a no-op cast because no bits change with this conversion.

https://llvm.org/docs/LangRef.html#bitcast-to-instruction

Maybe that has some specific definition in LLVM?

mlir/lib/Dialect/StandardOps/IR/Ops.cpp
491

Done

499

The llvm behavior is only different for non-elementwise bitcasts of vectors (which we're not doing here):

There is a caveat for bitcasts involving vector types in relation to endianess. For example bitcast <2 x i8> <value> to i16 puts element zero of the vector in the least significant bits of the i16 for little-endian while element zero ends up in the most significant bits for big-endian.

https://llvm.org/docs/LangRef.html#bitcast-to-instruction

It doesn't reference endianness otherwise. I think *usually* endianness shouldn't matter because the endianness of the source and target types would be the same, but I guess that wouldn't apply on weird architectures where floating-point is big-endian and integer is little-endian.

I think we've got two potential definitions here:

  1. The bitcast operation does a target-independent bitcast of the arbitrary-width value. I believe that's what these folders implement, assuming I understood how bitcastToAPInt works. It does mean that the statement that this is always a no-op is not true because a cast from f64 to i64 on a 32-bit system with big-endian floating point and little-endian integer would need to swap the words.
  2. The return value is target-dependent. The runtime no-op is still not true because of the float/int registers as Mehdi points out.

Given that we don't have target specifications at this level, I think I like (1) better. Then we have a clearly-defined target-independent operation for the compiler and an operation that is free on non-weird architectures, but needs special handling at the target-specific level for weird architectures.

It looks like the LLVM constant folding only handles casts to vectors: https://github.com/llvm/llvm-project/blob/18c19414eb70578d4c487d6f4b0f438aead71d6a/llvm/lib/Analysis/ConstantFolding.cpp#L149, so we can't really use that as a reference.

The other question would be how to lower to/from such an operation. For instance, I'm interested in lowering mhlo.bitcast, which really doesn't explain its view on endianness in much detail

so machines with different floating-point representations will give different results

https://www.tensorflow.org/xla/operation_semantics#bitcastconverttype

Implementation defined sounds like a free pass there :-P

Lowering to LLVM, would this op be able to be converted directly to an llvm bitcast? I think not, and that LLVM has the second semantics above. When lowering to LLVM, do we have the target information yet? Because then it would be easy to add a special case for the mixed-endian architectures.

538

Done.

GMNGeoffrey marked an inline comment as not done.Jul 13 2021, 4:06 PM
GMNGeoffrey added inline comments.
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
499

Or maybe not and we go with definition 2 to match LLVM, in which case I think constant folding is in-general impossible

GMNGeoffrey marked 2 inline comments as not done.Jul 20 2021, 11:29 AM
GMNGeoffrey added inline comments.
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
499

Ping on this?

mehdi_amini added inline comments.Jul 20 2021, 12:06 PM
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
499

Is "mixed endian" a thing?

GMNGeoffrey marked an inline comment as not done.Jul 20 2021, 12:09 PM
GMNGeoffrey added inline comments.
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
499

If by that you mean having different endianness for float and int, apparently yes. There are also other endianness formats that just kind of randomly strew the bytes about, apparently

GMNGeoffrey marked an inline comment as not done.Jul 20 2021, 12:12 PM

Ping on this? I can bring the question of exact bitcast semantics back to the RFC if you'd like

silvas added inline comments.Aug 2 2021, 2:49 PM
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
499

I was referring to "the return value is target dependent" case, which would mean that we could not fold without further information abotu the target. It seems that LLVM has a strong precedent for ignoring that issue, which is equivalent to defining the fp format in terms of the bit positions in the integer format. The critical code for the scalar case in LLVM is here: https://github.com/llvm/llvm-project/blob/18c19414eb70578d4c487d6f4b0f438aead71d6a/llvm/lib/IR/ConstantFold.cpp#L168 I would link to that code and explain that everything we are doing here is consistent with that. The DenseIntElementsAttr case seems like it requires further explanation though along with the DenseFPElementsAttr case.

  • Clarify op and folding semantics wrt endianness
GMNGeoffrey added inline comments.Aug 5 2021, 2:32 PM
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
499

I added a bunch more to operation semantics and some comments to the folding methods

silvas accepted this revision.Aug 5 2021, 2:35 PM
This revision is now accepted and ready to land.Aug 5 2021, 2:35 PM
rriddle added inline comments.Aug 5 2021, 2:41 PM
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
537

You should be able to move this to the folder, updating the root operation is fine you just need to return the existing result of the op for the OpFoldResult.

  • Move bitcast-of-bitcast simplification to folder
GMNGeoffrey marked 14 inline comments as done.Aug 5 2021, 3:49 PM
GMNGeoffrey added inline comments.
mlir/lib/Dialect/StandardOps/IR/Ops.cpp
537

Oh nice! Done.

This revision was automatically updated to reflect the committed changes.
GMNGeoffrey marked an inline comment as done.