offload/test/offloading/host_as_target.c

   1 // Check that specifying device as omp_get_initial_device():
   2 // - Doesn't cause the runtime to fail.
   3 // - Offloads code to the host.
   4 // - Doesn't transfer data.  In this case, just check that neither host data nor
   5 //   default device data are affected by the specified transfers.
   6 // - Works whether it's specified directly or as the default device.
   7
   8 // RUN: %libomptarget-compile-run-and-check-generic
   9
  10 #include <omp.h>
  11 #include <stdio.h>
  12
  13 static void check(char *X, int Dev) {
  14   printf("  host X = %c\n", *X);
  15   char DV = -1;
  16 #pragma omp target device(Dev) map(from : DV)
  17   DV = *X;
  18   printf("device X = %c\n", DV);
  19 }
  20
  21 #define CHECK_DATA() check(&X, DevDefault)
  22
  23 int main(void) {
  24   int DevDefault = omp_get_default_device();
  25   int DevInit = omp_get_initial_device();
  26
  27   //--------------------------------------------------
  28   // Initialize data on the host and default device.
  29   //--------------------------------------------------
  30
  31   //      CHECK:   host X = h
  32   // CHECK-NEXT: device X = d
  33   char X = 'd';
  34 #pragma omp target enter data map(to : X)
  35   X = 'h';
  36   CHECK_DATA();
  37
  38 //--------------------------------------------------
  39 // Check behavior when specifying host directly.
  40 //--------------------------------------------------
  41
  42 // CHECK-NEXT: omp_is_initial_device() = 1
  43 // CHECK-NEXT:   host X = h
  44 // CHECK-NEXT: device X = d
  45 #pragma omp target device(DevInit) map(always, tofrom : X)
  46   printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
  47   CHECK_DATA();
  48
  49 // CHECK-NEXT: omp_is_initial_device() = 1
  50 // CHECK-NEXT:   host X = h
  51 // CHECK-NEXT: device X = d
  52 #pragma omp target teams device(DevInit) num_teams(1) map(always, tofrom : X)
  53   printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
  54   CHECK_DATA();
  55
  56 // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
  57 // how to check that it actually pushes to the initial device.
  58 #pragma omp target teams device(DevInit) num_teams(1)
  59 #pragma omp distribute
  60   for (int i = 0; i < 2; ++i)
  61     ;
  62
  63 // CHECK-NEXT:   host X = h
  64 // CHECK-NEXT: device X = d
  65 #pragma omp target data device(DevInit) map(always, tofrom : X)
  66   ;
  67   CHECK_DATA();
  68
  69 // CHECK-NEXT:   host X = h
  70 // CHECK-NEXT: device X = d
  71 #pragma omp target enter data device(DevInit) map(always, to : X)
  72   ;
  73   CHECK_DATA();
  74
  75 // CHECK-NEXT:   host X = h
  76 // CHECK-NEXT: device X = d
  77 #pragma omp target exit data device(DevInit) map(always, from : X)
  78   ;
  79   CHECK_DATA();
  80
  81 // CHECK-NEXT:   host X = h
  82 // CHECK-NEXT: device X = d
  83 #pragma omp target update device(DevInit) to(X)
  84   ;
  85   CHECK_DATA();
  86
  87 // CHECK-NEXT:   host X = h
  88 // CHECK-NEXT: device X = d
  89 #pragma omp target update device(DevInit) from(X)
  90   ;
  91   CHECK_DATA();
  92
  93   //--------------------------------------------------
  94   // Check behavior when device defaults to host.
  95   //--------------------------------------------------
  96
  97   omp_set_default_device(DevInit);
  98
  99 // CHECK-NEXT: omp_is_initial_device() = 1
 100 // CHECK-NEXT:   host X = h
 101 // CHECK-NEXT: device X = d
 102 #pragma omp target map(always, tofrom : X)
 103   printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
 104   CHECK_DATA();
 105
 106 // CHECK-NEXT: omp_is_initial_device() = 1
 107 // CHECK-NEXT:   host X = h
 108 // CHECK-NEXT: device X = d
 109 #pragma omp target teams num_teams(1) map(always, tofrom : X)
 110   printf("omp_is_initial_device() = %d\n", omp_is_initial_device());
 111   CHECK_DATA();
 112
 113 // Check that __kmpc_push_target_tripcount_mapper doesn't fail. I'm not sure
 114 // how to check that it actually pushes to the initial device.
 115 #pragma omp target teams num_teams(1)
 116 #pragma omp distribute
 117   for (int i = 0; i < 2; ++i)
 118     ;
 119
 120 // CHECK-NEXT:   host X = h
 121 // CHECK-NEXT: device X = d
 122 #pragma omp target data map(always, tofrom : X)
 123   ;
 124   CHECK_DATA();
 125
 126 // CHECK-NEXT:   host X = h
 127 // CHECK-NEXT: device X = d
 128 #pragma omp target enter data map(always, to : X)
 129   ;
 130   CHECK_DATA();
 131
 132 // CHECK-NEXT:   host X = h
 133 // CHECK-NEXT: device X = d
 134 #pragma omp target exit data map(always, from : X)
 135   ;
 136   CHECK_DATA();
 137
 138 // CHECK-NEXT:   host X = h
 139 // CHECK-NEXT: device X = d
 140 #pragma omp target update to(X)
 141   ;
 142   CHECK_DATA();
 143
 144 // CHECK-NEXT:   host X = h
 145 // CHECK-NEXT: device X = d
 146 #pragma omp target update from(X)
 147   ;
 148   CHECK_DATA();
 149
 150   return 0;
 151 }