--- /dev/null
+// 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);
+}
--- /dev/null
+// 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);
+}
--- /dev/null
+// 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");
+}
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
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
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;
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;
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,
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,
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
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
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
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
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;
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;
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);
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) =
(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 =
(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 =
(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 =
(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 =