offload/test/offloading/d2d_memcpy.c

   1 // RUN: %libomptarget-compile-generic && \
   2 // RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \
   3 // RUN: %fcheck-generic -allow-empty
   4 // RUN: %libomptarget-compileopt-generic && \
   5 // RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \
   6 // RUN: %fcheck-generic -allow-empty
   7
   8 #include <assert.h>
   9 #include <omp.h>
  10 #include <stdio.h>
  11 #include <stdlib.h>
  12
  13 const int magic_num = 7;
  14
  15 int main(int argc, char *argv[]) {
  16   const int N = 128;
  17   const int num_devices = omp_get_num_devices();
  18
  19   // No target device, just return
  20   if (num_devices == 0) {
  21     printf("PASS\n");
  22     return 0;
  23   }
  24
  25   const int src_device = 0;
  26   int dst_device = num_devices - 1;
  27
  28   int length = N * sizeof(int);
  29   int *src_ptr = omp_target_alloc(length, src_device);
  30   int *dst_ptr = omp_target_alloc(length, dst_device);
  31
  32   assert(src_ptr && "src_ptr is NULL");
  33   assert(dst_ptr && "dst_ptr is NULL");
  34
  35 #pragma omp target teams distribute parallel for device(src_device)            \
  36     is_device_ptr(src_ptr)
  37   for (int i = 0; i < N; ++i) {
  38     src_ptr[i] = magic_num;
  39   }
  40
  41   int rc =
  42       omp_target_memcpy(dst_ptr, src_ptr, length, 0, 0, dst_device, src_device);
  43
  44   assert(rc == 0 && "error in omp_target_memcpy");
  45
  46   int *buffer = malloc(length);
  47
  48   assert(buffer && "failed to allocate host buffer");
  49
  50 #pragma omp target teams distribute parallel for device(dst_device)            \
  51     map(from : buffer[0 : N]) is_device_ptr(dst_ptr)
  52   for (int i = 0; i < N; ++i) {
  53     buffer[i] = dst_ptr[i] + magic_num;
  54   }
  55
  56   for (int i = 0; i < N; ++i)
  57     assert(buffer[i] == 2 * magic_num);
  58
  59   printf("PASS\n");
  60
  61   // Free host and device memory
  62   free(buffer);
  63   omp_target_free(src_ptr, src_device);
  64   omp_target_free(dst_ptr, dst_device);
  65
  66   return 0;
  67 }
  68
  69 // CHECK: PASS