diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 4ec0ee85ac4..d0235809fb3 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -25271,6 +25271,7 @@ c_parser_omp_dispatch_body (c_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INTEROP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) @@ -26963,6 +26964,9 @@ c_finish_omp_declare_variant (c_parser *parser, tree fndecl, tree parms) { error_at (c_parser_peek_token (parser)->location, "expected % or %"); + if (strcmp (p, "need_device_addr") == 0) + inform (c_parser_peek_token (parser)->location, + "% is not valid for C"); goto fail; } } diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 8a3472a4b34..15a5253b50d 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -49924,6 +49924,7 @@ cp_parser_omp_dispatch_body (cp_parser *parser) | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEPEND) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOVARIANTS) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOCONTEXT) \ + | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_HAS_DEVICE_ADDR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_INTEROP) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_IS_DEVICE_PTR) \ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NOWAIT)) diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 8495c45eddb..dd0d992a958 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -4124,27 +4124,39 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) arg_types = TREE_CHAIN (arg_types); bool need_device_ptr = false; - for (tree arg - = TREE_PURPOSE (TREE_VALUE (adjust_args_list)); - arg != NULL; arg = TREE_CHAIN (arg)) - { - if (TREE_VALUE (arg) - && TREE_CODE (TREE_VALUE (arg)) == INTEGER_CST - && wi::eq_p (i, wi::to_wide (TREE_VALUE (arg)))) - { - need_device_ptr = true; - break; - } - } + bool need_device_addr = false; + for (int need_addr = 0; need_addr <= 1; need_addr++) + for (tree arg = need_addr + ? TREE_VALUE (TREE_VALUE ( + adjust_args_list)) + : TREE_PURPOSE (TREE_VALUE ( + adjust_args_list)); + arg != NULL; arg = TREE_CHAIN (arg)) + { + if (TREE_VALUE (arg) + && TREE_CODE (TREE_VALUE (arg)) == INTEGER_CST + && wi::eq_p (i, wi::to_wide (TREE_VALUE (arg)))) + { + if (need_addr) + need_device_addr = true; + else + need_device_ptr = true; + break; + } + } - if (need_device_ptr) + if (need_device_ptr || need_device_addr) { bool is_device_ptr = false; + bool has_device_addr = false; + for (tree c = gimplify_omp_ctxp->clauses; c; c = TREE_CHAIN (c)) { - if (OMP_CLAUSE_CODE (c) - == OMP_CLAUSE_IS_DEVICE_PTR) + if ((OMP_CLAUSE_CODE (c) + == OMP_CLAUSE_IS_DEVICE_PTR) + || (OMP_CLAUSE_CODE (c) + == OMP_CLAUSE_HAS_DEVICE_ADDR)) { tree decl1 = DECL_NAME (OMP_CLAUSE_DECL (c)); tree decl2 @@ -4155,15 +4167,42 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value) || TREE_CODE (decl2) == PARM_DECL) { decl2 = DECL_NAME (decl2); - if (decl1 == decl2) - is_device_ptr = true; + if (decl1 == decl2 + && (OMP_CLAUSE_CODE (c) + == OMP_CLAUSE_IS_DEVICE_PTR)) + { + if (need_device_addr) + warning_at ( + OMP_CLAUSE_LOCATION (c), + OPT_Wopenmp, + "% for %qD does" + " not imply % " + "required for " + "%", + OMP_CLAUSE_DECL (c)); + is_device_ptr = true; + } + else if (decl1 == decl2) + { + if (need_device_ptr) + warning_at ( + OMP_CLAUSE_LOCATION (c), + OPT_Wopenmp, + "% for %qD does" + " not imply % " + "required for " + "%", + OMP_CLAUSE_DECL (c)); + has_device_addr = true; + } } } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE) device_num = OMP_CLAUSE_OPERAND (c, 0); } - if (!is_device_ptr) + if ((need_device_ptr && !is_device_ptr) + || (need_device_addr && !has_device_addr)) { if (device_num == NULL_TREE) { diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c new file mode 100644 index 00000000000..f62272cfb01 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-3.c @@ -0,0 +1,85 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +// Do diagnostic check / dump check only; +// Note: this test should work as run-test as well. + +#if 0 + #include +#else + #ifdef __cplusplus + extern "C" { + #endif + extern int omp_get_default_device (); + extern int omp_get_num_devices (); + #ifdef __cplusplus + } + #endif +#endif + + +void f(int *x, int *y); +#pragma omp declare variant(f) adjust_args(need_device_ptr: x, y) match(construct={dispatch}) +void g(int *x, int *y); + +void +sub (int *a, int *b) +{ + // The has_device_addr is a bit questionable as the caller is not actually + // passing a device address - but we cannot pass one because of the + // following: + // + // As for 'b' need_device_ptr has been specified and 'b' is not + // in the semantic requirement set 'is_device_ptr' (and only in 'has_device_addr') + // "the argument is converted in the same manner that a use_device_ptr clause + // on a target_data construct converts its pointer" + #pragma omp dispatch is_device_ptr(a), has_device_addr(b) /* { dg-warning "'has_device_addr' for 'b' does not imply 'is_device_ptr' required for 'need_device_ptr' \\\[-Wopenmp\\\]" } */ + g(a, b); +} + +void +f(int *from, int *to) +{ + static int cnt = 0; + cnt++; + if (cnt >= 3) + { + if (omp_get_default_device () != -1 + && omp_get_default_device () < omp_get_num_devices ()) + { + // On offload device but not mapped + if (from != (void *)0L) // Not mapped + __builtin_abort (); + } + else if (from[0] != 5) + __builtin_abort (); + return; + } + #pragma omp target is_device_ptr(from, to) + { + to[0] = from[0] * 10; + to[1] = from[1] * 10; + } +} + +int +main () +{ + int A[2], B[2] = {123, 456}, C[1] = {5}; + int *p = A; + #pragma omp target enter data map(A, B) + + /* Note: We don't add 'use_device_addr(B)' here; + if we do, it will fail with an illegal memory access (why?). */ + #pragma omp target data use_device_ptr(p) + { + sub(p, B); + sub(C, B); /* C is not mapped -> 'from' ptr == NULL */ + } + + #pragma omp target exit data map(A, B) +} + +// { dg-final { scan-tree-dump-times "#pragma omp dispatch has_device_addr\\(b\\) is_device_ptr\\(a\\)" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "__builtin_omp_get_mapped_ptr" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = __builtin_omp_get_mapped_ptr \\(b" 1 "gimple" } } +// { dg-final { scan-tree-dump-times "f \\(a, D\\.\[0-9\]+\\);" 1 "gimple" } } diff --git a/gcc/testsuite/gcc.dg/gomp/adjust-args-2.c b/gcc/testsuite/gcc.dg/gomp/adjust-args-2.c new file mode 100644 index 00000000000..ee4feffb2aa --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/adjust-args-2.c @@ -0,0 +1,5 @@ +void f(int *); +#pragma omp declare variant(f) adjust_args(need_device_addr: x) +/* { dg-error "expected 'nothing' or 'need_device_ptr'" "" { target *-*-* } .-1 } */ +/* { dg-note "'need_device_addr' is not valid for C" "" { target *-*-* } .-2 } */ +void g(int *x);