[OpenMP][libomptarget] Add support for close map modifier

[gtbercea]

Summary

This patch adds support for the close map modifier.

The close map modifier will overwrite the unified shared memory requirement and create a device copy of the data.

[Hahnfeld]

I think this mostly looks good, but depends on D65001, right?

[gtbercea]

In D65340#1604437, @Hahnfeld wrote:

I think this mostly looks good, but depends on D65001, right?

It does yes. I won't push it until the unified memory stuff is done.

[AlexEichenberger]

LGTM

[gtbercea]

• Add close modifier flag.
• Move test to appropriate folder.

[Hahnfeld]

I think this still needs some work.

libomptarget/src/device.cpp

166–167

Please update this comment.

206–207

Why do we need this check? If logic doesn't fail me, this should always be true with the condition above.

libomptarget/src/device.h

136

I'd call this HasCloseModifier (here and in other places)

libomptarget/src/omptarget.cpp

299

Please use IsCloseModifier / HasCloseModifier from above.

301

Can you remove the second Device.RTLRequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY? Either it's not set, or it's set and the condition needs to check if the mapping has the close modifier.

398

Same here, use IsCloseModifier / HasCloseModifier from above.

400

Same, logic can be simplified.

413–415

Can we have TgtPtrBegin == HstPtrBegin when the mapping has the close modifier?

libomptarget/test/unified_shared_memory/close_modifier.c

18

Please use int as in the other tests.

21

Please format.

26–27

Please don't cast to long long!

[gtbercea]

• Address comments.

[gtbercea]

• Update.

[Hahnfeld]

libomptarget/src/omptarget.cpp

301–302

Hmm, thinking about this again: Does close imply always? I don't find this in the spec, so I think HasCloseModifier should not be in the inner condition.

413

Same here.

libomptarget/test/unified_shared_memory/close_modifier.c

21

Ping, test is still not formatted.

[Hahnfeld]

I'd propose to mark the test close_modifier.c as unsupported with older compiler versions. Additionally, can you add a small, new test with manual calls to __tgt_target_data_begin / _end without relying on the compiler?

libomptarget/test/unified_shared_memory/close_modifier.c

29–31

Unfortunately, I don't think that's enough for this test to pass with older versions of Clang because it didn't support code generation for close. Did you test this?

[gtbercea]

libomptarget/src/omptarget.cpp

301–302

I think I see what you mean. The issue is not that this is in the innermost loop but that it may generate a copy to the device when IsNew is false.

[Hahnfeld]

libomptarget/src/omptarget.cpp

301–302

Yes, with "innermost condition" I meant the inner if statement, the one in the lines I marked.

[gtbercea]

• Format test.

[gtbercea]

libomptarget/test/unified_shared_memory/close_modifier.c

29–31

I don't think it will either but I'm not sure how we can actually make it work. I think there's not much more this test can do. To make this work probably requires a more forgiving check in the compiler.

[gtbercea]

• Add back close modifier flag check.

[gtbercea]

• Remove close flag from checks where it is not needed.

[gtbercea]

@Hahnfeld I have addressed your comments, please let me know if you have further comments.

[Hahnfeld]

The test will currently fail with older versions of Clang. It must at least be marked UNSUPPORTED for Clang versions older than what-will-be Clang 10.

Additionally, I'd still like to see a small test with manual calls to __tgt_target_data_begin / _end if possible. I'm thinking of the following:

int *a = // malloc
int *device;

__tgt_target_data_begin(...); // with close modifier for a
#pragma omp target data use_device_ptr(a)
{
  device = a;
}
__tgt_target_data_end(...); // correspondingly to the begin call above

// check...

Instead of the nested target data region, you could also use a simple target region in the outer data region, and have a check for use_device_ptr in another test.

[gtbercea]

In D65340#1616793, @Hahnfeld wrote:

The test will currently fail with older versions of Clang. It must at least be marked UNSUPPORTED for Clang versions older than what-will-be Clang 10.

Additionally, I'd still like to see a small test with manual calls to __tgt_target_data_begin / _end if possible. I'm thinking of the following:

int *a = // malloc
int *device;

__tgt_target_data_begin(...); // with close modifier for a
#pragma omp target data use_device_ptr(a)
{
  device = a;
}
__tgt_target_data_end(...); // correspondingly to the begin call above

// check...

Instead of the nested target data region, you could also use a simple target region in the outer data region, and have a check for use_device_ptr in another test.

