In D145831#4207338, @JonChesterfield wrote:

I'm really sure that locks at thread scope do not work on amdgpu or pre-volta nvptx. One of the threads wins the cas, all the others do not, and it immediately deadlocks.

Critical sections can be done by rewriting the cfg, general purpose locks can't.

What am I missing here?

@kevinsala if the issue is with returning the host thread a pointer that cannot be used on the host - except in a map clause and a target region - then perhaps we should not return the locked pointer at all. The lock function could operate in a "side effect" mode, where the user asks the runtime to lock memory and the ROCr keeps track of it. Does this make sense?

In D142512#4079226, @kevinsala wrote:

@carlo.bertolli Just tested the example below with the old plugins and it also segfaults:

#include <assert.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>

#include <omp.h>

void *llvm_omp_target_lock_mem(void *ptr, size_t size, int device_num);
void llvm_omp_target_unlock_mem(void *ptr, int device_num);

int main(int argc, char **argv) {
    const size_t size = 1000*sizeof(int);

    int *ptr1 = (int *) malloc(size);
    assert(ptr1 != NULL);

    memset(ptr1, 0, size);

    int *ptr2 = (int *) llvm_omp_target_lock_mem(ptr1, size, 0);
    assert(ptr2 != NULL);

    fprintf(stdout, "ptr1: %p, ptr2: %p\n", ptr1, ptr2);
    fprintf(stdout, "ptr1[0]: %d\n", ptr1[0]);
    fprintf(stdout, "ptr2[0]: %d\n", ptr2[0]);

    return 0;
}

When running the program, the access to ptr2[0] segfaults at the host side:

ptr1: 0x55b461f00ec0, ptr2: 0x7fc5576c8ec0
ptr1[0]: 0
Segmentation fault

So I don't think ptr2 is a locked host pointer. It's the pointer returned by hsa_amd_memory_lock as the agent_ptr output parameter. And probably it's a pointer that the GPU agents should use for asynchronous memory transfers involving the buffer [ptr1, size). I assumed the host locked pointer is ptr1 directly, after hsa_amd_memory_lock is executed. However, I'm not 100% sure since the documentation of the lock operation on HSA AMD is a bit confusing.

This test works correctly with the old plugin. A locked pointer is not a device pointer. It is a host pointer that has been pinned in memory and that has been made accessible by ROCr to the device agent that will be assigned the data transfers implementing the map clause. In abstract programming terms, the second feature is only an AMD one and it does not need to have a corresponding functionality on other targets.

In D139208#4059152, @kevinsala wrote:

@carlo.bertolli is there documentation regarding llvm_omp_target_lock_mem llvm_omp_target_unlock_mem? While reading the comments of this patch, I see that the API ensures that locked areas feature reference counting. But I have the following doubts. Should the API support locking memory buffers that were already locked (e.g., complete or partial overlapping)? What's the behavior in such a case?

[AMD Official Use Only - General]

Pushed a fix to the non-amdgpu build:
7928d9e12d47fcc226d0c6984e11f5f463670f4a
Tested on machine without ROCm installation or AMDGPU.

[OpenMP][libomptarget][AMDGPU] Address comments, including fix in error handling in is_locked function.

In D139208#4041170, @JonChesterfield wrote:

hsa_amd_memory_lock gives the corresponding region coarse semantics. No update visible until after a kernel exits. Is that right for this?

In D141227#4041529, @carlo.bertolli wrote:

In D141227#4041177, @JonChesterfield wrote:

What is locked memory, and could we define it in a comment/documentation somewhere if it isn't already?

Is it the same as mmap LOCKED? That plus extra?

Does it imply GPUs can read/write it? If so, can they read/write within a kernel execution, or can we implement locked memory by copy at the start and end of the execution?

In D141227#4041177, @JonChesterfield wrote:

What is locked memory, and could we define it in a comment/documentation somewhere if it isn't already?

Is it the same as mmap LOCKED? That plus extra?

Does it imply GPUs can read/write it? If so, can they read/write within a kernel execution, or can we implement locked memory by copy at the start and end of the execution?

In D139208#4041132, @ye-luo wrote:

lock_memory and unlock_memory not symmetric.

I mean
lock_memory(ptr) # lock happens
lock_memory(ptr) # no-op
unlock_memory(ptr) # unlock happens
unlock_memory(ptr) # no-op

I just feel unsafe. Feel more safe if
lock_memory(ptr) # lock happens
lock_memory(ptr) # no-op
unlock_memory(ptr) # no-op
unlock_memory(ptr) # unlock happens
but this requires doing reference counting.

In D139208#4041043, @ye-luo wrote:

I still don't feel safe about the amd plugin implementation

1. how is_locked handle error.

[OpenMP][libomptarget][AMDGPU] Addressed comments.

I changed the old plugin interface for tgt_rtl_data_lock to return an error code. It now returns the lockedptr as function argument. Let me know if this is not what was called for.
Thanks for this extension!

Thanks for being patience with this patch update. I hit a problem with the is_locked function, when the prelocked pointer is passed to the map clause: the address offset calculation was based on the system-allocator (e.g., malloc) pointer but it would not work if the agentBasePointer (locked) was passed in. Fixed now.

[OpenMP][libomptarget][AMDGPU] Address comments and use prelocked pointers in mep clause implementation.

[OpenMP][libomptarget][AMDGPU] Add lock/unlock to prevent races on Devices vector in plugin manager.

[OpenMP][libomptarget][AMDGPU] Apply requested changes and merge against trunk.

In D140077#3996828, @jdoerfert wrote:

Should this have a target? Do you want to check a and b after?

Updated test with checks of a and b pointers.

[OpenMP] Only rerun lit test for pinned memory, no need to rebuild.

[OpenMP] Add test for default allocator and move unlock call to appropriate site.

ping

ping

Apply comments

In D139208#3966980, @JonChesterfield wrote:

It's interesting that locking locked memory succeeds, but doesn't give you something that has to be unlocked twice. Not totally sure about the convenience vs error detection there. What does the proposed user facing interface look like?

Removed llvm_omp_target_host_mem_alloc as a case for pinning.

Pinning does not apply to device memory on AMDGPU target, only to host memory, as explained here:
https://github.com/RadeonOpenCompute/ROCR-Runtime/blob/master/src/inc/hsa_ext_amd.h#L1526

LGTM but @ABataev should review/approve this

Thanks for explaining. LGTM.

This looks good, but what happens when the user accidentally adds a nested parallel when this option is turned on? Do we get serial (correct) execution?

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++?