offload/test/mapping/auto_zero_copy.cpp

   1 // clang-format off
   2 // RUN: %libomptarget-compilexx-generic
   3 // RUN: env OMPX_APU_MAPS=1 HSA_XNACK=1 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \
   4 // RUN: | %fcheck-generic -check-prefix=INFO_ZERO -check-prefix=CHECK
   5
   6 // RUN: %libomptarget-compilexx-generic
   7 // RUN: env HSA_XNACK=0 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \
   8 // RUN: | %fcheck-generic -check-prefix=INFO_COPY -check-prefix=CHECK
   9
  10 // REQUIRES: amdgpu
  11 // REQUIRES: unified_shared_memory
  12
  13 // clang-format on
  14
  15 #include <cstdio>
  16
  17 int main() {
  18   int n = 1024;
  19
  20   // test various mapping types
  21   int *a = new int[n];
  22   int k = 3;
  23   int b[n];
  24
  25   for (int i = 0; i < n; i++)
  26     b[i] = i;
  27
  28     // clang-format off
  29   // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
  30   // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory
  31
  32   // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
  33   // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,
  34   // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
  35   // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)
  36 // clang-format on
  37 #pragma omp target teams distribute parallel for map(tofrom : a[ : n])         \
  38     map(to : b[ : n])
  39   for (int i = 0; i < n; i++)
  40     a[i] = i + b[i] + k;
  41
  42   int err = 0;
  43   for (int i = 0; i < n; i++)
  44     if (a[i] != i + b[i] + k)
  45       err++;
  46
  47   // CHECK: PASS
  48   if (err == 0)
  49     printf("PASS\n");
  50   return err;
  51 }