[OpenMP][libomptarget] Add support for unified memory for regular maps

[gtbercea]

Summary

This patch adds support for using unified memory in the case of regular maps that happen when a target region is offloaded to the device.

For cases where only a single version of the data is required then the host address can be used. When variables need to be privatized in any way or globalized, then the copy to the device is still required for correctness.

[gtbercea]

• Address comments.

[gtbercea]

libomptarget/src/api.cpp

118–123

I'm checking if the pointer is on the actual device. If unified memory is used then the pointers will match and the device present test will return false.
I have now refactored this check to make it more precise: if the host pointer is used then we have a flag for that.

libomptarget/src/device.cpp

292–293

Fixed.

[Hahnfeld]

I'd still like to see tests as mentioned in my last comment.

libomptarget/src/api.cpp

118–123

Okay, I think I get what you're saying: You want to check that this particular pointer has been mapped to that specific device_num, right?

But with unified shared memory, this shouldn't matter, right? Because all pointers can be accessed from all devices, no?

libomptarget/src/device.cpp

207–215

I still don't understand this.

[gtbercea]

libomptarget/src/api.cpp

118–123

Employing the API functions allows the user to bypass the unified memory behaviour and these functions allow the user to manage device pointers explicitly.

libomptarget/src/device.cpp

207–215

When use_device_ptr is employed the pointer is a true device pointer.

[gtbercea]

@Hahnfeld can you list the tests you would like to see please? And then I'll add them.

[Hahnfeld]

In D65001#1597305, @gtbercea wrote:

@Hahnfeld can you list the tests you would like to see please? And then I'll add them.

All new functionality should be tested unless there is a good reason it cannot:

In D65001#1594465, @Hahnfeld wrote:

Please add tests at least

• for the API, ie that omp_target_alloc returns distinct memory (I suggest to omp_target_memcpy from host to allocated memory to a different host memory and change the first host memory before copying back).
• that memory regions are indeed shared (in a target data region, the host should see updates from the host without an update and vice versa)

Additionally, I requested tests for omp_target_is_present inline.

[gtbercea]

• Add tests.

[gtbercea]

• Remove default device env var.

[gtbercea]

• Improve test.

[gtbercea]

@Hahnfeld I added several tests. Because these tests require unified memory to be supported by the underlying system I have added them as a new type of test: check-libomptarget-nvptx-unified (in addition to the check-libomptarget-nvptx one). These tests should only be run on platforms which support unified memory.

[Hahnfeld]

In D65001#1600995, @gtbercea wrote:

@Hahnfeld I added several tests. Because these tests require unified memory to be supported by the underlying system I have added them as a new type of test: check-libomptarget-nvptx-unified (in addition to the check-libomptarget-nvptx one). These tests should only be run on platforms which support unified memory.