I don't mind adding such a test but calling these functions manually is very tedious and error prone and hard to maintain. Are there any benefits I'm missing to this apart from running this test in isolation?

[Hahnfeld]

In D65340#1616943, @gtbercea wrote:

In D65340#1616793, @Hahnfeld wrote:

The test will currently fail with older versions of Clang. It must at least be marked UNSUPPORTED for Clang versions older than what-will-be Clang 10.

Additionally, I'd still like to see a small test with manual calls to __tgt_target_data_begin / _end if possible. I'm thinking of the following:

int *a = // malloc
int *device;

__tgt_target_data_begin(...); // with close modifier for a
#pragma omp target data use_device_ptr(a)
{
  device = a;
}
__tgt_target_data_end(...); // correspondingly to the begin call above

// check...

Instead of the nested target data region, you could also use a simple target region in the outer data region, and have a check for use_device_ptr in another test.

I don't mind adding such a test but calling these functions manually is very tedious and error prone and hard to maintain. Are there any benefits I'm missing to this apart from running this test in isolation?

It adds coverage with older compiler versions. If it's too complicated, I'm also fine with adding a check for use_device_ptr and restrict the test with close modifier to newer versions of Clang, up to you.

[gtbercea]

• Add unsupported clang versions to test.
• Add more tests.

[gtbercea]

• Clean-up.

[gtbercea]

• Clean test.

[gtbercea]

@Hahnfeld I have updated the tests, I added two new ones. Let me know if you have any further comments.

[gtbercea]

• Update.

[gtbercea]

@grokos I have updated this patch as suggested and have added new tests. Please let me know if you any further comments.

[gtbercea]

@Hahnfeld @grokos any further comments?

[grokos]

I don't have any comments about this patch other than the ones @Hahnfeld mentioned.

[gtbercea]

In D65340#1621626, @grokos wrote:

I don't have any comments about this patch other than the ones @Hahnfeld mentioned.

Great, @Hahnfeld please let me know if you have any other comments.

[Hahnfeld]

libomptarget/test/unified_shared_memory/close_enter_exit.c

7

I think you can safely remove clang-3 to clang-5 because CMake will block tests with anything before clang-6.

libomptarget/test/unified_shared_memory/close_manual.c

7

This completely misses the point: This test is supposed to work with older versions of clang, even those that don't support close.

47–53

Why do we need this?

libomptarget/test/unified_shared_memory/close_modifier.c

7

Same as above.

[gtbercea]

libomptarget/test/unified_shared_memory/close_manual.c

47–53

To initialize the device.

[gtbercea]

libomptarget/test/unified_shared_memory/close_manual.c

7

This is because of __tgt_register_requires not because of the other two methods.

[gtbercea]

• Remove requires manual call.

[gtbercea]

libomptarget/test/unified_shared_memory/close_manual.c

7

I believe clang-9 will work with unified shared memory. Because we are calling the tgt functions manually we can enable this test for clang-9 even if it tests the close functionality.

[gtbercea]

@Hahnfeld
I would like to clarify something about tgt_register_requires call.

This call is special and in a compiler that supports handling of requires clauses (like unified_shared_memory clause) ( i.e. Clang 9.0 onwards), this function is called before any main function code. Here is a printout with LIBOMPTARGET_DEBUG=1 to realize the difference:

===== WHERE COMPILER PLACES CALL TO __tgt_register_requires() ========
Libomptarget --> Loading RTLs... <<<< THIS IS THE VERY FIRST PRINT FROM THE RUNTIME. The call to __tgt_register_requires is just before that.
Libomptarget --> Loading library 'libomptarget.rtl.ppc64.so'...
Libomptarget --> Successfully loaded library 'libomptarget.rtl.ppc64.so'!
Libomptarget --> Registering RTL libomptarget.rtl.ppc64.so supporting 4 devices!
Libomptarget --> Loading library 'libomptarget.rtl.x86_64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.x86_64.so': libomptarget.rtl.x86_64.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.cuda.so'...
Target CUDA RTL --> Start initializing CUDA
Libomptarget --> Successfully loaded library 'libomptarget.rtl.cuda.so'!
Libomptarget --> Registering RTL libomptarget.rtl.cuda.so supporting 4 devices!
Libomptarget --> Loading library 'libomptarget.rtl.aarch64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.aarch64.so': libomptarget.rtl.aarch64.so: cannot open shared object file: No such file or directory!
Libomptarget --> RTLs loaded!
Libomptarget --> Image 0x0000000010020080 is compatible with RTL libomptarget.rtl.ppc64.so!
Libomptarget --> RTL 0x000000004747b3e0 has index 0!
Libomptarget --> Registering image 0x0000000010020080 with RTL libomptarget.rtl.ppc64.so!
Libomptarget --> Done registering entries!
===== MANUAL CALL TO __tgt_register_requires() IN MAIN (top of MAIN function!) ========

