52 lines · cpp
1// clang-format off2// RUN: %libomptarget-compilexx-generic3// RUN: env OMPX_APU_MAPS=1 HSA_XNACK=1 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \4// RUN: | %fcheck-generic -check-prefix=INFO_ZERO -check-prefix=CHECK5 6// RUN: %libomptarget-compilexx-generic7// RUN: env HSA_XNACK=0 LIBOMPTARGET_INFO=30 %libomptarget-run-generic 2>&1 \8// RUN: | %fcheck-generic -check-prefix=INFO_COPY -check-prefix=CHECK9 10// REQUIRES: amdgpu11// REQUIRES: unified_shared_memory12 13// clang-format on14 15#include <cstdio>16 17int main() {18 int n = 1024;19 20 // test various mapping types21 int *a = new int[n];22 int k = 3;23 int b[n];24 25 for (int i = 0; i < n; i++)26 b[i] = i;27 28 // clang-format off29 // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory30 // INFO_ZERO: Return HstPtrBegin 0x{{.*}} Size=4096 for unified shared memory31 32 // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,33 // INFO_COPY: Creating new map entry with HstPtrBase=0x{{.*}}, HstPtrBegin=0x{{.*}}, TgtAllocBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096,34 // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)35 // INFO_COPY: Mapping exists with HstPtrBegin=0x{{.*}}, TgtPtrBegin=0x{{.*}}, Size=4096, DynRefCount=1 (update suppressed)36// clang-format on37#pragma omp target teams distribute parallel for map(tofrom : a[ : n]) \38 map(to : b[ : n])39 for (int i = 0; i < n; i++)40 a[i] = i + b[i] + k;41 42 int err = 0;43 for (int i = 0; i < n; i++)44 if (a[i] != i + b[i] + k)45 err++;46 47 // CHECK: PASS48 if (err == 0)49 printf("PASS\n");50 return err;51}52