| /* { 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_realloc 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_alloc(2 * sizeof(int), ompx_gnu_managed_mem_alloc); | |
| if (!a) | |
| __builtin_abort (); | |
| a[0] = 42; | |
| a[1] = 43; | |
| /* Reallocate to larger size */ | |
| int *b = (int *) omp_realloc(a, 5 * sizeof(int), ompx_gnu_managed_mem_alloc, | |
| ompx_gnu_managed_mem_alloc); | |
| if (!b) | |
| __builtin_abort (); | |
| /* Check that original data is preserved */ | |
| if (b[0] != 42 || b[1] != 43) | |
| __builtin_abort (); | |
| b[2] = 44; | |
| b[3] = 45; | |
| b[4] = 46; | |
| uintptr_t b_p = (uintptr_t)b; | |
| #pragma omp target is_device_ptr(b) | |
| { | |
| if (b[0] != 42 || b[1] != 43 || b[2] != 44 || b[3] != 45 || b[4] != 46 | |
| || b_p != (uintptr_t)b) | |
| __builtin_abort (); | |
| } | |
| omp_free(b, ompx_gnu_managed_mem_alloc); | |
| return 0; | |
| } |