Manually calling __tgt_register_requires can never actually work correctly because any user space call to this function will be after the binary has been loaded and the flags already decided. The compiler just has to support the handling of the requires clauses.

[gtbercea]

libomptarget/test/unified_shared_memory/close_enter_exit.c

7

In this test, clang-9 needs to be unsupported because close is used explicitly in the code.

[Hahnfeld]

In D65340#1623236, @gtbercea wrote:

@Hahnfeld
I would like to clarify something about tgt_register_requires call.

This call is special and in a compiler that supports handling of requires clauses (like unified_shared_memory clause) ( i.e. Clang 9.0 onwards), this function is called before any main function code. Here is a printout with LIBOMPTARGET_DEBUG=1 to realize the difference:

===== WHERE COMPILER PLACES CALL TO __tgt_register_requires() ========
Libomptarget --> Loading RTLs... <<<< THIS IS THE VERY FIRST PRINT FROM THE RUNTIME. The call to __tgt_register_requires is just before that.
Libomptarget --> Loading library 'libomptarget.rtl.ppc64.so'...
Libomptarget --> Successfully loaded library 'libomptarget.rtl.ppc64.so'!
Libomptarget --> Registering RTL libomptarget.rtl.ppc64.so supporting 4 devices!
Libomptarget --> Loading library 'libomptarget.rtl.x86_64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.x86_64.so': libomptarget.rtl.x86_64.so: cannot open shared object file: No such file or directory!
Libomptarget --> Loading library 'libomptarget.rtl.cuda.so'...
Target CUDA RTL --> Start initializing CUDA
Libomptarget --> Successfully loaded library 'libomptarget.rtl.cuda.so'!
Libomptarget --> Registering RTL libomptarget.rtl.cuda.so supporting 4 devices!
Libomptarget --> Loading library 'libomptarget.rtl.aarch64.so'...
Libomptarget --> Unable to load library 'libomptarget.rtl.aarch64.so': libomptarget.rtl.aarch64.so: cannot open shared object file: No such file or directory!
Libomptarget --> RTLs loaded!
Libomptarget --> Image 0x0000000010020080 is compatible with RTL libomptarget.rtl.ppc64.so!
Libomptarget --> RTL 0x000000004747b3e0 has index 0!
Libomptarget --> Registering image 0x0000000010020080 with RTL libomptarget.rtl.ppc64.so!
Libomptarget --> Done registering entries!
===== MANUAL CALL TO __tgt_register_requires() IN MAIN (top of MAIN function!) ========

Manually calling __tgt_register_requires can never actually work correctly because any user space call to this function will be after the binary has been loaded and the flags already decided. The compiler just has to support the handling of the requires clauses.

I disagree, the information about requires directives only needs to be present when initializing the device, not when loading a library. I've posted D66019 to solve this and fix the test I mentioned in D65001.

libomptarget/test/unified_shared_memory/close_manual.c

47–53

This should happen in __tgt_target_data_begin which can be a legitimate first directive in a program.

[gtbercea]

In D65340#1623325, @Hahnfeld wrote:

I disagree, the information about requires directives only needs to be present when initializing the device, not when loading a library. I've posted D66019 to solve this and fix the test I mentioned in D65001.

Ok I'll add back the calls to the register_requires functions.

[gtbercea]

libomptarget/test/unified_shared_memory/close_manual.c

47–53

Yes target data can be a first directive but if there's no target region in the program then the runtime library won't be loaded/initialized.

[gtbercea]

• Add back requires registration call.

[Hahnfeld]

Can you please rebase on top of D66019? Sorry, that'll probably give you some conflicts :-/

libomptarget/test/unified_shared_memory/close_manual.c

47–53

Ah, got the intention. Can you please add a comment explaining that we need at least one target region to load a "binary" and make all of the data calls work? And I don't think we need to check that device_alloc is correctly transferred, do we? There are other tests for that (plus casting 10 to void * looks evil anyway).

