[Openmp-commits] [openmp] [OpenMP] Support for global variables when in auto zero-copy. (PR #80876)

[via Openmp-commits]

github-actions[bot] wrote:

<!--LLVM CODE FORMAT COMMENT: {clang-format}-->


:warning: C/C++ code formatter, clang-format found issues in your code. :warning:

<details>
<summary>
You can test this locally with the following command:
</summary>

``````````bash
git-clang-format --diff 93a2a8cb7f6ab815849e8320bff54c965edd09e7 8dcc66f6c7a7ceb0e8a1a7cf64ce82acc13b7429 -- openmp/libomptarget/test/mapping/auto_zero_copy_globals.cpp openmp/libomptarget/src/omptarget.cpp
``````````

</details>

<details>
<summary>
View the diff from clang-format here.
</summary>

``````````diff
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index 1a5ed7d663..821669d214 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -189,7 +189,7 @@ static int initLibrary(DeviceTy &Device) {
           // reference to the host global. We need to initialize the pointer on
           // the deive to point to the memory on the host.
           if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY) ||
-	      (PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {
+              (PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) {
             if (Device.RTL->data_submit(DeviceId, DeviceEntry.addr, Entry.addr,
                                         Entry.size) != OFFLOAD_SUCCESS)
               REPORT("Failed to write symbol for USM %s\n", Entry.name);
diff --git a/openmp/libomptarget/test/mapping/auto_zero_copy_globals.cpp b/openmp/libomptarget/test/mapping/auto_zero_copy_globals.cpp
index c6b7f80696..4a13d270ae 100644
--- a/openmp/libomptarget/test/mapping/auto_zero_copy_globals.cpp
+++ b/openmp/libomptarget/test/mapping/auto_zero_copy_globals.cpp
@@ -14,65 +14,72 @@
 
 // clang-format on
 
-#include <cstdio>
 #include <cstdint>
+#include <cstdio>
 
 /// Test for globals under automatic zero-copy.
 /// Because we are building without unified_shared_memory
 /// requirement pragma, all globals are allocated in the device
-/// memory of all used GPUs. To ensure those globals contain the intended values, we
-/// need to execute H2D and D2H memory copies even if we are running in automatic zero-copy.
-/// This only applies to globals. Local variables (their host pointers) are passed to the kernels by-value,
-/// according to the automatic zero-copy behavior.
+/// memory of all used GPUs. To ensure those globals contain the intended
+/// values, we need to execute H2D and D2H memory copies even if we are running
+/// in automatic zero-copy. This only applies to globals. Local variables (their
+/// host pointers) are passed to the kernels by-value, according to the
+/// automatic zero-copy behavior.
 
 #pragma omp begin declare target
-int32_t x;      // 4 bytes
-int32_t z[10];  // 40 bytes
-int32_t *k;     // 20 bytes
+int32_t x;     // 4 bytes
+int32_t z[10]; // 40 bytes
+int32_t *k;    // 20 bytes
 #pragma omp end declare target
 
 int main() {
   int32_t *dev_k = nullptr;
   x = 3;
   int32_t y = -1;
-  for(size_t t = 0; t < 10; t++)
+  for (size_t t = 0; t < 10; t++)
     z[t] = t;
   k = new int32_t[5];
 
   printf("Host pointer for k = %p\n", k);
-  for(size_t t = 0; t < 5; t++)
+  for (size_t t = 0; t < 5; t++)
     k[t] = -t;
 
-  /// target update to forces a copy between host and device global, which we must execute to
-  /// keep the two global copies consistent.
-  /// CHECK: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=40, Name=z
-  #pragma omp target update to(z[:10])
+/// target update to forces a copy between host and device global, which we must
+/// execute to keep the two global copies consistent. CHECK: Copying data from
+/// host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=40, Name=z
+#pragma omp target update to(z[ : 10])
 
-  /// target map with always modifier (for x) forces a copy between host and device global, which we must execute to
-  /// keep the two global copies consistent.
-  /// k's content (host address) is passed by-value to the kernel (Size=20 case).
-  /// y, being a local variable, is also passed by-value to the kernel (Size=4 case)
-  /// CHECK: Return HstPtrBegin {{.*}} Size=4 for unified shared memory
-  /// CHECK: Return HstPtrBegin {{.*}} Size=20 for unified shared memory
-  /// CHECK: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4, Name=x
-  #pragma omp target map(to:k[:5]) map(always, tofrom:x) map(tofrom:y) map(from:dev_k)
+/// target map with always modifier (for x) forces a copy between host and
+/// device global, which we must execute to keep the two global copies
+/// consistent. k's content (host address) is passed by-value to the kernel
+/// (Size=20 case). y, being a local variable, is also passed by-value to the
+/// kernel (Size=4 case) CHECK: Return HstPtrBegin {{.*}} Size=4 for unified
+/// shared memory CHECK: Return HstPtrBegin {{.*}} Size=20 for unified shared
+/// memory CHECK: Copying data from host to device, HstPtr={{.*}},
+/// TgtPtr={{.*}}, Size=4, Name=x
+#pragma omp target map(to : k[ : 5]) map(always, tofrom : x) map(tofrom : y)   \
+    map(from : dev_k)
   {
     x++;
     y++;
-    for(size_t t = 0; t < 10; t++)
+    for (size_t t = 0; t < 10; t++)
       z[t]++;
     dev_k = k;
   }
-  /// CHECK-NOT: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=20, Name=k
+/// CHECK-NOT: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}},
+/// Size=20, Name=k
 
-  /// CHECK: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4, Name=x
+/// CHECK: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}},
+/// Size=4, Name=x
 
-  /// CHECK: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=40, Name=z
-  #pragma omp target update from(z[:10])
+/// CHECK: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}},
+/// Size=40, Name=z
+#pragma omp target update from(z[ : 10])
 
   /// CHECK-NOT: k pointer not correctly passed to kernel
-  if (dev_k != k) printf("k pointer not correctly passed to kernel\n");
+  if (dev_k != k)
+    printf("k pointer not correctly passed to kernel\n");
 
-  delete [] k;
+  delete[] k;
   return 0;
 }

``````````

</details>


https://github.com/llvm/llvm-project/pull/80876