| /* { dg-do run } */ |
| /* { dg-require-effective-target omp_managedmem } */ |
| /* { dg-additional-options -foffload-options=amdgcn-amdhsa=-mxnack=on { target offload_target_amdgcn_with_xnack } } */ |
| |
| /* Check that omp_calloc can allocate Managed Memory, and that host and target |
| can see the data, at the same address, without a mapping. */ |
| |
| #include <omp.h> |
| #include <stdint.h> |
| |
| int |
| main () |
| { |
| int *a = (int *) omp_calloc(5, sizeof(int), ompx_gnu_managed_mem_alloc); |
| if (!a) |
| __builtin_abort (); |
| |
| /* Check that memory is zero-initialized */ |
| for (int i = 0; i < 5; i++) |
| if (a[i] != 0) |
| __builtin_abort (); |
| |
| a[0] = 42; |
| a[4] = 99; |
| uintptr_t a_p = (uintptr_t)a; |
| |
| #pragma omp target is_device_ptr(a) |
| { |
| if (a[0] != 42 || a[4] != 99 || a_p != (uintptr_t)a) |
| __builtin_abort (); |
| /* Check zero-initialization on device side */ |
| for (int i = 1; i < 4; i++) |
| if (a[i] != 0) |
| __builtin_abort (); |
| } |
| |
| omp_free(a, ompx_gnu_managed_mem_alloc); |
| return 0; |
| } |