[gtbercea]

• Rebase.

[gtbercea]

libomptarget/test/unified_shared_memory/close_manual.c

47–53

Of course, I'll add that.
Yeah I'll get rid of the check and the cast! :)

[Hahnfeld]

LG after addressing the last minor nits (comment + cast + check).

[gtbercea]

• Simplify target region.

libomptarget/include/omptarget.h

enum tgt_map_type {
   // return base device address of mapped data
   OMP_TGT_MAPTYPE_RETURN_PARAM    = 0x040,
   // private variable - not mapped
   OMP_TGT_MAPTYPE_PRIVATE         = 0x080,
   // copy by value - not mapped
   OMP_TGT_MAPTYPE_LITERAL         = 0x100,
   // mapping is implicit
   OMP_TGT_MAPTYPE_IMPLICIT        = 0x200,
+  // copy data to device
+  OMP_TGT_MAPTYPE_CLOSE           = 0x400,
   // member of struct, member given by [16 MSBs] - 1
   OMP_TGT_MAPTYPE_MEMBER_OF       = 0xffff000000000000
 };
 
 enum OpenMPOffloadingDeclareTargetFlags {
   /// Mark the entry as having a 'link' attribute.
   OMP_DECLARE_TARGET_LINK = 0x01,
   /// Mark the entry as being a global constructor.

libomptarget/src/device.h

DeviceTy& operator=(const DeviceTy &d) {
     LoopTripCnt = d.LoopTripCnt;
 
     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 &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true);
+      bool &IsNew, bool &IsHostPtr, bool IsImplicit, bool UpdateRefCount = true,
+      bool HasCloseModifier = false);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
       bool UpdateRefCount, bool &IsHostPtr);
-  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete);
+  int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
+                    bool HasCloseModifier = 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 &IsHostPtr, bool IsImplicit,
-    bool UpdateRefCount) {
+    bool UpdateRefCount, bool HasCloseModifier) {
   void *rc = NULL;
   IsHostPtr = false;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
 
   // Check if the pointer is contained.
   // If a variable is mapped to the device manually by the user - which would
   // lead to the IsContained flag to be true - then we must ensure that the
   // device address is returned even under unified memory conditions.
   if (lr.Flags.IsContained ||
       ((lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) && IsImplicit)) {
     auto &HT = *lr.Entry;
     IsNew = false;
 
void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
   } 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 unified shared memory is active, implicitly mapped variables that are not
     // privatized use host address. Any explicitly 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 (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
+    // In addition to the mapping rules above, the close map
+    // modifier forces the mapping of the variable to the device.
+    if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier) {
       DP("Return HstPtrBegin " DPxMOD " Size=%ld RefCount=%s\n",
          DPxPTR((uintptr_t)HstPtrBegin), Size, (UpdateRefCount ? " updated" : ""));
       IsHostPtr = true;
       rc = HstPtrBegin;
     } else {
       // 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);
       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;
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) {
-  if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)
+int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
+                            bool HasCloseModifier) {
+  if (RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier)
     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)

libomptarget/src/omptarget.cpp

if (member_of(arg_types[i]) < 0 && next_i < arg_num &&
       }
     }
 
     // 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;
+    // Force the creation of a device side copy of the data when:
+    // a close map modifier was associated with a map that contained a to.
+    bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
     // 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.
       Pointer_TgtPtrBegin = Device.getOrAllocTgtPtr(HstPtrBase, HstPtrBase,
-          sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef);
+          sizeof(void *), Pointer_IsNew, IsHostPtr, IsImplicit, UpdateRef,
+          HasCloseModifier);
       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, IsHostPtr, IsImplicit, UpdateRef);
+        data_size, IsNew, IsHostPtr, IsImplicit, UpdateRef, HasCloseModifier);
     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 (!(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)) {
+      if (!(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
+          HasCloseModifier) {
         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 (member_of(arg_types[i]) < 0 && next_i < arg_num &&
         data_size += padding;
       }
     }
 
     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;
+    bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
 
     // If PTR_AND_OBJ, HstPtrBegin is address of pointee
     void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, data_size, IsLast,
         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 (!(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY)) {
+        if (!(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
+            HasCloseModifier) {
           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) &&
             !(RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
               TgtPtrBegin == HstPtrBegin)) {
           DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",
               data_size, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));
           int rt = Device.data_retrieve(HstPtrBegin, TgtPtrBegin, data_size);
           if (rt != OFFLOAD_SUCCESS) {
             DP("Copying data from device failed.\n");
             return OFFLOAD_FAIL;
           }
         }
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,
+                                      HasCloseModifier);
         if (rt != OFFLOAD_SUCCESS) {
           DP("Deallocating data from device failed.\n");
           return OFFLOAD_FAIL;
         }
       }
     }
   }
 

