openmp/libomptarget/test/unified_shared_memory/close_modifier.c

   1 // RUN: %libomptarget-compile-run-and-check-generic
   2
   3 // REQUIRES: unified_shared_memory
   4 // UNSUPPORTED: clang-6, clang-7, clang-8, clang-9
   5
   6 // amdgpu runtime crash
   7 // UNSUPPORTED: amdgcn-amd-amdhsa
   8
   9 #include <omp.h>
  10 #include <stdio.h>
  11
  12 #pragma omp requires unified_shared_memory
  13
  14 #define N 1024
  15
  16 int main(int argc, char *argv[]) {
  17   int fails;
  18   void *host_alloc, *device_alloc;
  19   void *host_data, *device_data;
  20   int *alloc = (int *)malloc(N * sizeof(int));
  21   int data[N];
  22
  23   for (int i = 0; i < N; ++i) {
  24     alloc[i] = 10;
  25     data[i] = 1;
  26   }
  27
  28   host_data = &data[0];
  29   host_alloc = &alloc[0];
  30
  31 //
  32 // Test that updates on the device are not visible to host
  33 // when only a TO mapping is used.
  34 //
  35 #pragma omp target map(tofrom : device_data, device_alloc)                     \
  36     map(close, to : alloc[ : N], data[ : N])
  37   {
  38     device_data = &data[0];
  39     device_alloc = &alloc[0];
  40
  41     for (int i = 0; i < N; i++) {
  42       alloc[i] += 1;
  43       data[i] += 1;
  44     }
  45   }
  46
  47   // CHECK: Address of alloc on device different from host address.
  48   if (device_alloc != host_alloc)
  49     printf("Address of alloc on device different from host address.\n");
  50
  51   // CHECK: Address of data on device different from host address.
  52   if (device_data != host_data)
  53     printf("Address of data on device different from host address.\n");
  54
  55   // On the host, check that the arrays have been updated.
  56   // CHECK: Alloc host values not updated: Succeeded
  57   fails = 0;
  58   for (int i = 0; i < N; i++) {
  59     if (alloc[i] != 10)
  60       fails++;
  61   }
  62   printf("Alloc host values not updated: %s\n",
  63          (fails == 0) ? "Succeeded" : "Failed");
  64
  65   // CHECK: Data host values not updated: Succeeded
  66   fails = 0;
  67   for (int i = 0; i < N; i++) {
  68     if (data[i] != 1)
  69       fails++;
  70   }
  71   printf("Data host values not updated: %s\n",
  72          (fails == 0) ? "Succeeded" : "Failed");
  73
  74   //
  75   // Test that updates on the device are visible on host
  76   // when a from is used.
  77   //
  78
  79   for (int i = 0; i < N; i++) {
  80     alloc[i] += 1;
  81     data[i] += 1;
  82   }
  83
  84 #pragma omp target map(close, tofrom : alloc[ : N], data[ : N])
  85   {
  86     // CHECK: Alloc device values are correct: Succeeded
  87     fails = 0;
  88     for (int i = 0; i < N; i++) {
  89       if (alloc[i] != 11)
  90         fails++;
  91     }
  92     printf("Alloc device values are correct: %s\n",
  93            (fails == 0) ? "Succeeded" : "Failed");
  94     // CHECK: Data device values are correct: Succeeded
  95     fails = 0;
  96     for (int i = 0; i < N; i++) {
  97       if (data[i] != 2)
  98         fails++;
  99     }
 100     printf("Data device values are correct: %s\n",
 101            (fails == 0) ? "Succeeded" : "Failed");
 102
 103     // Update values on the device
 104     for (int i = 0; i < N; i++) {
 105       alloc[i] += 1;
 106       data[i] += 1;
 107     }
 108   }
 109
 110   // CHECK: Alloc host values updated: Succeeded
 111   fails = 0;
 112   for (int i = 0; i < N; i++) {
 113     if (alloc[i] != 12)
 114       fails++;
 115   }
 116   printf("Alloc host values updated: %s\n",
 117          (fails == 0) ? "Succeeded" : "Failed");
 118
 119   // CHECK: Data host values updated: Succeeded
 120   fails = 0;
 121   for (int i = 0; i < N; i++) {
 122     if (data[i] != 3)
 123       fails++;
 124   }
 125   printf("Data host values updated: %s\n",
 126          (fails == 0) ? "Succeeded" : "Failed");
 127
 128   free(alloc);
 129
 130   // CHECK: Done!
 131   printf("Done!\n");
 132
 133   return 0;
 134 }