70 lines · c
1// RUN: %libomptarget-compile-generic && \2// RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \3// RUN: %fcheck-generic -allow-empty4// RUN: %libomptarget-compileopt-generic && \5// RUN: env OMP_MAX_ACTIVE_LEVELS=2 %libomptarget-run-generic | \6// RUN: %fcheck-generic -allow-empty7 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 return20 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 memory62 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: PASS70