From 840c040498f3fad2875c72248862292796795dd1 Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Wed, 6 Apr 2022 15:21:50 -0400 Subject: [PATCH] [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 --- .../test/api/omp_device_managed_memory_alloc.c | 28 ++++++++++++++++++++++ openmp/libomptarget/test/api/omp_device_memory.c | 28 ++++++++++++++++++++++ .../test/api/omp_host_pinned_memory_alloc.c | 27 +++++++++++++++++++++ openmp/runtime/src/dllexports | 4 ++-- openmp/runtime/src/include/omp.h.var | 4 ---- openmp/runtime/src/include/omp_lib.f90.var | 2 -- openmp/runtime/src/include/omp_lib.h.var | 2 -- openmp/runtime/src/kmp.h | 2 -- openmp/runtime/src/kmp_alloc.cpp | 3 +-- openmp/runtime/src/kmp_global.cpp | 2 -- openmp/runtime/src/kmp_stub.cpp | 2 -- 11 files changed, 86 insertions(+), 18 deletions(-) create mode 100644 openmp/libomptarget/test/api/omp_device_managed_memory_alloc.c create mode 100644 openmp/libomptarget/test/api/omp_device_memory.c create mode 100644 openmp/libomptarget/test/api/omp_host_pinned_memory_alloc.c 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 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 +#include + +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 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 +#include + +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 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 +#include + +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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 = -- 2.7.4