xref: /llvm-project/offload/test/api/omp_device_managed_memory_alloc.c (revision 330d8983d25d08580fc1642fea48b2473f47a9da)
1*330d8983SJohannes Doerfert // RUN: %libomptarget-compile-run-and-check-generic
2*330d8983SJohannes Doerfert // RUN: %libomptarget-compileopt-run-and-check-generic
3*330d8983SJohannes Doerfert 
4*330d8983SJohannes Doerfert #include <omp.h>
5*330d8983SJohannes Doerfert #include <stdio.h>
6*330d8983SJohannes Doerfert 
main()7*330d8983SJohannes Doerfert int main() {
8*330d8983SJohannes Doerfert   const int N = 64;
9*330d8983SJohannes Doerfert 
10*330d8983SJohannes Doerfert   // Allocates device managed memory that is shared between the host and device.
11*330d8983SJohannes Doerfert   int *shared_ptr =
12*330d8983SJohannes Doerfert       omp_alloc(N * sizeof(int), llvm_omp_target_shared_mem_alloc);
13*330d8983SJohannes Doerfert 
14*330d8983SJohannes Doerfert #pragma omp target teams distribute parallel for is_device_ptr(shared_ptr)
15*330d8983SJohannes Doerfert   for (int i = 0; i < N; ++i) {
16*330d8983SJohannes Doerfert     shared_ptr[i] = 1;
17*330d8983SJohannes Doerfert   }
18*330d8983SJohannes Doerfert 
19*330d8983SJohannes Doerfert   int sum = 0;
20*330d8983SJohannes Doerfert   for (int i = 0; i < N; ++i)
21*330d8983SJohannes Doerfert     sum += shared_ptr[i];
22*330d8983SJohannes Doerfert 
23*330d8983SJohannes Doerfert   // CHECK: PASS
24*330d8983SJohannes Doerfert   if (sum == N)
25*330d8983SJohannes Doerfert     printf("PASS\n");
26*330d8983SJohannes Doerfert 
27*330d8983SJohannes Doerfert   omp_free(shared_ptr, llvm_omp_target_shared_mem_alloc);
28*330d8983SJohannes Doerfert }
29