[Openmp-commits] [openmp] 840c040 - [OpenMP] Change target memory tests to use allocators
Joseph Huber via Openmp-commits
openmp-commits at lists.llvm.org
Thu Apr 7 11:23:30 PDT 2022
Author: Joseph Huber
Date: 2022-04-07T14:23:14-04:00
New Revision: 840c040498f3fad2875c72248862292796795dd1
URL: https://github.com/llvm/llvm-project/commit/840c040498f3fad2875c72248862292796795dd1
DIFF: https://github.com/llvm/llvm-project/commit/840c040498f3fad2875c72248862292796795dd1.diff
LOG: [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
Added:
openmp/libomptarget/test/api/omp_device_managed_memory_alloc.c
openmp/libomptarget/test/api/omp_device_memory.c
openmp/libomptarget/test/api/omp_host_pinned_memory_alloc.c
Modified:
openmp/runtime/src/dllexports
openmp/runtime/src/include/omp.h.var
openmp/runtime/src/include/omp_lib.f90.var
openmp/runtime/src/include/omp_lib.h.var
openmp/runtime/src/kmp.h
openmp/runtime/src/kmp_alloc.cpp
openmp/runtime/src/kmp_global.cpp
openmp/runtime/src/kmp_stub.cpp
Removed:
################################################################################
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
index 0000000000000..d74ffebc472c8
--- /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
index 0000000000000..6bb02be7a50f5
--- /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
index 0000000000000..84b19fd2e184a
--- /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 90fb3c4e02b05..87989fe630929 100644
--- 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 8abd2428971b7..b3f9b67b12d31 100644
--- 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 12474b6c4bcb7..c72287422809a 100644
--- 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 a2094f45b0caa..9f5e58515e751 100644
--- 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 826a42407a338..fbc2eda04a8fb 100644
--- 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 0f76906714b1d..222637bb4eaee 100644
--- 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 fdabaad21f7b6..2a0378d4ce98c 100644
--- 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 da7340fd1404c..f25e24f09a03d 100644
--- 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 =
More information about the Openmp-commits
mailing list