[Patch v2] C, C++, Fortran, OpenMP: Add support for device-modifiers for 'omp target device'

Jakub Jelinek jakub@redhat.com
Mon Aug 23 17:47:33 GMT 2021


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



More information about the Gcc-patches mailing list