From bf14158b1a28c2c5b29c41071fa62c011d9f4f65 Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usmanov@samsung.com>
Date: Thu, 13 Feb 2014 15:58:28 +0400
Subject: [PATCH] OpenACC GENERIC nodes
---
gcc/doc/generic.texi | 45 ++++++++++++++++++
gcc/gimplify.c | 62 +++++++++++++++++++++++++
gcc/omp-low.c | 96 ++++++++++++++++++++++++++++++++------
gcc/tree-core.h | 61 ++++++++++++++++++++++---
gcc/tree-pretty-print.c | 119 ++++++++++++++++++++++++++++++++++++++++++++++++
gcc/tree.c | 44 +++++++++++++++++-
gcc/tree.def | 42 +++++++++++++++++
gcc/tree.h | 61 ++++++++++++++++++++++++-
8 files changed, 507 insertions(+), 23 deletions(-)
@@ -2052,6 +2052,15 @@ edge. Rethrowing the exception is represented using @code{RESX_EXPR}.
@node OpenMP
@subsection OpenMP
@tindex OACC_PARALLEL
+@tindex OACC_KERNELS
+@tindex OACC_DATA
+@tindex OACC_HOST_DATA
+@tindex OACC_DECLARE
+@tindex OACC_UPDATE
+@tindex OACC_ENTER_DATA
+@tindex OACC_EXIT_DATA
+@tindex OACC_WAIT
+@tindex OACC_CACHE
@tindex OMP_PARALLEL
@tindex OMP_FOR
@tindex OMP_SECTIONS
@@ -2073,6 +2082,42 @@ clauses used by the OpenMP API @w{@uref{http://www.openmp.org/}}.
Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
+@item OACC_KERNELS
+
+Represents @code{#pragma acc kernels [clause1 @dots{} clauseN]}.
+
+@item OACC_DATA
+
+Represents @code{#pragma acc data [clause1 @dots{} clauseN]}.
+
+@item OACC_HOST_DATA
+
+Represents @code{#pragma acc host_data [clause1 @dots{} clauseN]}.
+
+@item OACC_DECLARE
+
+Represents @code{#pragma acc declare [clause1 @dots{} clauseN]}.
+
+@item OACC_UPDATE
+
+Represents @code{#pragma acc update [clause1 @dots{} clauseN]}.
+
+@item OACC_ENTER_DATA
+
+Represents @code{#pragma acc enter data [clause1 @dots{} clauseN]}.
+
+@item OACC_EXIT_DATA
+
+Represents @code{#pragma acc exit data [clause1 @dots{} clauseN]}.
+
+@item OACC_WAIT
+
+Represents @code{#pragma acc wait [(num @dots{})]}.
+
+@item OACC_CACHE
+
+Represents @code{#pragma acc cache (var @dots{})}.
+
@item OMP_PARALLEL
Represents @code{#pragma omp parallel [clause1 @dots{} clauseN]}. It
@@ -4333,6 +4333,15 @@ is_gimple_stmt (tree t)
case ASM_EXPR:
case STATEMENT_LIST:
case OACC_PARALLEL:
+ case OACC_KERNELS:
+ case OACC_DATA:
+ case OACC_HOST_DATA:
+ case OACC_DECLARE:
+ case OACC_UPDATE:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OACC_WAIT:
+ case OACC_CACHE:
case OMP_PARALLEL:
case OMP_FOR:
case OMP_SIMD:
@@ -6157,6 +6166,23 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
+ case OMP_CLAUSE_HOST:
+ case OMP_CLAUSE_OACC_DEVICE:
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WAIT:
+ case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ remove = true;
+ break;
+
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_UNTIED:
@@ -6498,6 +6524,20 @@ gimplify_adjust_omp_clauses (tree *list_p)
case OMP_CLAUSE_DEPEND:
break;
+ case OMP_CLAUSE_HOST:
+ case OMP_CLAUSE_OACC_DEVICE:
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WAIT:
+ case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
default:
gcc_unreachable ();
}
@@ -7988,6 +8028,19 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
+ case OACC_KERNELS:
+ case OACC_DATA:
+ case OACC_HOST_DATA:
+ case OACC_DECLARE:
+ case OACC_UPDATE:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ case OACC_WAIT:
+ case OACC_CACHE:
+ sorry ("directive not yet implemented");
+ ret = GS_ALL_DONE;
+ break;
+
case OMP_PARALLEL:
gimplify_omp_parallel (expr_p, pre_p);
ret = GS_ALL_DONE;
@@ -8396,6 +8449,15 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
&& code != SWITCH_EXPR
&& code != TRY_FINALLY_EXPR
&& code != OACC_PARALLEL
+ && code != OACC_KERNELS
+ && code != OACC_DATA
+ && code != OACC_HOST_DATA
+ && code != OACC_DECLARE
+ && code != OACC_UPDATE
+ && code != OACC_ENTER_DATA
+ && code != OACC_EXIT_DATA
+ && code != OACC_WAIT
+ && code != OACC_CACHE
&& code != OMP_CRITICAL
&& code != OMP_FOR
&& code != OMP_MASTER
@@ -1491,6 +1491,18 @@ fixup_child_record_type (omp_context *ctx)
TREE_TYPE (ctx->receiver_decl) = build_pointer_type (type);
}
+static bool
+gimple_code_is_oacc (const_gimple g)
+{
+ switch (gimple_code (g))
+ {
+ case GIMPLE_OACC_PARALLEL:
+ return true;
+ default:
+ return false;
+ }
+}
+
/* Instantiate decls as necessary in CTX to satisfy the data sharing
specified by CLAUSES. */
@@ -1552,8 +1564,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_REDUCTION:
+ if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
case OMP_CLAUSE_LINEAR:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
do_private:
if (is_variable_sized (decl))
@@ -1591,7 +1608,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_COPYIN:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
by_ref = use_pointer_for_field (decl, NULL);
install_var_field (decl, by_ref, 3, ctx);
@@ -1602,8 +1619,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
break;
- case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_IF:
+ if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
+ case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
@@ -1611,14 +1633,14 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_SCHEDULE:
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEPEND:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
case OMP_CLAUSE_MAP:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
@@ -1641,7 +1663,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
#pragma omp target data, there is nothing to map for
those. */
- if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ if (!gimple_code_is_oacc (ctx->stmt)
&& gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
@@ -1710,17 +1732,34 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_MERGEABLE:
case OMP_CLAUSE_PROC_BIND:
case OMP_CLAUSE_SAFELEN:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
break;
case OMP_CLAUSE_ALIGNED:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
if (is_global_var (decl)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_HOST:
+ case OMP_CLAUSE_OACC_DEVICE:
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WAIT:
+ case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ sorry ("Clause not supported yet");
+ break;
+
default:
gcc_unreachable ();
}
@@ -1731,7 +1770,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_LASTPRIVATE:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
/* Let the corresponding firstprivate clause create
the variable. */
if (OMP_CLAUSE_LASTPRIVATE_GIMPLE_SEQ (c))
@@ -1743,8 +1782,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_FIRSTPRIVATE:
case OMP_CLAUSE_REDUCTION:
+ if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
case OMP_CLAUSE_LINEAR:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
decl = OMP_CLAUSE_DECL (c);
if (is_variable_sized (decl))
install_var_local (decl, ctx);
@@ -1757,7 +1801,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_SHARED:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
/* Ignore shared directives in teams construct. */
if (gimple_code (ctx->stmt) == GIMPLE_OMP_TEAMS)
break;
@@ -1767,7 +1811,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_MAP:
- if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ if (!gimple_code_is_oacc (ctx->stmt)
&& gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
break;
decl = OMP_CLAUSE_DECL (c);
@@ -1776,7 +1820,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& lookup_attribute ("omp declare target",
DECL_ATTRIBUTES (decl)))
{
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
break;
}
if (DECL_P (decl))
@@ -1804,10 +1848,15 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
break;
+ case OMP_CLAUSE_IF:
+ if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL)
+ {
+ sorry ("clause not supported yet");
+ break;
+ }
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_DEFAULT:
- case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
@@ -1827,7 +1876,24 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE__LOOPTEMP_:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ gcc_assert (!gimple_code_is_oacc (ctx->stmt));
+ break;
+
+ case OMP_CLAUSE_HOST:
+ case OMP_CLAUSE_OACC_DEVICE:
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WAIT:
+ case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE_INDEPENDENT:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ sorry ("Clause not supported yet");
break;
default:
@@ -213,19 +213,19 @@ enum omp_clause_code {
(c_parser_omp_variable_list). */
OMP_CLAUSE_ERROR = 0,
- /* OpenMP clause: private (variable_list). */
+ /* OpenMP/OpenACC clause: private (variable_list). */
OMP_CLAUSE_PRIVATE,
/* OpenMP clause: shared (variable_list). */
OMP_CLAUSE_SHARED,
- /* OpenMP clause: firstprivate (variable_list). */
+ /* OpenMP/OpenACC clause: firstprivate (variable_list). */
OMP_CLAUSE_FIRSTPRIVATE,
/* OpenMP clause: lastprivate (variable_list). */
OMP_CLAUSE_LASTPRIVATE,
- /* OpenMP clause: reduction (operator:variable_list).
+ /* OpenMP/OpenACC clause: reduction (operator:variable_list).
OMP_CLAUSE_REDUCTION_CODE: The tree_code of the operator.
Operand 1: OMP_CLAUSE_REDUCTION_INIT: Stmt-list to initialize the var.
Operand 2: OMP_CLAUSE_REDUCTION_MERGE: Stmt-list to merge private var
@@ -265,10 +265,37 @@ enum omp_clause_code {
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* OpenACC clause: host (variable_list). */
+ OMP_CLAUSE_HOST,
+
+ /* OpenACC clause: device (variable_list). */
+ OMP_CLAUSE_OACC_DEVICE,
+
+ /* OpenACC clause: device_resident (variable_list). */
+ OMP_CLAUSE_DEVICE_RESIDENT,
+
+ /* OpenACC clause: use_device (variable_list). */
+ OMP_CLAUSE_USE_DEVICE,
+
+ /* OpenACC clause: gang [(gang-argument-list)].
+ Where
+ gang-argument-list: [gang-argument-list, ] gang-argument
+ gang-argument: [num:] integer-expression
+ | static: size-expression
+ size-expression: * | integer-expression. */
+ OMP_CLAUSE_GANG,
+
+ /* OpenACC clause/directive: wait [(integer-expression-list)]. */
+ OMP_CLAUSE_WAIT,
+
+ /* Internal structure to hold OpenACC cache directive's variable-list.
+ #pragma acc cache (variable-list). */
+ OMP_NO_CLAUSE_CACHE,
+
/* Internal clause: temporary for combined loops expansion. */
OMP_CLAUSE__LOOPTEMP_,
- /* OpenMP clause: if (scalar-expression). */
+ /* OpenMP/OpenACC clause: if (scalar-expression). */
OMP_CLAUSE_IF,
/* OpenMP clause: num_threads (integer-expression). */
@@ -281,12 +308,13 @@ enum omp_clause_code {
OMP_CLAUSE_NOWAIT,
/* OpenMP clause: ordered. */
+ /* OpenACC clause: seq. */
OMP_CLAUSE_ORDERED,
/* OpenMP clause: default. */
OMP_CLAUSE_DEFAULT,
- /* OpenMP clause: collapse (constant-integer-expression). */
+ /* OpenMP/OpenACC clause: collapse (constant-integer-expression). */
OMP_CLAUSE_COLLAPSE,
/* OpenMP clause: untied. */
@@ -338,7 +366,28 @@ enum omp_clause_code {
OMP_CLAUSE_TASKGROUP,
/* Internally used only clause, holding SIMD uid. */
- OMP_CLAUSE__SIMDUID_
+ OMP_CLAUSE__SIMDUID_,
+
+ /* OpenACC clause: independent. */
+ OMP_CLAUSE_INDEPENDENT,
+
+ /* OpenACC clause: async [(integer-expression)]. */
+ OMP_CLAUSE_ASYNC,
+
+ /* OpenACC clause: worker [( [num:] integer-expression)]. */
+ OMP_CLAUSE_WORKER,
+
+ /* OpenACC clause: vector [( [length:] integer-expression)]. */
+ OMP_CLAUSE_VECTOR,
+
+ /* OpenACC clause: num_gangs (integer-expression). */
+ OMP_CLAUSE_NUM_GANGS,
+
+ /* OpenACC clause: num_workers (integer-expression). */
+ OMP_CLAUSE_NUM_WORKERS,
+
+ /* OpenACC clause: vector_length (integer-expression). */
+ OMP_CLAUSE_VECTOR_LENGTH
};
#undef DEFTREESTRUCT
@@ -326,6 +326,21 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE__LOOPTEMP_:
name = "_looptemp_";
goto print_remap;
+ case OMP_CLAUSE_HOST:
+ name = "host";
+ goto print_remap;
+ case OMP_CLAUSE_OACC_DEVICE:
+ name = "device";
+ goto print_remap;
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ name = "device_resident";
+ goto print_remap;
+ case OMP_CLAUSE_USE_DEVICE:
+ name = "use_device";
+ goto print_remap;
+ case OMP_NO_CLAUSE_CACHE:
+ name = "_cache_";
+ goto print_remap;
print_remap:
pp_string (buffer, name);
pp_left_paren (buffer);
@@ -634,6 +649,62 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
pp_right_paren (buffer);
break;
+ case OMP_CLAUSE_GANG:
+ pp_string (buffer, "gang(");
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_WAIT:
+ pp_string (buffer, "wait(");
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_ASYNC:
+ pp_string (buffer, "async");
+ if (OMP_CLAUSE_DECL (clause))
+ {
+ pp_character(buffer, '(');
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause),
+ spc, flags, false);
+ pp_character(buffer, ')');
+ }
+ break;
+
+ case OMP_CLAUSE_WORKER:
+ pp_string (buffer, "worker(");
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_VECTOR:
+ pp_string (buffer, "vector(");
+ dump_generic_node (buffer, OMP_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OMP_CLAUSE_NUM_GANGS:
+ pp_string (buffer, "num_gangs(");
+ dump_generic_node (buffer, OMP_CLAUSE_NUM_GANGS_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OMP_CLAUSE_NUM_WORKERS:
+ pp_string (buffer, "num_workers(");
+ dump_generic_node (buffer, OMP_CLAUSE_NUM_WORKERS_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OMP_CLAUSE_VECTOR_LENGTH:
+ pp_string (buffer, "vector_length(");
+ dump_generic_node (buffer, OMP_CLAUSE_VECTOR_LENGTH_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
case OMP_CLAUSE_INBRANCH:
pp_string (buffer, "inbranch");
break;
@@ -652,6 +723,9 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE_TASKGROUP:
pp_string (buffer, "taskgroup");
break;
+ case OMP_CLAUSE_INDEPENDENT:
+ pp_string (buffer, "independent");
+ break;
default:
/* Should never happen. */
@@ -2384,6 +2458,51 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
dump_omp_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags);
goto dump_omp_body;
+ case OACC_KERNELS:
+ pp_string (buffer, "#pragma acc kernels");
+ dump_omp_clauses (buffer, OACC_KERNELS_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_DATA:
+ pp_string (buffer, "#pragma acc data");
+ dump_omp_clauses (buffer, OACC_DATA_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_HOST_DATA:
+ pp_string (buffer, "#pragma acc host_data");
+ dump_omp_clauses (buffer, OACC_HOST_DATA_CLAUSES (node), spc, flags);
+ goto dump_omp_body;
+
+ case OACC_DECLARE:
+ pp_string (buffer, "#pragma acc declare");
+ dump_omp_clauses (buffer, OACC_DECLARE_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_UPDATE:
+ pp_string (buffer, "#pragma acc update");
+ dump_omp_clauses (buffer, OACC_UPDATE_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_ENTER_DATA:
+ pp_string (buffer, "#pragma acc enter data");
+ dump_omp_clauses (buffer, OACC_ENTER_DATA_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_EXIT_DATA:
+ pp_string (buffer, "#pragma acc exit data");
+ dump_omp_clauses (buffer, OACC_EXIT_DATA_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_WAIT:
+ pp_string (buffer, "#pragma acc wait");
+ dump_omp_clauses (buffer, OACC_WAIT_CLAUSES (node), spc, flags);
+ break;
+
+ case OACC_CACHE:
+ pp_string (buffer, "#pragma acc cache");
+ dump_omp_clauses (buffer, OACC_CACHE_CLAUSES(node), spc, flags);
+ break;
+
case OMP_PARALLEL:
pp_string (buffer, "#pragma omp parallel");
dump_omp_clauses (buffer, OMP_PARALLEL_CLAUSES (node), spc, flags);
@@ -259,6 +259,13 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 1, /* OMP_CLAUSE_HOST */
+ 1, /* OMP_CLAUSE_OACC_DEVICE */
+ 1, /* OMP_CLAUSE_DEVICE_RESIDENT */
+ 1, /* OMP_CLAUSE_USE_DEVICE */
+ 1, /* OMP_CLAUSE_GANG */
+ 1, /* OMP_CLAUSE_WAIT */
+ 1, /* OMP_NO_CLAUSE_CACHE */
1, /* OMP_CLAUSE__LOOPTEMP_ */
1, /* OMP_CLAUSE_IF */
1, /* OMP_CLAUSE_NUM_THREADS */
@@ -284,6 +291,13 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_SECTIONS */
0, /* OMP_CLAUSE_TASKGROUP */
1, /* OMP_CLAUSE__SIMDUID_ */
+ 0, /* OMP_CLAUSE_INDEPENDENT */
+ 1, /* OMP_CLAUSE_ASYNC */
+ 1, /* OMP_CLAUSE_WORKER */
+ 1, /* OMP_CLAUSE_VECTOR */
+ 1, /* OMP_CLAUSE_NUM_GANGS */
+ 1, /* OMP_CLAUSE_NUM_WORKERS */
+ 1, /* OMP_CLAUSE_VECTOR_LENGTH */
};
const char * const omp_clause_code_name[] =
@@ -303,6 +317,13 @@ const char * const omp_clause_code_name[] =
"from",
"to",
"map",
+ "host",
+ "device",
+ "device_resident",
+ "use_device",
+ "gang",
+ "wait",
+ "_cache_",
"_looptemp_",
"if",
"num_threads",
@@ -327,7 +348,14 @@ const char * const omp_clause_code_name[] =
"parallel",
"sections",
"taskgroup",
- "_simduid_"
+ "_simduid_",
+ "independent",
+ "async",
+ "worker",
+ "vector",
+ "num_gangs",
+ "num_workers",
+ "vector_length"
};
@@ -11034,6 +11062,19 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
switch (OMP_CLAUSE_CODE (*tp))
{
+ case OMP_CLAUSE_HOST:
+ case OMP_CLAUSE_OACC_DEVICE:
+ case OMP_CLAUSE_DEVICE_RESIDENT:
+ case OMP_CLAUSE_USE_DEVICE:
+ case OMP_CLAUSE_GANG:
+ case OMP_CLAUSE_WAIT:
+ case OMP_NO_CLAUSE_CACHE:
+ case OMP_CLAUSE_ASYNC:
+ case OMP_CLAUSE_WORKER:
+ case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NUM_GANGS:
+ case OMP_CLAUSE_NUM_WORKERS:
+ case OMP_CLAUSE_VECTOR_LENGTH:
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -11056,6 +11097,7 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0));
/* FALLTHRU */
+ case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_DEFAULT:
@@ -1017,6 +1017,24 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
+/* OpenACC - #pragma acc kernels [clause1 ... clauseN]
+ Operand 0: OACC_KERNELS_BODY: Sequence of kernels.
+ Operand 1: OACC_KERNELS_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
+
+/* OpenACC - #pragma acc data [clause1 ... clauseN]
+ Operand 0: OACC_DATA_BODY: Data construct body.
+ Operand 1: OACC_DATA_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2)
+
+/* OpenACC - #pragma acc host_data [clause1 ... clauseN]
+ Operand 0: OACC_HOST_DATA_BODY: Host_data construct body.
+ Operand 1: OACC_HOST_DATA_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_HOST_DATA, "oacc_host_data", tcc_statement, 2)
+
/* OpenMP - #pragma omp parallel [clause1 ... clauseN]
Operand 0: OMP_PARALLEL_BODY: Code to be executed by all threads.
Operand 1: OMP_PARALLEL_CLAUSES: List of clauses. */
@@ -1108,6 +1126,30 @@ DEFTREECODE (OMP_ORDERED, "omp_ordered", tcc_statement, 1)
Operand 1: OMP_CRITICAL_NAME: Identifier for critical section. */
DEFTREECODE (OMP_CRITICAL, "omp_critical", tcc_statement, 2)
+/* OpenACC - #pragma acc declare [clause1 ... clauseN]
+ Operand 0: OACC_DECLARE_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
+
+/* OpenACC - #pragma acc update [clause1 ... clauseN]
+ Operand 0: OACC_UPDATE_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_UPDATE, "oacc_update", tcc_statement, 1)
+
+/* OpenACC - #pragma acc enter data [clause1 ... clauseN]
+ Operand 0: OACC_ENTER_DATA_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1)
+
+/* OpenACC - #pragma acc exit data [clause1 ... clauseN]
+ Operand 0: OACC_EXIT_DATA_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_EXIT_DATA, "oacc_exit_data", tcc_statement, 1)
+
+/* OpenACC - #pragma acc wait [clause1 ... clauseN]
+ Operand 0: OACC_WAIT_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_WAIT, "oacc_wait", tcc_statement, 1)
+
+/* OpenACC - #pragma acc cache [clause1 ... clauseN]
+ Operand 0: OACC_CACHE_CLAUSES: List of clauses. */
+DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
+
/* OpenMP - #pragma omp target update [clause1 ... clauseN]
Operand 0: OMP_TARGET_UPDATE_CLAUSES: List of clauses. */
DEFTREECODE (OMP_TARGET_UPDATE, "omp_target_update", tcc_statement, 1)
@@ -1155,7 +1155,7 @@ extern void protected_set_expr_location (tree, location_t);
#define TRANSACTION_EXPR_RELAXED(NODE) \
(TRANSACTION_EXPR_CHECK (NODE)->base.public_flag)
-/* OpenMP directive and clause accessors. */
+/* OpenMP and OpenACC directive and clause accessors. */
#define OMP_BODY(NODE) \
TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_CRITICAL), 0)
@@ -1167,6 +1167,39 @@ extern void protected_set_expr_location (tree, location_t);
#define OACC_PARALLEL_CLAUSES(NODE) \
TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1)
+#define OACC_KERNELS_BODY(NODE) \
+ TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 0)
+#define OACC_KERNELS_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_KERNELS_CHECK(NODE), 1)
+
+#define OACC_DATA_BODY(NODE) \
+ TREE_OPERAND (OACC_DATA_CHECK (NODE), 0)
+#define OACC_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DATA_CHECK (NODE), 1)
+
+#define OACC_HOST_DATA_BODY(NODE) \
+ TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 0)
+#define OACC_HOST_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_HOST_DATA_CHECK (NODE), 1)
+
+#define OACC_DECLARE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DECLARE_CHECK (NODE), 0)
+
+#define OACC_ENTER_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_ENTER_DATA_CHECK (NODE), 0)
+
+#define OACC_EXIT_DATA_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_EXIT_DATA_CHECK (NODE), 0)
+
+#define OACC_UPDATE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_UPDATE_CHECK (NODE), 0)
+
+#define OACC_WAIT_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_WAIT_CHECK (NODE), 0)
+
+#define OACC_CACHE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_CACHE_CHECK (NODE), 0)
+
#define OMP_PARALLEL_BODY(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 0)
#define OMP_PARALLEL_CLAUSES(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 1)
@@ -1278,6 +1311,32 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_SCHEDULE_CHUNK_EXPR(NODE) \
OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_SCHEDULE), 0)
+/* OpenACC clause expressions */
+#define OMP_CLAUSE_GANG_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_GANG), 0)
+#define OMP_WAIT_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_WAIT), 0)
+#define OMP_CLAUSE_VECTOR_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0)
+#define OMP_CLAUSE_ASYNC_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_ASYNC), 0)
+#define OMP_CLAUSE_WORKER_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_WORKER), 0)
+#define OMP_CLAUSE_NUM_GANGS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_GANGS), 0)
+#define OMP_CLAUSE_NUM_WORKERS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_NUM_WORKERS), 0)
+#define OMP_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_VECTOR_LENGTH), 0)
+
#define OMP_CLAUSE_DEPEND_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEPEND)->omp_clause.subcode.depend_kind)
--
1.8.3.2