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
13const int magic_num = 7;
14
15int 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(format: "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(size: 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(format: "PASS\n");
60
61 // Free host and device memory
62 free(ptr: 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
70

source code of offload/test/offloading/d2d_memcpy.c