[OpenMP] Change target memory tests to use allocators

The target allocators have been supported for NVPTX offloading for
awhile. The tests should use the allocators instead of calling the
functions manually. Also the comments indicating these being a preview
should be removed.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D123242

diff --git a/openmp/libomptarget/test/api/omp_device_managed_memory_alloc.c b/openmp/libomptarget/test/api/omp_device_managed_memory_alloc.c
new file mode 100644 (file)
index 0000000..d74ffeb
--- /dev/null
+++ b/openmp/libomptarget/test/api/omp_device_managed_memory_alloc.c
@@ -0,0 +1,28 @@
+// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// REQUIRES: nvptx64-nvidia-cuda
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  const int N = 64;
+
+  // Allocates device managed memory that is shared between the host and device.
+  int *shared_ptr =
+      omp_alloc(N * sizeof(int), llvm_omp_target_shared_mem_alloc);
+
+#pragma omp target teams distribute parallel for is_device_ptr(shared_ptr)
+  for (int i = 0; i < N; ++i) {
+    shared_ptr[i] = 1;
+  }
+
+  int sum = 0;
+  for (int i = 0; i < N; ++i)
+    sum += shared_ptr[i];
+
+  // CHECK: PASS
+  if (sum == N)
+    printf("PASS\n");
+
+  omp_free(shared_ptr, llvm_omp_target_shared_mem_alloc);
+}

diff --git a/openmp/libomptarget/test/api/omp_device_memory.c b/openmp/libomptarget/test/api/omp_device_memory.c
new file mode 100644 (file)
index 0000000..6bb02be
--- /dev/null
+++ b/openmp/libomptarget/test/api/omp_device_memory.c
@@ -0,0 +1,28 @@
+// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// REQUIRES: nvptx64-nvidia-cuda
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  const int N = 64;
+
+  int *device_ptr =
+      omp_alloc(N * sizeof(int), llvm_omp_target_device_mem_alloc);
+
+#pragma omp target teams distribute parallel for is_device_ptr(device_ptr)
+  for (int i = 0; i < N; ++i) {
+    device_ptr[i] = 1;
+  }
+
+  int sum = 0;
+#pragma omp target reduction(+ : sum) is_device_ptr(device_ptr)
+  for (int i = 0; i < N; ++i)
+    sum += device_ptr[i];
+
+  // CHECK: PASS
+  if (sum == N)
+    printf("PASS\n");
+
+  omp_free(device_ptr, llvm_omp_target_device_mem_alloc);
+}

diff --git a/openmp/libomptarget/test/api/omp_host_pinned_memory_alloc.c b/openmp/libomptarget/test/api/omp_host_pinned_memory_alloc.c
new file mode 100644 (file)
index 0000000..84b19fd
--- /dev/null
+++ b/openmp/libomptarget/test/api/omp_host_pinned_memory_alloc.c
@@ -0,0 +1,27 @@
+// RUN: %libomptarget-compile-run-and-check-nvptx64-nvidia-cuda
+// REQUIRES: nvptx64-nvidia-cuda
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  const int N = 64;
+
+  int *hst_ptr = omp_alloc(N * sizeof(int), llvm_omp_target_host_mem_alloc);
+
+  for (int i = 0; i < N; ++i)
+    hst_ptr[i] = 2;
+
+#pragma omp target teams distribute parallel for map(tofrom : hst_ptr [0:N])
+  for (int i = 0; i < N; ++i)
+    hst_ptr[i] -= 1;
+
+  int sum = 0;
+  for (int i = 0; i < N; ++i)
+    sum += hst_ptr[i];
+
+  omp_free(hst_ptr, llvm_omp_target_shared_mem_alloc);
+  // CHECK: PASS
+  if (sum == N)
+    printf("PASS\n");
+}

