openmp/libomptarget/test/unified_shared_memory/shared_update.c

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