Improved the interface of methods commuting operands, improved X86-FMA3 mem-folding&coalescing.

[v_klochkov]

Summary

``The main goal of the change-set is to improve Memory-operand folding
and Coalescing optimizations performed for X86 FMA instructions
(Described in (1) below).

Unfortunately, that could not be done without interface changes
done in methods findCommutedOpIndices() and commuteInstruction().
(Described in (3) below).
The minor changes in non-X86 target sources: PowerPC, ARM and AMDGPU
were required by the new commuteInstruction() method with additional 2 operands
started being called from llvm/lib/CodeGen/* classes common for all targets.

The size of the fix is pretty big because of having (1) and (3) in one
change-set. The alternative to this change-set could be splitting of
the change-set into 2 parts:

• interface changes (described in (3) below)
• improvement of X86 FMA form selection (described in (1) below).

(1) Implemented optimal form selection (213/312/231) for X86 FMA instructions

   to improve Memory-Folding/Ciscization and Coalescing optimizations performed
   for FMAs. The change-set allows commuting any of FMA operands: 1st and 2nd, 
   1st and 3rd, 2nd and 3rd. 
   Previously, only 1st and 2nd operands could be commuted.
   
   Better Memory-folding and Coalescing optimizations help to reduce 
   registers pressure. Improvement from the changes can be shown on such 
   an example:

           for (int i = 0; i < N; i += 1) {
            val1 = _mm_and_pd(val1, val5);
            val2 = _mm_and_pd(val2, val6);
            val3 = _mm_and_pd(val3, val7);
            val4 = _mm_and_pd(val4, val8);
            val5 = _mm_xor_pd(val1, val5);
            val6 = _mm_xor_pd(val2, val6);
            val7 = _mm_xor_pd(val3, val7);
            val8 = _mm_xor_pd(val4, val8);

            v_accu1 = _mm_fmadd_pd(v_accu1, x1_arr[i], val1);
            v_accu2 = _mm_fmadd_pd(v_accu2, x2_arr[i], val2);
            v_accu3 = _mm_fmadd_pd(v_accu3, x3_arr[i], val3);
            v_accu4 = _mm_fmadd_pd(v_accu4, x4_arr[i], val4);
            v_accu5 = _mm_fmadd_pd(v_accu5, x5_arr[i], val5);
            v_accu6 = _mm_fmadd_pd(v_accu6, x6_arr[i], val6);
            v_accu7 = _mm_fmadd_pd(v_accu7, x7_arr[i], val7);
            v_accu8 = _mm_fmadd_pd(v_accu8, x8_arr[i], val8);
        }
		


    ASM code BEFORE the changes:
        .LBB1_2:                                # %for.body.6
                                        #   Parent Loop BB1_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
        vmovapd %xmm0, -56(%rsp)        # 16-byte Spill
        vandpd  %xmm7, %xmm3, %xmm7
        vandpd  %xmm5, %xmm12, %xmm5
        vandpd  %xmm6, %xmm9, %xmm6
        vmovapd -40(%rsp), %xmm10       # 16-byte Reload
        vandpd  %xmm10, %xmm13, %xmm10
        vmovapd %xmm10, -40(%rsp)       # 16-byte Spill
        vxorpd  %xmm7, %xmm3, %xmm3
        vxorpd  %xmm5, %xmm12, %xmm12
        vxorpd  %xmm6, %xmm9, %xmm9
        vxorpd  %xmm10, %xmm13, %xmm13
        vmovapd %xmm8, %xmm0
        vmovapd x1_arr+8192(%rcx), %xmm8
        vmovapd -24(%rsp), %xmm1        # 16-byte Reload
        vfmadd213pd     %xmm7, %xmm8, %xmm1
        vmovapd %xmm1, -24(%rsp)        # 16-byte Spill
        vmovapd %xmm0, %xmm8
        vmovapd x2_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm5, %xmm1, %xmm4
        vmovapd x3_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm6, %xmm1, %xmm8
        vmovapd x4_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm10, %xmm1, %xmm11
        vmovapd -56(%rsp), %xmm0        # 16-byte Reload
        vmovapd x5_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm3, %xmm1, %xmm15
        vmovapd x6_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm12, %xmm1, %xmm0
        vmovapd x7_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm9, %xmm1, %xmm2
        vmovapd x8_arr+8192(%rcx), %xmm1
        vfmadd213pd     %xmm13, %xmm1, %xmm14
        addq    $16, %rcx
        jne     .LBB1_2

        ASM code WITH the new changes (about 30% faster):
        .LBB1_2:                                # %for.body.6
                                        #   Parent Loop BB1_1 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
        vandpd  %xmm7, %xmm3, %xmm7
        vandpd  %xmm5, %xmm2, %xmm5
        vandpd  %xmm6, %xmm0, %xmm6
        vandpd  %xmm1, %xmm4, %xmm1
        vxorpd  %xmm7, %xmm3, %xmm3
        vxorpd  %xmm5, %xmm2, %xmm2
        vxorpd  %xmm6, %xmm0, %xmm0
        vfmadd132pd     x1_arr+8192(%rcx), %xmm7, %xmm15
        vfmadd132pd     x2_arr+8192(%rcx), %xmm5, %xmm8
        vfmadd132pd     x3_arr+8192(%rcx), %xmm6, %xmm9
        vfmadd132pd     x4_arr+8192(%rcx), %xmm1, %xmm10
        vfmadd132pd     x5_arr+8192(%rcx), %xmm3, %xmm14
        vfmadd132pd     x6_arr+8192(%rcx), %xmm2, %xmm11
        vfmadd132pd     x7_arr+8192(%rcx), %xmm0, %xmm12
        vxorpd  %xmm1, %xmm4, %xmm4
        vfmadd132pd     x8_arr+8192(%rcx), %xmm4, %xmm13
        addq    $16, %rcx
        jne     .LBB1_2

(2) Fixed a correctness problem caused by commuting 1st and 2nd operands of

   scalar FMAs generated for intrinsics. The problem is AUTOMATICALLY/for-free
   gets fixed by the proposed changes for (1).
   
   For FMA intrinsic call:

       __m128d foo(__m128d a, __m128d b, __m128d c) {
	     // must return XMM0={b[127:64], a[63:0]*b[63:0]+c[63:0]}
	     return _mm_fmadd_sd(b, a, c);
	   }

    The Coalescer/TwoAddressInstructionPass swapped the 1st and 2nd operands

of SCALAR FMA and invalidated the higher bits of the result returned
from foo().
The change-set fixes that and prohibits swapping 1st and 2nd operands
of scalar FMAs.

Swapping 1st and 2nd operands of scalar FMAs is possible and legal,
but only after special analysis of FMA users. Such optimization/analysis
can be implemented separately.

(3) The changes performed for (1) and (2) could not be implemented without

   interface change in 2 methods of TargetInstrInfo class and it's child classes:
       
	   bool TargetInstrInfo::findCommutedOpIndices(MachineInstr *MI,
                                                   unsigned &SrcOpIdx1,
                                                   unsigned &SrcOpIdx2) const;

       The operands SrcOpIdx1 and SrcOpIdx2 used only for OUTPUT from
	   the method previously.
	   Now they are INPUT and OUTPUT arguments.
	   INPUT values specify the indices of operands that are wanted to be swapped.
	   The input value ~0U gives the findCommutedOpIndices freedom to pick 
	   any commutable operand (i.e. defines the _old_ behaviour of the method).
	   
	   MachineInstr *TargetInstrInfo::commuteInstruction(MachineInstr *MI,
                                                         bool NewMI,
                                                         unsigned Idx1,
                                                         unsigned Idx2) const;

       Two arguments Idx1 and Idx2 were added to the method; they specify 
	   the operands to be swapped/commuted.
	    
    The old commuteInstruction() method did not let you to ask to commute

1st and 3rd operands or 2nd and 3rd operands.

The changes in TwoAddressInstructionPass.cpp show how the updated methods
can be used to fix missing optimization opportunities that could happen

        previously.

Readability and risky assumptions.
Previously, something similar to this sequence was used in several places:

	    unsigned Idx1;
		unsigned Idx2;
	    if (findCommutedOpIndices(MI, Idx1, Idx2) &&
		    (Idx1 == 1 || Idx2 == 1)) {
	      commuteInstrction(MI, false); //!!! how can we know that Idx1
		                                //    and Idx2 are commuted here?
	      <do something with Idx1 and Idx2 operands here>
		}

The updated functions allow to write more clear, safe and readable code:

        unsigned Idx1 = 1 /*want to commute 1st operand*/;
	    unsigned Idx2 = ~0U /*Don't care, choose any other operand*/;
	        if (findCommutedOpIndices(MI, Idx1, Idx2)) {
	          commuteInstrction(MI, false, Idx1, Idx2);
	          <do something with Idx1 and Idx2 operands here>
			}

The old method commuteInstruction() not specifying the commuted operands
was not removed as removing it would require restructuring/improvements
in some other places (similar to what was done in TwoAddressInstructionPass),
which cannot be done in this change-set as the current version of
the change-set is already too big.
``

[delena]

Hi Slava,

I added some comments, but I'm definitely not a reviewer. I suggest to send a mail to LLVM-dev and tell that you want to improve FMA code and ask people who can review the patch.
Do you have "OK" for the up-streaming?

llvm/lib/Target/X86/X86InstrInfo.cpp

2931 ↗

