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