In D95062#2510863, @jdoerfert wrote:

In D95062#2510856, @tianshilei1992 wrote:

None of the three atomic operations provided by OpenMP, read, write, and update, sound like similar semantics as the four operations here.

Please finish reading the section about atomics (https://www.openmp.org/spec-html/5.1/openmpsu105.html#x138-1480002.19.7). All of the above operations are expressible, as far as I can tell.
add and inc are atomic update, exchange is atomic capture, max and cas are atomic compare, all with appropriate expression statement following them.

In D94972#2507666, @echristo wrote:

Hmm?

In D94972#2507385, @echristo wrote:

Hard to tell because there aren't any comments around any of the code that changed in either patch. Can you fix that in a separate change and actually document what everything is doing? :)

In D94972#2507344, @echristo wrote:

I think this would be more clear as a combined patch rather than 3 separate and some reverted patches. Would it make more sense to revert the chain and put them all back together in a single review?

In D94972#2507235, @echristo wrote:

Is the summary correct? The patch that this is fixing isn't reverted...

It's not clear what this is fixing and why it's fixing it.

Thanks!

-eric

In D93967#2496809, @echristo wrote:

If you changed the existing patch it won't be clear that this is something that needs review. Can you make a new review with the bug fix?

Thanks!

-eric

In D94961#2506460, @JonChesterfield wrote:

This patch was written, roughly, by:

• copying the known-working openmp driver from rocm into the trunk source tree
• deleting lots of stuff that didn't look necessary
• deleting some stuff that is broadly necessary, but the specifics are up for debate

The idea is to move language-independent but amdgcn-specific code into ROCMToolChain. Some has already gone in, others (like computeMSVCVersion) will likely move too.

Regarding the rest of the end to end stack:

• host plugin works, same code in trunk / rocm / aomp
• device plugin will work once it's building as openmp, modulo printf and malloc
• compiler backend will work for spmd kernels today, will work for generic kernels after D94648 or equivalent lands

Regarding tests (which need the unimplemented bit filled in with the next patch):

• Runtime tests (for spmd kernels) are working locally with a jury rigged devicertl
• Codegen tests are proving awkward to update. Dropping line number would help, but there's still a difference in addrspace cast distribution. I'm hoping the scripts involved in generating the nvptx cases can be adapted.

LGTM

Bug fix.

In D94315#2487643, @jdoerfert wrote:

What we should do, as we move to the OpenMPIRBuilder, is to use runtime interfaces that match OpenMP directives.
Here is the omp parallel one for the host: D94332
The device one should look the same, potentially after extending it, as we want to apply the same logic regardless of the device.
This will also make things like parallel region merging (@ggeorgakoudis) much easier.

I would prefer to fix PR48686 with this patch though until we switch to a new interface. The part where I remove the fn argument
meddling is on it's own valuable. We should not (need to) overwrite the function attributes for performance reasons given that we
have not inserted them. The new noinline will establish the invariant D94332 is going to have as well and which makes OpenMPOpt
much simpler.

Rework, bug fixes.

In D94315#2487257, @jdoerfert wrote:

In D94315#2487244, @ABataev wrote:

In D94315#2487242, @jdoerfert wrote:

I don't understand where inaccessiblemem_or_argmemonly is coming from. This prevents us from inlining the outlined parallel region.

kmpc_serialized_parallel and kmpc_end_serialized_parallel functions are marked with this attribute and I think OpenMPOpt optimizes read_only functions because of this. I.e. it does not consider kmpc_serialized_parallel and kmpc_end_serialized_paralle function calls as opimization barriers.

Can you provide an example why this is a problem?

In D94315#2487242, @jdoerfert wrote:

In D94315#2487164, @ABataev wrote:

In D94315#2487150, @JonChesterfield wrote:

I'm guessing we're using the function boundary as a compiler barrier. That seems fragile in the face of improving cross-function optimisation.

Looks like applying inaccessiblemem_or_argmemonly attribute to the OpenMP functions is way too optimistic since we still can access this (inaccessible) memory using other OpenMP functions. Not sure about the semantics of this attribute, though.

I don't understand where inaccessiblemem_or_argmemonly is coming from. This prevents us from inlining the outlined parallel region.

In D94315#2487150, @JonChesterfield wrote:

I'm guessing we're using the function boundary as a compiler barrier. That seems fragile in the face of improving cross-function optimisation.

Bug fixes

In D93967#2486129, @mstorsjo wrote:

This caused misoptimizations for armv7, where code that previously worked correctly now produce different results. (The code is clean under ubsan, so it shouldn't be relying on anything undefined.)

The issue appears with https://martin.st/temp/interplayvideo-preproc.c, compiled with clang -target armv7-linux-gnueabihf -O2.

The diff in generated code, before/after, looks like this:

        vmov.32 d16[0], lr
        vmov.32 d16[1], r2
 .LBB27_3:                               @ %if.end
                                         @   in Loop: Header=BB27_4 Depth=1
        vmov.32 r2, d16[1]
        add     r3, r3, #1
        vmov.32 r5, d16[0]
        cmp     r3, #8
+       vdup.32 d16, d16[0]
        vmov.16 d17[1], r2
        vmov.16 d18[0], r5
        vdup.16 d21, d17[1]
        vdup.16 d20, d18[0]
        vst1.16 {d20, d21}, [r1], r12
        beq     .LBB27_9
 .LBB27_4:                               @ %for.body
                                         @ =>This Inner Loop Header: Depth=1
        tst     r3, #3
        bne     .LBB27_3

If it loops back to .LBB27_3, the vector element d16[1] no longer has the value it was expected to have.

LG

Removed outdated comment.

In D92668#2473257, @rupprecht wrote:

I think I'm seeing a miscompile as a result of the patch. The issue is not the transformation itself, but a dependent instruction is reading the wrong value.

Before this patch, we have this:

  %171 = load <2 x i32>, <2 x i32>* %114, align 8
  %reorder_shuffle = shufflevector <2 x i32> %171, <2 x i32> undef, <2 x i32> <i32 1, i32 0>
  %shuffle = shufflevector <2 x i32> %reorder_shuffle, <2 x i32> undef, <4 x i32> <i32 0, i32 1, i32 0, i32 1>
...
  %186 = add nsw <2 x i32> %reorder_shuffle, <i32 -1, i32 -1>

After this patch the same snippet looks like:

  %171 = load <2 x i32>, <2 x i32>* %114, align 8
  %shuffle = shufflevector <2 x i32> %171, <2 x i32> undef, <4 x i32> <i32 1, i32 0, i32 1, i32 0>
...
  %186 = add nsw <2 x i32> %171, <i32 -1, i32 -1>

IIUC, if we suppose %171 loads the values (a, b), then:

• In the original version, %reorder_shuffle is (b, a), so %186 is (b - 1, a - 1)
• In the modified version, %186 reads directly from %171, so the result is (a - 1, b - 1), which is swapped

Do you know what the issue might be? I'm trying to reduce the test case so I can share more details.

In D93502#2463936, @dblaikie wrote:

In D93502#2463383, @ABataev wrote:

What I saw with the nvcc that it generated the debug info just like in the tests before. So, I just tried to mimic this behavior, since ptxas is very fragile and can be broken very easily. I would try to test it using ptxas/cuda-gdb to check that it does not break anything.

Oh, huh - interesting. Is it possible for me to reproduce/play around with this behavior in some way? (I've never used nvcc, not sure where to download/install it from, etc)

What I saw with the nvcc that it generated the debug info just like in the tests before. So, I just tried to mimic this behavior, since ptxas is very fragile and can be broken very easily. I would try to test it using ptxas/cuda-gdb to check that it does not break anything.

In D93068#2457946, @yaxunl wrote:

@ABataev Is this patch OK for OpenMP? It is NFC for OpenMP toolchain but affects using clang-offload-bundler as a standalone tool. Thanks.

LG

Restore original formatting in test cases that were not directly affected by the patch. Also, I would start with a single kind of expression rather than trying to cover as many kinds of expressions as possible. It makes it easier to understand and to review it.

In D57059#2443492, @dtemirbulatov wrote:

Do you mean compile time increasing? With this patch?

no, just compile-time error.

In D57059#2443350, @dtemirbulatov wrote:

While reviewing the latest update, I think I spotted SLP compile-time failure in SingleSource/Benchmarks/Misc/oourafft.c, here is the reduced testcase to reporduce:
source_filename = "/home/dtemirbulatov/llvm/test-suite/SingleSource/Benchmarks/Misc/oourafft.c"
target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
target triple = "x86_64-unknown-linux-gnu"

define dso_local fastcc void @cft1st(double* %a) unnamed_addr #0 {
entry:

%0 = or i64 16, 2
%arrayidx107 = getelementptr inbounds double, double* %a, i64 %0
%1 = or i64 16, 3
%arrayidx114 = getelementptr inbounds double, double* %a, i64 %1
%2 = or i64 16, 4
%arrayidx131 = getelementptr inbounds double, double* %a, i64 %2
%3 = or i64 16, 6
%arrayidx134 = getelementptr inbounds double, double* %a, i64 %3
%4 = load double, double* %arrayidx134, align 8
%5 = or i64 16, 5
%arrayidx138 = getelementptr inbounds double, double* %a, i64 %5
%6 = or i64 16, 7
%arrayidx141 = getelementptr inbounds double, double* %a, i64 %6
%7 = load double, double* %arrayidx141, align 8
%sub149 = fsub double undef, %4
%sub156 = fsub double undef, %7
store double undef, double* %arrayidx131, align 8
store double undef, double* %arrayidx138, align 8
%sub178 = fsub double undef, %sub156
%add179 = fadd double undef, %sub149
%mul180 = fmul double undef, %sub178
%sub182 = fsub double %mul180, undef
store double %sub182, double* %arrayidx107, align 8
%mul186 = fmul double undef, %add179
%add188 = fadd double %mul186, undef
store double %add188, double* %arrayidx114, align 8
unreachable

}

attributes #0 = { "target-features"="+avx,+avx2,+bmi,+bmi2,+cx16,+cx8,+f16c,+fma,+fsgsbase,+fxsr,+invpcid,+lzcnt,+mmx,+movbe,+pclmul,+popcnt,+rdrnd,+sahf,+sse,+sse2,+sse3,+sse4.1,+sse4.2,+ssse3,+x87,+xsave,+xsaveopt" }

!llvm.ident = !{!0}

!0 = !{!"clang version 12.0.0 (https://github.com/llvm/llvm-project.git aaa925795f93c389a96ee01bab73bc2b6b771cbb)"}

In D57059#2442542, @nikic wrote:

In D57059#2442527, @anton-afanasyev wrote:

In D57059#2442286, @RKSimon wrote:

AFAICT the only outstanding question is whether the compile time increase is acceptable?

I'd agree that geomean = 0.2% is acceptable for the change with such awesome performance impact, just noted that changed time compilation is significant in comparision with other changes. Generally it looks good to me apart from one minor unaddressed comment.

Could the summary of the revision be updated with the performance data? The thread is very long and I didn't spot where the measurements are, so it's hard to say what we're trading off here...

Rebase

General approach to fix the analysis of pointers.

Fix analysis for overlapping pointers.

In D86119#2437676, @cchen wrote:

Just tried this patch and found that below test abort inside generateInfoCapture.

#include <stdio.h>
#define LEN 100
int buf[LEN];

int main()
{
    int i;
    int *p = buf;

    for (i = 0; i < LEN; i++) {
        p[i] = 0;
    }

#pragma omp target map(p) map(p[:100])
    {
        *(p+5) = 1;
    }

}

Replaced += with just =

Address comments

Updated and rebased

Rebased ands reworked. Excluded SLM as its cost model is not quite correct.
Gives perf improvement for cpu2017.511.povray_r ~3% and improves performance up to 8% in some other cases.

LG

Rebase

I would just remove this extra symbol.

LG

Rebase