In D132306#4207335, @josemonsalve2 wrote:

Adding another limitation to this approach. In the following code:

void aaaa(int b) {
    // TARGET A
    #pragma omp target
    {}

    // TARGET B
    #pragma omp target nowait
    {}

    // TARGET C
    #pragma omp target
    {}
}

The task A->B dependency should be respected, since A is originally synchronous, it should have been executed before B. However, B will be spawned in a different thread and A can potentially execute after B.

This message could be printed when LIBOMPTARGET_INFO="something". I don't think it should be printed by default as it could be false positive.

I am not sure if we should generate this message. The spec says "Each list item must have a corresponding list item in the device data environment or be accessible on the target device."
The pointer may be in the unified_shared_memory.

In D134565#3812623, @jhuber6 wrote:

In D134565#3812616, @tianshilei1992 wrote:

We currently have a new plugin interface https://reviews.llvm.org/D134396, and the use of libelf is expected to go away.

The libelf dependency has already been removed in LLVM, we only use LLVM's libraries now.

LGTM

LGTM

LGTM

Thanks Andrey for the fix

LGTM

You should not be using kmpc in libomptarget

If the current behavior is only the pointer field of the struct and not other members of the struct are updated after the array is allocated and copied.

Are you concerned about a misbehavior here?

I am just answering the question George asked about the expected behavior.

#pragma omp target map(s1.p[0:10])

Shilei,

How much time do  you think you  need to resolve or conclude to revert or disable with macros in 12.0

Some would like to stabilize their performance numbers and would like to do it as early as possible.
Thanks
Ravi

Ron,

Prefer if you  remove this from the mail " [AMD Official Use Only - Internal Distribution Only]"

Thanks
Ravi

Reasonable to drop the device number for alloc_host as all devices should be able to access host memory.

This is where we are heading towards in OpenMP spec
The mandatory value specifies that program execution is terminated if a device construct or
device memory routine is encountered and the device is not available or is not supported by the
implementation. or if omp_get_num_devices() would return 0.

If OMP_TARGET_OFFLOAD=mandatory it will not fail but just return

In D96999#2573081, @grokos wrote:

Is this the missing behavior from D83061? Looks good, thanks!

In D95483#2525182, @JonChesterfield wrote:

There are a bunch of side channels for communicating between the compiler and the runtime already. Compiler emitted magic symbols, a msgpack blob for amdgpu, writing extra stuff into the elf in general etc.

What information do you intend to pass along through this new channel, why yaml, and why for all plugins?

LGTM

LGTM

Can you also remove the definition in the header file include/omptarget.h

In D84422#2173372, @jdenny wrote:

In D84422#2172898, @jdenny wrote:

Has anyone clarified the motivation for this behavior?

I meant, is there any insight into why the spec specifies this behavior?

In D84422#2172926, @grokos wrote:

Instead of introducing new API functions and making all these changes in all these files, wouldn't it be easier if we just unset the PRESENT flag from arg_types in clang when we generate the call to __tgt_target_data_end_* if we are exiting from a scoped environment?

Ah, that does sound simpler. Thanks. I'll look into it.

Suppressing the presence check on exit from omp target would require a runtime change in addition to the Clang change you suggest for omp target data. However, I've so far failed to formulate a reasonable test case. Specifically, I don't yet see a way to guarantee that the data will definitely be present at the start of omp target but might not be present by the end. Is it possible? If not, then maybe we should leave the check in place for omp target.

In D84422#2170667, @grokos wrote:

So is the test case that motivated this patch illegal OpenMP code?

#pragma omp target enter data map(alloc:i)
#pragma omp target data map(present, alloc: i)
{
  #pragma omp target exit data map(delete:i) // you cannot delete that object in the scope, illegal code?
} // fails presence check here

In D84422#2170285, @grokos wrote:

What confuses me about this interpretation of the standard is the inconsistency at data exit. So if we have an explicit omp target exit data map(present...) then we should respect the "present" semantics, whereas when we have a scoped data exit:

#pragma omp target data map(present,...)
{
  ...
} // implicit "exit data" here

then "present" should be ignored.

I agree that the paragraph from the standard leaves little room for other interpretations, I'd just like to point out that it looks inconsistent - at least to me.

How do you plan to support
#pragma omp target update to (arr[1:2][1:2][0:2], x, b[1:5][0:2])
Are you going to split this into 3 updates since your are using the arg fields.

I don't recall saying needs changes.

LGTM

Can you enable LIBOMPTARGET_DEBUG and see if this is printed from libomptarget to confirm libomptarget is loaded and its api is used
DP("Call to omp_get_num_devices returning %zd\n", Devices_size);

Ok.
Thanks.

I did not see the code where check is done if Nestingdepth is 0 at end of compilation.

We should just go with generating an error if the DeclareTargetNestingLevel is not 0 at the end of compilation unit.
Hard to detect if user accidentally forgot to have end declare in header file and had it in the include file or it was intentional.

Is there a way to tell if the header files have matching omp declare target/omp end declare target.
The reason is if one of the header files is missing a matching omp end declare target all files which include it will end up having everything marked with declare target

If the objects which are copied are not compatible between host and target due to use of pointers which have different size in a structure then it is users responsibility.
I am more interested in where you allocate memory on device you get back an opaque object of size 64. The opaque object may be 64bit device pointer or a high 32bit pointer to device memory and low 32bit offset from that memory or it can be a pointer to a structure which has other information. The idea having libomptarget is to provide a thin layer between compiler and device specific plugin and we need to support different types of plugin.

Provide additional APIs in plugins if necessary to accomplish the arithmetic operations.
By doing arithmetic for target address in libomptarget you are assuming the pointer size is 64.

Alex: If you feel strongly about your request, I will be happy to move the code into CheckDeviceAndCtors.
It is just a suggestion. Either way is fine.

It may be named libOMPtarget and is in openmp projects but we want to use the offloading library with other runtime libraries.

We should not be having calls to _-kmpc_omp_taskwait in libomptarget.
The wait should be done in the user code and then libomptarget routine should be called.

DEFAULT implies MANDATORY or DISABLED. Don't need to communicate to libomp

Openmp has defined OMP_TARGET_OFFLOAD to control offloading, We should be using this

OpenMP says : " If the corresponding list item is not present in the device data environment then no assignment occurs to or from the original list item."