libomptarget/test/unified_shared_memory/close_enter_exit.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc = 0, *device_alloc = 0;
+  int *a = (int *)malloc(N * sizeof(int));
+
+  // Init
+  for (int i = 0; i < N; ++i) {
+    a[i] = 10;
+  }
+  host_alloc = &a[0];
+
+  //
+  // map + target no close
+  //
+#pragma omp target data map(tofrom : a[ : N]) map(tofrom : device_alloc)
+  {
+#pragma omp target map(tofrom : device_alloc)
+    { device_alloc = &a[0]; }
+  }
+
+  // CHECK: a used from unified memory.
+  if (device_alloc == host_alloc)
+    printf("a used from unified memory.\n");
+
+  //
+  // map + target with close
+  //
+  device_alloc = 0;
+#pragma omp target data map(close, tofrom : a[ : N]) map(tofrom : device_alloc)
+  {
+#pragma omp target map(tofrom : device_alloc)
+    { device_alloc = &a[0]; }
+  }
+  // CHECK: a copied to device.
+  if (device_alloc != host_alloc)
+    printf("a copied to device.\n");
+
+  //
+  // map + use_device_ptr no close
+  //
+  device_alloc = 0;
+#pragma omp target data map(tofrom : a[ : N]) use_device_ptr(a)
+  { device_alloc = &a[0]; }
+
+  // CHECK: a used from unified memory with use_device_ptr.
+  if (device_alloc == host_alloc)
+    printf("a used from unified memory with use_device_ptr.\n");
+
+  //
+  // map + use_device_ptr close
+  //
+  device_alloc = 0;
+#pragma omp target data map(close, tofrom : a[ : N]) use_device_ptr(a)
+  { device_alloc = &a[0]; }
+
+  // CHECK: a used from device memory with use_device_ptr.
+  if (device_alloc != host_alloc)
+    printf("a used from device memory with use_device_ptr.\n");
+
+  //
+  // map enter/exit + close
+  //
+  device_alloc = 0;
+#pragma omp target enter data map(close, to : a[ : N])
+
+#pragma omp target map(from : device_alloc)
+  { device_alloc = &a[0]; }
+
+#pragma omp target exit data map(from : a[ : N])
+
+  // CHECK: a has been mapped to the device.
+  if (device_alloc != host_alloc)
+    printf("a has been mapped to the device.\n");
+
+  free(a);
+
+  // CHECK: Done!
+  printf("Done!\n");
+
+  return 0;
+}

libomptarget/test/unified_shared_memory/close_manual.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+#include <omp.h>
+#include <stdio.h>
+
+// ---------------------------------------------------------------------------
+// Various definitions copied from OpenMP RTL
+
+extern void __tgt_register_requires(int64_t);
+
+extern void __tgt_target_data_begin(int64_t device_id, int32_t arg_num,
+                                    void **args_base, void **args,
+                                    int64_t *arg_sizes, int64_t *arg_types);
+
+extern void __tgt_target_data_end(int64_t device_id, int32_t arg_num,
+                                  void **args_base, void **args,
+                                  int64_t *arg_sizes, int64_t *arg_types);
+
+// End of definitions copied from OpenMP RTL.
+// ---------------------------------------------------------------------------
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc = 0, *device_alloc = 0;
+  int *a = (int *)malloc(N * sizeof(int));
+
+  // Manual registration of requires flags for Clang versions
+  // that do not support requires.
+  __tgt_register_requires(8);
+
+  // Init
+  for (int i = 0; i < N; ++i) {
+    a[i] = 10;
+  }
+  host_alloc = &a[0];
+
+// Dummy target region that ensures the runtime library is loaded when
+// the target data begin/end functions are manually called below.
+#pragma omp target
+  {}
+
+  // Manual calls
+  int device_id = omp_get_default_device();
+  int arg_num = 1;
+  void **args_base = (void **)&a;
+  void **args = (void **)&a;
+  int64_t arg_sizes[arg_num];
+
+  arg_sizes[0] = sizeof(int) * N;
+
+  int64_t arg_types[arg_num];
+
+  // Ox400 enables the CLOSE map type in the runtime:
+  // OMP_TGT_MAPTYPE_CLOSE = 0x400
+  // OMP_TGT_MAPTYPE_TO    = 0x001
+  arg_types[0] = 0x400 | 0x001;
+
+  device_alloc = host_alloc;
+
+  __tgt_target_data_begin(device_id, arg_num, args_base, args, arg_sizes,
+                          arg_types);
+
+#pragma omp target data use_device_ptr(a)
+  { device_alloc = a; }
+
+  __tgt_target_data_end(device_id, arg_num, args_base, args, arg_sizes,
+                        arg_types);
+
+  // CHECK: a was copied to the device
+  if (device_alloc != host_alloc)
+    printf("a was copied to the device\n");
+
+  free(a);
+
+  // CHECK: Done!
+  printf("Done!\n");
+
+  return 0;
+}

