offload/test/offloading/d2d_memcpy_sync.c

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