forked from OSchip/llvm-project
64 lines
1.8 KiB
C++
64 lines
1.8 KiB
C++
// RUN: %libomptarget-compilexx-run-and-check-aarch64-unknown-linux-gnu
|
|
// RUN: %libomptarget-compilexx-run-and-check-powerpc64-ibm-linux-gnu
|
|
// RUN: %libomptarget-compilexx-run-and-check-powerpc64le-ibm-linux-gnu
|
|
// RUN: %libomptarget-compilexx-run-and-check-x86_64-pc-linux-gnu
|
|
// RUN: %libomptarget-compilexx-run-and-check-nvptx64-nvidia-cuda
|
|
|
|
#include <omp.h>
|
|
#include <stdio.h>
|
|
|
|
#define N 1024
|
|
|
|
int A[N];
|
|
int B[N];
|
|
int C[N];
|
|
int main() {
|
|
for (int i = 0; i < N; i++)
|
|
A[i] = B[i] = i;
|
|
|
|
#pragma omp parallel num_threads(2)
|
|
{
|
|
if (omp_get_thread_num() == 1) {
|
|
// map data A & B and move to
|
|
#pragma omp target enter data map(to : A, B) depend(out : A[0]) nowait
|
|
|
|
// no data move since already mapped
|
|
#pragma omp target map(A, B) depend(out : A[0]) nowait
|
|
{
|
|
for (int i = 0; i < N; i++)
|
|
++A[i];
|
|
for (int i = 0; i < N; i++)
|
|
++B[i];
|
|
}
|
|
|
|
// no data move since already mapped
|
|
#pragma omp target teams num_teams(1) map(A, B) depend(out : A[0]) nowait
|
|
{
|
|
for (int i = 0; i < N; i++)
|
|
++A[i];
|
|
for (int i = 0; i < N; i++)
|
|
++B[i];
|
|
}
|
|
|
|
// A updated via update
|
|
#pragma omp target update from(A) depend(out : A[0]) nowait
|
|
|
|
// B updated via exit, A just released
|
|
#pragma omp target exit data map(release \
|
|
: A) map(from \
|
|
: B) depend(out \
|
|
: A[0]) nowait
|
|
} // if
|
|
} // parallel
|
|
|
|
int Sum = 0;
|
|
for (int i = 0; i < N; i++)
|
|
Sum += A[i] + B[i];
|
|
// Sum is 2 * N * (2 + N - 1 + 2) / 2
|
|
// CHECK: Sum = 1051648.
|
|
printf("Sum = %d.\n", Sum);
|
|
|
|
return Sum != 2 * N * (2 + N - 1 + 2) / 2;
|
|
}
|
|
|