offload/test/mapping/target_uses_allocator.c

   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("FAILE\n");
  58   else
  59     // CHECK: PASSED
  60     printf("PASSED\n");
  61 }