Can we instead have the tests work with the host plugins? That would be much easier. (Sorry, don't have the time to look at the tests right now)

[gtbercea]

• Add tests.

[gtbercea]

In D65001#1601037, @Hahnfeld wrote:

In D65001#1600995, @gtbercea wrote:

@Hahnfeld I added several tests. Because these tests require unified memory to be supported by the underlying system I have added them as a new type of test: check-libomptarget-nvptx-unified (in addition to the check-libomptarget-nvptx one). These tests should only be run on platforms which support unified memory.

Can we instead have the tests work with the host plugins? That would be much easier. (Sorry, don't have the time to look at the tests right now)

Done!

[Hahnfeld]

libomptarget/test/api/api.c

9 ↗

(On Diff #211963)

Please add a manual call to the runtime function such that this test works with older versions of Clang.

22 ↗

(On Diff #211963)

Is there a particular reason to use the second device? Can we just take omp_get_default_device?

69 ↗

(On Diff #211963)

Can we change the type to int to avoid any problems with floating point numbers?

77–79 ↗

(On Diff #211963)

I'd suggest to make this different tests, and make clear that the test exercises unified_shared_memory (maybe a new directory)

82–88 ↗

(On Diff #211963)

And I still don't understand why we need omp_target_is_present to return false here. What would be the disadvantage of saying "all data is present on all devices"?

98 ↗

(On Diff #211963)

because this sounds really odd: A can be accessed on the device, but is not "present"?

libomptarget/test/offloading/requires_unified_shared_memory_local.c

17–18 ↗

(On Diff #211963)

Can we make this int?

25–26 ↗

(On Diff #211963)

Why are the pointers cast to long long? Can we just compare them as void *?

[Hahnfeld]

Oh, and there's still no test for use_device_ptr

[gtbercea]

libomptarget/test/api/api.c

22 ↗

(On Diff #211963)

No there isn't and yes we can.

69 ↗

(On Diff #211963)

Sure.

82–88 ↗

(On Diff #211963)

I said it and I'll say it again: when the user manually requests data to be allocated on the device, data must be allocated on the device. In this case the use of the omp_target_alloc will allocate data on the device even if unified shared memory is active.

98 ↗

(On Diff #211963)

This is just a presence test using the OpenMP API. If A is not associated to a device instance then it won't be considered present. Again, if the user wants to handle all of this manually the option is there and unified memory can't make this not an option for the user.

libomptarget/test/offloading/requires_unified_shared_memory_local.c

17–18 ↗

(On Diff #211963)

Yes

25–26 ↗

(On Diff #211963)

Is there a problem I'm missing if they are cast to long long?

[Hahnfeld]

libomptarget/test/api/api.c

82–88 ↗

(On Diff #211963)

I don't see the relation to what I asked: Sure, omp_target_alloc has to return fresh memory. But why does this matter for omp_target_is_present?

98 ↗

(On Diff #211963)

If A is not associated to a device instance then it won't be considered present.

Can you link me to a paragraph in the spec that mandates this implication? I only know of the following from the Effect of omp_target_is_present:

This routine returns non-zero if the specified pointer would be found present on device device_num by a map clause; otherwise, it returns zero.

I think there's nothing that prevents a non-zero return value because in fact all pointers would be found present by a map clause under unified shared memory.

libomptarget/test/offloading/requires_unified_shared_memory_local.c

25–26 ↗

(On Diff #211963)

Yes, casting to long long might truncate, or at least I'm not aware that it guarantees to be of the same size as pointers.

[gtbercea]

libomptarget/test/api/api.c

98 ↗

(On Diff #211963)

You would be right to think so if this was not in combination with omp_target_alloc. The fact that it is it means that the pointer was not officially mapped (explicitely or implicitly).

[Hahnfeld]

libomptarget/test/api/api.c

98 ↗

(On Diff #211963)

So what would be a full motivating example where libomptarget can't return true for any pointer passed to omp_target_is_present? Is it in this test?

[grokos]

libomptarget/src/device.cpp

160–161

IsHostPtr is never initialized in this function, I think it's better to set it to false explicitly instead of relying on the assumption that the caller has set it to false. You can init it right after rc.

166–177

This comma is confusing here, can you move it after "active"?

If unified shared memory is active, implicitly mapped variables that are not privatized use host address.

167

explicitely --> explicitly

174–176

I'm confused with this condition. What do we want to test?

The condition can be simplified to:

IsImplicit || !lr.Flags.IsContained || Size

because if Size==0, then lookupMapping cannot return ExtendsBefore/After, it may only return IsContained. So we want a mapping that's either

1. implicit or
2. explicit and it is either not contained or has Size>0.

The latter doesn't make sense to me... For instance, what if Size>0 and the mapping explicitly extends after? The condition will evaluate to true but this is invalid use no matter whether we use unified shared memory or not.

178–181

This evaluates to just HstPtrBegin, the two HT.HstPtrBegin cancel out, so you can skip defining tp altogether.

237–245

This init can be moved outside the locked-mutex region.

248

Same here, just use HstPtrBegin instead of tp.

libomptarget/src/omptarget.cpp

247–249

This is correct, with one little exception. Although the OpenMP standard does not mandate it, upstream clang supports use_device_ptr on pointers which are struct members. Because they are struct members, they are not marked with TARGET_PARAM (only the combined entry is considered a target parameter, not the individual members). On the other hand, they are marked with PTR_AND_OBJ and have some value in the MEMBER_OF bits.

Once again, it's a non-standard extension so we are free to decide whether to support it or not in the unified shared memory scenario.

259

ne --> be

379–381

Same here (if we decide to support the struct member case).

507

unified shared memory

541

unified shared memory

553–578

I think we can skip the whole shadow pointer loop when we use unified shared memory.

libomptarget/test/offloading/requires_unified_shared_memory_local.c

25–26 ↗

(On Diff #211963)

You can cast them to uintptr_t, it is a datatype guaranteed to be long enough to represent a pointer.

[Hahnfeld]

libomptarget/src/omptarget.cpp

247–249

Can we please first answer my question why we need to care about the existence of use_device_ptr? Why does it make a difference for unified shared memory?

libomptarget/test/offloading/requires_unified_shared_memory_local.c

25–26 ↗

(On Diff #211963)

Again: Why do we need to cast? We can just compare void *!

[gtbercea]

libomptarget/src/omptarget.cpp

247–249

Because use_device_ptr implies the use of a true device pointer and again that needs to be respected even unified memory is used.

[gtbercea]

libomptarget/test/api/api.c

98 ↗

(On Diff #211963)

Wait, this is not about a motivating example, it's about correctness. The function should return false if A is not "present" on the device. So I think the question is what is the definition of present in this case? As stated before I think in this case A is not actually mapped, perhaps the pointer to A may be implicitly mapped but the actual data belonging to A has not been associated with the pointer.

[Hahnfeld]

libomptarget/src/omptarget.cpp

247–249

Where is that in the spec?

libomptarget/test/api/api.c

98 ↗

(On Diff #211963)

But the data of A will not copied to the device by a map clause, right? So it's already present from that perspective.

[gtbercea]

libomptarget/src/omptarget.cpp

247–249

The pointers returned by omp_target_alloc and accessed through use_device_ptr are guaranteed to be pointer values that can support pointer arithmetic while still being native device pointers. (Section 2.4 page 61 bottom).

[Hahnfeld]

libomptarget/src/omptarget.cpp

247–249

With unified_shared_memory we get: (Section 2.4, page 62, lines 4:7)

Additionally, memory in the device data environment of any device visible to OpenMP, including but not limited to the host, is considered part of the device data environment of all devices accessible through OpenMP except as noted below. Every device address allocated through OpenMP device memory routines is a valid host pointer.

From my understanding, this implies all shared memory can be referenced by a host pointer. This guarantees pointer arithmetic per the C / C++ standard.

[AlexEichenberger]

libomptarget/src/omptarget.cpp

247–249

Experience from our users show a significant performance degradation when every data is mapped to host memory under unified. Overwhelmingly, users want specific data to be allocated on device & mapped, while at the same time they want generic memory used by the device to be host memory (e.g. for linked list structures).

Recognizing this important use case, the specs has introduced the "close" modifier to the map for users to indicate that the data being mapped should be set "close" to the device, aka for us allocated on the device. Similarly, if the user went through the additional exercise of using target alloc / free, we want to respect this explicit request of the user by allocating the data on the device.

This is why, even under unified memory model, some data is still mapped with duplicated copies between host and device. Thus requiring the maps & use_device_ptr to be tested and enforced when data was actually mapped.

Fortunately for the overhead, we know that only a very small subset of all data accessed will be labeled with "close" or allocated with target_alloc, so the cost of performing this mapping will be very significantly reduced.

[Hahnfeld]

libomptarget/src/omptarget.cpp

247–249

Alex, I fully agree with all of your points: Yes, we need means to allocate memory on the device (omp_target_alloc and close) and this must be reflected by maps and use_device_ptr. I'm not arguing against this, the behavior of the two methods is pretty clear in the spec.

What I'm asking about is something different, though related: If we have "generic" memory (as you call it, so normal memory on the host that is "shared" with all devices), do we need to allocate device memory if the user specifies use_device_ptr? Because that is what the current patch will do in getOrAllocTgtPtr if IsInUseDevicePtrClause = true.

Let's discuss a concrete example: (assume that the TU has a requires unified_shared_memory)

int A[N];
// init ...
int *Aptr = &A[0];

#pragma omp target data use_device_ptr(A)
{
  // What is the value of A?
}

I think we should have A == Aptr, but this patch will allocate disjoint device memory. (And while writing this, I wonder how libomptarget knows about the size of A? According to the spec, the clause expects a list of pointers, not array sections. So how will this work with opaque pointers passed to an orphaned function?)

[gtbercea]

libomptarget/src/omptarget.cpp

247–249

Your program returns the following errors:

test_maps.c:49:42: error: expected pointer or reference to pointer in 'use_device_ptr' clause

#pragma omp target data use_device_ptr(A)
                                       ^

test_maps.c:49:3: error: expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'

#pragma omp target data use_device_ptr(A)

[Hahnfeld]

libomptarget/src/omptarget.cpp

247–249

That happens if I don't compile test the small snippets I post :-/ either make A a heap pointer or use Aptr in the clause, that should work (though I can't test right now)

[gtbercea]

• Remove use device ptr flag.

[gtbercea]

• Address comments.

libomptarget/src/omptarget.cpp

247–249

So I think I had misunderstood the spec the first time around, I have now removed the special treatment of the use_device_ptr. The host address should be preserved now.

[gtbercea]

libomptarget/src/device.cpp

178–181

I tried and it leads to failing tests when handling structs.

[gtbercea]

• Move tests to new folder.
• Fix tests.
• Move requires test.

[gtbercea]

@Hahnfeld @grokos new version of the patch with addressed comments.

@Hahnfeld I cannot replace the usage of the requires pragma with the call to the registration function directly because the registration function is the first function that gets called so if I explicitely invoke it will only be the 2nd call to that function. A subsequent call with a different set of flags will lead to a mismatch in requires clauses error. (first implicit call is without unified_shared_memory and the second call is with unified_shared_memory).

[Hahnfeld]

In D65001#1607199, @gtbercea wrote:

@Hahnfeld I cannot replace the usage of the requires pragma with the call to the registration function directly because the registration function is the first function that gets called so if I explicitely invoke it will only be the 2nd call to that function. A subsequent call with a different set of flags will lead to a mismatch in requires clauses error. (first implicit call is without unified_shared_memory and the second call is with unified_shared_memory).

Yes, but you can have both, the requires directive and the manual call. I understand that this is not needed for newer Clang versions, but right now the test will fail for older releases that don't support requires.

---

For the tests, I think we should keep requires.c in test/offloading/ because it's not specific to unified_shared_memory. Can you also rename requires_unified_shared_memory_local.c to shared_update.c (or something like that) to make the name more concise?

libomptarget/src/api.cpp

118–123

I'm still not sure about this, what do others think?

libomptarget/src/device.cpp

174–176

I agree with @grokos here.

178–181

Again I agree with the (static) analysis by @grokos. If that leads to failing tests, we need to understand why. What was the idea behind the calculation of tp?

207–213

use_device_ptr is no more, so this change shouldn't be needed anymore?

292–293

Is deallocTgtPtr still called under unified_shared_memory?

libomptarget/src/omptarget.cpp

418–427

We can just move the error check into the branch, or even better move the check for unified_shared_memory to the outer condition.

521–522

Do we need the shadow pointers under unified_shared_memory and if TgtPtrBegin == HstPtrBegin? If not, we can just move the check to the outer condition.

576–577

As above, but I think this needs to check TgtPtrBegin == HstPtrBegin.

I think we can just do the following around line 490:

} else if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
                TgtPtrBegin == HstPtrBegin) {
  DP("hst data:" DPxMOD " unified and shared, becomes a noop\n", DPxPTR(HstPtrBegin));
  continue;
}

699–709

Can move the check Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && TgtPtrBegin == HstPtrBegin to line 690, in a similar else branch behind !Pointer_TgtPtrBegin.

libomptarget/test/unified_shared_memory/api.c

14 ↗

(On Diff #212436)

Please format the test.

[gtbercea]

libomptarget/src/api.cpp

118–123

See Alex's answer below!

[gtbercea]

libomptarget/src/device.cpp

178–181

The answer is simple: this is how it was done before. Look at the existing pointer calculation I see no reason why it should have to be simplified. The simplification you propose has nothing to do with unified memory so if you think it needs to be done then you should open it up as a separate issue in a separate patch.

[Hahnfeld]

libomptarget/src/api.cpp

118–123

This doesn't even mention omp_target_is_present, so this is no answer.

[gtbercea]

• Rename test.
• Revert file move. Simplify condition.
• Address comments.

libomptarget/src/device.cpp

178–181

I misread the comment. I think it's no problem doing the simplification.

[gtbercea]

libomptarget/src/api.cpp

118–123

So the problem is that if I return true all the time for the presence test when unified memory is on, then there's no way to distinguish the case where I use omp_target_alloc() but I haven't associated the host pointer with the device pointer yet.

So in other words the presence test is always an "honest" device presence test: it will only return true when there is a host version and corresponding device duplicate version. I think that case is the more meaningful to cover since for the other cases you don't really need a presence test.

[gtbercea]

libomptarget/src/api.cpp

118–123

I think your confusion is based on the fact that you assume that the presence test is for checking if a future usage of a variable in a device target region leads to a seg fault or not. And I agree that in the non-unified case this intent coincides with that of a map existing.
In the unified case that's not necessarily the case: you can use a variable on the device by referring to its host version without the need of a map existing. The presence test in that case will return false and that's fine.

[gtbercea]

@Hahnfeld I hope I have answered all your concerns. Please let me know if you have any other comments.

[Hahnfeld]

There's still no call to __tgt_register_requires in the two tests, so I guess they won't pass with older versions of Clang.

libomptarget/src/device.cpp

258

Hehe, casting a void * to uintptr_t and back to void *? Please remove :D

libomptarget/src/omptarget.cpp

566–572

This can never happen, the loop will continue above (currently line 493).

libomptarget/test/unified_shared_memory/shared_update.c

20 ↗

(On Diff #212622)

This test isn't formatted either.

[gtbercea]

In D65001#1610445, @Hahnfeld wrote:

There's still no call to __tgt_register_requires in the two tests, so I guess they won't pass with older versions of Clang.

ah I knew I was forgetting something :)

[gtbercea]

• Add manual for registering the requires flag.
• Address comments.

[gtbercea]

• Clean-up.

[Hahnfeld]

I'm fine with the last update.

[Hahnfeld]

Oh, but the tests again look odd. Please run them through clang-format.

[grokos]

libomptarget/src/device.cpp

189

I still don't understand this condition (even in the simplified form). Can you explain what we are testing?

[gtbercea]

• Fix condition.

[gtbercea]

libomptarget/src/device.cpp

189

The condition comes from the 1st and 3rd if conditions. I have updated the condition to its more explicit version.

But, from what I can tell, IsContained can never be true in this case, ExtendesBefore/After I think cannot be true either.

IsImplicit can vary, Size also.

So I would reduce the condition to:

(IsImplicit || Size)

Does this now make sense?

[gtbercea]

libomptarget/src/device.cpp

189

Actually it should just be Size I think. Testing it now.

[gtbercea]

• Fix condition again.

[gtbercea]

libomptarget/src/device.cpp

189

Condition has now been fixed. This is it in its simplest form.

[gtbercea]

• Fix condition to exclude already mapped device data.

[gtbercea]

• Format tests.

[gtbercea]

@grokos I have updated the condition. I have also amended the comment to better explain the conditions.

[gtbercea]

@grokos any thoughts on the updated condition?

[grokos]

libomptarget/src/device.cpp

189

I still think the condition is not correct - it can lead to invalid mappings escaping the runtime check. E.g. what if we use unified memory and we try to explicitly "map more"? The condition will evaluate to true (since we try to extend the mapping, isContained will be false and Size will be > 0), but the mapping is illegal and should be caught by the runtime.

I believe the whole if branch this patch introduces should be removed. Instead, we must let the original code perform the checks and in the last else-branch (line 185 in the original code "else if (Size)" we should check whether or not we use unified shared memory; if yes, then simply return the host address, if no, then proceed with the allocation of memory on the device. In other words, let this function find out whether the requested mapping already exists on the device; if it does we are done, if it doesn't then find out whether we should allocate device memory or use the host version of data.

[gtbercea]

• Move condition.

[gtbercea]

libomptarget/src/device.cpp

189

Moved it.

[gtbercea]

• Clean up conditions.

[grokos]

libomptarget/src/device.cpp

189

OK, it looks good now. Make sure to rebase the next patch (the close modifier) and extend this condition to check for !close before returning the host address (RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier).

[gtbercea]

libomptarget/src/device.cpp

189

Already done :)

[Hahnfeld]

openmp/trunk/libomptarget/test/unified_shared_memory/shared_update.c

28–30 ↗

(On Diff #213945)

Apparently, this does not work: The generated code will call __tgt_register_lib first which caches the global RequiresFlags in Device.RTLRequiresFlags. Because __tgt_register_requires has not been called yet, the value is still 0 so the new code won't be executed. Please fix and test on your end that it works with older versions of the compiler!

[gtbercea]

openmp/trunk/libomptarget/test/unified_shared_memory/shared_update.c

28–30 ↗

(On Diff #213945)

@Hahnfeld this test works correctly for compiler versions which support unified shared memory.

My proposed fix is to remove all the manual calls and restrict the test to new versions of Clang and do that for every test here and in the close modifier patches. None of the close or unified memory pieces of functionality need to be tested with older clang versions because they are not supported on those versions.

[Hahnfeld]

openmp/trunk/libomptarget/test/unified_shared_memory/shared_update.c

28–30 ↗

(On Diff #213945)

This will lose a lot of test coverage for the runtime library because the tests can only be run with a not-even-released versions of the compiler.

[gtbercea]

openmp/trunk/libomptarget/test/unified_shared_memory/shared_update.c

28–30 ↗

(On Diff #213945)

New features are only added to the newest version of the compiler and it will work with that and I think that's where the tests need to pass to make sure that incoming functionality doesn't break existing one.

I don't see any reason to test with older versions of Clang that can't even use these features. Am I missing something? What's the gain?

[tianshilei1992]

I never get those unified shared memory tests work on NVIDIA GPUs. Is it expected?

libomptarget/deviceRTLs/nvptx/CMakeLists.txt

foreach(sm ${nvptx_sm_list})
                          $<TARGET_FILE_DIR:omptarget-nvptx>)
 
       # Install bitcode library under the lib destination folder.
       install(FILES ${CMAKE_CURRENT_BINARY_DIR}/libomptarget-nvptx-sm_${sm}.bc DESTINATION "${OPENMP_INSTALL_LIBDIR}")
     endforeach()
   endif()
 
   add_subdirectory(test)
+  add_subdirectory(test-unified)
 else()
   libomptarget_say("Not building CUDA offloading device RTL: CUDA tools not found in the system.")
 endif()

libomptarget/deviceRTLs/nvptx/test-unified/CMakeLists.txt

+if(NOT OPENMP_TEST_COMPILER_ID STREQUAL "Clang")
+  # Silently return, no need to annoy the user.
+  return()
+endif()
+
+set(deps omptarget-nvptx omptarget omp)
+if(LIBOMPTARGET_NVPTX_ENABLE_BCLIB)
+  set(deps ${deps} omptarget-nvptx-bc)
+endif()
+
+# Don't run by default.
+set(EXCLUDE_FROM_ALL True)
+# Run with only one thread to only launch one application to the GPU at a time.
+add_openmp_testsuite(check-libomptarget-nvptx-unified
+    "Running libomptarget-nvptx tests for unified memory" ${CMAKE_CURRENT_BINARY_DIR}
+    DEPENDS ${deps} ARGS -j1)
+
+set(LIBOMPTARGET_NVPTX_TEST_FLAGS "" CACHE STRING
+    "Extra compiler flags to send to the test compiler.")
+set(LIBOMPTARGET_NVPTX_TEST_OPENMP_FLAGS
+    "-fopenmp -fopenmp-targets=nvptx64-nvidia-cuda" CACHE STRING
+    "OpenMP compiler flags to use for testing libomptarget-nvptx.")
+
+# Configure the lit.site.cfg.in file
+set(AUTO_GEN_COMMENT "## Autogenerated by libomptarget-nvptx configuration.\n# Do not edit!")
+configure_file(lit.site.cfg.in lit.site.cfg @ONLY)

libomptarget/deviceRTLs/nvptx/test-unified/api/api.c

+// RUN: %compile-run-and-check
+
+#include <stdio.h>
+#include <omp.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+void init(double A[], double B[], double C[]) {
+  for(int i=0; i<N; ++i) {
+    A[i] = 0;
+    B[i] = 1;
+    C[i] = i;
+  }
+}
+
+int main(int argc, char *argv[]) {
+  const int device = 1;
+
+  // CHECK: Initial device: -10
+  printf("Initial device: %d\n", omp_get_initial_device());
+
+  //
+  // Target alloc & target memcpy
+  //
+  double A[N], B[N], C[N];
+
+  // Init
+  init(A,B,C);
+
+  double *pA, *pB, *pC;
+
+  // map ptrs
+  pA = &A[0];
+  pB = &B[0];
+  pC = &C[0];
+
+  double *d_A = (double *)omp_target_alloc(N*sizeof(double), device);
+  double *d_B = (double *)omp_target_alloc(N*sizeof(double), device);
+  double *d_C = (double *)omp_target_alloc(N*sizeof(double), device);
+
+  // CHECK: omp_target_alloc succeeded
+  printf("omp_target_alloc %s\n", d_A && d_B && d_C ?
+      "succeeded" : "failed");
+
+  omp_target_memcpy(d_B, pB, N*sizeof(double), 0, 0,
+      device, omp_get_initial_device());
+  omp_target_memcpy(d_C, pC, N*sizeof(double), 0, 0,
+      device, omp_get_initial_device());
+
+  #pragma omp target is_device_ptr(d_A, d_B, d_C) device(device)
+  {
+    #pragma omp parallel for schedule(static,1)
+    for (int i = 0; i < N; i++) {
+      d_A[i] = d_B[i] + d_C[i] + 1;
+    }
+  }
+
+  omp_target_memcpy(pA, d_A, N*sizeof(double), 0, 0,
+      omp_get_initial_device(), device);
+
+  // CHECK: Test omp_target_memcpy: Succeeded
+  int fail = 0;
+  for(int i=0; i<N; ++i) {
+    if (A[i] != (double)(i+2)) fail++;
+  }
+  if (fail) {
+    printf ("Test omp_target_memcpy: Failed\n");
+  } else {
+    printf ("Test omp_target_memcpy: Succeeded\n");
+  }
+
+  //
+  // target_is_present and target_associate/disassociate_ptr
+  //
+  init(A,B,C);
+
+  // CHECK: B is not present, associating it...
+  // CHECK: omp_target_associate_ptr B succeeded
+  if (!omp_target_is_present(B, device)) {
+    printf("B is not present, associating it...\n");
+    int rc = omp_target_associate_ptr(B, d_B, N*sizeof(double), 0, device);
+    printf("omp_target_associate_ptr B %s\n", !rc ? "succeeded" : "failed");
+  }
+
+  // CHECK: C is not present, associating it...
+  // CHECK: omp_target_associate_ptr C succeeded
+  if (!omp_target_is_present(C, device)) {
+    printf("C is not present, associating it...\n");
+    int rc = omp_target_associate_ptr(C, d_C, N*sizeof(double), 0, device);
+    printf("omp_target_associate_ptr C %s\n", !rc ? "succeeded" : "failed");
+  }
+
+  // CHECK: Inside target data: A is not present
+  // CHECK: Inside target data: B is present
+  // CHECK: Inside target data: C is present
+  #pragma omp target data map(from: B, C) device(device)
+  {
+    printf("Inside target data: A is%s present\n",
+        omp_target_is_present(A, device) ? "" : " not");
+    printf("Inside target data: B is%s present\n",
+        omp_target_is_present(B, device) ? "" : " not");
+    printf("Inside target data: C is%s present\n",
+        omp_target_is_present(C, device) ? "" : " not");
+
+    #pragma omp target map(from: A) device(device)
+    {
+      #pragma omp parallel for schedule(static,1)
+      for (int i = 0; i < N; i++)
+        A[i] = B[i] + C[i] + 1;
+    }
+  }
+
+  // CHECK: B is present, disassociating it...
+  // CHECK: omp_target_disassociate_ptr B succeeded
+  // CHECK: C is present, disassociating it...
+  // CHECK: omp_target_disassociate_ptr C succeeded
+  if (omp_target_is_present(B, device)) {
+    printf("B is present, disassociating it...\n");
+    int rc = omp_target_disassociate_ptr(B, device);
+    printf("omp_target_disassociate_ptr B %s\n", !rc ? "succeeded" : "failed");
+  }
+  if (omp_target_is_present(C, device)) {
+    printf("C is present, disassociating it...\n");
+    int rc = omp_target_disassociate_ptr(C, device);
+    printf("omp_target_disassociate_ptr C %s\n", !rc ? "succeeded" : "failed");
+  }
+
+  // CHECK: Test omp_target_associate_ptr: Succeeded
+  fail = 0;
+  for(int i=0; i<N; ++i) {
+    if (A[i] != (double)(i+2)) fail++;
+  }
+  if (fail) {
+    printf ("Test omp_target_associate_ptr: Failed\n");
+  } else {
+    printf ("Test omp_target_associate_ptr: Succeeded\n");
+  }
+
+  omp_target_free(d_A, device);
+  omp_target_free(d_B, device);
+  omp_target_free(d_C, device);
+
+  return 0;
+}

libomptarget/deviceRTLs/nvptx/test-unified/lit.cfg

+# -*- Python -*- vim: set ft=python ts=4 sw=4 expandtab tw=79:
+# Configuration file for the 'lit' test runner.
+
+import os
+import lit.formats
+
+# Tell pylint that we know config and lit_config exist somewhere.
+if 'PYLINT_IMPORT' in os.environ:
+    config = object()
+    lit_config = object()
+
+def prepend_library_path(name, value, sep):
+    if name in config.environment:
+        config.environment[name] = value + sep + config.environment[name]
+    else:
+        config.environment[name] = value
+
+# name: The name of this test suite.
+config.name = 'libomptarget-nvptx'
+
+# suffixes: A list of file extensions to treat as test files.
+config.suffixes = ['.c', '.cpp', '.cc']
+
+# test_source_root: The root path where tests are located.
+config.test_source_root = os.path.dirname(__file__)
+
+# test_exec_root: The root object directory where output is placed
+config.test_exec_root = config.binary_dir
+
+# test format
+config.test_format = lit.formats.ShTest()
+
+# compiler flags
+config.test_flags = " -I " + config.omp_header_directory + \
+    " -L " + config.library_dir + \
+    " --libomptarget-nvptx-path=" + config.library_dir;
+
+if config.omp_host_rtl_directory:
+    config.test_flags = config.test_flags + \
+        " -L " + config.omp_host_rtl_directory
+
+config.test_flags = config.test_flags + " " + config.test_extra_flags
+
+# Setup environment to find dynamic library at runtime.
+prepend_library_path('LD_LIBRARY_PATH', config.library_dir, ":")
+prepend_library_path('LD_LIBRARY_PATH', config.omp_host_rtl_directory, ":")
+
+# Forbid fallback to host.
+config.environment["OMP_TARGET_OFFLOAD"] = "MANDATORY"
+
+config.environment["OMP_DEFAULT_DEVICE"] = "1"
+
+# substitutions
+config.substitutions.append(("%compilexx-run-and-check",
+    "%compilexx-and-run | " + config.libomptarget_filecheck + " %s"))
+config.substitutions.append(("%compile-run-and-check",
+    "%compile-and-run | " + config.libomptarget_filecheck + " %s"))
+config.substitutions.append(("%compilexx-and-run", "%compilexx && %run"))
+config.substitutions.append(("%compile-and-run", "%compile && %run"))
+
+config.substitutions.append(("%compilexx",
+    "%clangxx %openmp_flags %flags %s -o %t"))
+config.substitutions.append(("%compile",
+    "%clang %openmp_flags %flags %s -o %t"))
+
+config.substitutions.append(("%clangxx", config.test_cxx_compiler))
+config.substitutions.append(("%clang", config.test_c_compiler))
+config.substitutions.append(("%openmp_flags", config.test_openmp_flags))
+config.substitutions.append(("%flags", config.test_flags))
+
+config.substitutions.append(("%run", "%t"))

libomptarget/deviceRTLs/nvptx/test-unified/lit.site.cfg.in

+@AUTO_GEN_COMMENT@
+
+config.test_c_compiler = "@OPENMP_TEST_C_COMPILER@"
+config.test_cxx_compiler = "@OPENMP_TEST_CXX_COMPILER@"
+config.test_openmp_flags = "@LIBOMPTARGET_NVPTX_TEST_OPENMP_FLAGS@"
+config.test_extra_flags = "@LIBOMPTARGET_NVPTX_TEST_FLAGS@"
+config.binary_dir = "@CMAKE_CURRENT_BINARY_DIR@"
+config.library_dir = "@LIBOMPTARGET_LIBRARY_DIR@"
+config.omp_header_directory = "@LIBOMPTARGET_OPENMP_HEADER_FOLDER@"
+config.omp_host_rtl_directory = "@LIBOMPTARGET_OPENMP_HOST_RTL_FOLDER@"
+config.libomptarget_filecheck = "@OPENMP_FILECHECK_EXECUTABLE@"
+
+# Let the main config do the real work.
+lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")

libomptarget/deviceRTLs/nvptx/test-unified/offloading/requires_unified_shared_memory_link.c

+// RUN: %compile-run-and-check
+
+#include <stdio.h>
+#include <omp.h>
+
+#pragma omp requires unified_shared_memory
+
+double VAR = 10.0;
+
+#pragma omp declare target link(VAR)
+
+int main(int argc, char *argv[]) {
+  // This also runs reg requires for the first time.
+  long long device_addr, host_addr;
+
+#pragma omp target map(tofrom: device_addr)
+  {
+    VAR += 1;
+    device_addr = (long long) &VAR;
+  }
+
+  host_addr = (long long) &VAR;
+
+  // CHECK: VAR = 11
+  printf("VAR = %f\n", VAR);
+
+  // CHECK: DEVICE address is HOST address.
+  if (device_addr == host_addr)
+    printf("DEVICE address is HOST address.\n");
+
+  printf("DEVICE_ADDR = %lld\n", device_addr);
+  printf("HOST_ADDR = %lld\n", host_addr);
+  return 0;
+}

libomptarget/deviceRTLs/nvptx/test-unified/offloading/requires_unified_shared_memory_local.c

+// RUN: %compile-run-and-check
+
+#include <stdio.h>
+#include <omp.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  long long host_alloc, device_alloc;
+  long long host_data, device_data;
+  double *alloc = (double *)malloc(N*sizeof(double));
+  double data[N];
+
+  for(int i=0; i<N; ++i){
+    alloc[i] = 10.0;
+    data[i] = 1.0;
+  }
+
+  host_data = (long long) &data[0];
+  host_alloc = (long long) &alloc[0];
+
+  // implicit mapping of data
+  #pragma omp target device(1) map(tofrom: device_data, device_alloc)
+  {
+    device_data = (long long) &data[0];
+    device_alloc = (long long) &alloc[0];
+
+    for(int i=0; i<N; i++) {
+      alloc[i] += 1.0;
+      data[i] += 1.0;
+    }
+  }
+
+  // CHECK: Address of alloc on device matches host address.
+  if (device_alloc == host_alloc)
+    printf("Address of alloc on device matches host address.\n");
+
+  // CHECK: Address of data on device matches host address.
+  if (device_data == host_data)
+    printf("Address of data on device matches host address.\n");
+
+  // On the host, check that the arrays have been updated.
+  // CHECK: Alloc values updated: Succeeded
+  fails = 0;
+  for(int i=0; i<N; i++) {
+    if (alloc[i] != 11.0)
+      fails++;
+  }
+  printf("Alloc values updated: %s\n", (fails == 0) ? "Succeeded" : "Failed");
+
+  // CHECK: Data values updated: Succeeded
+  fails = 0;
+  for(int i=0; i<N; i++) {
+    if (data[i] != 2.0)
+      fails++;
+  }
+  printf("Data values updated: %s\n", (fails == 0) ? "Succeeded" : "Failed");
+
+  free(alloc);
+
+  printf("Done!\n");
+
+  return 0;
+}

libomptarget/deviceRTLs/nvptx/test-unified/offloading/requires_unified_shared_memory_to.c

+// RUN: %compile-run-and-check
+
+#include <stdio.h>
+#include <omp.h>
+
+#pragma omp requires unified_shared_memory
+
+double VAR_TO = 10.0;
+#pragma omp declare target to(VAR_TO)
+
+#pragma omp declare target
+double VAR_DT = 20.0;
+#pragma omp end declare target
+
+int main(int argc, char *argv[]) {
+  // This also runs reg requires for the first time.
+  long long device_addr_to, host_addr_to;
+  long long device_addr_dt, host_addr_dt;
+
+#pragma omp target map(tofrom: device_addr_to, device_addr_dt)
+  {
+    VAR_TO += 1;
+    VAR_DT += 1;
+    device_addr_to = (long long) &VAR_TO;
+    device_addr_dt = (long long) &VAR_DT;
+  }
+
+  host_addr_to = (long long) &VAR_TO;
+  host_addr_dt = (long long) &VAR_DT;
+
+  // CHECK: VAR_TO = 11
+  printf("VAR_TO = %f\n", VAR_TO);
+
+  // CHECK: VAR_DT = 21
+  printf("VAR_DT = %f\n", VAR_DT);
+
+  // CHECK: DEVICE_TO address is HOST_TO address.
+  if (device_addr_to == host_addr_to)
+    printf("DEVICE_TO address is HOST_TO address.\n");
+
+  // CHECK: DEVICE_DT address is HOST_DT address.
+  if (device_addr_dt == host_addr_dt)
+    printf("DEVICE_DT address is HOST_DT address.\n");
+
+  printf("DEVICE_ADDR = %lld\n", device_addr_to);
+  printf("HOST_ADDR = %lld\n", host_addr_to);
+
+  printf("DEVICE_ADDR = %lld\n", device_addr_dt);
+  printf("HOST_ADDR = %lld\n", host_addr_dt);
+
+  return 0;
+}

libomptarget/src/api.cpp

EXTERN int omp_target_is_present(void *ptr, int device_num) {
   if (Devices_size <= (size_t)device_num) {
     DP("Call to omp_target_is_present with invalid device ID, returning "
         "false\n");
     return false;
   }
 
   DeviceTy& Device = Devices[device_num];
   bool IsLast; // not used
-  int rc = (Device.getTgtPtrBegin(ptr, 0, IsLast, false) != NULL);
+  bool IsHostPtr;
+  void *TgtPtr = Device.getTgtPtrBegin(ptr, 0, IsLast, false, IsHostPtr);
+  int rc = (TgtPtr != NULL);
+  // Under unified memory the host pointer can be returned by the
+  // getTgtPtrBegin() function which means that there is no device
+  // corresponding point for ptr. This function should return false
+  // in that situation.
+  if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)
+    rc = !IsHostPtr;
   DP("Call to omp_target_is_present returns %d\n", rc);
   return rc;
 }
 
 EXTERN int omp_target_memcpy(void *dst, void *src, size_t length,
     size_t dst_offset, size_t src_offset, int dst_device, int src_device) {
   DP("Call to omp_target_memcpy, dst device %d, src device %d, "
       "dst addr " DPxMOD ", src addr " DPxMOD ", dst offset %zu, "

libomptarget/src/device.h

DeviceTy& operator=(const DeviceTy &d) {
     RTLRequiresFlags = d.RTLRequiresFlags;
 
     return *this;
   }
 
   long getMapEntryRefCnt(void *HstPtrBegin);
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
-      bool &IsNew, bool IsImplicit, bool UpdateRefCount = true);
+      bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true,
+      bool IsInUseDevicePtrClause = false);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-      bool UpdateRefCount);
-  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete);
+      bool UpdateRefCount, bool &IsHostPtr);
+  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
+                    bool IsInUseDevicePtrClause = false);
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);
 
   // calls to RTL
   int32_t initOnce();
   __tgt_target_table *load_binary(void *Img);
 
   int32_t data_submit(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size);

libomptarget/src/device.cpp

 
 // Used by target_data_begin
 // Return the target pointer begin (where the data will be moved).
 // Allocate memory if this is the first occurrence of this mapping.
 // Increment the reference counter.
 // If NULL is returned, then either data allocation failed or the user tried
 // to do an illegal mapping.
 void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
-    int64_t Size, bool &IsNew, bool IsImplicit, bool UpdateRefCount) {
+    int64_t Size, bool &IsNew, bool &IsHostPtr, bool IsImplicit,
+    bool UpdateRefCount, bool IsInUseDevicePtrClause) {
   void *rc = NULL;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
 
+  // If unified shared memory is active implicitly mapped variables that are not
+  // privatized, use host address. Any explicitely mapped variables also use
+  // host address where correctness is not impeded. In all other cases
+  // maps are respected.
+  // TODO: In addition to the mapping rules above, when the close map
+  // modifier is implemented, foce the mapping of the variable to the device.
+  if (RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+      !IsInUseDevicePtrClause &&
+      (IsImplicit || !((lr.Flags.IsContained ||
+                        lr.Flags.ExtendsBefore ||
+                        lr.Flags.ExtendsAfter) && !Size))) {
+    auto &HT = *lr.Entry;
+    uintptr_t tp = HT.HstPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+    DP("Return HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n", DPxPTR(tp),
+       Size, (UpdateRefCount ? " updated" : ""));
+    IsHostPtr = true;
+    rc = (void *)tp;
+  } else {
     // Check if the pointer is contained.
     if (lr.Flags.IsContained ||
         ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && IsImplicit)) {
       auto &HT = *lr.Entry;
       IsNew = false;
 
       if (UpdateRefCount)
         ++HT.RefCount;
 
       uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
       DP("Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
           "Size=%ld,%s RefCount=%s\n", (IsImplicit ? " (implicit)" : ""),
           DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
           (UpdateRefCount ? " updated" : ""),
           (CONSIDERED_INF(HT.RefCount)) ? "INF" :
               std::to_string(HT.RefCount).c_str());
       rc = (void *)tp;
     } else if ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && !IsImplicit) {
       // Explicit extension of mapped data - not allowed.
       DP("Explicit extension of mapping is not allowed.\n");
     } else if (Size) {
       // If it is not contained and Size > 0 we should create a new entry for it.
       IsNew = true;
-    uintptr_t tp = (uintptr_t)RTL->data_alloc(RTLDeviceID, Size, HstPtrBegin);
+      uintptr_t tp = (uintptr_t)NULL;
+      // If the use_device_ptr clause is used, we create a copy of the variable
+      // on the device even in the unified shared memory case.
+      // The explicit usage of the use_device_ptr clause forces the variable to
+      // have a device side version.
+      if (!(RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
+          (RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+           IsInUseDevicePtrClause))
+        tp = (uintptr_t)RTL->data_alloc(RTLDeviceID, Size, HstPtrBegin);
       DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", "
           "HstEnd=" DPxMOD ", TgtBegin=" DPxMOD "\n", DPxPTR(HstPtrBase),
           DPxPTR(HstPtrBegin), DPxPTR((uintptr_t)HstPtrBegin + Size), DPxPTR(tp));
       HostDataToTargetMap.push_front(HostDataToTargetTy((uintptr_t)HstPtrBase,
           (uintptr_t)HstPtrBegin, (uintptr_t)HstPtrBegin + Size, tp));
       rc = (void *)tp;
     }
+  }
 
   DataMapMtx.unlock();
   return rc;
 }
 
 // Used by target_data_begin, target_data_end, target_data_update and target.
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from target_data_end.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-    bool UpdateRefCount) {
+    bool UpdateRefCount, bool &IsHostPtr) {
   void *rc = NULL;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
+  IsHostPtr = false;
 
+  // If the value isn't found in the mapping and unified shared memory
+  // is on then it means we have stumbled upon an implcitly mapped value
+  // which we need to use directly from the host.
+  if (RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+      !(lr.Flags.IsContained ||
+        lr.Flags.ExtendsBefore ||
+        lr.Flags.ExtendsAfter)) {
+    auto &HT = *lr.Entry;
+    IsLast = false;
+    uintptr_t tp = HT.HstPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
+    DP("Get HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n", DPxPTR(tp),
+       Size, (UpdateRefCount ? " updated" : ""));
+    IsHostPtr = true;
+    rc = (void *)tp;
+  } else {
     if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
       auto &HT = *lr.Entry;
       IsLast = !(HT.RefCount > 1);
 
       if (HT.RefCount > 1 && UpdateRefCount)
         --HT.RefCount;
 
       uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
       DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
           "Size=%ld,%s RefCount=%s\n", DPxPTR(HstPtrBegin), DPxPTR(tp), Size,
           (UpdateRefCount ? " updated" : ""),
           (CONSIDERED_INF(HT.RefCount)) ? "INF" :
               std::to_string(HT.RefCount).c_str());
       rc = (void *)tp;
     } else {
       IsLast = false;
     }
+  }
 
   DataMapMtx.unlock();
   return rc;
 }
 
 // Return the target pointer begin (where the data will be moved).
 // Lock-free version called when loading global symbols from the fat binary.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size) {
   uintptr_t hp = (uintptr_t)HstPtrBegin;
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     uintptr_t tp = HT.TgtPtrBegin + (hp - HT.HstPtrBegin);
     return (void *)tp;
   }
 
   return NULL;
 }
 
-int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete) {
+int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
+                            bool IsInUseDevicePtrClause) {
+  if (RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+      !IsInUseDevicePtrClause)
+    return OFFLOAD_SUCCESS;
   // Check if the pointer is contained in any sub-nodes.
   int rc;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     if (ForceDelete)
       HT.RefCount = 1;

libomptarget/src/omptarget.cpp

if (member_of(arg_types[i]) < 0 && next_i < arg_num &&
         HstPtrBegin = (char *) HstPtrBegin - padding;
         data_size += padding;
       }
     }
 
     // Address of pointer on the host and device, respectively.
     void *Pointer_HstPtrBegin, *Pointer_TgtPtrBegin;
     bool IsNew, Pointer_IsNew;
+    bool IsHostPtr = false;
     bool IsImplicit = arg_types[i] & OMP_TGT_MAPTYPE_IMPLICIT;
+    // TODO: Check if this is correct
+    bool IsInUseDevicePtrClause = arg_types[i] & OMP_TGT_MAPTYPE_TARGET_PARAM &&
+        arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM;
     // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
     // have reached this point via __tgt_target_data_begin and not __tgt_target
     // then no argument is marked as TARGET_PARAM ("omp target data map" is not
     // associated with a target region, so there are no target parameters). This
     // may be considered a hack, we could revise the scheme in the future.
     bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF);
     if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       DP("Has a pointer entry: \n");
       // base is address of pointer.
+      // TODO: Check if IsInUseDevicePtrClause needs to ne passed.
       Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase,
-          sizeof(void *), Pointer_IsNew, IsImplicit, UpdateRef);
+          sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef,
+          IsInUseDevicePtrClause);
       if (!Pointer_TgtPtrBegin) {
         DP("Call to getOrAllocTgtPtr returned null pointer (device failure or "
             "illegal mapping).\n");
         return OFFLOAD_FAIL;
       }
       DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"
           "\n", sizeof(void *), DPxPTR(Pointer_TgtPtrBegin),
           (Pointer_IsNew ? "" : " not"));
       Pointer_HstPtrBegin = HstPtrBase;
       // modify current entry.
       HstPtrBase = *(void **)HstPtrBase;
       UpdateRef = true; // subsequently update ref count of pointee
     }
 
     void *TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBegin, HstPtrBase,
-        data_size, IsNew, IsImplicit, UpdateRef);
+        data_size, IsNew, IsHostPtr, IsImplicit, UpdateRef,
+        IsInUseDevicePtrClause);
     if (!TgtPtrBegin && data_size) {
       // If data_size==0, then the argument could be a zero-length pointer to
       // NULL, so getOrAlloc() returning NULL is not an error.
       DP("Call to getOrAllocTgtPtr returned null pointer (device failure or "
           "illegal mapping).\n");
     }
     DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
         " - is%s new\n", data_size, DPxPTR(TgtPtrBegin),
         (IsNew ? "" : " not"));
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
       uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
       void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
       DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
       args_base[i] = TgtPtrBase;
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
       bool copy = false;
+      if (!(Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)) {
         if (IsNew || (arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS)) {
           copy = true;
         } else if (arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) {
           // Copy data only if the "parent" struct has RefCount==1.
           int32_t parent_idx = member_of(arg_types[i]);
           long parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
           assert(parent_rc > 0 && "parent struct not found");
           if (parent_rc == 1) {
             copy = true;
           }
         }
+      }
 
-      if (copy) {
+      if (copy && !IsHostPtr) {
         DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",
             data_size, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
         int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, data_size);
         if (rt != OFFLOAD_SUCCESS) {
           DP("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
     }
 
-    if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
+    if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {
       DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n",
           DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin));
       uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
       void *TgtPtrBase = (void *)((uint64_t)TgtPtrBegin - Delta);
       int rt = Device.data_submit(Pointer_TgtPtrBegin, &TgtPtrBase,
           sizeof(void *));
       if (rt != OFFLOAD_SUCCESS) {
         DP("Copying data to device failed.\n");
if (member_of(arg_types[i]) < 0 && next_i < arg_num &&
       if (padding) {
         DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD
             "\n", padding, DPxPTR(HstPtrBegin));
         HstPtrBegin = (char *) HstPtrBegin - padding;
         data_size += padding;
       }
     }
 
-    bool IsLast;
+    bool IsLast, IsHostPtr;
     bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) ||
         (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ);
     bool ForceDelete = arg_types[i] & OMP_TGT_MAPTYPE_DELETE;
+    // TODO: Check if this is correct
+    bool IsInUseDevicePtrClause = arg_types[i] & OMP_TGT_MAPTYPE_TARGET_PARAM &&
+        arg_types[i] & OMP_TGT_MAPTYPE_RETURN_PARAM;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast,
-        UpdateRef);
+        UpdateRef, IsHostPtr);
     DP("There are %" PRId64 " bytes allocated at target address " DPxMOD
         " - is%s last\n", data_size, DPxPTR(TgtPtrBegin),
         (IsLast ? "" : " not"));
 
     bool DelEntry = IsLast || ForceDelete;
 
     if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
         !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
       DelEntry = false; // protect parent struct from being deallocated
     }
 
     if ((arg_types[i] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
       // Move data back to the host
       if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
         bool Always = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS;
         bool CopyMember = false;
+        if (!(Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)) {
           if ((arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF) &&
               !(arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {
             // Copy data only if the "parent" struct has RefCount==1.
             int32_t parent_idx = member_of(arg_types[i]);
             long parent_rc = Device.getMapEntryRefCnt(args[parent_idx]);
             assert(parent_rc > 0 && "parent struct not found");
             if (parent_rc == 1) {
               CopyMember = true;
             }
           }
+        }
 
         if (DelEntry || Always || CopyMember) {
           DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
               data_size, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
-          int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, data_size);
+          int rt;
+          if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+              TgtPtrBegin == HstPtrBegin)
+            rt = OFFLOAD_SUCCESS;
+          else
+            rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, data_size);
           if (rt != OFFLOAD_SUCCESS) {
             DP("Copying data from device failed.\n");
             return OFFLOAD_FAIL;
           }
         }
       }
 
       // If we copied back to the host a struct/array containing pointers, we
       // need to restore the original host pointer values from their shadow
       // copies. If the struct is going to be deallocated, remove any remaining
       // shadow pointer entries for this struct.
       uintptr_t lb = (uintptr_t) HstPtrBegin;
