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