PING for review

Besides the unpromotable alloca issue due to indirect accesses, such coercion to GLOBAL pointer directly is not safe as, in HIP/CUDA, both CONSTANT and GLOBAL pointers would be passed as the kernel arguments. Without introducing a new address space combing GLOBAL/CONSTANT, such coercion would be unsafe.

Fix coding style following clang-tidy.

Add amdgpu-kernel-arg-pointer-type.cu back and revise its checks.

Revise the comment and point the safety issue by coercing the kernel argument
from a generic pointer to a global one.

In D89980#2348339, @tra wrote:

Are there any tests to illustrate what this change does to IR or generated code?

Test case is enhanced to check that no kernel argument type is coerced.

Fix more regression tests due to the enhanced AMDGPU AA.

Remove unordered check.

Rebase

In D89447#2351887, @dmgreen wrote:

Has the SystremZ Prefetch issues been resolved?

Rebase

PING for review. As a similar patch was approved, shall I just commit it again with the compilation time issue is addressed?

Just kingly PING for review.

Add an option to turn on/off the use of AA during codegen.

This patch is still required as MMO's alignment is calculated based on the offset from the base alignment. As the base alignment is the alignment from the pointer in the IR, it cannot be modified. We need extra logic to re-align MMO operand if we widen the original one. For instance of a 16-bit load from ptr has an alignment of 2, if ptr is equivalent to base - 2 and base's alignment is 4, we could widen that 16-bit load to 32-bit load from ptr - 2with an alignment 4. But, as we cannot change IR in MMO, we need extra stuff to in the new MMO could assume that new alignment.

Rebase and revise.

PING for review.

In D89525#2336485, @rampitec wrote:

In D89525#2336479, @hliao wrote:

In D89525#2336008, @rampitec wrote:

In D89525#2336002, @yaxunl wrote:

In D89525#2335991, @rampitec wrote:

In D89525#2335989, @yaxunl wrote:

In D89525#2335972, @rampitec wrote:

In D89525#2335949, @yaxunl wrote:

I think they are correct for OpenCL, since in OpenCL shared var can only be declared in kernel function or passed by kernel arg.

However I am not sure whether a constant pointer can pointer to shared memory, i.e, whether the address of a shared variable is compile time constant, or whether the following is valid code:

__shared__ int a;

__constant__ int *b = &a;

Currently clang allows it but nvcc does not https://godbolt.org/z/9W8vee

I tends to agree with nvcc's treatment since this allows more flexible way of implementing shared variable supports in backend. @tra for advice

But you are not checking for a constant pointer here!

In HIP __constant__ is a variable attribute, not the address space of the pointee. __constant__ int * means a pointer itself in constant address space and pointing to generic/flat address space.