if ((arg_types[i] & OMP_TGT_MAPTYPE_FROM) || DelEntry) {
         } else {
           ++it;
         }
       }
       Device.ShadowMtx.unlock();
 
       // Deallocate map
       if (DelEntry) {
-        int rt = Device.deallocTgtPtr(HstPtrBegin, data_size, ForceDelete);
+        int rt = Device.deallocTgtPtr(HstPtrBegin, data_size, ForceDelete,
+                                      IsInUseDevicePtrClause);
         if (rt != OFFLOAD_SUCCESS) {
           DP("Deallocating data from device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
     }
   }
 
   return OFFLOAD_SUCCESS;
 }
 
 /// Internal function to pass data to/from the target.
 int target_data_update(DeviceTy &Device, int32_t arg_num,
     void **args_base, void **args, int64_t *arg_sizes, int64_t *arg_types) {
   // process each input.
   for (int32_t i = 0; i < arg_num; ++i) {
     if ((arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) ||
         (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE))
       continue;
 
     void *HstPtrBegin = args[i];
     int64_t MapSize = arg_sizes[i];
-    bool IsLast;
+    bool IsLast, IsHostPtr;
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, MapSize, IsLast,
-        false);
+        false, IsHostPtr);
     if (!TgtPtrBegin) {
       DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
       continue;
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_FROM) {
       DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
           arg_sizes[i], DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
-      int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, MapSize);
+      int rt;
+      // Act as if a copy to device was successful in the case of
+      // unified memory where only a host version of the data exists.
+      if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+          TgtPtrBegin == HstPtrBegin)
+        rt = OFFLOAD_SUCCESS;
+      else
+        rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, MapSize);
       if (rt != OFFLOAD_SUCCESS) {
         DP("Copying data from device failed.\n");
         return OFFLOAD_FAIL;
       }
 
       uintptr_t lb = (uintptr_t) HstPtrBegin;
       uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize;
       Device.ShadowMtx.lock();
       for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
           it != Device.ShadowPtrMap.end(); ++it) {
         void **ShadowHstPtrAddr = (void**) it->first;
         if ((uintptr_t) ShadowHstPtrAddr < lb)
           continue;
         if ((uintptr_t) ShadowHstPtrAddr >= ub)
           break;
         DP("Restoring original host pointer value " DPxMOD " for host pointer "
             DPxMOD "\n", DPxPTR(it->second.HstPtrVal),
             DPxPTR(ShadowHstPtrAddr));
         *ShadowHstPtrAddr = it->second.HstPtrVal;
       }
       Device.ShadowMtx.unlock();
     }
 
     if (arg_types[i] & OMP_TGT_MAPTYPE_TO) {
       DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",
           arg_sizes[i], DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));
