1// RUN: %libomptarget-compile-run-and-check-generic
2
3// FIXME: https://github.com/llvm/llvm-project/issues/77841
4// UNSUPPORTED: amdgcn-amd-amdhsa
5// UNSUPPORTED: nvptx64-nvidia-cuda
6// UNSUPPORTED: nvptx64-nvidia-cuda-LTO
7
8#include <omp.h>
9#include <stdio.h>
10
11#define N 1024
12
13int test_omp_aligned_alloc_on_device() {
14 int errors = 0;
15
16 omp_memspace_handle_t memspace = omp_default_mem_space;
17 omp_alloctrait_t traits[2] = {{omp_atk_alignment, 64}, {omp_atk_access, 64}};
18 omp_allocator_handle_t alloc =
19 omp_init_allocator(omp_default_mem_space, 1, traits);
20
21#pragma omp target map(tofrom : errors) uses_allocators(alloc(traits))
22 {
23 int *x;
24 int not_correct_array_values = 0;
25
26 x = (int *)omp_aligned_alloc(64, N * sizeof(int), alloc);
27 if (x == NULL) {
28 errors++;
29 } else {
30#pragma omp parallel for simd simdlen(16) aligned(x : 64)
31 for (int i = 0; i < N; i++) {
32 x[i] = i;
33 }
34
35#pragma omp parallel for simd simdlen(16) aligned(x : 64)
36 for (int i = 0; i < N; i++) {
37 if (x[i] != i) {
38#pragma omp atomic write
39 not_correct_array_values = 1;
40 }
41 }
42 if (not_correct_array_values) {
43 errors++;
44 }
45 omp_free(x, alloc);
46 }
47 }
48
49 omp_destroy_allocator(alloc);
50
51 return errors;
52}
53
54int main() {
55 int errors = 0;
56 if (test_omp_aligned_alloc_on_device())
57 printf(format: "FAILED\n");
58 else
59 // CHECK: PASSED
60 printf(format: "PASSED\n");
61}
62

source code of offload/test/mapping/target_uses_allocator.c