(On Diff #30193)

I suggest to change from "Do not call" to "If you call"

2940 ↗

(On Diff #30193)

Please remove one empty line.

3169 ↗

(On Diff #30193)

Looks huge. I'm not sure but may be these enums in ABC order and we can compare against first and last?
Or auto-generate something?

3378 ↗

(On Diff #30193)

This method may be static. Right?

3387 ↗

(On Diff #30193)

I suggest to separate scalar from vector. You handle them separately, right?
I also think that you can put all FMA tables in a separate header file.

3546 ↗

(On Diff #30193)

I think this interface is inconvenient. I suggest to separate input and output. You can put ~0U as default value of input.

[MatzeB]

Some comments on the proposed interfaces:

llvm/include/llvm/Target/TargetInstrInfo.h

97–109 ↗

(On Diff #30193)

This is just a helper used to implement some of the other commuting methods. It should not be public and virtual, it doesn't even access any state and could be static or a helper function elsewhere.

274–277 ↗

(On Diff #30193)

I think just adding Idx1,Idx2 to the existing method should be enough. The two operand commutable case will also work with that and it less confusing than having two overloaded variants around.

289–312 ↗

(On Diff #30193)

This interface looks complicated, two functions for querying commutability with different "query" styles. I think this should be simplified to a function that simply returns all operands that are commutable. I would imagine something like:

bool findCommutedOpIndicess(MachineInstr *MI, SmallVectorImpl<unsigned> &CommutableOperandNums) const;

or

ArrayRef<unsigned> findCommutedOpIndices(MachineInstr *MI) const;

[qcolombet]

Hi Vyacheslav,

I second Matthias' comments.

Just a few more things, see the inline comments.

Thanks,
-Quentin

llvm/include/llvm/Target/TargetInstrInfo.h

100 ↗

(On Diff #30193)

Please define a static constant for this magic value.
I do not like magic values wandering around without context.
I do not have a good naming right now though, maybe UndefinedIndex?

289–312 ↗

(On Diff #30193)

I think for the general case instead of having a list of unsigned, we would need a list of pair of unsigned.

[v_klochkov]

`Elena, Matthias, Quentin,

Thank you for the code-review and comments.
The interface of the functions findCommutedOpIndices() and areOpsCommutable()
caused the biggest questions.

Combining Elena's idea with some of Quentin's comments regarding
the magic consts can give this first alternative:

const unsigned CommuteAnyOperandIndex = ~0U;
bool findCommutedOpIndices(MachineInstr *MI, 
                           unsigned &OpIdx1,
                           unsigned &OpIdx2,
                           unsigned MustCommuteIdx1 = CommuteAnyOperandIndex,
                           unsigned MustCommuteIdx2 = CommuteAnyOperandIndex);

The second alternative is to have a method returning all pairs
of commutable operands:

bool findCommutedOpIndices(MachineInstr *MI, 
                           SmallVectorImpl<_some_type_containing_a_pair_of_unsigned_indices_>);

(I am not sure if I should define some new class/struct or reuse some existing type for <pair of indices>,
but that ok, let this question wait until we can decide what alternative seems better).

I really like Elena's idea; Please let me explain why.
The reason (a) below is the most important one.

a) In many cases we only need to know if some known operand is commutable with others.
   For example, RegisterCoalescer.cpp wants to know if 'UseOpIdx' can be swapped with something else;
   TwoAddressInstructionPass.cpp wants to know if 'BaseOpIdx' is commutable with something else;
   In some other cases we may even know the operands that need to be commuted and do not want
   to know about other operands commutativity.
   
   So, the first alternative provides a flexible instrument that helps to know just what we need to know,
   while the second alternative makes findCommutaleOpIndices() to collect information that often is
   not needed later.
   
   For example, I want to know if 2nd and 3rd operands of FMA are commutable.
   The 1st alternative just gives me the answer:
     if (findCommutedOpIndices(MI, Idx1, Idx2, 2, 3)) {}
   The 2nd alternative would probably return 3 pairs: <1,2>, <1,3>, <2,3>.
   I would not only need to find the desired <2,3> in the set of 3 pairs, but I would
   also ask findCommutedOpIndices() to do potentially expensive analysis regarding commutativity of
   the 1st operand (if FMA is scalar, then 1st operand is commutable only if users use only the lowest 
   element of XMM).
   
b) The 1st alternative helps to make the change-set a bit more compact than it is now,
   while the 2nd would require additional changes that would additionally complicate the places 
   where findCOmmutedOpIndices() is called now.

c) It would be very simple to remove the method areOpsCommutable() if it seems redundant.
   The calls of that method could be easily replaced with the calls of findCommutedOpIndices().
   For example:
     if (areOpsCommutable(MI, 1, 2)) {}
   ->
     if (findCommutedOpIndices(MI, Idx1, Idx2, 1, 2)) {}

Please let me know if you agree with the reasoning and if you are OK with Elena's idea.

Thanks,
Slava`

llvm/lib/Target/X86/X86InstrInfo.cpp

2931 ↗

(On Diff #30193)

This "Do not call" comment was moved to here from the old version of include/llvm/Target/TargetInstrInfo.h
The TargetInstrInfo::commuteInstruction() has assert verifying that MI is commutable.
After taking that assert into account this comment seems quite precise.

2940 ↗

(On Diff #30193)

Ok, removed it, the updated version of the change-set will have this fix.

3169 ↗

(On Diff #30193)

Unfortunately it is huge, I agree.
Comparing against the first and last or having some assumptions about how and in which order the opcodes were defined
seems a very risky approach causing unexpected effects/errors in future. I am pretty sure that we should not go this way.

I considered the idea of having a special bit for FMAs (something similar to the fields defined in llvm/include/llvm/Target/Target.td:
isReturn,isBitcast,etc). Adding isFMA3 to there would be inappropriate as FMA3 is meaningful only for X86, while all other 32 1-bit fields defined there are quite generic and usable for all targets. Also, adding even 1 bit to there will increase the size of IR.
Unfortunately, I could not find anything similar but for X86 platform only.

3378 ↗

(On Diff #30193)

Yes, this method could be static and be similar to existing methods like "static bool isFrameLoadOpcode(int Opcode)", etc.

The reason why I passed 'MachineInstruction' argument instead of 'Opcode' to this function and why this method is not static now,
is that I wanted to reserve the opportunity to handle SCALAR FMAs and their 1st operand more optimistically later (when additional analysis of scalar FMA users would be implemented); please see the FIXME comment at the line 3487.

3387 ↗

(On Diff #30193)

There is one loop handling all vector and scalar FMAs below, I did not handle them separately.
The 'IsScalar' field was needed only to handle the 1st operand with extra carefulness as commuting 1st operand of scalar FMA requires some additional analysis.

Regarding the separating FMA tables into a separate header file...

Separating it to a header file makes sense only when it would be used by something else, i.e. not only by one method. Otherwise, it is more convenient to have this array definition closer to the function/method using that table.
Also, In my opinion the function local/static array OpcodeAlts is written using the same style that was used in several other places in this file (Please see the definition of MemoryFoldTable2Addr, MemoryFoldTable0, etc). Moving all similar static arrays of structures to a header file deserves a special/separate change-set.

3546 ↗

(On Diff #30193)

Special thank you for this comment! Separating INPUT and OUTPUT arguments seems very reasonable. I like this idea.

In my opinion both approaches have right to live though.
Before adding 2 additional arguments to findCommutedOpIndices() and fixing other places I would wait for more comments from reviewers.

[qcolombet]

So, the first alternative provides a flexible instrument that helps to know just what we need to know, while the second alternative makes findCommutaleOpIndices() to collect information that often is not needed later.

Right, but the first method does not tell the consumer that there are other alternatives for a given operand. The second method is much more general.

That being said, I believe you are right that most users of this method do not care about the alternatives, at least for now, so this is fine to have them iterate on the other indexes to check if there are alternatives. Exactly like you do in the two address pass.

Now, regarding the API, I would be in favor for something simpler, i.e.:

• Kill areOpsCommutable.
• Keep findCommutableOps with its current signature.

— Just make the two unsigned input/output parameters, like you did.
— Do not add two extra unsigned input parameters.

llvm/lib/CodeGen/TargetInstrInfo.cpp

208 ↗

(On Diff #30193)

Shouldn’t we just need two default arguments ~0U, instead of duplicating the prototype?

llvm/lib/Target/X86/X86InstrInfo.cpp

3546 ↗

(On Diff #30193)

I do not see why it is better to separate the input and output parameters here.
As long as the parameter will have the value: CommuteAnyOperandIndex, we know how to make the distinction.

[v_klochkov]

Thank you for the answer, Quentin.

I agree with your new comments and started preparing an updated change-set.
Regarding the default ~0U values for arguments of commuteInstruction() method, I still need your opinion - please see my answer to your question ('Inline Comment') for details.

Summarizing the planned additional changes:

• remove areOpsCommutable()
• fix ~0U magic const
• change fixCommutedOpindices() to a helper method.
• ? remove old commuteInstruction() method and duplicate code handling ~0U args 5 times..., but ONLY if you recommend doing that.

I would also add that if eventually we need a method that would return ALL pairs of commutable operands,
then it could be a NEW method that would have the word 'All' in its name:
... findAllCommutedOpIndices();
Such method (if needed) should be added in a separate change-set though.

Thank you,
Slava

llvm/include/llvm/Target/TargetInstrInfo.h

97–109 ↗

(On Diff #30193)

Yes, I agree, good catch. This will be fixed.

100 ↗

(On Diff #30193)

I agree, we need a const for this value. How about 'CommuteAnyOperandIndex'
or 'AnyCommutableOperandIndex'?

274–277 ↗

(On Diff #30193)

Matthias, can you please explain your idea?
In particular, what do you mean by saying

'...should be enough. The two operand commutable case will also work with that'?

I really want to just remove the old method that does not specify the commuted operands (i.e. to do exactly what you recommended here and to just add two operands to existing method). Unfortunately, I cannot do that without rewriting several places in LLVM.

This change-set replaces some calls of commuteInstruction(MI) calls with calls of commuteInstruction(MI,false,Idx1,Idx2). That made the changes in those places more clear.

There are though 7 or 8 places where the old style method is called and rewriting those places would make this change-set significantly bigger.
For example, CodeGen/MachineCSE.cpp has 2 calls of old commuteInstruction() method.
That place is obviously has opportunities for improvement, but that should be done in a separate change-set as the current change-set is already very big.

So, the only reason for having two variants of commuteInstruction() method is the need to limit the size of the change-set, to limit the efforts needed for code-review, and to eliminate the risks (correctness, etc.) that are brought by too huge change-sets.

llvm/lib/CodeGen/TargetInstrInfo.cpp

208 ↗

(On Diff #30193)

That could be done this way (i.e. use ~0U value as default values for the indices).
This way is good as there will be only one commuteInstruction() method.

The only big disadvantage on that way is that I'll need to duplicate the code of this method 5 times:

1. in TargetInstrInfo.cpp (in commuteInstruction()
2. in X86 specific implementation of commuteInstruction() method
3. in PowerPC specific implementation
4. in AMDGPU specific implementation
5. in ARM specific implementation.

The duplicated code would be something like this:

if (Idx1 == ~0U || Idx2 == ~0U) {
  if (!findCommutedOpIndices(MI, Idx1, Idx2)) {
    asssert(...);
    return nullptr;
  }
}

After comparing the advantages and disadvantages... I would prefer to have this code in one place (i.e. as it is implemented now in the 1st change-set), but this question is not something very important for me
and I will do as you recommend to do.

[v_klochkov]

Quentin,

I updated the change-set accordingly to the comments and recommendations from reviewers.

In this change-set I did not remove the old commuteInstruction() method as I am waiting for comments from you regarding my comment that says that this method can be removed only with the cost of duplication of the method code 5 times.

Thank you,
Slava

[qcolombet]

Hi Slava,

In this change-set I did not remove the old commuteInstruction() method as I am waiting for comments from you regarding my comment that says that this method can be removed only with the cost of duplication of the method code 5 times.

I think it makes sense to only expose one API and have one overridable internal API. Then put the boring code shared in the public function.
E.g., something like:
class stuff {
protected:

virtual ty PublicAPIImpl(ty2);

public:

/* not virtual */ ty PublicAPI(ty2) {
  /* boring code */
 return PublicAPIImpl(ty2);
}

};

Thanks,
-Quentin

llvm/lib/Target/X86/X86InstrInfo.cpp

3499 ↗

(On Diff #31721)

My understanding is that you are address this point here:

(2) Fixed a correctness problem caused by commuting 1st and 2nd operands of scalar FMAs generated for intrinsics.

Most of the time I think we do not care about the high level bits of the value (which is what you are fixing here). Therefore, I wonder if we are not being pessimistic on the commutation opportunities.
I agree we should seek correctness first, but I wonder how often that high level setting is actually expected… We had this bug forever and apparently nobody noticed it.

Anyway, what is your plan to get us the performance back?

[v_klochkov]

This change-set (3rd revision) is done accordingly to Quentin's suggestion to have 'protected virtual commuteInstructionImpl()' method
and to have the other method 'commuteInstruction()' non-virtual. The last one can accept CommuteAnyOperandIndex arguments.

This solution gives us only one public commuteInstruction() method instead of 2 methods with different interfaces/prototypes as it was in 2nd revision.

[v_klochkov]

Hi Quentin,

Thank you for the good idea (commuteInstructionImpl()). Hopefully, I understood it right.
I uploaded a new change-set. Would you please review the additional changes?

Also, I answered your question regarding the stability/correctness fix - please see my answer right after your inline comment/question. In my opinion, stability/correctness has priority over performance in this particular question.
There are at least 2 way how to fix the conservative code-gen.

Thank you,
Slava

llvm/lib/Target/X86/X86InstrInfo.cpp

3499 ↗

(On Diff #31721)

That correctness problem exists for FMAs and does not exist for ADD/MUL operations.
Also, FMAs are relatively new instructions.

For example, if you compile the test:

#include <immintrin.h>
double func(double y, double x) {
  return y + x;
}
__m128d funcx(__m128d y, __m128d x) {
  return _mm_add_sd(x, y);
}

then you'll see that only 1 instruction is generated for func() and 2 instructions for funcx().
func() just ignores the upper bits of returned XMM and funcx() correctly handles the upper bits of returned XMM value.

The difference in IR is: ADDSDrr opcode is used in func(), ADDSDrr_Int opcode is used in funcx().

So, one of possible solutions could be to add *_Int opcodes for FMA operations like it was done for ADD and MUL operations, and be more conservative for *FMA*_Int opcodes only.

Another solution is mentioned in FIXME comment above, i.e. to implement functionality that can tell if only the lowest element of the result of scalar FMA is used.

In my opinion, these 2 solutions do not exclude each other; they both should be implemented.

Currently, we do not have *FMA*_Int opcodes, that is why it would be better to be more conservative and correct. This patch might make the code a little bit worse/conservative on some corner cases, but it also improves code-gen for many other cases, for example, for those cases where the 1st or 2nd operand can be swapped with 3rd operand when it helps to do memory-op-folding optimization.

[qcolombet]

Hi Slava,

I think we are getting close to the final approval :).
Thanks for working on this.

I haven’t looked at the X86 specific part yet, as I would like the patch to be split in two:

• One patch for the change in API and NFC for all the backends.
• One patch to extend the commute code to handle the FMA cases.

Also, a few general comments:

• Some formatting look strange to me, please use clang-format on the patch.
• Please do not insert blank comment lines around comments, i.e.,

<— Remove those
<Some comment>
// <— Remove those

• Do not repeat the comment from the header in the cpp files.

Thanks,
-Quentin

llvm/include/llvm/Target/TargetInstrInfo.h

299 ↗

(On Diff #32646)

We can’t override this one anymore, please update the comment.

302 ↗

(On Diff #32646)

It would be useful to specify what is the behavior when both commute indices are CommuteAnyOperandIndex.

322 ↗

(On Diff #32646)

Same thing: what happens if both indices are set to CommuteAnyOperandIndex?

llvm/lib/CodeGen/TargetInstrInfo.cpp

124 ↗

(On Diff #32646)

We usually do not repeat the doxygen comment when they are set in the header.
By doing so, we risk to have them being out-of-sync. Therefore, please keep only the one in the header.
(Ditto for the other functions you’ve modified.)

llvm/lib/CodeGen/TwoAddressInstructionPass.cpp

1167 ↗

(On Diff #32646)

Please document the meaning of the other arguments as well, like Dist, BaseOpKilled, etc.

1182 ↗

(On Diff #32646)

Invert the condition and use “continue”. (Per LLVM coding standard.)

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

179 ↗

(On Diff #32646)

It is not handled at all, right?

Shouldn’t findCommutedOpIndices return only register operands?

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

888 ↗

(On Diff #32646)

I am guessing that you shouldn’t be the one doing this can of changes. Or at least, it should be a separate patch.

[v_klochkov]

I removed the X86 FMA specific changes and left here only the interface changes
for commuteInstruction() and findCommutedOpIndices() methods accordingly to Quentin's request.

Also, removed blank comment lines and duplicated methods' description/comments.

[v_klochkov]

`Hi Quentin,

I appreciate the time you spend to this code-review request.

In this change-set I removed X86 FMA specific changes - that will be a separate change-set
as you asked. This totally excluded the changes from the files:

llvm\test\CodeGen\X86\fma-commute-x86.ll
llvm\test\CodeGen\X86\fma_patterns.ll 
llvm\lib\Target\X86\X86InstrFMA.td

I removed the duplicated comments for methods (header vs cpp)
even though I personally like such duplication.
It is good to have function description in header file, but it is also
so convenient to have it in .cpp when you're looking at the method implementation;
you can see what arguments mean, etc. without need to take a look at *.h file.

AMDGPU changes:

I really did NOT want to do those non-obvious changes for AMDGPU.
Avoiding the changes would let me to avoid questions and to simplify the review/approve process.

Unfortunately, I had to do those changes. 
Also, those changes cannot be separated from the interface changes I did for
commuteInstruction() and findCommutedOpIndices().

The problem I met there can be explained this way:
1) There are some places like this:
     if (MI->IsCommutable() && TII->commuteInstruction(MI)) {}
    I.e. commuteInstruction() was called without preceding call of findCommutedOpIndices().
2) commuteInstruction() implementation for AMDGPU can commute Reg and Imm operands.

So, the solution is:
a) To allow AMDGPU implementation of findCommutedOpIndices() to return true
    when the second operand is Imm.
    Otherwise, all calls mentioned in problem (1) above would not commute instructions.
b) To fix SIFoldOperands.cpp/tryAddToFoldList() and add additional check there because
     Imm operand is not wanted there.
     I updated the FIXME comment there to make it more informative.

Thank you,
Slava
`

[qcolombet]

Hi Slava,

I still do not get the AMDGPU changes.

Unfortunately, I had to do those changes.
Also, those changes cannot be separated from the interface changes I did for
commuteInstruction() and findCommutedOpIndices().

We didn't change the basic semantic of those APIs, just making them more powerful, right?
So, why do we need to change the way we use them for this target.
Indeed, " if (MI->IsCommutable() && TII->commuteInstruction(MI)) {}" seems like a reasonable pattern to me.

What am I missing?

Thanks,
-Quentin

llvm/lib/CodeGen/TargetInstrInfo.cpp

212 ↗

(On Diff #33366)

The formatting still looks suspicious to me.
I expect the "else" to be on the same line as the '}'.
Have you run clang-format?

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

186 ↗

(On Diff #33366)

Since you said the target can commute imm and reg operand, why do we need this change?

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

796 ↗

(On Diff #33366)

I still don't get why we need to turn this if into an assert.

[v_klochkov]

The code 'if (MI->IsCommutable() && TII->commuteInstruction(MI))' works differently before and after the changes in commuteInstruction() method .

BEFORE:

If MI has the attribute 'MCID::Commutable' set to true, then 
    1) go to AMDGPU specific implementation of commuteInstruction()
        and commute operands (even if those are Reg and Imm).

AFTER:

If MI has the attribute 'MCID::Commutable' set to true, then 
  1) go to TargetInstrInfo::commuteInstruction() and try to commute operands
      1a) call AMDGPU specific implementation of findCommutedOpIndices()
      1b) if could find commutable operands, then call AMDGPU specific implementation of commuteInstruction()

The step 1a) returned false and no operands commute happened, which caused LIT tests fails.
So, I just synchronized understanding of commutable operands in AMDGPU methods findCommutedOpIndices() and commuteInstruction().
The updated findCommutedOpIndices() returns true for Reg and Imm operands if such can be commuted by commuteInstruction().

The fix in SIFoldOperands.cpp was needed because the updated findCommutedOpIndices() may return true for commutable Reg and Imm operands.
In such cases the index of Imm operand is stored into 'FoldList' std::vector object.
Later (downstream of the optimization) the elements of returned FoldList are treated as if they are all REG operands, which causes assert violations for Imm operands stored to 'FoldList'.

I'll fix the mentioned formatting issues, and will try to run format-clang tools.

Thanks,
Slava

llvm/include/llvm/Target/TargetInstrInfo.h

287–288 ↗

(On Diff #33366)

Fixed.

304–331 ↗

(On Diff #33366)

Fixed.

llvm/lib/CodeGen/TargetInstrInfo.cpp

124 ↗

(On Diff #33366)

Ok, removed the duplicated comments/descriptions.

llvm/lib/CodeGen/TwoAddressInstructionPass.cpp

1167 ↗

(On Diff #33366)

Fixed.

1182 ↗

(On Diff #33366)

Fixed.

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

179 ↗

(On Diff #33366)

findCommutedOpIndices() returns commutable operands. They can be Imm operands for AMDGPU target.
I also updated the FIXME comment to make it more informative.

[v_klochkov]

Fixed coding style/standard violations such as too long lines, indentations, etc.
using the recommendations from 'clang-format-diff.py' tool.

[v_klochkov]

Please see my answer regarding the new assert in AMDGPU version of commuteInstructionImpl() method.
Thanks,
Slava

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

796 ↗

(On Diff #33366)

(Src1Idx == -1) is impossible here as this place is reachable only after findCommutedOpIndices() call which filters out such situations.
This assert for Src1Idx is semantically and stylistically equivalent to the assert at the line 788 (assert for Src0Idx).

[qcolombet]

Hi Slava,

LGTM.

The code 'if (MI->IsCommutable() && TII->commuteInstruction(MI))' works differently before and after the changes in commuteInstruction() method .

Thanks for the explanation, now I see the difference in semantic. You’ll want to send an email to give a heads-up to out of tree targets for this change.

Also, before landing the patch, ping Tom ([e-mail], code owner of the AMDGPU), to see if the changes in AMDGPU also look good to him.

Cheers,
-Quentin

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

178 ↗

(On Diff #33495)

The comment is just fine without the FIXME in front of it.
Just remove it, i.e., I do not think there is anything to fix here in the end.

[v_klochkov]

Mr. Stellard,
Please review and approve the changes in 3 files owned by AMDGPU target:

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp
llvm/lib/Target/AMDGPU/SIInstrInfo.h
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

All the changes (14 files, including AMDGPU) have been reviewed by Quentin Colombet.

Also, below I attached the e-mails I sent you to your e-mail at amd.com
Please see more details in it.

Thank you,
Vyacheslav Klochkov

From: Klochkov, Vyacheslav N  
Sent: Monday, September 14, 2015 4:46 PM
To: '<EDITED>@amd.com'
Cc: Klochkov, Vyacheslav N
Subject: RE: LLVM code-review: http://reviews.llvm.org/D11370
 
Mr. Stellard,
 
In this e-mail I am asking you for approval for AMDGMU specific changes
that are a small part of changes improving methods that commute operands of Machine Instructions.
 
Please see the details below.
Also, I will be more than happy to answer your questions if you have any.
 
Thank you,
Vyacheslav Klochkov
------------------------------------------------------------------------------------
From: Klochkov, Vyacheslav N 
Sent: Tuesday, September 8, 2015 4:05 PM
To: <EDITED>@amd.com
Cc: Klochkov, Vyacheslav N
Subject: LLVM code-review: http://reviews.llvm.org/D11370
 
Dear Mr. Stellard,
 
Would you please approve the AMDGPU specific changes in this code-review tracker:
http://reviews.llvm.org/D11370
 
This change-set was reviewed and approved by several people:
-        Quentin Colombet (official code-reviewer).
-        David Kreitzer (code-review before submitting changes to Open Source community);
-        Michael Kuperstein (code-review before submitting changes to community);
-        Matthias Braun and Elena Demikhovsky (not official reviewers, sent some comments).
Quentin recommended to ask for your approval for AMDGPU changes.
 
This change-set includes changes in 14 files, 3 of 14 are AMDGPU specific:
llvm/lib/Target/AMDGPU/SIFoldOperands.cpp
llvm/lib/Target/AMDGPU/SIInstrInfo.h
llvm/lib/Target/AMDGPU/SIInstrInfo.cpp
 
 
The change-set improves the interface of the findCommutedOpIndices() and commuteInstruction() methods.
It did not change the code-generation for AMDGPU, all LIT tests passed.
 
The old commuteInstruction() method did not specify the indices of the operands to be commuted.
The main idea of the change-set is that commuteInstruction() must be able to have the commuted operands be specified explicitly.
That is needed when a caller of commuteInstruction() method knows what operands must be commuted.
 
For example, if 1st and 2nd ops are commutable and 1st and 3rd operands are commutable too,
the old commuteInstruction() method did not allow to do the second commute transformation.
The new method fixes that problem.
 
It is still possible not to specify the operands to be commuted, but in such cases the operands to be commuted must be found by the method findCommutedOpIndices()
(Please see the commuteInstruction() method in llvm/lib/CodeGen/TargetInstrInfo.cpp).
 
This interface change in llvm/lib/CodeGen/TargetInstrInfo.cpp caused the need to do minor changes in target specific implementation of
findCommutedOpIndices() and commuteInstruction().
 
The changes in ARM, X86, PowerPC went very smooth. The changes in AMDGPU required me to do a little bit bigger changes.
 
For AMDGPU I tuned the findCommutedOpIndices() such a way that now it returns TRUE for commutable Reg and Imm operands
as Reg and Imm operands can be commuted in AMDGPU specific implementation of commuteInstruction() method.
So, findCommutedOpIndices() and commuteInstruction() are in sync now, i.e. the 1st returns true when the 2nd can do the commute.
 
The change in findCommutedOpIndices() allowed me to add an assert on one of operands in commuteInstructionImpl()  AMDGPU method.
Also, I needed to add a simple check to SIFoldOperands.cpp to avoid Imm operands in the transformation where only Reg operands are expected.
 
Thank you,
Vyacheslav Klochkov

[arsenm]

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

874 ↗

(On Diff #33495)

Typo: immeditate

885–887 ↗

(On Diff #33495)

It's not entirely accurate to use isVOP2 / isVOP3 for checking the number of operands. VOP* instructions are always available in a VOP3 encoding, but will still have < 3 operands. Checking if AMDGPU::OpName::src2 is a valid operand is a more reliable check. SALU instructions with an immediate can also be commuted, although there is less reason to do so other than canonicalization. Although it looks like isVOP2/isVOP3 is what commuteInstruction already checks so I guess this is OK for now.

[arsenm]

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

798–799 ↗

(On Diff #33495)

I would prefer to deMorgan's law this and distribute the !

[v_klochkov]

Thank you for the comments.
I fixed the "immeditate" misprint,
replaced the if statement: if (!(A && B) && !(C && D)) --> if ((!A || !B) && (!C || !D))
and removed the 'FIXME' word in SIFoldOperands.cpp (accordingly to recommendation from Quentin Colombet).

llvm/lib/Target/AMDGPU/SIFoldOperands.cpp

178 ↗

(On Diff #33495)

Fixed. I removed the "FIXME" word.

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

798–799 ↗

(On Diff #33495)

Ok, Fixed.

874 ↗

(On Diff #33495)

Fixed: immeditate -> immediate.

885–887 ↗

(On Diff #33495)

Thank you for the explanations.
I am quite happy that you are Ok with the current version of the changes as the fixing of such subtle things should be done by AMDGPU experts. In this change-set I just synchronized the checks in findOpIndicesToCommute() and commuteInstructionImpl() (i.e. re-used the checks from commuteInstructionImpl() in findOpIndicesToCommute()).

[arsenm]

LGTM

llvm/lib/Target/AMDGPU/SIInstrInfo.cpp

779–780 ↗

(On Diff #33495)

I just removed this check a few days ago, so you probably will have a conflict when you apply this to trunk

800–801 ↗

(On Diff #33495)

These asserts can be removed. I decided that there's no point to checking if src0/src1 are valid operands

llvm/trunk/include/llvm/Target/TargetInstrInfo.h

protected:
   /// than producing a value, or if it requres any address registers that are
   /// not always available.
   /// Requirements must be check as stated in isTriviallyReMaterializable() .
   virtual bool isReallyTriviallyReMaterializable(const MachineInstr *MI,
                                                  AliasAnalysis *AA) const {
     return false;
   }
 
+  /// This method commutes the operands of the given machine instruction MI.
+  /// The operands to be commuted are specified by their indices OpIdx1 and
+  /// OpIdx2.
+  ///
+  /// If a target has any instructions that are commutable but require
+  /// converting to different instructions or making non-trivial changes
+  /// to commute them, this method can be overloaded to do that.
+  /// The default implementation simply swaps the commutable operands.
+  ///
+  /// If NewMI is false, MI is modified in place and returned; otherwise, a
+  /// new machine instruction is created and returned.
+  ///
+  /// Do not call this method for a non-commutable instruction.
+  /// Even though the instruction is commutable, the method may still
+  /// fail to commute the operands, null pointer is returned in such cases.
+  virtual MachineInstr *commuteInstructionImpl(MachineInstr *MI,
+                                               bool NewMI,
+                                               unsigned OpIdx1,
+                                               unsigned OpIdx2) const;
+
+  /// Assigns the (CommutableOpIdx1, CommutableOpIdx2) pair of commutable
+  /// operand indices to (ResultIdx1, ResultIdx2).
+  /// One or both input values of the pair: (ResultIdx1, ResultIdx2) may be
+  /// predefined to some indices or be undefined (designated by the special
+  /// value 'CommuteAnyOperandIndex').
+  /// The predefined result indices cannot be re-defined.
+  /// The function returns true iff after the result pair redefinition
+  /// the fixed result pair is equal to or equivalent to the source pair of
+  /// indices: (CommutableOpIdx1, CommutableOpIdx2). It is assumed here that
+  /// the pairs (x,y) and (y,x) are equivalent.
+  static bool fixCommutedOpIndices(unsigned &ResultIdx1,
+                                   unsigned &ResultIdx2,
+                                   unsigned CommutableOpIdx1,
+                                   unsigned CommutableOpIdx2);
+
 private:
   /// For instructions with opcodes for which the M_REMATERIALIZABLE flag is
   /// set and the target hook isReallyTriviallyReMaterializable returns false,
   /// this function does target-independent tests to determine if the
   /// instruction is really trivially rematerializable.
   bool isReallyTriviallyReMaterializableGeneric(const MachineInstr *MI,
                                                 AliasAnalysis *AA) const;
 
public:
   /// performed, otherwise it returns the last new instruction.
   ///
   virtual MachineInstr *
   convertToThreeAddress(MachineFunction::iterator &MFI,
                    MachineBasicBlock::iterator &MBBI, LiveVariables *LV) const {
     return nullptr;
   }
 
-  /// If a target has any instructions that are commutable but require
-  /// converting to different instructions or making non-trivial changes to
-  /// commute them, this method can overloaded to do that.
-  /// The default implementation simply swaps the commutable operands.
+  // This constant can be used as an input value of operand index passed to
+  // the method findCommutedOpIndices() to tell the method that the
+  // corresponding operand index is not pre-defined and that the method
+  // can pick any commutable operand.
+  static const unsigned CommuteAnyOperandIndex = ~0U;
+
+  /// This method commutes the operands of the given machine instruction MI.
+  ///
+  /// The operands to be commuted are specified by their indices OpIdx1 and
+  /// OpIdx2. OpIdx1 and OpIdx2 arguments may be set to a special value
+  /// 'CommuteAnyOperandIndex', which means that the method is free to choose
+  /// any arbitrarily chosen commutable operand. If both arguments are set to
+  /// 'CommuteAnyOperandIndex' then the method looks for 2 different commutable
+  /// operands; then commutes them if such operands could be found.
+  ///
   /// If NewMI is false, MI is modified in place and returned; otherwise, a
-  /// new machine instruction is created and returned.  Do not call this
-  /// method for a non-commutable instruction, but there may be some cases
-  /// where this method fails and returns null.
-  virtual MachineInstr *commuteInstruction(MachineInstr *MI,
-                                           bool NewMI = false) const;
-
-  /// If specified MI is commutable, return the two operand indices that would
-  /// swap value. Return false if the instruction
-  /// is not in a form which this routine understands.
+  /// new machine instruction is created and returned.
+  ///
+  /// Do not call this method for a non-commutable instruction or
+  /// for non-commuable operands.
+  /// Even though the instruction is commutable, the method may still
+  /// fail to commute the operands, null pointer is returned in such cases.
+  MachineInstr *
+  commuteInstruction(MachineInstr *MI,
+                     bool NewMI = false,
+                     unsigned OpIdx1 = CommuteAnyOperandIndex,
+                     unsigned OpIdx2 = CommuteAnyOperandIndex) const;
+
+  /// Returns true iff the routine could find two commutable operands in the
+  /// given machine instruction.
+  /// The 'SrcOpIdx1' and 'SrcOpIdx2' are INPUT and OUTPUT arguments.
+  /// If any of the INPUT values is set to the special value
+  /// 'CommuteAnyOperandIndex' then the method arbitrarily picks a commutable
+  /// operand, then returns its index in the corresponding argument.
+  /// If both of INPUT values are set to 'CommuteAnyOperandIndex' then method
+  /// looks for 2 commutable operands.
+  /// If INPUT values refer to some operands of MI, then the method simply
+  /// returns true if the corresponding operands are commutable and returns
+  /// false otherwise.
+  ///
+  /// For example, calling this method this way:
+  ///     unsigned Op1 = 1, Op2 = CommuteAnyOperandIndex;
+  ///     findCommutedOpIndices(MI, Op1, Op2);
+  /// can be interpreted as a query asking to find an operand that would be
+  /// commutable with the operand#1.
   virtual bool findCommutedOpIndices(MachineInstr *MI, unsigned &SrcOpIdx1,
                                      unsigned &SrcOpIdx2) const;
 
   /// A pair composed of a register and a sub-register index.
   /// Used to give some type checking when modeling Reg:SubReg.
   struct RegSubRegPair {
     unsigned Reg;
     unsigned SubReg;

llvm/trunk/lib/CodeGen/RegisterCoalescer.cpp

if (!DefMI->isCommutable())
     return false;
   // If DefMI is a two-address instruction then commuting it will change the
   // destination register.
   int DefIdx = DefMI->findRegisterDefOperandIdx(IntA.reg);
   assert(DefIdx != -1);
   unsigned UseOpIdx;
   if (!DefMI->isRegTiedToUseOperand(DefIdx, &UseOpIdx))
     return false;
-  unsigned Op1, Op2, NewDstIdx;
-  if (!TII->findCommutedOpIndices(DefMI, Op1, Op2))
-    return false;
-  if (Op1 == UseOpIdx)
-    NewDstIdx = Op2;
-  else if (Op2 == UseOpIdx)
-    NewDstIdx = Op1;
-  else
+
+  // FIXME: The code below tries to commute 'UseOpIdx' operand with some other
+  // commutable operand which is expressed by 'CommuteAnyOperandIndex'value
+  // passed to the method. That _other_ operand is chosen by
+  // the findCommutedOpIndices() method.
+  //
+  // That is obviously an area for improvement in case of instructions having
+  // more than 2 operands. For example, if some instruction has 3 commutable
+  // operands then all possible variants (i.e. op#1<->op#2, op#1<->op#3,
+  // op#2<->op#3) of commute transformation should be considered/tried here.
+  unsigned NewDstIdx = TargetInstrInfo::CommuteAnyOperandIndex;
+  if (!TII->findCommutedOpIndices(DefMI, UseOpIdx, NewDstIdx))
     return false;
 
   MachineOperand &NewDstMO = DefMI->getOperand(NewDstIdx);
   unsigned NewReg = NewDstMO.getReg();
   if (NewReg != IntB.reg || !IntB.Query(AValNo->def).isKill())
     return false;
 
   // Make sure there are no other definitions of IntB that would reach the
bool RegisterCoalescer::removeCopyByCommutingDef(const CoalescerPair &CP,
   }
 
   DEBUG(dbgs() << "\tremoveCopyByCommutingDef: " << AValNo->def << '\t'
                << *DefMI);
 
   // At this point we have decided that it is legal to do this
   // transformation.  Start by commuting the instruction.
   MachineBasicBlock *MBB = DefMI->getParent();
-  MachineInstr *NewMI = TII->commuteInstruction(DefMI);
+  MachineInstr *NewMI =
+      TII->commuteInstruction(DefMI, false, UseOpIdx, NewDstIdx);
   if (!NewMI)
     return false;
   if (TargetRegisterInfo::isVirtualRegister(IntA.reg) &&
       TargetRegisterInfo::isVirtualRegister(IntB.reg) &&
       !MRI->constrainRegClass(IntB.reg, MRI->getRegClass(IntA.reg)))
     return false;
   if (NewMI != DefMI) {
     LIS->ReplaceMachineInstrInMaps(DefMI, NewMI);

llvm/trunk/lib/CodeGen/TargetInstrInfo.cpp

TargetInstrInfo::ReplaceTailWithBranchTo(MachineBasicBlock::iterator Tail,
 
   // If MBB isn't immediately before MBB, insert a branch to it.
   if (++MachineFunction::iterator(MBB) != MachineFunction::iterator(NewDest))
     InsertBranch(*MBB, NewDest, nullptr, SmallVector<MachineOperand, 0>(),
                  Tail->getDebugLoc());
   MBB->addSuccessor(NewDest);
 }
 
-// commuteInstruction - The default implementation of this method just exchanges
-// the two operands returned by findCommutedOpIndices.
-MachineInstr *TargetInstrInfo::commuteInstruction(MachineInstr *MI,
-                                                  bool NewMI) const {
+MachineInstr *TargetInstrInfo::commuteInstructionImpl(MachineInstr *MI,
+                                                      bool NewMI,
+                                                      unsigned Idx1,
+                                                      unsigned Idx2) const {
   const MCInstrDesc &MCID = MI->getDesc();
   bool HasDef = MCID.getNumDefs();
   if (HasDef && !MI->getOperand(0).isReg())
     // No idea how to commute this instruction. Target should implement its own.
     return nullptr;
-  unsigned Idx1, Idx2;
-  if (!findCommutedOpIndices(MI, Idx1, Idx2)) {
-    assert(MI->isCommutable() && "Precondition violation: MI must be commutable.");
-    return nullptr;
-  }
 
+  unsigned CommutableOpIdx1 = Idx1, CommutableOpIdx2 = Idx2;
+  assert(findCommutedOpIndices(MI, CommutableOpIdx1, CommutableOpIdx2) &&
+         CommutableOpIdx1 == Idx1 && CommutableOpIdx2 == Idx2 &&
+         "TargetInstrInfo::CommuteInstructionImpl(): not commutable operands.");
   assert(MI->getOperand(Idx1).isReg() && MI->getOperand(Idx2).isReg() &&
          "This only knows how to commute register operands so far");
+
   unsigned Reg0 = HasDef ? MI->getOperand(0).getReg() : 0;
   unsigned Reg1 = MI->getOperand(Idx1).getReg();
   unsigned Reg2 = MI->getOperand(Idx2).getReg();
   unsigned SubReg0 = HasDef ? MI->getOperand(0).getSubReg() : 0;
   unsigned SubReg1 = MI->getOperand(Idx1).getSubReg();
   unsigned SubReg2 = MI->getOperand(Idx2).getSubReg();
   bool Reg1IsKill = MI->getOperand(Idx1).isKill();
   bool Reg2IsKill = MI->getOperand(Idx2).isKill();
MachineInstr *TargetInstrInfo::commuteInstructionImpl(MachineInstr *MI,
   MI->getOperand(Idx1).setIsKill(Reg2IsKill);
   MI->getOperand(Idx2).setIsUndef(Reg1IsUndef);
   MI->getOperand(Idx1).setIsUndef(Reg2IsUndef);
   MI->getOperand(Idx2).setIsInternalRead(Reg1IsInternal);
   MI->getOperand(Idx1).setIsInternalRead(Reg2IsInternal);
   return MI;
 }
 
-/// findCommutedOpIndices - If specified MI is commutable, return the two
-/// operand indices that would swap value. Return true if the instruction
-/// is not in a form which this routine understands.
+MachineInstr *TargetInstrInfo::commuteInstruction(MachineInstr *MI,
+                                                  bool NewMI,
+                                                  unsigned OpIdx1,
+                                                  unsigned OpIdx2) const {
+  // If OpIdx1 or OpIdx2 is not specified, then this method is free to choose
+  // any commutable operand, which is done in findCommutedOpIndices() method
+  // called below.
+  if ((OpIdx1 == CommuteAnyOperandIndex || OpIdx2 == CommuteAnyOperandIndex) &&
+      !findCommutedOpIndices(MI, OpIdx1, OpIdx2)) {
+    assert(MI->isCommutable() &&
+           "Precondition violation: MI must be commutable.");
+    return nullptr;
+  }
+  return commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
+}
+
+bool TargetInstrInfo::fixCommutedOpIndices(unsigned &ResultIdx1,
+                                           unsigned &ResultIdx2,
+                                           unsigned CommutableOpIdx1,
+                                           unsigned CommutableOpIdx2) {
+  if (ResultIdx1 == CommuteAnyOperandIndex &&
+      ResultIdx2 == CommuteAnyOperandIndex) {
+    ResultIdx1 = CommutableOpIdx1;
+    ResultIdx2 = CommutableOpIdx2;
+  } else if (ResultIdx1 == CommuteAnyOperandIndex) {
+    if (ResultIdx2 == CommutableOpIdx1)
+      ResultIdx1 = CommutableOpIdx2;
+    else if (ResultIdx2 == CommutableOpIdx2)
+      ResultIdx1 = CommutableOpIdx1;
+    else
+      return false;
+  } else if (ResultIdx2 == CommuteAnyOperandIndex) {
+    if (ResultIdx1 == CommutableOpIdx1)
+      ResultIdx2 = CommutableOpIdx2;
+    else if (ResultIdx1 == CommutableOpIdx2)
+      ResultIdx2 = CommutableOpIdx1;
+    else
+      return false;
+  } else
+    // Check that the result operand indices match the given commutable
+    // operand indices.
+    return (ResultIdx1 == CommutableOpIdx1 && ResultIdx2 == CommutableOpIdx2) ||
+           (ResultIdx1 == CommutableOpIdx2 && ResultIdx2 == CommutableOpIdx1);
+
+  return true;
+}
+
 bool TargetInstrInfo::findCommutedOpIndices(MachineInstr *MI,
                                             unsigned &SrcOpIdx1,
                                             unsigned &SrcOpIdx2) const {
   assert(!MI->isBundle() &&
          "TargetInstrInfo::findCommutedOpIndices() can't handle bundles");
 
   const MCInstrDesc &MCID = MI->getDesc();
   if (!MCID.isCommutable())
     return false;
+
   // This assumes v0 = op v1, v2 and commuting would swap v1 and v2. If this
   // is not true, then the target must implement this.
-  SrcOpIdx1 = MCID.getNumDefs();
-  SrcOpIdx2 = SrcOpIdx1 + 1;
+  unsigned CommutableOpIdx1 = MCID.getNumDefs();
+  unsigned CommutableOpIdx2 = CommutableOpIdx1 + 1;
+  if (!fixCommutedOpIndices(SrcOpIdx1, SrcOpIdx2,
+                            CommutableOpIdx1, CommutableOpIdx2))
+    return false;
+
   if (!MI->getOperand(SrcOpIdx1).isReg() ||
       !MI->getOperand(SrcOpIdx2).isReg())
     // No idea.
     return false;
   return true;
 }
 
-
 bool
 TargetInstrInfo::isUnpredicatedTerminator(const MachineInstr *MI) const {
   if (!MI->isTerminator()) return false;
 
   // Conditional branch is a special case.
   if (MI->isBranch() && !MI->isBarrier())
     return true;
   if (!MI->isPredicable())

llvm/trunk/lib/CodeGen/TwoAddressInstructionPass.cpp

class TwoAddressInstructionPass : public MachineFunctionPass {
 
   bool isRevCopyChain(unsigned FromReg, unsigned ToReg, int Maxlen);
 
   bool noUseAfterLastDef(unsigned Reg, unsigned Dist, unsigned &LastDef);
 
   bool isProfitableToCommute(unsigned regA, unsigned regB, unsigned regC,
                              MachineInstr *MI, unsigned Dist);
 
-  bool commuteInstruction(MachineBasicBlock::iterator &mi,
-                          unsigned RegB, unsigned RegC, unsigned Dist);
+  bool commuteInstruction(MachineInstr *MI,
+                          unsigned RegBIdx, unsigned RegCIdx, unsigned Dist);
 
   bool isProfitableToConv3Addr(unsigned RegA, unsigned RegB);
 
   bool convertInstTo3Addr(MachineBasicBlock::iterator &mi,
                           MachineBasicBlock::iterator &nmi,
                           unsigned RegA, unsigned RegB, unsigned Dist);
 
   bool isDefTooClose(unsigned Reg, unsigned Dist, MachineInstr *MI);
 
   bool rescheduleMIBelowKill(MachineBasicBlock::iterator &mi,
                              MachineBasicBlock::iterator &nmi,
                              unsigned Reg);
   bool rescheduleKillAboveMI(MachineBasicBlock::iterator &mi,
                              MachineBasicBlock::iterator &nmi,
                              unsigned Reg);
 
   bool tryInstructionTransform(MachineBasicBlock::iterator &mi,
                                MachineBasicBlock::iterator &nmi,
                                unsigned SrcIdx, unsigned DstIdx,
                                unsigned Dist, bool shouldOnlyCommute);
 
+  bool tryInstructionCommute(MachineInstr *MI,
+                             unsigned DstOpIdx,
+                             unsigned BaseOpIdx,
+                             bool BaseOpKilled,
+                             unsigned Dist);
   void scanUses(unsigned DstReg);
 
   void processCopy(MachineInstr *MI);
 
   typedef SmallVector<std::pair<unsigned, unsigned>, 4> TiedPairList;
   typedef SmallDenseMap<unsigned, TiedPairList> TiedOperandMap;
   bool collectTiedOperands(MachineInstr *MI, TiedOperandMap&);
   void processTiedPairs(MachineInstr *MI, TiedPairList&, unsigned &Dist);
isProfitableToCommute(unsigned regA, unsigned regB, unsigned regC,
   // Since there are no intervening uses for both registers, then commute
   // if the def of regC is closer. Its live interval is shorter.
   return LastDefB && LastDefC && LastDefC > LastDefB;
 }
 
 /// commuteInstruction - Commute a two-address instruction and update the basic
 /// block, distance map, and live variables if needed. Return true if it is
 /// successful.
-bool TwoAddressInstructionPass::
-commuteInstruction(MachineBasicBlock::iterator &mi,
-                   unsigned RegB, unsigned RegC, unsigned Dist) {
-  MachineInstr *MI = mi;
+bool TwoAddressInstructionPass::commuteInstruction(MachineInstr *MI,
+                                                   unsigned RegBIdx,
+                                                   unsigned RegCIdx,
+                                                   unsigned Dist) {
+  unsigned RegC = MI->getOperand(RegCIdx).getReg();
   DEBUG(dbgs() << "2addr: COMMUTING  : " << *MI);
-  MachineInstr *NewMI = TII->commuteInstruction(MI);
+  MachineInstr *NewMI = TII->commuteInstruction(MI, false, RegBIdx, RegCIdx);
 
   if (NewMI == nullptr) {
     DEBUG(dbgs() << "2addr: COMMUTING FAILED!\n");
     return false;
   }
 
   DEBUG(dbgs() << "2addr: COMMUTED TO: " << *NewMI);
   assert(NewMI == MI &&
if (LIS) {
     LV->removeVirtualRegisterKilled(Reg, KillMI);
     LV->addVirtualRegisterKilled(Reg, MI);
   }
 
   DEBUG(dbgs() << "\trescheduled kill: " << *KillMI);
   return true;
 }
 
+/// Tries to commute the operand 'BaseOpIdx' and some other operand in the
+/// given machine instruction to improve opportunities for coalescing and
+/// elimination of a register to register copy.
+///
+/// 'DstOpIdx' specifies the index of MI def operand.
+/// 'BaseOpKilled' specifies if the register associated with 'BaseOpIdx'
+/// operand is killed by the given instruction.
+/// The 'Dist' arguments provides the distance of MI from the start of the
+/// current basic block and it is used to determine if it is profitable
+/// to commute operands in the instruction.
+///
+/// Returns true if the transformation happened. Otherwise, returns false.
+bool TwoAddressInstructionPass::tryInstructionCommute(MachineInstr *MI,
+                                                      unsigned DstOpIdx,
+                                                      unsigned BaseOpIdx,
+                                                      bool BaseOpKilled,
+                                                      unsigned Dist) {
+  unsigned DstOpReg = MI->getOperand(DstOpIdx).getReg();
+  unsigned BaseOpReg = MI->getOperand(BaseOpIdx).getReg();
+  unsigned OpsNum = MI->getDesc().getNumOperands();
+  unsigned OtherOpIdx = MI->getDesc().getNumDefs();
+  for (; OtherOpIdx < OpsNum; OtherOpIdx++) {
+    // The call of findCommutedOpIndices below only checks if BaseOpIdx
+    // and OtherOpIdx are commutable, it does not really searches for
+    // other commutable operands and does not change the values of passed
+    // variables.
+    if (OtherOpIdx == BaseOpIdx ||
+        !TII->findCommutedOpIndices(MI, BaseOpIdx, OtherOpIdx))
+      continue;
+
+    unsigned OtherOpReg = MI->getOperand(OtherOpIdx).getReg();
+    bool AggressiveCommute = false;
+
+    // If OtherOp dies but BaseOp does not, swap the OtherOp and BaseOp
+    // operands. This makes the live ranges of DstOp and OtherOp joinable.
+    bool DoCommute =
+        !BaseOpKilled && isKilled(*MI, OtherOpReg, MRI, TII, LIS, false);
+
+    if (!DoCommute &&
+        isProfitableToCommute(DstOpReg, BaseOpReg, OtherOpReg, MI, Dist)) {
+      DoCommute = true;
+      AggressiveCommute = true;
+    }
+
+    // If it's profitable to commute, try to do so.
+    if (DoCommute && commuteInstruction(MI, BaseOpIdx, OtherOpIdx, Dist)) {
+      ++NumCommuted;
+      if (AggressiveCommute)
+        ++NumAggrCommuted;
+      return true;
+    }
+  }
+  return false;
+}
+
 /// tryInstructionTransform - For the case where an instruction has a single
 /// pair of tied register operands, attempt some transformations that may
 /// either eliminate the tied operands or improve the opportunities for
 /// coalescing away the register copy.  Returns true if no copy needs to be
 /// inserted to untie mi's operands (either because they were untied, or
 /// because mi was rescheduled, and will be visited again later). If the
 /// shouldOnlyCommute flag is true, only instruction commutation is attempted.
 bool TwoAddressInstructionPass::
tryInstructionTransform(MachineBasicBlock::iterator &mi,
 
   assert(TargetRegisterInfo::isVirtualRegister(regB) &&
          "cannot make instruction into two-address form");
   bool regBKilled = isKilled(MI, regB, MRI, TII, LIS, true);
 
   if (TargetRegisterInfo::isVirtualRegister(regA))
     scanUses(regA);
 
-  // Check if it is profitable to commute the operands.
-  unsigned SrcOp1, SrcOp2;
-  unsigned regC = 0;
-  unsigned regCIdx = ~0U;
-  bool TryCommute = false;
-  bool AggressiveCommute = false;
-  if (MI.isCommutable() && MI.getNumOperands() >= 3 &&
-      TII->findCommutedOpIndices(&MI, SrcOp1, SrcOp2)) {
-    if (SrcIdx == SrcOp1)
-      regCIdx = SrcOp2;
-    else if (SrcIdx == SrcOp2)
-      regCIdx = SrcOp1;
-
-    if (regCIdx != ~0U) {
-      regC = MI.getOperand(regCIdx).getReg();
-      if (!regBKilled && isKilled(MI, regC, MRI, TII, LIS, false))
-        // If C dies but B does not, swap the B and C operands.
-        // This makes the live ranges of A and C joinable.
-        TryCommute = true;
-      else if (isProfitableToCommute(regA, regB, regC, &MI, Dist)) {
-        TryCommute = true;
-        AggressiveCommute = true;
-      }
-    }
-  }
+  bool Commuted = tryInstructionCommute(&MI, DstIdx, SrcIdx, regBKilled, Dist);
 
   // If the instruction is convertible to 3 Addr, instead
   // of returning try 3 Addr transformation aggresively and
   // use this variable to check later. Because it might be better.
   // For example, we can just use `leal (%rsi,%rdi), %eax` and `ret`
   // instead of the following code.
   //   addl     %esi, %edi
   //   movl     %edi, %eax
   //   ret
-  bool Commuted = false;
-
-  // If it's profitable to commute, try to do so.
-  if (TryCommute && commuteInstruction(mi, regB, regC, Dist)) {
-    Commuted = true;
-    ++NumCommuted;
-    if (AggressiveCommute)
-      ++NumAggrCommuted;
-    if (!MI.isConvertibleTo3Addr())
-      return false;
-  }
+  if (Commuted && !MI.isConvertibleTo3Addr())
 
   if (shouldOnlyCommute)
     return false;
 
   // If there is one more use of regB later in the same MBB, consider
   // re-schedule this MI below it.
   if (!Commuted && EnableRescheduling && rescheduleMIBelowKill(mi, nmi, regB)) {
     ++NumReSchedDowns;

llvm/trunk/lib/Target/AMDGPU/SIFoldOperands.cpp

if (!TII->isOperandLegal(MI, OpNo, OpToFold)) {
     // If we are already folding into another operand of MI, then
     // we can't commute the instruction, otherwise we risk making the
     // other fold illegal.
     if (isUseMIInFoldList(FoldList, MI))
       return false;
 
     // Operand is not legal, so try to commute the instruction to
     // see if this makes it possible to fold.
-    unsigned CommuteIdx0;
-    unsigned CommuteIdx1;
+    unsigned CommuteIdx0 = TargetInstrInfo::CommuteAnyOperandIndex;
+    unsigned CommuteIdx1 = TargetInstrInfo::CommuteAnyOperandIndex;
     bool CanCommute = TII->findCommutedOpIndices(MI, CommuteIdx0, CommuteIdx1);
 
     if (CanCommute) {
       if (CommuteIdx0 == OpNo)
         OpNo = CommuteIdx1;
       else if (CommuteIdx1 == OpNo)
         OpNo = CommuteIdx0;
     }
 
-    if (!CanCommute || !TII->commuteInstruction(MI))
+    // One of operands might be an Imm operand, and OpNo may refer to it after
+    // the call of commuteInstruction() below. Such situations are avoided
+    // here explicitly as OpNo must be a register operand to be a candidate
+    // for memory folding.
+    if (CanCommute && (!MI->getOperand(CommuteIdx0).isReg() ||
+                       !MI->getOperand(CommuteIdx1).isReg()))
+      return false;
+
+    if (!CanCommute ||
+        !TII->commuteInstruction(MI, false, CommuteIdx0, CommuteIdx1))
       return false;
 
     if (!TII->isOperandLegal(MI, OpNo, OpToFold))
       return false;
   }
 
   FoldList.push_back(FoldCandidate(MI, OpNo, OpToFold));
   return true;

llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.h

void addUsersToMoveToVALUWorklist(
     unsigned Reg, MachineRegisterInfo &MRI,
     SmallVectorImpl<MachineInstr *> &Worklist) const;
 
   bool checkInstOffsetsDoNotOverlap(MachineInstr *MIa,
                                     MachineInstr *MIb) const;
 
   unsigned findUsedSGPR(const MachineInstr *MI, int OpIndices[3]) const;
 
+protected:
+  MachineInstr *commuteInstructionImpl(MachineInstr *MI,
+                                       bool NewMI,
+                                       unsigned OpIdx0,
+                                       unsigned OpIdx1) const override;
+
 public:
   explicit SIInstrInfo(const AMDGPUSubtarget &st);
 
   const SIRegisterInfo &getRegisterInfo() const override {
     return RI;
   }
 
   bool isReallyTriviallyReMaterializable(const MachineInstr *MI,
public:
   // \brief Returns an opcode that can be used to move a value to a \p DstRC
   // register.  If there is no hardware instruction that can store to \p
   // DstRC, then AMDGPU::COPY is returned.
   unsigned getMovOpcode(const TargetRegisterClass *DstRC) const;
 
   LLVM_READONLY
   int commuteOpcode(const MachineInstr &MI) const;
 
-  MachineInstr *commuteInstruction(MachineInstr *MI,
-                                   bool NewMI = false) const override;
   bool findCommutedOpIndices(MachineInstr *MI,
                              unsigned &SrcOpIdx1,
                              unsigned &SrcOpIdx2) const override;
 
   bool areMemAccessesTriviallyDisjoint(
     MachineInstr *MIa, MachineInstr *MIb,
     AliasAnalysis *AA = nullptr) const override;
 

llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.cpp

BuildMI(MBB, MI, DL, get(AMDGPU::V_CNDMASK_B32_e64), DstHi)
         .addOperand(SrcCond);
     MI->eraseFromParent();
     break;
   }
   }
   return true;
 }
 
-MachineInstr *SIInstrInfo::commuteInstruction(MachineInstr *MI,
-                                              bool NewMI) const {
+/// Commutes the operands in the given instruction.
+/// The commutable operands are specified by their indices OpIdx0 and OpIdx1.
+///
+/// Do not call this method for a non-commutable instruction or for
+/// non-commutable pair of operand indices OpIdx0 and OpIdx1.
+/// Even though the instruction is commutable, the method may still
+/// fail to commute the operands, null pointer is returned in such cases.
+MachineInstr *SIInstrInfo::commuteInstructionImpl(MachineInstr *MI,
+                                                  bool NewMI,
+                                                  unsigned OpIdx0,
+                                                  unsigned OpIdx1) const {
   int CommutedOpcode = commuteOpcode(*MI);
   if (CommutedOpcode == -1)
     return nullptr;
 
   int Src0Idx = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
                                            AMDGPU::OpName::src0);
   MachineOperand &Src0 = MI->getOperand(Src0Idx);
   if (!Src0.isReg())
     return nullptr;
 
   int Src1Idx = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
                                            AMDGPU::OpName::src1);
+
+  if ((OpIdx0 != static_cast<unsigned>(Src0Idx) ||
+       OpIdx1 != static_cast<unsigned>(Src1Idx)) &&
+      (OpIdx0 != static_cast<unsigned>(Src1Idx) ||
+       OpIdx1 != static_cast<unsigned>(Src0Idx)))
+    return nullptr;
+
   MachineOperand &Src1 = MI->getOperand(Src1Idx);
 
   // Make sure it's legal to commute operands for VOP2.
   if (isVOP2(MI->getOpcode()) &&
       (!isOperandLegal(MI, Src0Idx, &Src1) ||
        !isOperandLegal(MI, Src1Idx, &Src0))) {
     return nullptr;
   }
if (!Src1.isReg()) {
     if (Src1.isImm())
       Src0.ChangeToImmediate(Src1.getImm());
     else
       llvm_unreachable("Should only have immediates");
 
     Src1.ChangeToRegister(Reg, false);
     Src1.setSubReg(SubReg);
   } else {
-    MI = TargetInstrInfo::commuteInstruction(MI, NewMI);
+    MI = TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx0, OpIdx1);
   }
 
   if (MI)
     MI->setDesc(get(CommutedOpcode));
 
   return MI;
 }
 
 // This needs to be implemented because the source modifiers may be inserted
 // between the true commutable operands, and the base
 // TargetInstrInfo::commuteInstruction uses it.
 bool SIInstrInfo::findCommutedOpIndices(MachineInstr *MI,
-                                        unsigned &SrcOpIdx1,
-                                        unsigned &SrcOpIdx2) const {
+                                        unsigned &SrcOpIdx0,
+                                        unsigned &SrcOpIdx1) const {
   const MCInstrDesc &MCID = MI->getDesc();
   if (!MCID.isCommutable())
     return false;
 
   unsigned Opc = MI->getOpcode();
   int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
   if (Src0Idx == -1)
     return false;
 
   // FIXME: Workaround TargetInstrInfo::commuteInstruction asserting on
-  // immediate.
+  // immediate. Also, immediate src0 operand is not handled in
+  // SIInstrInfo::commuteInstruction();
   if (!MI->getOperand(Src0Idx).isReg())
     return false;
 
   int Src1Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src1);
   if (Src1Idx == -1)
     return false;
 
-  if (!MI->getOperand(Src1Idx).isReg())
+  MachineOperand &Src1 = MI->getOperand(Src1Idx);
+  if (Src1.isImm()) {
+    // SIInstrInfo::commuteInstruction() does support commuting the immediate
+    // operand src1 in 2 and 3 operand instructions.
+    if (!isVOP2(MI->getOpcode()) && !isVOP3(MI->getOpcode()))
       return false;
-
+  } else if (Src1.isReg()) {
     // If any source modifiers are set, the generic instruction commuting won't
     // understand how to copy the source modifiers.
     if (hasModifiersSet(*MI, AMDGPU::OpName::src0_modifiers) ||
         hasModifiersSet(*MI, AMDGPU::OpName::src1_modifiers))
       return false;
+  } else
+    return false;
 
-  SrcOpIdx1 = Src0Idx;
-  SrcOpIdx2 = Src1Idx;
-  return true;
+  return fixCommutedOpIndices(SrcOpIdx0, SrcOpIdx1, Src0Idx, Src1Idx);
 }
 
 MachineInstr *SIInstrInfo::buildMovInstr(MachineBasicBlock *MBB,
                                          MachineBasicBlock::iterator I,
                                          unsigned DstReg,
                                          unsigned SrcReg) const {
   return BuildMI(*MBB, I, MBB->findDebugLoc(I), get(AMDGPU::V_MOV_B32_e32),
                  DstReg) .addReg(SrcReg);

llvm/trunk/lib/Target/ARM/ARMBaseInstrInfo.h

protected:
   /// with the pair \p MI, \p DefIdx. False otherwise.
   ///
   /// \pre MI.isInsertSubregLike().
   bool
   getInsertSubregLikeInputs(const MachineInstr &MI, unsigned DefIdx,
                             RegSubRegPair &BaseReg,
                             RegSubRegPairAndIdx &InsertedReg) const override;
 
+  /// Commutes the operands in the given instruction.
+  /// The commutable operands are specified by their indices OpIdx1 and OpIdx2.
+  ///
+  /// Do not call this method for a non-commutable instruction or for
+  /// non-commutable pair of operand indices OpIdx1 and OpIdx2.
+  /// Even though the instruction is commutable, the method may still
+  /// fail to commute the operands, null pointer is returned in such cases.
+  MachineInstr *commuteInstructionImpl(MachineInstr *MI,
+                                       bool NewMI,
+                                       unsigned OpIdx1,
+                                       unsigned OpIdx2) const override;
+
 public:
   // Return whether the target has an explicit NOP encoding.
   bool hasNOP() const;
 
   // Return the non-pre/post incrementing version of 'Opc'. Return 0
   // if there is not such an opcode.
   virtual unsigned getUnindexedOpcode(unsigned Opc) const =0;
 
public:
   void reMaterialize(MachineBasicBlock &MBB, MachineBasicBlock::iterator MI,
                      unsigned DestReg, unsigned SubIdx,
                      const MachineInstr *Orig,
                      const TargetRegisterInfo &TRI) const override;
 
   MachineInstr *duplicate(MachineInstr *Orig,
                           MachineFunction &MF) const override;
 
-  MachineInstr *commuteInstruction(MachineInstr*,
-                                   bool=false) const override;
-
   const MachineInstrBuilder &AddDReg(MachineInstrBuilder &MIB, unsigned Reg,
                                      unsigned SubIdx, unsigned State,
                                      const TargetRegisterInfo *TRI) const;
 
   bool produceSameValue(const MachineInstr *MI0, const MachineInstr *MI1,
                         const MachineRegisterInfo *MRI) const override;
 
   /// areLoadsFromSameBasePtr - This is used by the pre-regalloc scheduler to

llvm/trunk/lib/Target/ARM/ARMBaseInstrInfo.cpp

unsigned llvm::getMatchingCondBranchOpcode(unsigned Opc) {
   if (Opc == ARM::tB)
     return ARM::tBcc;
   if (Opc == ARM::t2B)
     return ARM::t2Bcc;
 
   llvm_unreachable("Unknown unconditional branch opcode!");
 }
 
-/// commuteInstruction - Handle commutable instructions.
-MachineInstr *
-ARMBaseInstrInfo::commuteInstruction(MachineInstr *MI, bool NewMI) const {
+MachineInstr *ARMBaseInstrInfo::commuteInstructionImpl(MachineInstr *MI,
+                                                       bool NewMI,
+                                                       unsigned OpIdx1,
+                                                       unsigned OpIdx2) const {
   switch (MI->getOpcode()) {
   case ARM::MOVCCr:
   case ARM::t2MOVCCr: {
     // MOVCC can be commuted by inverting the condition.
     unsigned PredReg = 0;
     ARMCC::CondCodes CC = getInstrPredicate(MI, PredReg);
     // MOVCC AL can't be inverted. Shouldn't happen.
     if (CC == ARMCC::AL || PredReg != ARM::CPSR)
       return nullptr;
-    MI = TargetInstrInfo::commuteInstruction(MI, NewMI);
+    MI = TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
     if (!MI)
       return nullptr;
     // After swapping the MOVCC operands, also invert the condition.
     MI->getOperand(MI->findFirstPredOperandIdx())
       .setImm(ARMCC::getOppositeCondition(CC));
     return MI;
   }
   }
-  return TargetInstrInfo::commuteInstruction(MI, NewMI);
+  return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
 }
 
 /// Identify instructions that can be folded into a MOVCC instruction, and
 /// return the defining instruction.
 static MachineInstr *canFoldIntoMOVCC(unsigned Reg,
                                       const MachineRegisterInfo &MRI,
                                       const TargetInstrInfo *TII) {
   if (!TargetRegisterInfo::isVirtualRegister(Reg))

llvm/trunk/lib/Target/ARM/Thumb2SizeReduction.cpp

if (Reg0 != Reg2) {
         return false;
       // Try to commute the operands to make it a 2-address instruction.
       MachineInstr *CommutedMI = TII->commuteInstruction(MI);
       if (!CommutedMI)
         return false;
     }
   } else if (Reg0 != Reg1) {
     // Try to commute the operands to make it a 2-address instruction.
-    unsigned CommOpIdx1, CommOpIdx2;
+    unsigned CommOpIdx1 = 1;
+    unsigned CommOpIdx2 = TargetInstrInfo::CommuteAnyOperandIndex;
     if (!TII->findCommutedOpIndices(MI, CommOpIdx1, CommOpIdx2) ||
-        CommOpIdx1 != 1 || MI->getOperand(CommOpIdx2).getReg() != Reg0)
+        MI->getOperand(CommOpIdx2).getReg() != Reg0)
       return false;
-    MachineInstr *CommutedMI = TII->commuteInstruction(MI);
+    MachineInstr *CommutedMI =
+        TII->commuteInstruction(MI, false, CommOpIdx1, CommOpIdx2);
     if (!CommutedMI)
       return false;
   }
   if (Entry.LowRegs2 && !isARMLowRegister(Reg0))
     return false;
   if (Entry.Imm2Limit) {
     unsigned Imm = MI->getOperand(2).getImm();
     unsigned Limit = (1 << Entry.Imm2Limit) - 1;

llvm/trunk/lib/Target/PowerPC/PPCInstrInfo.h

bool StoreRegToStackSlot(MachineFunction &MF,
                            SmallVectorImpl<MachineInstr*> &NewMIs,
                            bool &NonRI, bool &SpillsVRS) const;
   bool LoadRegFromStackSlot(MachineFunction &MF, DebugLoc DL,
                             unsigned DestReg, int FrameIdx,
                             const TargetRegisterClass *RC,
                             SmallVectorImpl<MachineInstr*> &NewMIs,
                             bool &NonRI, bool &SpillsVRS) const;
   virtual void anchor();
+
+protected:
+  /// Commutes the operands in the given instruction.
+  /// The commutable operands are specified by their indices OpIdx1 and OpIdx2.
+  ///
+  /// Do not call this method for a non-commutable instruction or for
+  /// non-commutable pair of operand indices OpIdx1 and OpIdx2.
+  /// Even though the instruction is commutable, the method may still
+  /// fail to commute the operands, null pointer is returned in such cases.
+  ///
+  /// For example, we can commute rlwimi instructions, but only if the
+  /// rotate amt is zero.  We also have to munge the immediates a bit.
+  MachineInstr *commuteInstructionImpl(MachineInstr *MI,
+                                       bool NewMI,
+                                       unsigned OpIdx1,
+                                       unsigned OpIdx2) const override;
+
 public:
   explicit PPCInstrInfo(PPCSubtarget &STI);
 
   /// getRegisterInfo - TargetInstrInfo is a superset of MRegister info.  As
   /// such, whenever a client has an instance of instruction info, it should
   /// always be able to get register info as well (through this method).
   ///
   const PPCRegisterInfo &getRegisterInfo() const { return RI; }
public:
   bool isCoalescableExtInstr(const MachineInstr &MI,
                              unsigned &SrcReg, unsigned &DstReg,
                              unsigned &SubIdx) const override;
   unsigned isLoadFromStackSlot(const MachineInstr *MI,
                                int &FrameIndex) const override;
   unsigned isStoreToStackSlot(const MachineInstr *MI,
                               int &FrameIndex) const override;
 
-  // commuteInstruction - We can commute rlwimi instructions, but only if the
-  // rotate amt is zero.  We also have to munge the immediates a bit.
-  MachineInstr *commuteInstruction(MachineInstr *MI, bool NewMI) const override;
-
   bool findCommutedOpIndices(MachineInstr *MI, unsigned &SrcOpIdx1,
                              unsigned &SrcOpIdx2) const override;
 
   void insertNoop(MachineBasicBlock &MBB,
                   MachineBasicBlock::iterator MI) const override;
 
 
   // Branch analysis.

llvm/trunk/lib/Target/PowerPC/PPCInstrInfo.cpp

if (MI->getOperand(1).isImm() && !MI->getOperand(1).getImm() &&
       FrameIndex = MI->getOperand(2).getIndex();
       return MI->getOperand(0).getReg();
     }
     break;
   }
   return 0;
 }
 
-// commuteInstruction - We can commute rlwimi instructions, but only if the
-// rotate amt is zero.  We also have to munge the immediates a bit.
-MachineInstr *
-PPCInstrInfo::commuteInstruction(MachineInstr *MI, bool NewMI) const {
+MachineInstr *PPCInstrInfo::commuteInstructionImpl(MachineInstr *MI,
+                                                   bool NewMI,
+                                                   unsigned OpIdx1,
+                                                   unsigned OpIdx2) const {
   MachineFunction &MF = *MI->getParent()->getParent();
 
   // Normal instructions can be commuted the obvious way.
   if (MI->getOpcode() != PPC::RLWIMI &&
       MI->getOpcode() != PPC::RLWIMIo)
-    return TargetInstrInfo::commuteInstruction(MI, NewMI);
+    return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
   // Note that RLWIMI can be commuted as a 32-bit instruction, but not as a
   // 64-bit instruction (so we don't handle PPC::RLWIMI8 here), because
   // changing the relative order of the mask operands might change what happens
   // to the high-bits of the mask (and, thus, the result).
 
   // Cannot commute if it has a non-zero rotate count.
   if (MI->getOperand(3).getImm() != 0)
     return nullptr;
 
   // If we have a zero rotate count, we have:
   //   M = mask(MB,ME)
   //   Op0 = (Op1 & ~M) | (Op2 & M)
   // Change this to:
   //   M = mask((ME+1)&31, (MB-1)&31)
   //   Op0 = (Op2 & ~M) | (Op1 & M)
 
   // Swap op1/op2
+  assert(((OpIdx1 == 1 && OpIdx2 == 2) || (OpIdx1 == 2 && OpIdx2 == 1)) &&
+         "Only the operands 1 and 2 can be swapped in RLSIMI/RLWIMIo.");
   unsigned Reg0 = MI->getOperand(0).getReg();
   unsigned Reg1 = MI->getOperand(1).getReg();
   unsigned Reg2 = MI->getOperand(2).getReg();
   unsigned SubReg1 = MI->getOperand(1).getSubReg();
   unsigned SubReg2 = MI->getOperand(2).getSubReg();
   bool Reg1IsKill = MI->getOperand(1).isKill();
   bool Reg2IsKill = MI->getOperand(2).isKill();
   bool ChangeReg0 = false;
bool PPCInstrInfo::findCommutedOpIndices(MachineInstr *MI, unsigned &SrcOpIdx1,
   // For VSX A-Type FMA instructions, it is the first two operands that can be
   // commuted, however, because the non-encoded tied input operand is listed
   // first, the operands to swap are actually the second and third.
 
   int AltOpc = PPC::getAltVSXFMAOpcode(MI->getOpcode());
   if (AltOpc == -1)
     return TargetInstrInfo::findCommutedOpIndices(MI, SrcOpIdx1, SrcOpIdx2);
 
-  SrcOpIdx1 = 2;
-  SrcOpIdx2 = 3;
-  return true;
+  // The commutable operand indices are 2 and 3. Return them in SrcOpIdx1
+  // and SrcOpIdx2.
+  return fixCommutedOpIndices(SrcOpIdx1, SrcOpIdx2, 2, 3);
 }
 
 void PPCInstrInfo::insertNoop(MachineBasicBlock &MBB,
                               MachineBasicBlock::iterator MI) const {
   // This function is used for scheduling, and the nop wanted here is the type
   // that terminates dispatch groups on the POWER cores.
   unsigned Directive = Subtarget.getDarwinDirective();
   unsigned Opcode;

llvm/trunk/lib/Target/X86/X86InstrInfo.h

public:
   ///
   /// This method returns a null pointer if the transformation cannot be
   /// performed, otherwise it returns the new instruction.
   ///
   MachineInstr *convertToThreeAddress(MachineFunction::iterator &MFI,
                                       MachineBasicBlock::iterator &MBBI,
                                       LiveVariables *LV) const override;
 
-  /// commuteInstruction - We have a few instructions that must be hacked on to
-  /// commute them.
+  /// Returns true iff the routine could find two commutable operands in the
+  /// given machine instruction.
+  /// The 'SrcOpIdx1' and 'SrcOpIdx2' are INPUT and OUTPUT arguments. Their
+  /// input values can be re-defined in this method only if the input values
+  /// are not pre-defined, which is designated by the special value
+  /// 'CommuteAnyOperandIndex' assigned to it.
+  /// If both of indices are pre-defined and refer to some operands, then the
+  /// method simply returns true if the corresponding operands are commutable
+  /// and returns false otherwise.
   ///
-  MachineInstr *commuteInstruction(MachineInstr *MI, bool NewMI) const override;
-
+  /// For example, calling this method this way:
+  ///     unsigned Op1 = 1, Op2 = CommuteAnyOperandIndex;
+  ///     findCommutedOpIndices(MI, Op1, Op2);
+  /// can be interpreted as a query asking to find an operand that would be
+  /// commutable with the operand#1.
   bool findCommutedOpIndices(MachineInstr *MI, unsigned &SrcOpIdx1,
                              unsigned &SrcOpIdx2) const override;
 
   // Branch analysis.
   bool isUnpredicatedTerminator(const MachineInstr* MI) const override;
   bool AnalyzeBranch(MachineBasicBlock &MBB, MachineBasicBlock *&TBB,
                      MachineBasicBlock *&FBB,
                      SmallVectorImpl<MachineOperand> &Cond,
MachineInstr* optimizeLoadInstr(MachineInstr *MI,
                                   MachineInstr *&DefMI) const override;
 
   std::pair<unsigned, unsigned>
   decomposeMachineOperandsTargetFlags(unsigned TF) const override;
 
   ArrayRef<std::pair<unsigned, const char *>>
   getSerializableDirectMachineOperandTargetFlags() const override;
 
+protected:
+  /// Commutes the operands in the given instruction by changing the operands
+  /// order and/or changing the instruction's opcode and/or the immediate value
+  /// operand.
+  ///
+  /// The arguments 'CommuteOpIdx1' and 'CommuteOpIdx2' specify the operands
+  /// to be commuted.
+  ///
+  /// Do not call this method for a non-commutable instruction or
+  /// non-commutable operands.
+  /// Even though the instruction is commutable, the method may still
+  /// fail to commute the operands, null pointer is returned in such cases.
+  MachineInstr *commuteInstructionImpl(MachineInstr *MI, bool NewMI,
+                                       unsigned CommuteOpIdx1,
+                                       unsigned CommuteOpIdx2) const override;
+
 private:
   MachineInstr * convertToThreeAddressWithLEA(unsigned MIOpc,
                                               MachineFunction::iterator &MFI,
                                               MachineBasicBlock::iterator &MBBI,
                                               LiveVariables *LV) const;
 
   /// isFrameOperand - Return true and the FrameIndex if the specified
   /// operand and follow operands form a reference to the stack frame.
   bool isFrameOperand(const MachineInstr *MI, unsigned int Op,
                       int &FrameIndex) const;
 };
 
 } // End llvm namespace
 
 #endif

llvm/trunk/lib/Target/X86/X86InstrInfo.cpp

if (LV) {  // Update live variables
     if (Dest.isDead())
       LV->replaceKillInstruction(Dest.getReg(), MI, NewMI);
   }
 
   MFI->insert(MBBI, NewMI);          // Insert the new inst
   return NewMI;
 }
 
-/// We have a few instructions that must be hacked on to commute them.
-///
-MachineInstr *
-X86InstrInfo::commuteInstruction(MachineInstr *MI, bool NewMI) const {
+MachineInstr *X86InstrInfo::commuteInstructionImpl(MachineInstr *MI,
+                                                   bool NewMI,
+                                                   unsigned OpIdx1,
+                                                   unsigned OpIdx2) const {
   switch (MI->getOpcode()) {
   case X86::SHRD16rri8: // A = SHRD16rri8 B, C, I -> A = SHLD16rri8 C, B, (16-I)
   case X86::SHLD16rri8: // A = SHLD16rri8 B, C, I -> A = SHRD16rri8 C, B, (16-I)
   case X86::SHRD32rri8: // A = SHRD32rri8 B, C, I -> A = SHLD32rri8 C, B, (32-I)
   case X86::SHLD32rri8: // A = SHLD32rri8 B, C, I -> A = SHRD32rri8 C, B, (32-I)
   case X86::SHRD64rri8: // A = SHRD64rri8 B, C, I -> A = SHLD64rri8 C, B, (64-I)
   case X86::SHLD64rri8:{// A = SHLD64rri8 B, C, I -> A = SHRD64rri8 C, B, (64-I)
     unsigned Opc;
case X86::SHLD64rri8:{// A = SHLD64rri8 B, C, I -> A = SHRD64rri8 C, B, (64-I)
     unsigned Amt = MI->getOperand(3).getImm();
     if (NewMI) {
       MachineFunction &MF = *MI->getParent()->getParent();
       MI = MF.CloneMachineInstr(MI);
       NewMI = false;
     }
     MI->setDesc(get(Opc));
     MI->getOperand(3).setImm(Size-Amt);
-    return TargetInstrInfo::commuteInstruction(MI, NewMI);
+    return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
   }
   case X86::BLENDPDrri:
   case X86::BLENDPSrri:
   case X86::PBLENDWrri:
   case X86::VBLENDPDrri:
   case X86::VBLENDPSrri:
   case X86::VBLENDPDYrri:
   case X86::VBLENDPSYrri:
case X86::VPBLENDWYrri:{
     // Only the least significant bits of Imm are used.
     unsigned Imm = MI->getOperand(3).getImm() & Mask;
     if (NewMI) {
       MachineFunction &MF = *MI->getParent()->getParent();
       MI = MF.CloneMachineInstr(MI);
       NewMI = false;
     }
     MI->getOperand(3).setImm(Mask ^ Imm);
-    return TargetInstrInfo::commuteInstruction(MI, NewMI);
+    return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
   }
   case X86::PCLMULQDQrr:
   case X86::VPCLMULQDQrr:{
     // SRC1 64bits = Imm[0] ? SRC1[127:64] : SRC1[63:0]
     // SRC2 64bits = Imm[4] ? SRC2[127:64] : SRC2[63:0]
     unsigned Imm = MI->getOperand(3).getImm();
     unsigned Src1Hi = Imm & 0x01;
     unsigned Src2Hi = Imm & 0x10;
     if (NewMI) {
       MachineFunction &MF = *MI->getParent()->getParent();
       MI = MF.CloneMachineInstr(MI);
       NewMI = false;
     }
     MI->getOperand(3).setImm((Src1Hi << 4) | (Src2Hi >> 4));
-    return TargetInstrInfo::commuteInstruction(MI, NewMI);
+    return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
   }
   case X86::CMPPDrri:
   case X86::CMPPSrri:
   case X86::VCMPPDrri:
   case X86::VCMPPSrri:
   case X86::VCMPPDYrri:
   case X86::VCMPPSYrri: {
     // Float comparison can be safely commuted for
     // Ordered/Unordered/Equal/NotEqual tests
     unsigned Imm = MI->getOperand(3).getImm() & 0x7;
     switch (Imm) {
     case 0x00: // EQUAL
     case 0x03: // UNORDERED
     case 0x04: // NOT EQUAL
     case 0x07: // ORDERED
       if (NewMI) {
         MachineFunction &MF = *MI->getParent()->getParent();
         MI = MF.CloneMachineInstr(MI);
         NewMI = false;
       }
-      return TargetInstrInfo::commuteInstruction(MI, NewMI);
+      return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
     default:
       return nullptr;
     }
   }
   case X86::VPCOMBri: case X86::VPCOMUBri:
   case X86::VPCOMDri: case X86::VPCOMUDri:
   case X86::VPCOMQri: case X86::VPCOMUQri:
   case X86::VPCOMWri: case X86::VPCOMUWri: {
default:
       break;
     }
     if (NewMI) {
       MachineFunction &MF = *MI->getParent()->getParent();
       MI = MF.CloneMachineInstr(MI);
       NewMI = false;
     }
     MI->getOperand(3).setImm(Imm);
-    return TargetInstrInfo::commuteInstruction(MI, NewMI);
+    return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
   }
   case X86::CMOVB16rr:  case X86::CMOVB32rr:  case X86::CMOVB64rr:
   case X86::CMOVAE16rr: case X86::CMOVAE32rr: case X86::CMOVAE64rr:
   case X86::CMOVE16rr:  case X86::CMOVE32rr:  case X86::CMOVE64rr:
   case X86::CMOVNE16rr: case X86::CMOVNE32rr: case X86::CMOVNE64rr:
   case X86::CMOVBE16rr: case X86::CMOVBE32rr: case X86::CMOVBE64rr:
   case X86::CMOVA16rr:  case X86::CMOVA32rr:  case X86::CMOVA64rr:
   case X86::CMOVL16rr:  case X86::CMOVL32rr:  case X86::CMOVL64rr:
if (NewMI) {
       MachineFunction &MF = *MI->getParent()->getParent();
       MI = MF.CloneMachineInstr(MI);
       NewMI = false;
     }
     MI->setDesc(get(Opc));
     // Fallthrough intended.
   }
   default:
-    return TargetInstrInfo::commuteInstruction(MI, NewMI);
+    return TargetInstrInfo::commuteInstructionImpl(MI, NewMI, OpIdx1, OpIdx2);
   }
 }
 
-bool X86InstrInfo::findCommutedOpIndices(MachineInstr *MI, unsigned &SrcOpIdx1,
+bool X86InstrInfo::findCommutedOpIndices(MachineInstr *MI,
+                                         unsigned &SrcOpIdx1,
                                          unsigned &SrcOpIdx2) const {
   switch (MI->getOpcode()) {
     case X86::CMPPDrri:
     case X86::CMPPSrri:
     case X86::VCMPPDrri:
     case X86::VCMPPSrri:
     case X86::VCMPPDYrri:
     case X86::VCMPPSYrri: {
       // Float comparison can be safely commuted for
       // Ordered/Unordered/Equal/NotEqual tests
       unsigned Imm = MI->getOperand(3).getImm() & 0x7;
       switch (Imm) {
         case 0x00: // EQUAL
         case 0x03: // UNORDERED
         case 0x04: // NOT EQUAL
         case 0x07: // ORDERED
-        SrcOpIdx1 = 1;
-        SrcOpIdx2 = 2;
-        return true;
+          // The indices of the commutable operands are 1 and 2.
+          // Assign them to the returned operand indices here.
+          return fixCommutedOpIndices(SrcOpIdx1, SrcOpIdx2, 1, 2);
       }
       return false;
     }
     case X86::VFMADDPDr231r:
     case X86::VFMADDPSr231r:
     case X86::VFMADDSDr231r:
     case X86::VFMADDSSr231r:
     case X86::VFMSUBPDr231r:
switch (MI->getOpcode()) {
     case X86::VFMADDPDr231rY:
     case X86::VFMADDPSr231rY:
     case X86::VFMSUBPDr231rY:
     case X86::VFMSUBPSr231rY:
     case X86::VFNMADDPDr231rY:
     case X86::VFNMADDPSr231rY:
     case X86::VFNMSUBPDr231rY:
     case X86::VFNMSUBPSr231rY:
-      SrcOpIdx1 = 2;
-      SrcOpIdx2 = 3;
-      return true;
+          // The indices of the commutable operands are 2 and 3.
+          // Assign them to the returned operand indices here.
+          return fixCommutedOpIndices(SrcOpIdx1, SrcOpIdx2, 2, 3);
     default:
       return TargetInstrInfo::findCommutedOpIndices(MI, SrcOpIdx1, SrcOpIdx2);
   }
+  return false;
 }
 
 static X86::CondCode getCondFromBranchOpc(unsigned BrOpc) {
   switch (BrOpc) {
   default: return X86::COND_INVALID;
   case X86::JE_1:  return X86::COND_E;
   case X86::JNE_1: return X86::COND_NE;
   case X86::JL_1:  return X86::COND_L;
if (I != OpcodeTablePtr->end()) {
       }
       return NewMI;
     }
   }
 
   // If the instruction and target operand are commutable, commute the
   // instruction and try again.
   if (AllowCommute) {
-    unsigned OriginalOpIdx = OpNum, CommuteOpIdx1, CommuteOpIdx2;
+    unsigned CommuteOpIdx1 = OpNum, CommuteOpIdx2 = CommuteAnyOperandIndex;
     if (findCommutedOpIndices(MI, CommuteOpIdx1, CommuteOpIdx2)) {
       bool HasDef = MI->getDesc().getNumDefs();
       unsigned Reg0 = HasDef ? MI->getOperand(0).getReg() : 0;
       unsigned Reg1 = MI->getOperand(CommuteOpIdx1).getReg();
       unsigned Reg2 = MI->getOperand(CommuteOpIdx2).getReg();
-      bool Tied0 =
-          0 == MI->getDesc().getOperandConstraint(CommuteOpIdx1, MCOI::TIED_TO);
       bool Tied1 =
+          0 == MI->getDesc().getOperandConstraint(CommuteOpIdx1, MCOI::TIED_TO);
+      bool Tied2 =
           0 == MI->getDesc().getOperandConstraint(CommuteOpIdx2, MCOI::TIED_TO);
 
       // If either of the commutable operands are tied to the destination
       // then we can not commute + fold.
-      if ((HasDef && Reg0 == Reg1 && Tied0) ||
-          (HasDef && Reg0 == Reg2 && Tied1))
+      if ((HasDef && Reg0 == Reg1 && Tied1) ||
+          (HasDef && Reg0 == Reg2 && Tied2))
         return nullptr;
 
-      if ((CommuteOpIdx1 == OriginalOpIdx) ||
-          (CommuteOpIdx2 == OriginalOpIdx)) {
-        MachineInstr *CommutedMI = commuteInstruction(MI, false);
-        if (!CommutedMI) {
-          // Unable to commute.
-          return nullptr;
-        }
-        if (CommutedMI != MI) {
-          // New instruction. We can't fold from this.
-          CommutedMI->eraseFromParent();
-          return nullptr;
-        }
+      MachineInstr *CommutedMI =
+          commuteInstruction(MI, false, CommuteOpIdx1, CommuteOpIdx2);
 
-        // Attempt to fold with the commuted version of the instruction.
-        unsigned CommuteOp =
-            (CommuteOpIdx1 == OriginalOpIdx ? CommuteOpIdx2 : CommuteOpIdx1);
-        NewMI =
-            foldMemoryOperandImpl(MF, MI, CommuteOp, MOs, InsertPt, Size, Align,
-                                  /*AllowCommute=*/false);
-        if (NewMI)
-          return NewMI;
+      NewMI = foldMemoryOperandImpl(MF, MI, CommuteOpIdx2, MOs, InsertPt,
+                                    Size, Align, /*AllowCommute=*/false);
 
-        // Folding failed again - undo the commute before returning.
-        MachineInstr *UncommutedMI = commuteInstruction(MI, false);
-        if (!UncommutedMI) {
-          // Unable to commute.
-          return nullptr;
-        }
-        if (UncommutedMI != MI) {
-          // New instruction. It doesn't need to be kept.
-          UncommutedMI->eraseFromParent();
-          return nullptr;
-        }
+      MachineInstr *UncommutedMI =
+          commuteInstruction(MI, false, CommuteOpIdx1, CommuteOpIdx2);
 
-        // Return here to prevent duplicate fuse failure report.
-        return nullptr;
-      }
-    }
-  }
 
   // No fusion
   if (PrintFailedFusing && !MI->isCopy())
     dbgs() << "We failed to fuse operand " << OpNum << " in " << *MI;
   return nullptr;
 }
 
 /// Return true for all instructions that only update