On Fri, Aug 20, 2021 at 09:18:32PM +0200, Marcel Vollweiler wrote:
> --- a/gcc/c/c-parser.c
> +++ b/gcc/c/c-parser.c
> @@ -15864,37 +15864,81 @@ c_parser_omp_clause_map (c_parser *parser, tree
> list)
> }
>
> /* OpenMP 4.0:
> - device ( expression ) */
> +>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>>
> device ( expression )
Please remove all the >>>>>s.
> +
> + OpenMP 5.0:
> + device ( [device-modifier :] integer-expression )
> +
> + device-modifier:
> + ancestor | device_num */
>
> + /* A requires directive with the reverse_offload clause must be
> + specified. */
> + if ((omp_requires_mask & OMP_REQUIRES_REVERSE_OFFLOAD) == 0)
> + {
> + c_parser_error (parser, "a %<requires%> directive with the "
> + "%<reverse_offload%> clause must be "
> + "specified");
[BI think this diagnostics is confusing, it tells the user that it has to
do something but doesn't tell why. It is also not a parser error.
So I think it should be instead
error_at (tok->location, "%<ancestor%> device modifier not "
"preceded by %<requires%> directive "
"with %<reverse_offload%> clause");
> + parens.skip_until_found_close (parser);
> + return list;
> + }
> + ancestor = true;
> + }
> + if (!INTEGRAL_TYPE_P (TREE_TYPE (t)))
> + {
> + c_parser_error (parser, "expected integer expression");
> + return list;
> }
>
> + check_no_duplicate_clause (list, OMP_CLAUSE_DEVICE, "device");
> +
> + c = build_omp_clause (clause_loc, OMP_CLAUSE_DEVICE);
> +
> + OMP_CLAUSE_DEVICE_ID (c) = t;
> + OMP_CLAUSE_CHAIN (c) = list;
> + OMP_CLAUSE_DEVICE_ANCESTOR (c) = ancestor;
> +
> + list = c;
> return list;
> }
>
> diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
> index 5349ef1..b4d8d81 100644
> --- a/gcc/c/c-typeck.c
> +++ b/gcc/c/c-typeck.c
> @@ -15139,6 +15139,22 @@ c_finish_omp_clauses (tree clauses, enum
> c_omp_region_type ort)
> case OMP_CLAUSE_COLLAPSE:
> case OMP_CLAUSE_FINAL:
> case OMP_CLAUSE_DEVICE:
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE
> + && OMP_CLAUSE_DEVICE_ANCESTOR (c))
> + {
> + t = OMP_CLAUSE_DEVICE_ID (c);
> + if (TREE_CODE (t) == INTEGER_CST
> + && wi::to_widest (t) != 1)
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + "the %<device%> clause expression must evaluate to "
> + "%<1%>");
> + remove = true;
> + break;
> + }
> + }
> + /* FALLTHRU */
For the C FE, I'd suggest to move this to the c_parser_omp_clause_device
routine like other similar checking is done there too. And you can use
if (TREE_CODE (t) == INTEGER_CST && !integer_onep (t))
> + error_at (tok->location, "a %<requires%> directive with the "
> + "%<reverse_offload%> clause must be "
> + "specified");
See above.
> @@ -38562,6 +38601,7 @@ cp_parser_omp_clause_device (cp_parser *parser, tree
> list,
> c = build_omp_clause (location, OMP_CLAUSE_DEVICE);
> OMP_CLAUSE_DEVICE_ID (c) = t;
> OMP_CLAUSE_CHAIN (c) = list;
> + OMP_CLAUSE_DEVICE_ANCESTOR (c) = ancestor;
But in C++ the INTEGER_CST checking shouldn't be done here, because
the argument could be type or value dependent.
> --- a/gcc/cp/semantics.c
> +++ b/gcc/cp/semantics.c
> @@ -7334,6 +7334,15 @@ finish_omp_clauses (tree clauses, enum
> c_omp_region_type ort)
> "%<device%> id must be integral");
> remove = true;
> }
> + else if (OMP_CLAUSE_DEVICE_ANCESTOR (c)
> + && TREE_CODE (t) == INTEGER_CST
> + && wi::to_widest (t) != 1)
!integer_onep (t)
> + if (!(gfc_current_ns->omp_requires & OMP_REQ_REVERSE_OFFLOAD))
> + {
> + gfc_error ("a %<requires%> directive with the "
> + "%<reverse_offload%> clause must be "
> + "specified at %C");
See above.
> + else if (gfc_match ("%e )", &c->device) == MATCH_YES)
> + {
> + }
> + else
Better != MATCH_YES and drop the {} else ?
> + {
> + gfc_error ("Expected integer expression or a single device-"
> + "modifier %<device_num%> or %<ancestor%> at %C");
> + break;
> + }
> + continue;
> + }
> if ((mask & OMP_CLAUSE_DEVICE)
> && openacc
> && gfc_match ("device ( ") == MATCH_YES
> + if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEVICE
> + && OMP_CLAUSE_DEVICE_ANCESTOR (c))
> + {
> + if (code != OMP_TARGET)
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + "%<device%> clause with %<ancestor%> is only "
> + "allowed on %<target%> construct");
> + remove = true;
> + }
Formatting, {/} are correctly indented, but error_at and remove should be
indented 2 columns to the right from that, not 4 columns.
Also, it should have break; there too, so that it doesn't fallthrough
to the next one:
> +
> + tree clauses = *orig_list_p;
> + for (; clauses ; clauses = OMP_CLAUSE_CHAIN (clauses))
> + if (OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_DEVICE
> + && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_FIRSTPRIVATE
> + && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_PRIVATE
> + && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_DEFAULTMAP
> + && OMP_CLAUSE_CODE (clauses) != OMP_CLAUSE_MAP
> + )
> + {
> + error_at (OMP_CLAUSE_LOCATION (c),
> + "with %<ancestor%>, only the %<device%>, "
> + "%<firstprivate%>, %<private%>, %<defaultmap%>, "
> + "and %<map%> clauses may appear on the "
> + "construct");
> + remove = true;
> + }
> + }
> + /* Fall through. */
> +
> @@ -4001,6 +4011,20 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool
> *handled_ops_p,
> "OpenMP runtime API call %qD in a region with "
> "%<order(concurrent)%> clause", fndecl);
> }
> + if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
> + && gimple_omp_target_kind (ctx->stmt) ==
> + GF_OMP_TARGET_KIND_REGION)
Formatting. Neither == nor = should be at the end of lines. So,
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
&& (gimple_omp_target_kind (ctx->stmt)
== GF_OMP_TARGET_KIND_REGION))
> + {
> + tree c =
> + omp_find_clause (gimple_omp_target_clauses (ctx->stmt),
> + OMP_CLAUSE_DEVICE);
And probably use a tree tgt_clauses = gimple_omp_target_clauses (ctx->stmt);
temporary to make tree c = omp_find_clause (tgt_clauses, OMP_CLAUSE_DEVICE);
fit nicely.
> + if (c && OMP_CLAUSE_DEVICE_ANCESTOR (c))
> + {
> + error_at (gimple_location (stmt),
> + "OpenMP runtime API call %qD in a region with "
> + "%<device(ancestor)%> clause", fndecl);
> + }
Single statement in if body shouldn't be wrapped with {}s.
> + }
> }
> }
> }
> diff --git a/gcc/testsuite/c-c++-common/gomp/target-device-1.c
> b/gcc/testsuite/c-c++-common/gomp/target-device-1.c
> new file mode 100644
> index 0000000..dafa643
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/gomp/target-device-1.c
> @@ -0,0 +1,34 @@
> +/* { dg-do compile } */
> +
> +void
> +foo (void)
> +{
> + /* Test to ensure that 'device_num' is parsed correctly in device clauses.
> */
> +
> + int n;
Better use
foo (int n)
such that it isn't an uninitialized use. Or initialize n to something.
> +
> + #pragma omp target device (1)
> + ;
> +
> + #pragma omp target device (n)
> + ;
> +
> + #pragma omp target device (n + 1)
> + ;
> +
> +end
> \ No newline at end of file
Please avoid these in all the tests.
Otherwise LGTM.
Jakub