diff --git a/openmp/runtime/src/dllexports b/openmp/runtime/src/dllexports
index 90fb3c4..87989fe 100644 (file)
--- a/openmp/runtime/src/dllexports
+++ b/openmp/runtime/src/dllexports
@@ -567,7 +567,7 @@ kmp_set_disp_num_buffers                    890
     omp_cgroup_mem_alloc                   DATA
     omp_pteam_mem_alloc                    DATA
     omp_thread_mem_alloc                   DATA
-    # Preview of target memory support
+
     llvm_omp_target_host_mem_alloc         DATA
     llvm_omp_target_shared_mem_alloc       DATA
     llvm_omp_target_device_mem_alloc       DATA
@@ -577,7 +577,7 @@ kmp_set_disp_num_buffers                    890
     omp_const_mem_space                    DATA
     omp_high_bw_mem_space                  DATA
     omp_low_lat_mem_space                  DATA
-    # Preview of target memory support
+
     llvm_omp_target_host_mem_space         DATA
     llvm_omp_target_shared_mem_space       DATA
     llvm_omp_target_device_mem_space       DATA

diff --git a/openmp/runtime/src/include/omp.h.var b/openmp/runtime/src/include/omp.h.var
index 8abd242..b3f9b67 100644 (file)
--- a/openmp/runtime/src/include/omp.h.var
+++ b/openmp/runtime/src/include/omp.h.var
@@ -368,7 +368,6 @@
     extern __KMP_IMP omp_allocator_handle_t const omp_cgroup_mem_alloc;
     extern __KMP_IMP omp_allocator_handle_t const omp_pteam_mem_alloc;
     extern __KMP_IMP omp_allocator_handle_t const omp_thread_mem_alloc;
-    /* Preview of target memory support */
     extern __KMP_IMP omp_allocator_handle_t const llvm_omp_target_host_mem_alloc;
     extern __KMP_IMP omp_allocator_handle_t const llvm_omp_target_shared_mem_alloc;
     extern __KMP_IMP omp_allocator_handle_t const llvm_omp_target_device_mem_alloc;
@@ -379,7 +378,6 @@
     extern __KMP_IMP omp_memspace_handle_t const omp_const_mem_space;
     extern __KMP_IMP omp_memspace_handle_t const omp_high_bw_mem_space;
     extern __KMP_IMP omp_memspace_handle_t const omp_low_lat_mem_space;
-    /* Preview of target memory support */
     extern __KMP_IMP omp_memspace_handle_t const llvm_omp_target_host_mem_space;
     extern __KMP_IMP omp_memspace_handle_t const llvm_omp_target_shared_mem_space;
     extern __KMP_IMP omp_memspace_handle_t const llvm_omp_target_device_mem_space;
@@ -399,7 +397,6 @@
       omp_cgroup_mem_alloc = 6,
       omp_pteam_mem_alloc = 7,
       omp_thread_mem_alloc = 8,
-      /* Preview of target memory support */
       llvm_omp_target_host_mem_alloc = 100,
       llvm_omp_target_shared_mem_alloc = 101,
       llvm_omp_target_device_mem_alloc = 102,
@@ -416,7 +413,6 @@
       omp_const_mem_space = 2,
       omp_high_bw_mem_space = 3,
       omp_low_lat_mem_space = 4,
-      /* Preview of target memory support */
       llvm_omp_target_host_mem_space = 100,
       llvm_omp_target_shared_mem_space = 101,
       llvm_omp_target_device_mem_space = 102,

diff --git a/openmp/runtime/src/include/omp_lib.f90.var b/openmp/runtime/src/include/omp_lib.f90.var
index 12474b6..c722874 100644 (file)
--- a/openmp/runtime/src/include/omp_lib.f90.var
+++ b/openmp/runtime/src/include/omp_lib.f90.var
@@ -139,7 +139,6 @@
         integer (kind=omp_allocator_handle_kind), parameter :: omp_cgroup_mem_alloc = 6
         integer (kind=omp_allocator_handle_kind), parameter :: omp_pteam_mem_alloc = 7
         integer (kind=omp_allocator_handle_kind), parameter :: omp_thread_mem_alloc = 8
