openmp/libomptarget/test/unified_shared_memory/api.c

   1 // RUN: %libomptarget-compile-run-and-check-generic
   2 // XFAIL: nvptx64-nvidia-cuda
   3 // XFAIL: nvptx64-nvidia-cuda-LTO
   4
   5 // Fails on amdgpu with error: GPU Memory Error
   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 void init(int A[], int B[], int C[]) {
  24   for (int i = 0; i < N; ++i) {
  25     A[i] = 0;
  26     B[i] = 1;
  27     C[i] = i;
  28   }
  29 }
  30
  31 int main(int argc, char *argv[]) {
  32   const int device = omp_get_default_device();
  33
  34   // Manual registration of requires flags for Clang versions
  35   // that do not support requires.
  36   __tgt_register_requires(8);
  37
  38   // CHECK: Initial device: [[INITIAL_DEVICE:[0-9]+]]
  39   printf("Initial device: %d\n", omp_get_initial_device());
  40   // CHECK: Num devices: [[INITIAL_DEVICE]]
  41   printf("Num devices: %d\n", omp_get_num_devices());
  42
  43   //
  44   // Target alloc & target memcpy
  45   //
  46   int A[N], B[N], C[N];
  47
  48   // Init
  49   init(A, B, C);
  50
  51   int *pA, *pB, *pC;
  52
  53   // map ptrs
  54   pA = &A[0];
  55   pB = &B[0];
  56   pC = &C[0];
  57
  58   int *d_A = (int *)omp_target_alloc(N * sizeof(int), device);
  59   int *d_B = (int *)omp_target_alloc(N * sizeof(int), device);
  60   int *d_C = (int *)omp_target_alloc(N * sizeof(int), device);
  61
  62   // CHECK: omp_target_alloc succeeded
  63   printf("omp_target_alloc %s\n", d_A && d_B && d_C ? "succeeded" : "failed");
  64
  65   omp_target_memcpy(d_B, pB, N * sizeof(int), 0, 0, device,
  66                     omp_get_initial_device());
  67   omp_target_memcpy(d_C, pC, N * sizeof(int), 0, 0, device,
  68                     omp_get_initial_device());
  69
  70 #pragma omp target is_device_ptr(d_A, d_B, d_C) device(device)
  71   {
  72 #pragma omp parallel for schedule(static, 1)
  73     for (int i = 0; i < N; i++) {
  74       d_A[i] = d_B[i] + d_C[i] + 1;
  75     }
  76   }
  77
  78   omp_target_memcpy(pA, d_A, N * sizeof(int), 0, 0, omp_get_initial_device(),
  79                     device);
  80
  81   // CHECK: Test omp_target_memcpy: Succeeded
  82   int fail = 0;
  83   for (int i = 0; i < N; ++i) {
  84     if (A[i] != i + 2)
  85       fail++;
  86   }
  87   if (fail) {
  88     printf("Test omp_target_memcpy: Failed\n");
  89   } else {
  90     printf("Test omp_target_memcpy: Succeeded\n");
  91   }
  92
  93   //
  94   // target_is_present and target_associate/disassociate_ptr
  95   //
  96   init(A, B, C);
  97
  98   // CHECK: B is not present, associating it...
  99   // CHECK: omp_target_associate_ptr B succeeded
 100   if (!omp_target_is_present(B, device)) {
 101     printf("B is not present, associating it...\n");
 102     int rc = omp_target_associate_ptr(B, d_B, N * sizeof(int), 0, device);
 103     printf("omp_target_associate_ptr B %s\n", !rc ? "succeeded" : "failed");
 104   }
 105
 106   // CHECK: C is not present, associating it...
 107   // CHECK: omp_target_associate_ptr C succeeded
 108   if (!omp_target_is_present(C, device)) {
 109     printf("C is not present, associating it...\n");
 110     int rc = omp_target_associate_ptr(C, d_C, N * sizeof(int), 0, device);
 111     printf("omp_target_associate_ptr C %s\n", !rc ? "succeeded" : "failed");
 112   }
 113
 114 // CHECK: Inside target data: A is not present
 115 // CHECK: Inside target data: B is present
 116 // CHECK: Inside target data: C is present
 117 #pragma omp target data map(from : B, C) device(device)
 118   {
 119     printf("Inside target data: A is%s present\n",
 120            omp_target_is_present(A, device) ? "" : " not");
 121     printf("Inside target data: B is%s present\n",
 122            omp_target_is_present(B, device) ? "" : " not");
 123     printf("Inside target data: C is%s present\n",
 124            omp_target_is_present(C, device) ? "" : " not");
 125
 126 #pragma omp target map(from : A) device(device)
 127     {
 128 #pragma omp parallel for schedule(static, 1)
 129       for (int i = 0; i < N; i++)
 130         A[i] = B[i] + C[i] + 1;
 131     }
 132   }
 133
 134   // CHECK: B is present, disassociating it...
 135   // CHECK: omp_target_disassociate_ptr B succeeded
 136   // CHECK: C is present, disassociating it...
 137   // CHECK: omp_target_disassociate_ptr C succeeded
 138   if (omp_target_is_present(B, device)) {
 139     printf("B is present, disassociating it...\n");
 140     int rc = omp_target_disassociate_ptr(B, device);
 141     printf("omp_target_disassociate_ptr B %s\n", !rc ? "succeeded" : "failed");
 142   }
 143   if (omp_target_is_present(C, device)) {
 144     printf("C is present, disassociating it...\n");
 145     int rc = omp_target_disassociate_ptr(C, device);
 146     printf("omp_target_disassociate_ptr C %s\n", !rc ? "succeeded" : "failed");
 147   }
 148
 149   // CHECK: Test omp_target_associate_ptr: Succeeded
 150   fail = 0;
 151   for (int i = 0; i < N; ++i) {
 152     if (A[i] != i + 2)
 153       fail++;
 154   }
 155   if (fail) {
 156     printf("Test omp_target_associate_ptr: Failed\n");
 157   } else {
 158     printf("Test omp_target_associate_ptr: Succeeded\n");
 159   }
 160
 161   omp_target_free(d_A, device);
 162   omp_target_free(d_B, device);
 163   omp_target_free(d_C, device);
 164
 165   printf("Done!\n");
 166
 167   return 0;
 168 }