libomptarget/test/unified_shared_memory/close_modifier.c

+// RUN: %libomptarget-compile-run-and-check-aarch64-unknown-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-powerpc64le-ibm-linux-gnu
+// RUN: %libomptarget-compile-run-and-check-x86_64-pc-linux-gnu
+
+// UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
+
+#include <omp.h>
+#include <stdio.h>
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int fails;
+  void *host_alloc, *device_alloc;
+  void *host_data, *device_data;
+  int *alloc = (int *)malloc(N * sizeof(int));
+  int data[N];
+
+  for (int i = 0; i < N; ++i) {
+    alloc[i] = 10;
+    data[i] = 1;
+  }
+
+  host_data = &data[0];
+  host_alloc = &alloc[0];
+
+//
+// Test that updates on the device are not visible to host
+// when only a TO mapping is used.
+//
+#pragma omp target map(tofrom                                                  \
+                       : device_data, device_alloc) map(close, to              \
+                                                        : alloc[:N], data      \
+                                                        [:N])
+  {
+    device_data = &data[0];
+    device_alloc = &alloc[0];
+
+    for (int i = 0; i < N; i++) {
+      alloc[i] += 1;
+      data[i] += 1;
+    }
+  }
+
+  // CHECK: Address of alloc on device different from host address.
+  if (device_alloc != host_alloc)
+    printf("Address of alloc on device different from host address.\n");
+
+  // CHECK: Address of data on device different from host address.
+  if (device_data != host_data)
+    printf("Address of data on device different from host address.\n");
+
+  // On the host, check that the arrays have been updated.
+  // CHECK: Alloc host values not updated: Succeeded
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (alloc[i] != 10)
+      fails++;
+  }
+  printf("Alloc host values not updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  // CHECK: Data host values not updated: Succeeded
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (data[i] != 1)
+      fails++;
+  }
+  printf("Data host values not updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  //
+  // Test that updates on the device are visible on host
+  // when a from is used.
+  //
+
+  for (int i = 0; i < N; i++) {
+    alloc[i] += 1;
+    data[i] += 1;
+  }
+
+#pragma omp target map(close, tofrom : alloc[:N], data[:N])
+  {
+    // CHECK: Alloc device values are correct: Succeeded
+    fails = 0;
+    for (int i = 0; i < N; i++) {
+      if (alloc[i] != 11)
+        fails++;
+    }
+    printf("Alloc device values are correct: %s\n",
+           (fails == 0) ? "Succeeded" : "Failed");
+    // CHECK: Data device values are correct: Succeeded
+    fails = 0;
+    for (int i = 0; i < N; i++) {
+      if (data[i] != 2)
+        fails++;
+    }
+    printf("Data device values are correct: %s\n",
+           (fails == 0) ? "Succeeded" : "Failed");
+
+    // Update values on the device
+    for (int i = 0; i < N; i++) {
+      alloc[i] += 1;
+      data[i] += 1;
+    }
+  }
+
+  // CHECK: Alloc host values updated: Succeeded
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (alloc[i] != 12)
+      fails++;
+  }
+  printf("Alloc host values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  // CHECK: Data host values updated: Succeeded
+  fails = 0;
+  for (int i = 0; i < N; i++) {
+    if (data[i] != 3)
+      fails++;
+  }
+  printf("Data host values updated: %s\n",
+         (fails == 0) ? "Succeeded" : "Failed");
+
+  free(alloc);
+
+  // CHECK: Done!
+  printf("Done!\n");
+
+  return 0;
+}