-        ! Preview of target memory support
         integer (kind=omp_allocator_handle_kind), parameter :: llvm_omp_target_host_mem_alloc = 100
         integer (kind=omp_allocator_handle_kind), parameter :: llvm_omp_target_shared_mem_alloc = 101
         integer (kind=omp_allocator_handle_kind), parameter :: llvm_omp_target_device_mem_alloc = 102
@@ -149,7 +148,6 @@
         integer (kind=omp_memspace_handle_kind), parameter :: omp_const_mem_space = 2
         integer (kind=omp_memspace_handle_kind), parameter :: omp_high_bw_mem_space = 3
         integer (kind=omp_memspace_handle_kind), parameter :: omp_low_lat_mem_space = 4
-        ! Preview of target memory support
         integer (kind=omp_memspace_handle_kind), parameter :: llvm_omp_target_host_mem_space = 100
         integer (kind=omp_memspace_handle_kind), parameter :: llvm_omp_target_shared_mem_space = 101
         integer (kind=omp_memspace_handle_kind), parameter :: llvm_omp_target_device_mem_space = 102

diff --git a/openmp/runtime/src/include/omp_lib.h.var b/openmp/runtime/src/include/omp_lib.h.var
index a2094f4..9f5e585 100644 (file)
--- a/openmp/runtime/src/include/omp_lib.h.var
+++ b/openmp/runtime/src/include/omp_lib.h.var
@@ -218,7 +218,6 @@
       parameter(omp_pteam_mem_alloc=7)
       integer(kind=omp_allocator_handle_kind)omp_thread_mem_alloc
       parameter(omp_thread_mem_alloc=8)
-      ! Preview of target memory support
       integer(omp_allocator_handle_kind)llvm_omp_target_host_mem_alloc
       parameter(llvm_omp_target_host_mem_alloc=100)
       integer(omp_allocator_handle_kind)llvm_omp_target_shared_mem_alloc
@@ -236,7 +235,6 @@
       parameter(omp_high_bw_mem_space=3)
       integer(kind=omp_memspace_handle_kind)omp_low_lat_mem_space
       parameter(omp_low_lat_mem_space=4)
-      ! Preview of target memory support
       integer(omp_memspace_handle_kind)llvm_omp_target_host_mem_space
       parameter(llvm_omp_target_host_mem_space=100)
       integer(omp_memspace_handle_kind)llvm_omp_target_shared_mem_space

diff --git a/openmp/runtime/src/kmp.h b/openmp/runtime/src/kmp.h
index 826a424..fbc2eda 100644 (file)
--- a/openmp/runtime/src/kmp.h
+++ b/openmp/runtime/src/kmp.h
@@ -967,7 +967,6 @@ extern omp_memspace_handle_t const omp_large_cap_mem_space;
 extern omp_memspace_handle_t const omp_const_mem_space;
 extern omp_memspace_handle_t const omp_high_bw_mem_space;
 extern omp_memspace_handle_t const omp_low_lat_mem_space;
-// Preview of target memory support
 extern omp_memspace_handle_t const llvm_omp_target_host_mem_space;
 extern omp_memspace_handle_t const llvm_omp_target_shared_mem_space;
 extern omp_memspace_handle_t const llvm_omp_target_device_mem_space;
@@ -987,7 +986,6 @@ extern omp_allocator_handle_t const omp_low_lat_mem_alloc;
 extern omp_allocator_handle_t const omp_cgroup_mem_alloc;
 extern omp_allocator_handle_t const omp_pteam_mem_alloc;
 extern omp_allocator_handle_t const omp_thread_mem_alloc;
-// Preview of target memory support
 extern omp_allocator_handle_t const llvm_omp_target_host_mem_alloc;
 extern omp_allocator_handle_t const llvm_omp_target_shared_mem_alloc;
 extern omp_allocator_handle_t const llvm_omp_target_device_mem_alloc;

