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 --- /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 --- /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 --- /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 --- a/openmp/runtime/src/dllexports +++ b/openmp/runtime/src/dllexports @@ -567,7 +567,7 @@ 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 @@ 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 --- a/openmp/runtime/src/include/omp.h.var +++ b/openmp/runtime/src/include/omp.h.var @@ -368,7 +368,7 @@ 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 +379,7 @@ 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 +399,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 +415,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.h.var b/openmp/runtime/src/include/omp_lib.h.var --- a/openmp/runtime/src/include/omp_lib.h.var +++ b/openmp/runtime/src/include/omp_lib.h.var @@ -218,7 +218,7 @@ 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 +236,7 @@ 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/include/omp_lib.f90.var b/openmp/runtime/src/include/omp_lib.f90.var --- a/openmp/runtime/src/include/omp_lib.f90.var +++ b/openmp/runtime/src/include/omp_lib.f90.var @@ -139,7 +139,7 @@ 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 +149,7 @@ 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/kmp.h b/openmp/runtime/src/kmp.h --- a/openmp/runtime/src/kmp.h +++ b/openmp/runtime/src/kmp.h @@ -967,7 +967,6 @@ 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_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 --- a/openmp/runtime/src/kmp_alloc.cpp +++ b/openmp/runtime/src/kmp_alloc.cpp @@ -1242,7 +1242,6 @@ 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 @@ 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 --- a/openmp/runtime/src/kmp_global.cpp +++ b/openmp/runtime/src/kmp_global.cpp @@ -316,7 +316,7 @@ (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 +337,7 @@ (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 --- a/openmp/runtime/src/kmp_stub.cpp +++ b/openmp/runtime/src/kmp_stub.cpp @@ -350,7 +350,6 @@ (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)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 =