[gomp4] Initial support for accelerator support parsing

Hi!

This patch contains C++ parser changes etc. to handle
#pragma omp {teams,target {,data,update},distribute} parsing
all the way through till omp lowering (it bombs badly in omp expansion,
but already omp lowering will need to be tought out).
Things not handled yet are #pragma omp declare target/#pragma omp end
declare target (I assume we can handle it by automatically adding
"omp target" attribute to vars/functions in there) and there is no support
for array sections yet (also relevant for OMP_CLAUSE_DEPEND).

Say:

void baz (float *, float *, int);

float
foo (int x)
{
  float b[1024], c[1024], s = 0;
  int i;
  baz (b, c, x);        
  #pragma omp target map(to: b, c)
    #pragma omp parallel for reduction(+:s)
      for (i = 0; i < x; i++)
        s += b[i] * c[i];
  return s;
}

float
bar (int x, int y, int z)
{
  float b[1024], c[1024], s = 0;
  int i, j;
  baz (b, c, x);
  #pragma omp target data map(to: b)
  {
    #pragma omp target map(to: c)
      #pragma omp teams num_teams(y) num_threads(z) reduction(+:s)
        #pragma omp distribute dist_schedule(static, 4) collapse(1)
          for (j=0; j < x; j += y)
            #pragma omp parallel for reduction(+:s)
              for (i = j; i < j + y; i++)
                s += b[i] * c[i];
    #pragma omp target update from(b, c)
  }
  return s;
}

now parses with g++ -S -fopenmp -fdump-tree-all and shows dump up to
*.cfg (crash during ompexp).

Comments on this before I commit it to gomp-4_0-branch?

