In D151423#4387264, @foad wrote:

@qcolombet I was planning to go back to the previous diff which used auto for the return types: https://reviews.llvm.org/D151423?vs=on&id=525565. Do you have a preference?

Could you re-upload the patch with the full context please?

I like the improved readability.
Thanks!

Ping, @nikic maybe?

LGTM

Right, it does complain about the LiveIntervals, but I thought that indicated that the LiveIntervals were not computed correctly, not that the underlying MIR was malformed.

I need some help understanding the background here. Why does LiveIntervals::computeVirtRegs only call splitSeparateComponents if it detected dead defs in a live interval? Is this some artifact of the way MIR is typically constructed by instruction selection?

In D147457#4365865, @pawosm01 wrote:

In D147457#4348901, @qcolombet wrote:

I was hoping this patch could address my problem reported here: https://github.com/llvm/llvm-project/issues/62736
Sadly, it isn't. Maybe it's worth to extend it?

Let me take a look. Thanks for the pointer.

Hello @qcolombet

I'd like to kindly ask you whether you had a chance to look at it and to have any conclusions.

Given what you observed and this limit is pretty much never hit in practice, I'm fine with bumping it.

I was hoping this patch could address my problem reported here: https://github.com/llvm/llvm-project/issues/62736
Sadly, it isn't. Maybe it's worth to extend it?

On a second thought the ABI concerns are probably fine because we can get the information directly from the LLVM IR.

Ping^5

I'm still wondering if somehow adding more semantic operands to FP operations would be less intrusive.

I can help get the patch reviewed if it is still relevant.
Just one high-level remark: the changes in the peephole optimizer should be a separate patch, right?

Adding more reviews.

Sure!
I'll probably do that beginning of next week.

Thanks for the fix

Given the patch does not modify generic parts anymore, I'm happy with it.
I leave the final decision to @arsenm since AMDGPU is outside my domain of expertise.

Ping^4

Ping^3

LGTM, please add a test for the lowering to "simple" pad.

LGTM

Ping^2

Do we want to support MachineFunction/Info cloning?

This has to work and needs to be serializable.

• Use all_of and llvm::seq where appropriate
• Use getMixedDimensions to future proof the code w.r.t. dynamic dimensions
• Early-exit to avoid else indentation

Forgot to submit these comments!

LGTM

Ping

Nice catch!

Sorry for the delay I don't have an answer, but I have a couple of questions:

1. Would it work to attach the delegate in MachineFunction::initTargetMachineFunctionInfo?
2. Do we want to support MachineFunction/Info cloning?

I don't think this is correct.
Here we are merging the main live-ranges, so if we set the join subrange flag we may not check the right thing. In particular, the LaneBitmask in this call is incorrect, unless I am mistaken, for a subrange check.

LGTM but I think we're missing a test for populateResolveExtractStridedMetadataPatterns.
@nicolasvasilache should we use transform dialect for the test?

> I think the break statement cannot be removed, as it would crash if the last instruction is a debug instruction. Please correct me if I am missing something here.

• Fix bazel build

I only had a cursory look but the overall direction seems good.

• Remove the test pass
• Add a tranform op instead
• Add a comment for strides when creating the subview
• Add a test for non-1 strides (only for vector.transfer_write, @nicolasvasilache let me know if you want more)
• Report failures with notifyMatchFailure

• Add missing vector dialect dependency to the test-extract-address-computations pass

• Add support for vector.transfer_read/write
• Move the populate function in its own header (we'll want to move others as well in an NFC patch)
• Fix the dropping of the nontemporal flag (and add a test for that)
• Refactor getSrcMemRef to not choke on tensor type
• Refactor getViewSizeForEachDim to use the same implementation for ldmatrix, vector.transfer_read/write
• Add tests for vector.transfer_read/write

Abandoning that one for now. I think it would make sense to introduce some interface for this kind of rewriting.

Nice clean-up!

Food for thoughts:
Since the canonicalization is going to undo the decomposition, should we rewrite the expression (e.g., as part of the canonicalization) to the ordering we want?
The downside is we would still rely on the backend to do the hoisting of the loop invariant stuff.