offload/test/offloading/force-usm.cpp

   1 // clang-format off
   2 // RUN: %libomptarget-compilexx-generic
   3 // RUN: env LIBOMPTARGET_INFO=32 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=NO-USM
   4 //
   5 // RUN: %libomptarget-compilexxx-generic-force-usm
   6 // RUN: env HSA_XNACK=1 LIBOMPTARGET_INFO=32 \
   7 // RUN:       %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefix=FORCE-USM
   8 //
   9 // REQUIRES: unified_shared_memory
  10 //
  11 // UNSUPPORTED: nvptx64-nvidia-cuda
  12 // UNSUPPORTED: nvptx64-nvidia-cuda-LTO
  13 // clang-format on
  14
  15 #include <cassert>
  16 #include <cstdio>
  17 #include <cstdlib>
  18
  19 int GI;
  20 #pragma omp declare target
  21 int *pGI;
  22 #pragma omp end declare target
  23
  24 int main(void) {
  25
  26   GI = 0;
  27   // Implicit mappings
  28   int alpha = 1;
  29   int beta[3] = {2, 5, 8};
  30
  31   // Require map clauses for non-USM execution
  32   pGI = (int *)malloc(sizeof(int));
  33   *pGI = 42;
  34
  35 #pragma omp target map(pGI[ : 1], GI)
  36   {
  37     GI = 1 * alpha;
  38     *pGI = 2 * beta[1];
  39   }
  40
  41   assert(GI == 1);
  42   assert(*pGI == 10);
  43
  44   printf("SUCCESS\n");
  45
  46   return 0;
  47 }
  48
  49 // clang-format off
  50 // NO-USM: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
  51 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=12
  52 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=4
  53 // NO-USM-NEXT: omptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=8, Name=pGI
  54 // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
  55 // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=12
  56 // NO-USM-NEXT: omptarget device 0 info: Copying data from device to host, TgtPtr={{.*}}, HstPtr={{.*}}, Size=4
  57 // NO-USM-NEXT: SUCCESS
  58
  59 // FORCE-USM: SUCCESS
  60 //
  61 // clang-format on