-      int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, MapSize);
+      int rt;
+      // No need to actually copy any data to device. The data is available
+      // on the host and can be accessed by the device via unified memory.
+      // When host and device pointers are equal it means that this is called
+      // internally by the runtime and not directly by the user via the API.
+      if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+          TgtPtrBegin == HstPtrBegin)
+        rt = OFFLOAD_SUCCESS;
+      else
+        rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, MapSize);
       if (rt != OFFLOAD_SUCCESS) {
         DP("Copying data to device failed.\n");
         return OFFLOAD_FAIL;
       }
       uintptr_t lb = (uintptr_t) HstPtrBegin;
       uintptr_t ub = (uintptr_t) HstPtrBegin + MapSize;
       Device.ShadowMtx.lock();
       for (ShadowPtrListTy::iterator it = Device.ShadowPtrMap.begin();
           it != Device.ShadowPtrMap.end(); ++it) {
         void **ShadowHstPtrAddr = (void**) it->first;
         if ((uintptr_t) ShadowHstPtrAddr < lb)
           continue;
         if ((uintptr_t) ShadowHstPtrAddr >= ub)
           break;
         DP("Restoring original target pointer value " DPxMOD " for target "
             "pointer " DPxMOD "\n", DPxPTR(it->second.TgtPtrVal),
             DPxPTR(it->second.TgtPtrAddr));
+        if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+            TgtPtrBegin == HstPtrBegin)
+          rt = OFFLOAD_SUCCESS;
+        else
           rt = Device.data_submit(it->second.TgtPtrAddr,
               &it->second.TgtPtrVal, sizeof(void *));
         if (rt != OFFLOAD_SUCCESS) {
           DP("Copying data to device failed.\n");
           Device.ShadowMtx.unlock();
           return OFFLOAD_FAIL;
         }
       }
       Device.ShadowMtx.unlock();
     }
   }
   return OFFLOAD_SUCCESS;
 }
 
 static const unsigned LambdaMapping = OMP_TGT_MAPTYPE_PTR_AND_OBJ |
                                       OMP_TGT_MAPTYPE_LITERAL |
                                       OMP_TGT_MAPTYPE_IMPLICIT;
