LGTM, thanks for this.

Replace list of shadow pointer to be restored with simple synchronization.

It looks like https://reviews.llvm.org/D115823 will take a bit longer to be merged. I will merge this in to fix the API error, as it can go in first.

Add context.

LGTM

Trying rebase again.

rebase onto main

[OpenMP][NFC] Apply requested change: remove "nowait" from function name.
Add context (sorry about that!)

[OpenMP] Add missing hsa declarations/definitions when building runtime without rocr (or hsa library) installed on the system

In D115279#3179975, @ye-luo wrote:

I know what happened to my machine. Some CMake change caused offload plugins are not compiled. Sign. broken upstream.
My intention is to check pinned status. Before the first lock(not pinned), after the first lock(pinned), after the offload region(pinned), after the unlock(unpinned).
Could you also verify with rocprof hsa trace that the lock and unlock are both called twice?

In D115279#3179916, @ye-luo wrote:

I failed to verify your first lock behaves as intended.

#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <omp.h>
#include <stdio.h>

#define N 100293

int checkLocked(void *ptr) {
  hsa_amd_pointer_info_t info;
  hsa_status_t herr;

  herr = hsa_amd_pointer_info(ptr, &info, NULL, NULL, NULL);
  if (herr != HSA_STATUS_SUCCESS) {
    printf("  hsa_amd_pointer_info failed\n");
    return 1;
  }

  if (info.type != HSA_EXT_POINTER_TYPE_LOCKED) {
    printf("  pointer is noooooooooooot locked\n");
    return 1;
  } else
    printf("  pointer is locked\n");

  return 0;
}

int main() {
  int n = N;
  int *a = new int[n];
  for (int i = 0; i < n; i++)
    a[i] = 0;

  int *a_locked = nullptr;
  hsa_status_t herr =
      hsa_amd_memory_lock(a, n * sizeof(int), nullptr, 0, (void **)&a_locked);
  if (herr != HSA_STATUS_SUCCESS) {
    printf("Locking failed\n");
    return 1;
  }

  checkLocked(a);

#pragma omp target parallel for map(tofrom : a_locked[:n])
  for (int i = 0; i < n; i++)
    a_locked[i] = i;

  herr = hsa_amd_memory_unlock(a);
  if (herr != HSA_STATUS_SUCCESS) {
    printf("Unlocking failed\n");
    return 1;
  }

  int err = 0;
  for (int i = 0; i < n; i++)
    if (a[i] != i) {
      err++;
      printf("Err at %d, got %d expected %d\n", i, a[i], i);
      if (err > 10)
        break;
    }

  delete[] a;

  return err;
}

I got failure at the first check with "hsa_amd_pointer_info failed". Could you take a look?

In D115279#3179687, @ye-luo wrote:

I have not tried with memory that has already been locked, but I will. In any case, with this patch, if locking fails, then we revert to malloc+lock+unlock+free. This is not ideal, and this case is added for other reasons, but it should be supporting the case.

It will be better skipping lock/free if the memory is known to HSA already. I think IBM XL skips its pinned memory optimization when it sees the pointer pinned already for CUDA.
I have code managing lock/unlock via HIP. Even if a lock call from the plugin succeeds, and then a plugin unlock call succeeds, the user unlock call fails.
For this reason, check memory info is required.
fallback to "malloc+lock+unlock+free" is the worst option.

In D115279#3178711, @JonChesterfield wrote:

Yes, should revert and update the dlopen HSA source. Apologies for not checking this builds before committing it.

I don't know whether there's more ritual to to around whether it's already pinned - @carlo.bertolli did you look into testing if the memory is already pinned before doing so? Particularly interested if already pinned is a reason for lock to fail

@JonChesterfield thanks for the quick review. Would you mind merging this for me?

format

@JonChesterfield this looks more readable as a first step.

Abandoning this one for readability issues and moving to two separate patches.
First one:
https://reviews.llvm.org/D115267

