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 | |
13 | int 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 | |
54 | int main() { |
55 | int errors = 0; |
56 | if (test_omp_aligned_alloc_on_device()) |
57 | printf(format: "FAILE\n" ); |
58 | else |
59 | // CHECK: PASSED |
60 | printf(format: "PASSED\n" ); |
61 | } |
62 | |