2013-05-24  Jakub Jelinek  <jakub@redhat.com>

	* tree.def (OMP_TEAMS, OMP_TARGET_DATA, OMP_TARGET,
	OMP_TARGET_UPDATE): New tree codes.
	* tree-cfg.c (make_edges): Handle GIMPLE_OMP_TARGET
	and GIMPLE_OMP_TEAMS.
	* omp-low.c (scan_sharing_clauses): Handle OMP_CLAUSE_DIST_SCHEDULE.
	* gimple-low.c (lower_stmt): Handle GIMPLE_OMP_TARGET
	and GIMPLE_OMP_TEAMS.
	* tree.h (OMP_TEAMS_BODY, OMP_TEAMS_CLAUSES, OMP_TARGET_DATA_BODY,
	OMP_TARGET_DATA_CLAUSES, OMP_TARGET_BODY, OMP_TARGET_CLAUSES,
	OMP_TARGET_UPDATE_CLAUSES): Define.
 	* tree-nested.c (convert_nonlocal_reference_stmt,
	convert_local_reference_stmt, convert_gimple_call): Handle
	GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* tree-inline.c (estimate_num_insns): Likewise.
	(remap_gimple_stmt): Likewise.  Adjust gimple_build_omp_for
	caller.
	* gimple.def: Adjust comments describing OMP_CLAUSEs.
	(GIMPLE_OMP_TARGET, GIMPLE_OMP_TEAMS): New GIMPLE stmts.
	* tree-parloops.c (create_parallel_loop): Adjust gimple_build_omp_for
	caller.
	* tree-pretty-print.c (dump_generic_node): Handle OMP_TEAMS,
	OMP_TARGET, OMP_TARGET_DATA and OMP_TARGET_UPDATE.
	* gimple.h (GF_OMP_TARGET_KIND_MASK, GF_OMP_TARGET_KIND_REGION,
	GF_OMP_TARGET_KIND_DATA, GF_OMP_TARGET_KIND_UPDATE): New.
	(gimple_build_omp_for): Add kind argument to prototype.
	(gimple_build_omp_target, gimple_build_omp_teams): New prototypes.
	(gimple_has_substatements): Handle GIMPLE_OMP_TARGET and
	GIMPLE_OMP_TEAMS.
	(gimple_omp_subcode): Change GIMPLE_OMP_SINGLE to GIMPLE_OMP_TEAMS.
	(gimple_omp_target_clauses, gimple_omp_target_clauses_ptr,
	gimple_omp_target_set_clauses, gimple_omp_target_kind,
	gimple_omp_target_set_kind, gimple_omp_teams_clauses,
	gimple_omp_teams_clauses_ptr, gimple_omp_teams_set_clauses): New
	inline functions.
	(gimple_return_set_retval): Handle GIMPLE_OMP_TARGET and
	GIMPLE_OMP_TEAMS.
	* gimple.c (gimple_build_omp_for): Add kind argument, call
	gimple_omp_for_set_kind.
	(gimple_build_omp_target, gimple_build_omp_teams): New functions.
	(walk_gimple_op, walk_gimple_stmt, gimple_copy): Handle
	GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* gimple-pretty-print.c (dump_gimple_omp_target,
	dump_gimple_omp_teams): New functions.
	(pp_gimple_stmt_1): Handle GIMPLE_OMP_TARGET and GIMPLE_OMP_TEAMS.
	* gimplify.c (enum gimplify_omp_var_data): Add GOVD_MAP.
	(enum omp_region_type): Add ORT_TEAMS, ORT_TARGET and ORT_TARGET_DATA.
	(omp_add_variable): Add temporary assertions.
	(omp_notice_threadprivate_variable): Complain if threadprivate vars
	appear in target region.
	(omp_notice_variable): ORT_TARGET, ORT_TARGET_DATA and ORT_TEAMS
	handling.
	(omp_check_private): Ignore ORT_TARGET and ORT_TARGET_DATA regions.
	(gimplify_scan_omp_clauses): Handle OMP_CLAUSE_MAP, OMP_CLAUSE_TO,
	OMP_CLAUSE_FROM, OMP_CLAUSE_NUM_TEAMS, OMP_CLAUSE_DIST_SCHEDULE
	and OMP_CLAUSE_DEVICE.
	(gimplify_adjust_omp_clauses): Likewise.
	(gimplify_adjust_omp_clauses_1): Handle GOVD_MAP.  Fix up
	check for privatization by also testing for GOVD_LINEAR.
	(gimplify_omp_for): Adjust gimple_build_omp_for caller.
	Clear *expr_p.
	(gimplify_omp_workshare): Handle also OMP_TARGET, OMP_TARGET_DATA
	and OMP_TEAMS.  Clear *expr_p.
	(gimplify_omp_target_update): New function.
	(gimplify_expr): Handle OMP_TARGET, OMP_TARGET_DATA, OMP_TARGET_UPDATE
	and OMP_TEAMS.
cp/
	* parser.c (cp_parser_omp_clause_cancelkind): Remove diagnostics.
	(cp_parser_omp_all_clauses): Require that OMP_CLAUSE_{TO,FROM}
	and OMP_CLAUSE_{PARALLEL,FOR,SECTIONS,TASKGROUP} must be first in
	the list of clauses.
	(OMP_TEAMS_CLAUSE_MASK, OMP_TARGET_CLAUSE_MASK,
	OMP_TARGET_DATA_CLAUSE_MASK, OMP_TARGET_UPDATE_CLAUSE_MASK,
	OMP_DISTRIBUTE_CLAUSE_MASK): Define.
	(cp_parser_omp_teams, cp_parser_omp_target, cp_parser_omp_target_data,
	cp_parser_omp_target_update, cp_parser_omp_distribute): New functions.
	(cp_parser_omp_construct): Handle PRAGMA_OMP_DISTRIBUTE and
	PRAGMA_OMP_TEAMS.
	(cp_parser_pragma): Handle PRAGMA_OMP_DISTRIBUTE, PRAGMA_OMP_TEAMS
	and PRAGMA_OMP_TARGET.
	* pt.c (tsubst_expr): Handle OMP_TEAMS, OMP_TARGET, OMP_TARGET_DATA
	and OMP_TARGET_UPDATE.

	Jakub

[gomp4] Initial support for accelerator support parsing

Commit Message

Patch