llvm/offload/test/offloading/d2d_memcpy.c

// RUN: %libomptarget-compile-generic && \
// RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \
// RUN: %fcheck-generic -allow-empty
// RUN: %libomptarget-compileopt-generic && \
// RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \
// RUN: %fcheck-generic -allow-empty

#include <assert.h>
#include <omp.h>
#include <stdio.h>
#include <stdlib.h>

const int magic_num = 7;

int main(int argc, char *argv[]) {
  const int N = 128;
  const int num_devices = omp_get_num_devices();

  // No target device, just return
  if (num_devices == 0) {
    printf("PASS\n");
    return 0;
  }

  const int src_device = 0;
  int dst_device = num_devices - 1;

  int length = N * sizeof(int);
  int *src_ptr = omp_target_alloc(length, src_device);
  int *dst_ptr = omp_target_alloc(length, dst_device);

  assert(src_ptr && "src_ptr is NULL");
  assert(dst_ptr && "dst_ptr is NULL");

#pragma omp target teams distribute parallel for device(src_device)            \
    is_device_ptr(src_ptr)
  for (int i = 0; i < N; ++i) {
    src_ptr[i] = magic_num;
  }

  int rc =
      omp_target_memcpy(dst_ptr, src_ptr, length, 0, 0, dst_device, src_device);

  assert(rc == 0 && "error in omp_target_memcpy");

  int *buffer = malloc(length);

  assert(buffer && "failed to allocate host buffer");

#pragma omp target teams distribute parallel for device(dst_device)            \
    map(from : buffer[0 : N]) is_device_ptr(dst_ptr)
  for (int i = 0; i < N; ++i) {
    buffer[i] = dst_ptr[i] + magic_num;
  }

  for (int i = 0; i < N; ++i)
    assert(buffer[i] == 2 * magic_num);

  printf("PASS\n");

  // Free host and device memory
  free(buffer);
  omp_target_free(src_ptr, src_device);
  omp_target_free(dst_ptr, dst_device);

  return 0;
}

// CHECK: PASS