Where do you check for this specifically in this block:

} else if (const Argument *Arg = dyn_cast<Argument>(ObjA)) {
   const Function *F = Arg->getParent();
   switch (F->getCallingConv()) {
   case CallingConv::AMDGPU_KERNEL:
     // In the kernel function, kernel arguments won't alias to (local)
     // variables in shared or private address space.
     return NoAlias;

I was talking about semantic check in language. Here is the IR. In IR a kernel arg can pointing to constant or global addr due to promotion. Originally all kernel arg of HIP points to generic addr space only.

But not in OpenCL.

For OpenCL, since it won't allow generic pointer as kernel function arguments, there never be such a case, a generic pointer argument.

OK, that makes sense. Now we only need to make sure we never compile anything but HIP and OpenCL <= 2.0. For example that we will not support fortran maybe? And never extend any of the languages to allow different address space arguments.

My point here is that it is language specific, but there is nothing language specific on the AA implementation. A possible solution is to define some attributes or metadata saying that a certain situation may not ever happen and then check it in AA. That property has to ve set by the language though as AA may not know the source semantics.

In D89447#2336699, @jonpa wrote:

I removed that check that one of those two MMOs needs to be store. No change to test now. But, that sounds a little weird that one instruction claims mayLoad and mayStore but only has isLoad MMO. Maybe we need to change that MMO in prefetch to be both isLoad and isStore.

The SystemZ Prefetch Data (PFD) instruction has a flag bit to control if the prefetch is read or write, so in a way those two flags make sense. However, since the PFD does not in fact clobber any memory I think the idea behind the mayLoad and mayStore flags is to keep the instruction in place in the block as much as possible. It shouldn't necessarily matter, but generally it is probably better spread out the memory accesses / prefetches rather than having them all happen at once, I would think. If the block is big with many prefetches they shouldn't all end up in the end of the block...

If you think that check is valuable, then maybe we could try to find another way to keep the PFD instructions in their places. I tried adding hasSideEffects = 1 on PFD/PFDRL instructions instead of removing the check on the MMOs, but that did not seem to be NFC on benchmarks...

In D89525#2336008, @rampitec wrote:

In D89525#2336002, @yaxunl wrote:

In D89525#2335991, @rampitec wrote:

In D89525#2335989, @yaxunl wrote:

In D89525#2335972, @rampitec wrote:

In D89525#2335949, @yaxunl wrote:

I think they are correct for OpenCL, since in OpenCL shared var can only be declared in kernel function or passed by kernel arg.

However I am not sure whether a constant pointer can pointer to shared memory, i.e, whether the address of a shared variable is compile time constant, or whether the following is valid code:

__shared__ int a;

__constant__ int *b = &a;

Currently clang allows it but nvcc does not https://godbolt.org/z/9W8vee

I tends to agree with nvcc's treatment since this allows more flexible way of implementing shared variable supports in backend. @tra for advice

But you are not checking for a constant pointer here!

In HIP __constant__ is a variable attribute, not the address space of the pointee. __constant__ int * means a pointer itself in constant address space and pointing to generic/flat address space.

Where do you check for this specifically in this block:

} else if (const Argument *Arg = dyn_cast<Argument>(ObjA)) {
   const Function *F = Arg->getParent();
   switch (F->getCallingConv()) {
   case CallingConv::AMDGPU_KERNEL:
     // In the kernel function, kernel arguments won't alias to (local)
     // variables in shared or private address space.
     return NoAlias;

I was talking about semantic check in language. Here is the IR. In IR a kernel arg can pointing to constant or global addr due to promotion. Originally all kernel arg of HIP points to generic addr space only.

But not in OpenCL.

In D89525#2335949, @yaxunl wrote:

In D89525#2333924, @hliao wrote:

@yaxunl could you double-check that OpenCL also follows that rule.
@nhaehnle could you check whether that potentially breaks graphics.

I think they are correct for OpenCL, since in OpenCL shared var can only be declared in kernel function or passed by kernel arg.

However I am not sure whether a constant pointer can pointer to shared memory, i.e, whether the address of a shared variable is compile time constant, or whether the following is valid code:

__shared__ int a;

__constant__ int *b = &a;

Currently clang allows it but nvcc does not https://godbolt.org/z/9W8vee

I tends to agree with nvcc's treatment since this allows more flexible way of implementing shared variable supports in backend. @tra for advice

In D89447#2334424, @jonpa wrote:

In D89447#2332645, @uweigand wrote:

@jonpa can you check whether the SystemZ test case you added still checks what it was intended to check here?

I actually see a problem with this patch: The SystemZ test case passes the original test by not producing any scalar load + vector element insertions instructions. However now I see that all PFD (prefetch) instructions end up last in the block, whereas before they were not. This may or not be good if the block is huge and the prefetching is therefore not done too many iterations ahead. Maybe this should be checked with benchmarks to make sure the current prefetching tuning does not loose by this.

The reason seems to be that the SystemZ::PFD instruction is marked both with mayLoad and mayStore, but now this patch looks at the *memory operand* and figures out that it is only loading and therefore does not alias with another load. This check was previously only done with the MI flags. The post-RA scheduler now puts all of them at the end.

Remove the MMO non-store check.

In D89525#2334916, @foad wrote:

In D89525#2333924, @hliao wrote:

@yaxunl could you double-check that OpenCL also follows that rule.
@nhaehnle could you check whether that potentially breaks graphics.

FYI I've done some Vulkan CTS testing using LLPC with this patch and didn't notice any problems.

Fix typos and revise the coding style following clang-tidy.

@yaxunl could you double-check that OpenCL also follows that rule.
@nhaehnle could you check whether that potentially breaks graphics.

In D89447#2333315, @hliao wrote:

Add limit on memory operand AA check.

Add limit on memory operand AA check.

In D89447#2331842, @Kayjukh wrote:

This looks like the same as my old patch (https://reviews.llvm.org/D80161). There are additional details available on the commit's Phab page. To put it in a nutshell, some codes can trigger a very large amount of calls to the aliasing check (see the repro provided by @nemanjai: https://pastebin.com/tRtTQdSa), which results in a very large increase in compilation time.

Bounding the number of checks may be a good solution, even though it would be nicer to have something more clever that could allow all the operands to be checked. Not sure how feasible this would be though.

In D89447#2331842, @Kayjukh wrote:

This looks like the same as my old patch (https://reviews.llvm.org/D80161). There are additional details available on the commit's Phab page. To put it in a nutshell, some codes can trigger a very large amount of calls to the aliasing check (see the repro provided by @nemanjai: https://pastebin.com/tRtTQdSa), which results in a very large increase in compilation time.

Bounding the number of checks may be a good solution, even though it would be nicer to have something more clever that could allow all the operands to be checked. Not sure how feasible this would be though.

In D89447#2331763, @dmgreen wrote:

This looks familiar. It is the same as https://reviews.llvm.org/D80161? That patch was apparently reverted in 65cd2c7a8015577fea15c861f41d2e4b5768961f, because it was hitting timeouts on certain targets.

I'm not sure if @Kayjukh remembers more? But it may be worth putting a limit on the total number of alias checks it can perform.

The test changes seem fine to me.

That change triggers a few regression test failures. All of them are due to different code schedule due to memory instructions with multiple memory operands. Please help me double-check that's the case and whether that sounds a better code sequence. Thanks.

Register pressure tests have to disable AMDGPU AA to pass the test; otherwise, the register pressure is reduced after using AA.

Kindly PING for review.

This change is reverted as, on hosts without LBR supported but with LIBPFM installed and used, this change makes llvm/test/tools/llvm-exegesis/X86/lbr/mov-add.s failed. On that host, perf_event_open fails with EOPNOTSUPP on LBR config. That change's basic assumption

Rebase

In D88303#2298259, @bader wrote:

Thanks!

Split the original into 2. This is the first part, which add
correctly-rounded-device-sqrt-fp-math for OpenCL only. The second part will
remove that attribute annotating completely.

Remove the irrelevant change on .clang-format.

In D87939#2285876, @qcolombet wrote:

The comment seems outdated, if my understanding is right, and even the original code cannot perform that change since, once the 2nd COPY with same source is found in L1407, the check @ L1419 just skips that earlier as the 1st COPY has no subreg and the 2nd COPY has sub1.

Good point!

Now, I am wondering why is this change not just NFC then?

In D87939#2285776, @qcolombet wrote:

Hi @hliao,

I must be missing something, but it feels to me that this patch is actually making the situation worse.

Could you look at my example inlined below and explain how it would still work with this patch?

Cheers,
-Quentin

In D87972#2284096, @lebedev.ri wrote:

In D87972#2284064, @xbolva00 wrote:

In D87972#2284060, @MaskRay wrote:

I have tested this patch internally and seen gains and losses. On one document search related benchmark 3~5% improvement. One zippy (snappy) there is 3~5% regression. Perhaps we do need a conditional extra SROA run.

Snappy - you mean public https://github.com/google/snappy?

Well, it should be possible to analyze it...

@lebedev.ri any perf data from testsuite/rawspeed?

I did look.

sroa-after-unroll.rsbench.txt68 KBDownload

This suggests that geomean is -0.8% runtime improvement,
with ups&downs.

But as i have said in the patch's description, i stumbled into this when writing new code, where the effect is much larger.

Rebase.

This patch enhances the peephole-opt to fix the redundant copy issues once to be fixed in D87556. With the enhancement, we could remove that redundant COPY locally. Test cases are revised due to the code quality improvement or change. Fortunately, AMDGPU and ARM tests need addressing that difference.

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947880, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947851, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947724, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947723, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947722, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947689, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947688, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947687, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947684, @arsenm wrote:

Can you revert this? I think that this transform itself is a workaround, and even if it were a good idea, I think it doesn't belong in another loop over the function in finalizeLowering

Could you elaborate on why that would be a workaround? Basically, after instruction selection, the COPY from SGPR to VGPR should be lowered to a native instruction.

Because this should be done after register allocation like is already done. Replacing a copy with something else should only interfere with generic optimizations

That would be too late for MachineCSE and other optimization to remove the redundant COPYs and reduce the register usage. Moving that after RA won't reduce register pressure.

So you are working around something MahcineCSE isn't doing on copies. You could just have MachineCSE do this for these copies

MachineCSE and other optimizations are designed not to handle that target-independent COPY or like, which is added with the intention that, potentially, the source and destination operands are coalesced and that COPY is removed finally. As SGPR and VGPR are different register banks and won't be coalesced anyway, native instruction should be used instead.

This does not make sense, COPY is what generic optimizations do understand. If it's useful to CSE cross bank copies, MachineCSE should handle them. As a representational choice, not-copy is worse than copy

A target-independent cross-bank COPY is definitely useful but the current COPY should not be used for that purpose considering how it's used in RA-related passes, especially only architectural constraints are changed between the source and destination operands where the propagation should be stopped.

This is exactly what copy is for. It wouldn't make sense to have a separate copy for this

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947724, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947723, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947722, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947689, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947688, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947687, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947684, @arsenm wrote:

Can you revert this? I think that this transform itself is a workaround, and even if it were a good idea, I think it doesn't belong in another loop over the function in finalizeLowering

Could you elaborate on why that would be a workaround? Basically, after instruction selection, the COPY from SGPR to VGPR should be lowered to a native instruction.

Because this should be done after register allocation like is already done. Replacing a copy with something else should only interfere with generic optimizations

That would be too late for MachineCSE and other optimization to remove the redundant COPYs and reduce the register usage. Moving that after RA won't reduce register pressure.

So you are working around something MahcineCSE isn't doing on copies. You could just have MachineCSE do this for these copies

MachineCSE and other optimizations are designed not to handle that target-independent COPY or like, which is added with the intention that, potentially, the source and destination operands are coalesced and that COPY is removed finally. As SGPR and VGPR are different register banks and won't be coalesced anyway, native instruction should be used instead.

This does not make sense, COPY is what generic optimizations do understand. If it's useful to CSE cross bank copies, MachineCSE should handle them. As a representational choice, not-copy is worse than copy

Revise formatting following the clang-format suggestion.

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947722, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947689, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947688, @arsenm wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947687, @hliao wrote:

In rGc3492a1aa1b98c8d81b0969d52cea7681f0624c2#947684, @arsenm wrote:

Can you revert this? I think that this transform itself is a workaround, and even if it were a good idea, I think it doesn't belong in another loop over the function in finalizeLowering

Could you elaborate on why that would be a workaround? Basically, after instruction selection, the COPY from SGPR to VGPR should be lowered to a native instruction.

Because this should be done after register allocation like is already done. Replacing a copy with something else should only interfere with generic optimizations

That would be too late for MachineCSE and other optimization to remove the redundant COPYs and reduce the register usage. Moving that after RA won't reduce register pressure.

So you are working around something MahcineCSE isn't doing on copies. You could just have MachineCSE do this for these copies