if (!(arg_types[i] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {
         unsigned idx = member_of(arg_types[i]);
         int tgtIdx = tgtArgsPositions[idx];
         assert(tgtIdx != -1 && "Base address must be translated already.");
         // The parent lambda must be processed already and it must be the last
         // in tgt_args and tgt_offsets arrays.
         void *HstPtrVal = args[i];
         void *HstPtrBegin = args_base[i];
         void *HstPtrBase = args[idx];
-        bool IsLast; // unused.
+        bool IsLast, IsHostPtr; // unused.
         void *TgtPtrBase =
             (void *)((intptr_t)tgt_args[tgtIdx] + tgt_offsets[tgtIdx]);
         DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase));
         uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
         void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
         void *Pointer_TgtPtrBegin =
-            Device.getTgtPtrBegin(HstPtrVal, arg_sizes[i], IsLast, false);
+            Device.getTgtPtrBegin(HstPtrVal, arg_sizes[i], IsLast, false,
+                                  IsHostPtr);
         if (!Pointer_TgtPtrBegin) {
           DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
              DPxPTR(HstPtrVal));
           continue;
         }
         DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n",
            DPxPTR(Pointer_TgtPtrBegin), DPxPTR(TgtPtrBegin));
-        int rt = Device.data_submit(TgtPtrBegin, &Pointer_TgtPtrBegin,
-                                    sizeof(void *));
+        int rt;
+        if (Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+            TgtPtrBegin == HstPtrBegin)
+          rt = OFFLOAD_SUCCESS;
+        else
+          rt = Device.data_submit(TgtPtrBegin, &Pointer_TgtPtrBegin,
         if (rt != OFFLOAD_SUCCESS) {
           DP("Copying data to device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
       continue;
     }
     void *HstPtrBegin = args[i];
     void *HstPtrBase = args_base[i];
     void *TgtPtrBegin;
     ptrdiff_t TgtBaseOffset;
-    bool IsLast; // unused.
+    bool IsLast, IsHostPtr; // unused.
     if (arg_types[i] & OMP_TGT_MAPTYPE_LITERAL) {
       DP("Forwarding first-private value " DPxMOD " to the target construct\n",
           DPxPTR(HstPtrBase));
       TgtPtrBegin = HstPtrBase;
       TgtBaseOffset = 0;
     } else if (arg_types[i] & OMP_TGT_MAPTYPE_PRIVATE) {
       // Allocate memory for (first-)private array
       TgtPtrBegin = Device.RTL->data_alloc(Device.RTLDeviceID,
#endif
         int rt = Device.data_submit(TgtPtrBegin, HstPtrBegin, arg_sizes[i]);
         if (rt != OFFLOAD_SUCCESS) {
           DP ("Copying data to device failed, failed.\n");
           return OFFLOAD_FAIL;
         }
       }
     } else if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {
       TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast,
-          false);
+          false, IsHostPtr);
       TgtBaseOffset = 0; // no offset for ptrs.
       DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD " to "
          "object " DPxMOD "\n", DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBase),
          DPxPTR(HstPtrBase));
     } else {
       TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast,
-          false);
+          false, IsHostPtr);
       TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
 #ifdef OMPTARGET_DEBUG
       void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);
       DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",
           DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));
 #endif
     }
     tgtArgsPositions[i] = tgt_args.size();