diff --git a/openmp/runtime/src/kmp_alloc.cpp b/openmp/runtime/src/kmp_alloc.cpp
index 0f76906..222637b 100644 (file)
--- a/openmp/runtime/src/kmp_alloc.cpp
+++ b/openmp/runtime/src/kmp_alloc.cpp
@@ -1242,7 +1242,6 @@ static void **mk_hbw_preferred_hugetlb;
 static void **mk_dax_kmem;
 static void **mk_dax_kmem_all;
 static void **mk_dax_kmem_preferred;
-// Preview of target memory support
 static void *(*kmp_target_alloc_host)(size_t size, int device);
 static void *(*kmp_target_alloc_shared)(size_t size, int device);
 static void *(*kmp_target_alloc_device)(size_t size, int device);
@@ -1352,7 +1351,7 @@ void __kmp_fini_memkind() {
   mk_dax_kmem_preferred = NULL;
 #endif
 }
-// Preview of target memory support
+
 void __kmp_init_target_mem() {
   *(void **)(&kmp_target_alloc_host) = KMP_DLSYM("llvm_omp_target_alloc_host");
   *(void **)(&kmp_target_alloc_shared) =

diff --git a/openmp/runtime/src/kmp_global.cpp b/openmp/runtime/src/kmp_global.cpp
index fdabaad..2a0378d 100644 (file)
--- a/openmp/runtime/src/kmp_global.cpp
+++ b/openmp/runtime/src/kmp_global.cpp
@@ -316,7 +316,6 @@ omp_allocator_handle_t const omp_pteam_mem_alloc =
     (omp_allocator_handle_t const)7;
 omp_allocator_handle_t const omp_thread_mem_alloc =
     (omp_allocator_handle_t const)8;
-// Preview of target memory support
 omp_allocator_handle_t const llvm_omp_target_host_mem_alloc =
     (omp_allocator_handle_t const)100;
 omp_allocator_handle_t const llvm_omp_target_shared_mem_alloc =
@@ -337,7 +336,6 @@ omp_memspace_handle_t const omp_high_bw_mem_space =
     (omp_memspace_handle_t const)3;
 omp_memspace_handle_t const omp_low_lat_mem_space =
     (omp_memspace_handle_t const)4;
-// Preview of target memory support
 omp_memspace_handle_t const llvm_omp_target_host_mem_space =
     (omp_memspace_handle_t const)100;
 omp_memspace_handle_t const llvm_omp_target_shared_mem_space =

diff --git a/openmp/runtime/src/kmp_stub.cpp b/openmp/runtime/src/kmp_stub.cpp
index da7340f..f25e24f 100644 (file)
--- a/openmp/runtime/src/kmp_stub.cpp
+++ b/openmp/runtime/src/kmp_stub.cpp
@@ -350,7 +350,6 @@ omp_allocator_handle_t const omp_pteam_mem_alloc =
     (omp_allocator_handle_t const)7;
 omp_allocator_handle_t const omp_thread_mem_alloc =
     (omp_allocator_handle_t const)8;
-// Preview of target memory support
 omp_allocator_handle_t const llvm_omp_target_host_mem_alloc =
     (omp_allocator_handle_t const)100;
 omp_allocator_handle_t const llvm_omp_target_shared_mem_alloc =
@@ -368,7 +367,6 @@ omp_memspace_handle_t const omp_high_bw_mem_space =
     (omp_memspace_handle_t const)3;
 omp_memspace_handle_t const omp_low_lat_mem_space =
     (omp_memspace_handle_t const)4;
-// Preview of target memory support
 omp_memspace_handle_t const llvm_omp_target_host_mem_space =
     (omp_memspace_handle_t const)100;
 omp_memspace_handle_t const llvm_omp_target_shared_mem_space =