/* { dg-do run } */ #include #include int v = 6; void bar (long *x, long *y) { *x += 2; *y += 3; } int baz (void) { return 5; } #pragma omp declare target to (bar, baz, v) __attribute__((noinline, noclone)) void foo (int a, int b, long c, long d) { int err; if (omp_get_num_teams () != 1) abort (); /* The OpenMP 4.5 spec says that these expressions are evaluated before target region on combined target teams, so those cases are always fine. */ #pragma omp target map(from: err) err = omp_get_num_teams () != 1; if (err) abort (); #pragma omp target map(from: err) #pragma omp teams err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1; if (err) abort (); #pragma omp target teams map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1; if (err) abort (); #pragma omp target map(from: err) #pragma omp teams num_teams (4) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > 4; if (err) abort (); #pragma omp target teams num_teams (4) map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > 4; if (err) abort (); #pragma omp target map(from: err) #pragma omp teams thread_limit (7) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_thread_limit () > 7; if (err) abort (); #pragma omp target teams thread_limit (7) map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_thread_limit () > 7; if (err) abort (); #pragma omp target map(from: err) #pragma omp teams num_teams (4) thread_limit (8) { { err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8; } } if (err) abort (); #pragma omp target teams num_teams (4) thread_limit (8) map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > 4 || omp_get_thread_limit () > 8; if (err) abort (); #pragma omp target map(from: err) #pragma omp teams num_teams (a) thread_limit (b) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > a || omp_get_thread_limit () > b; if (err) abort (); #pragma omp target teams num_teams (a) thread_limit (b) map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > a || omp_get_thread_limit () > b; if (err) abort (); #pragma omp target map(from: err) #pragma omp teams num_teams (c + 1) thread_limit (d - 1) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1; if (err) abort (); #pragma omp target teams num_teams (c + 1) thread_limit (d - 1) map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1; if (err) abort (); #pragma omp target map (always, to: c, d) map(from: err) #pragma omp teams num_teams (c + 1) thread_limit (d - 1) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1; if (err) abort (); #pragma omp target data map (to: c, d) { #pragma omp target defaultmap (tofrom: scalar) bar (&c, &d); /* This is one of the cases which can't be generally optimized, the c and d are (or could be) already mapped and whether their device and original values match is unclear. */ #pragma omp target map (to: c, d) map(from: err) #pragma omp teams num_teams (c + 1) thread_limit (d - 1) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > c + 1 || omp_get_thread_limit () > d - 1; if (err) abort (); } /* This can't be optimized, there are function calls inside of target involved. */ #pragma omp target map(from: err) #pragma omp teams num_teams (baz () + 1) thread_limit (baz () - 1) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1; if (err) abort (); #pragma omp target teams num_teams (baz () + 1) thread_limit (baz () - 1) map(from: err) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > baz () + 1 || omp_get_thread_limit () > baz () - 1; if (err) abort (); /* This one can't be optimized, as v might have different value between host and target. */ #pragma omp target map(from: err) #pragma omp teams num_teams (v + 1) thread_limit (v - 1) err = omp_get_num_teams () < 1 || omp_get_thread_limit () < 1 || omp_get_num_teams () > v + 1 || omp_get_thread_limit () > v - 1; if (err) abort (); } int main () { foo (3, 5, 7, 9); return 0; }