62 lines · c
1// RUN: %libomptarget-compile-run-and-check-generic2 3// FIXME: https://github.com/llvm/llvm-project/issues/778414// UNSUPPORTED: amdgcn-amd-amdhsa5// UNSUPPORTED: nvptx64-nvidia-cuda6// UNSUPPORTED: nvptx64-nvidia-cuda-LTO7 8#include <omp.h>9#include <stdio.h>10 11#define N 102412 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 write39 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("FAILED\n");58 else59 // CHECK: PASSED60 printf("PASSED\n");61}62