openmp/libomptarget/test/mapping/target_uses_allocator.c

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