| /* { dg-additional-options "-fdump-tree-gimple" } */ |
| |
| #include <stdint.h> |
| #include <omp.h> |
| |
| int |
| main () |
| { |
| int x, *xbuf[10]; |
| omp_allocator_handle_t my_alloc; |
| const omp_alloctrait_t trait[1]= {{omp_atk_alignment,128}}; |
| |
| #pragma omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none) |
| #pragma omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf) |
| { |
| if ((uintptr_t) &x % 128 != 0) |
| __builtin_abort (); |
| if ((uintptr_t) xbuf % 128 != 0) |
| __builtin_abort (); |
| } |
| |
| my_alloc = (omp_allocator_handle_t) 0xABCD; |
| |
| #pragma omp target uses_allocators(traits(trait): my_alloc) defaultmap(none) map(tofrom: x, xbuf) |
| #pragma omp parallel allocate(allocator(my_alloc): x, xbuf) if(0) firstprivate(x, xbuf) |
| { |
| if ((uintptr_t) &x % 128 != 0) |
| __builtin_abort (); |
| if ((uintptr_t) xbuf % 128 != 0) |
| __builtin_abort (); |
| } |
| |
| if (my_alloc != (omp_allocator_handle_t) 0xABCD) |
| __builtin_abort (); |
| |
| /* The following creates an allocator with empty traits + default mem space. */ |
| #pragma omp target uses_allocators(my_alloc) map(tofrom: x, xbuf) defaultmap(none) |
| #pragma omp parallel allocate(allocator(my_alloc), align(128): x, xbuf) if(0) firstprivate(x, xbuf) |
| { |
| if ((uintptr_t) &x % 128 != 0) |
| __builtin_abort (); |
| if ((uintptr_t) xbuf % 128 != 0) |
| __builtin_abort (); |
| } |
| |
| if (my_alloc != (omp_allocator_handle_t) 0xABCD) |
| __builtin_abort (); |
| |
| return 0; |
| } |
| |
| /* { dg-final { scan-tree-dump-times "#pragma omp target .* uses_allocators\\(memspace\\(\\), traits\\(trait\\) : my_alloc\\)" 1 "gimple" } } */ |
| /* { dg-final { scan-tree-dump-times "#pragma omp target .* uses_allocators\\(memspace\\(\\), traits\\(\\) : my_alloc\\)" 1 "gimple" } } */ |