This is already a lot of code with parse+sema. I wonder if we should split the patch into two, i.e. 1. parse+sema; 2. code gen? @ABataev ?
It should simplify maintenance of the patch and allow time to extend the OpenMP IR builder.

LGTM

In D111280#3085993, @jdoerfert wrote:

In D111280#3048057, @carlo.bertolli wrote:

Thanks for this. Do you think this could use a test?
I assume a test would be easier to write once this information is used somewhere, but it's hard to tell without context.

How would a test look like? And what do you mean by "information is used somewhere"?
If we land this code it is "used" and will avoid execution of kernels with an associated loop trip count of 0.

Thanks for fixing this. It would be nice if you could add a comment where it is used at the beginning of parallel_51, stating what are the initial value of parallelLevel (kmpc_parallel_level) and those of kmpc_is_spmd_exec_mode in the two cases of generic and spmd modes (assuming generic spmd mode is handled as a spmd). Something like:

Thanks for this. Do you think this could use a test?
I assume a test would be easier to write once this information is used somewhere, but it's hard to tell without context.

LGTM

Updated patch based on comments: add new map table entry field to track USM maps; use it to determine behavior in getTargetPointer and deallocTgtPtr; update var name and add explaining comments.

In D102449#2977711, @tianshilei1992 wrote:

Another thing is how we deal with a corner case. Say the OpenMP code is written in the following way:

#pragma omp atomic compare
  x = e < x ? e : x;

That's how OpenMP spec defines the atomic operation. x is always in "else statement" of a conditional statement.

Now we need to lower it to LLVM IR, which is atomicrmw operation. Based on the LLVM IR reference, it only supports the atomic operations that x is in the "then statement". For example: x = x > e ? x : e. See the x here is before :. In order to lower the OpenMP statement, we have to do a transformation. In order to swap e and x, we need to transform it to x = e >= x : x : e, a.k.a. x = x <= e : x : e. However, we don't have an atomic operation for <=. We only have <. So if x != e, the result is good.

The incorrectness happens if x == e. Recall at the OpenMP statement, when x == e, the result should be x = x. But if we look at our lowered LLVM IR, x = x < e : x : e, when x == e, it becomes x = e, which doesn't conform with OpenMP spec.

What should we do here?

Thanks for uploading this.

Update environment variable name to reflect comments and intended name

I have very small comments. LGTM.

@ABataev can you please merge this for me? I still have to ask for commit privileges.

In D103899#2817081, @JonChesterfield wrote:

Nice, thanks. I guess we already have an equivalent test in c++?

[OpenMP] Add test for C language of target uses_allocators clause

Ah! I forgot to add a reference to a filed bugzilla:
https://bugs.llvm.org/show_bug.cgi?id=50619

[OpenMP] Add regression test: this verifies that only one call to the runtime for reduction is done in a combined construct with distribute, without simd, and with at least one pragma accepting reduction as a clause.

That's correct, debug printfs were reporting zero threads to be used in SPMD region instead of the actual number used.

Committed revision 321759.

Committed revision 319696.

Committed revision 318692.

Committed revision 317893.

Committed revision 314905.

Committed revision 302031.

Re-Committed revision 301340.

Committed revision 301223.

[OpenMP] Refactor after split of original patch in two patches (sema and cg). Transform lambdas to static functions, where possible, define default constructors, format, and update against trunk.

Committed revision 300802.

Also reported by bbot slave:

Committed revision r295497

Address comments on initial patch after splitting. Added a dependency with sema patch.

About this:

I will split this patch into smaller patches, but I will keep this open in case I need to refer back to it: some of the changes that can go into independent patches require a justification and this patch shows it.
Thanks for your comments!

[OpenMP] Updating the patch to show full context

Committed revision r294195

Committed to r294026

Commited r290898.

Committed revision 275282.

Committed revision 275271.

[OpenMP] Add regression test for ast, remove 'brief's, and rebase.

[OpenMP] Add regression test to check correctness of ast building, remove 'brief's, and rebase.

[OpenMP] Remove generation of capture (move to codegen) and rebase.