LGTM

fp64-atomics-gfx90a.ll is failing.

The code as written does not guarantee it will increase the occupancy if a def was sunk. Moreover it will sink defs even in regions with low RP. Essentially it may easily pessimise the code without improving anything. You really need to check that sinking is beneficial.

• Check for calls along with inline asm not relying on function attribute.
• Moved -global-isel flag in tests.

Any tests?

LGTM

Same thing is possible even with calls but requires a full call graph examination. If annotate kernel features pass will move to attributer it shall be possible to analyze inline asm in calls and propagate this info.

In D117189#3241074, @arsenm wrote:

Testcase?

Fixed tuple parsing as well.
Removed w/a in 2 places which were using it.

In D117057#3236020, @arsenm wrote:

In D117057#3235661, @rampitec wrote:

In D117057#3235626, @rampitec wrote:

I guess I can run this code in the AMDGPUDAGToDAGISel::runOnMachineFunction or even SIMachineFunctionInfo constructor and skip the attribution.

The downside is that I have to scan all instructions for that, and AMDGPUAnnotateKernelFeatures already does that. Where are you planning to move its code?

AMDGPUAttributor. If this were to be an attribute, which is pretty ugly, it should be the inverse. Why can't you just select to AGPRs, and later we can adjust the register classes if necessary?

In D117057#3235626, @rampitec wrote:

I guess I can run this code in the AMDGPUDAGToDAGISel::runOnMachineFunction or even SIMachineFunctionInfo constructor and skip the attribution.

I guess I can run this code in the AMDGPUDAGToDAGISel::runOnMachineFunction or even SIMachineFunctionInfo constructor and skip the attribution.

In D117057#3235597, @arsenm wrote:

Why does this specifically need to check for inline asm? Is this only checked in the kernel or functions too?

I'm trying to delete this pass and don't want to add more stuff to it. If it's just for kernels can't you check for AGPR physreg uses later?

In D116819#3231986, @arsenm wrote:

I think this is an extreme interpretation of optnone. This is a minor optimization which happens as part of lowering. The fact we do this as a separate step is just an artifact of how we happen to lower control flow. It's not strictly true that no optimizations occur at -O0, especially if they are cheap and provide benefit

In D114126#3227404, @spatel wrote:

If we removed some strictness about uses, we probably wouldn't require this patch or the follow-ons for complicated bitwise logic optimizations. Those are only needed because we miss the simpler, intermediate folds in complex expressions with extra uses. I suppose we'd be more at risk for infinite combine loops if we allowed creating extra instructions, but it might be worth experimenting with removing some use checks on bitwise logic folds to see how that plays out.

Rebased.

It seems to cause a real regression, block-should-not-be-in-alive-blocks.mir fails.

In D116270#3224818, @alex-t wrote:

In D116270#3215217, @rampitec wrote:

In D116270#3209307, @alex-t wrote:

This looks like a regression in xnor.ll :

	s_not_b32 s0, s0                        	v_not_b32_e32 v0, v0
	v_xor_b32_e32 v0, s0, v0                        v_xor_b32_e32 v0, s4, v0

but it is not really. All the nodes in the example are divergent and the divergent ( xor, x -1) is selected to V_NOT_B32 as of https://reviews.llvm.org/D115884 has been committed.
S_NOT_B32 appears at the left because of the custom optimization that converts S_XNOR_B32 back to NOT (XOR) for the targets which have no V_XNOR. This optimization relies on the fact that if the NOT operand is SGPR and V_XOR_B32_e32 can accept SGPR as a first source operand.
I am not sure if it is always safe. The VALU instructions execution is controlled by the EXEC mask but SALU is not.

This is indeed a regression. It is always safe to keep s_not_b32 on SALU. Also note this effectively makes SIInstrInfo::lowerScalarXnor() useless. This is why XNOR was left behind by the D111907.

SIInstrInfo::lowerScalarXnor() is exactly the part of the "manual" SALU to VALU lowering that I am trying to get rid of.
The divergent "not" must be selected to the "V_NOT_B32_e32/64" otherwise we still have illegal VGPR to SGPR copies.
This happens because the divergent "not" node has divergent operands and their result will be likely in VGPR.
Also, we should select everything correctly first and can apply some peephole optimizations after.
In other words: we should not "cheat ourselves" during the selection. The selection should be done fairly corresponding to the node divergence bit.
Then we can apply the optimization in case it is safe.
Note that this is not the only case when we would like to further optimize the code after selection.
I'm planning to further add a separate pass for that.

We cannot solve the problem in the custom selection procedure because NOT node operand has not yet been selected and we do not know if it is SGPR or VGPR.
The only way, for now, is to post-process not(xor)/xor(not) in SIFixSGPRCopies. This may be considered a temporary hack until we have no proper pass for that.

In D116231#3225213, @spatel wrote:

In D116231#3224577, @foad wrote:

I don't really have an opinion on the patch, but I'm curious.

Given (~a & ~b & c) | (~a & ~c & b), do we make any attempt to pull out the ~a like this: ~a & ((~b & c) | (~c & b)) ? If not, why not? Would that be a good thing to address? Then it's just a case of simplifying the two-variable (~b & c) | (~c & b) -> b ^ c which is easy.

InstCombine does try that transform via InstCombinerImpl::SimplifyAssociativeOrCommutative() and/or InstCombinerImpl::SimplifyUsingDistributiveLaws(). And that works as expected - for example if we alter the 1st modified test in this patch to be like this:

define i32 @not_and_not_and_and_or(i32 %a, i32 %b, i32 %c) {
  %nota = xor i32 %a, -1
  %notb = xor i32 %b, -1
  %and1 = and i32 %nota, %c
  %and2 = and i32 %and1, %notb
  %or1 = or i32 %a, %c
  %not1 = xor i32 %or1, -1
  %and3 = and i32 %not1, %b
  %or3 = or i32 %and2, %and3
  call void @use(i32 %nota)
  call void @use(i32 %notb)
  ret i32 %or3

Then it reduces because we find the 'b' and '~b' values directly in the operands of the 'and' instructions. So this might be a question for "-reassociate" - can we get that pass to arrange the operands such that -instcombine can fold this (without breaking some other pattern)?

In D116194#3225474, @spatel wrote:

In D116194#3223711, @rampitec wrote:

In D116194#3223354, @spatel wrote:

Making the caller pass a dummy seems awkward.
I don't have any suggestions to improve it, but adding some more potential reviewers.

I can pass a pointer with default null. Is that better?

Wouldn't it be a pointer to a pointer in that case?

In D116194#3223354, @spatel wrote:

Making the caller pass a dummy seems awkward.
I don't have any suggestions to improve it, but adding some more potential reviewers.

ping

ping

ping

As discussed it would be worth to do a separate change to skip the optimization at -O0 alltogether.

LGTM

In D116270#3209307, @alex-t wrote:

This looks like a regression in xnor.ll :

	s_not_b32 s0, s0                        	v_not_b32_e32 v0, v0
	v_xor_b32_e32 v0, s0, v0                        v_xor_b32_e32 v0, s4, v0

but it is not really. All the nodes in the example are divergent and the divergent ( xor, x -1) is selected to V_NOT_B32 as of https://reviews.llvm.org/D115884 has been committed.
S_NOT_B32 appears at the left because of the custom optimization that converts S_XNOR_B32 back to NOT (XOR) for the targets which have no V_XNOR. This optimization relies on the fact that if the NOT operand is SGPR and V_XOR_B32_e32 can accept SGPR as a first source operand.
I am not sure if it is always safe. The VALU instructions execution is controlled by the EXEC mask but SALU is not.

LGTM

A side note: this small change fully or partially simplifies 28 out of original 255 ternary cases.

Fixed formatting.
It will go under the same LHS as above after D116231 (plus it will need some more tests in this case).

Thanks! LGTM.

LGTM

LGTM given the new tests.

I will precommit tests if that is generally OK.

In D116053#3205477, @vangthao wrote:

IR is essentially a single thread representation. The implicit exec use is our way to model mutithreaded divergence. Consider this transformation which shall now become legal:

int lid = get_local_id(0);      int lid = get_local_id(0);
int i = 0;                      int i = 0;
x = def();                      do {
do {                        =>    x = def();
  use1(x);                        use1(x);
} while(i++ < lid);             } while(i++ < lid);
use2(x);                        use2(x);

def dominates use2 in both cases, but in the second case not with every lane. All lanes except first will use an undef.

We will not sink into a loop if the def is outside of the loop. In the test case loop_sink_fmac, the def was already in a loop and was why it was able to be sinked.

In D116053#3205007, @vangthao wrote:

In D116053#3203538, @rampitec wrote:

I am not sure I can prove to myself this is legal. For example you are sinking a def into a loop with divergent condition and this def is used after the loop. Can this happen? If so a def might be done with an exec smaller than a use which creates an undef. Hoisting was OK because def was moved into a direction where exec is strictly not less than before. Did you run PSDB on it?

I believe the requirement for defs to dominate all uses prevents this from happening but I can add more test cases to check for this. This passed PSDB.

There are still 2 comments around tests.

I am not sure I can prove to myself this is legal. For example you are sinking a def into a loop with divergent condition and this def is used after the loop. Can this happen? If so a def might be done with an exec smaller than a use which creates an undef. Hoisting was OK because def was moved into a direction where exec is strictly not less than before. Did you run PSDB on it?

In D115996#3203263, @bcahoon wrote:

In D115996#3203119, @rampitec wrote:

Will the slot remain in the final allocation?

Yes. At least in the example I'm seeing.

Will the slot remain in the final allocation?

In D115996#3202932, @arsenm wrote:

Can we rip out this special "spill" handling since D109301 @cdevadas

In D115880#3199545, @abinavpp wrote:

In D115880#3198240, @rampitec wrote:

I do not believe really needed for D115881.

This is to show the change in SIInsertWaitcnts for local-atomics-fp.ll in D115881 at the moment. If we fix the problems in D115881, hopefully the regression should go away. In that case, this change is a mere test generation.

LGTM

LGTM

It also obscures what is being tested and I do not believe really needed for D115881.

Rebased.
Added handling of newly added case (~A & B & C) | ~(A | B) --> (C | ~B) & ~A.

LGTM, thank you!

Ping

In D115413#3190617, @JonChesterfield wrote:

I am doubtful that this would be the correct strategy for initializers of LDS - I suspect that's going to have to be handled by injecting initialisation code into the prologue of kernels and changing the initialisers to undef - but even if it is, we can retrieve it from git at will. Right now it makes callgraph crawling to avoid unnecessary allocation of LDS much more complicated (or imprecise).

I definitely agree that LDS handling is under-tested at present but there's no value I can see to these specific tests. They'll need to be changed to be more useful at which point we're better off deciding what it is we want to check and writing tests that hit that.

LGTM

LGTM