This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[PATCH 4/6] [GOMP4] OpenACC 1.0+ support in fortran front-end
- From: Ilmir Usmanov <i dot usmanov at samsung dot com>
- To: gcc-patches at gcc dot gnu dot org, Thomas Schwinge <thomas at codesourcery dot com>, jakub at redhat dot com
- Cc: Evgeny Gavrin <e dot gavrin at samsung dot com>, GarbuzovViacheslav <v dot garbuzov at samsung dot com>, Dmitri Botcharnikov <dmitry dot b at samsung dot com>
- Date: Thu, 23 Jan 2014 22:04:45 +0400
- Subject: [PATCH 4/6] [GOMP4] OpenACC 1.0+ support in fortran front-end
- Authentication-results: sourceware.org; auth=none
- References: <52E158EF dot 9050009 at samsung dot com> <52E1595D dot 9000007 at samsung dot com> <52E1597E dot 8070809 at samsung dot com> <52E15999 dot 6010505 at samsung dot com>
>From 37806068fffcab95a21b51829c900d49be14961d Mon Sep 17 00:00:00 2001
From: Ilmir Usmanov <i.usmanov@samsung.com>
Date: Thu, 23 Jan 2014 21:08:05 +0400
Subject: [PATCH 4/6] OpenACC GENERIC nodes
---
gcc/gimplify.c | 73 ++++++++++++
gcc/omp-low.c | 84 ++++++++++++--
gcc/tree-core.h | 106 +++++++++++++++++-
gcc/tree-pretty-print.c | 286 +++++++++++++++++++++++++++++++++++++++++++++++-
gcc/tree-pretty-print.h | 7 ++
gcc/tree.c | 96 ++++++++++++++--
gcc/tree.def | 50 +++++++++
gcc/tree.h | 95 +++++++++++++++-
8 files changed, 770 insertions(+), 27 deletions(-)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index e45bed2..8af4368 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -4274,6 +4274,15 @@ is_gimple_stmt (tree t)
case ASM_EXPR:
case STATEMENT_LIST:
case OACC_PARALLEL:
+ case OACC_KERNELS:
+ case OACC_DATA:
+ case OACC_CACHE:
+ case OACC_WAIT:
+ case OACC_HOST_DATA:
+ case OACC_DECLARE:
+ case OACC_UPDATE:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
case OMP_PARALLEL:
case OMP_FOR:
case OMP_SIMD:
@@ -6025,6 +6034,32 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
+ case OACC_CLAUSE_COPY:
+ case OACC_CLAUSE_COPYOUT:
+ case OACC_CLAUSE_CREATE:
+ case OACC_CLAUSE_PRESENT:
+ case OACC_CLAUSE_PRESENT_OR_COPY:
+ case OACC_CLAUSE_PRESENT_OR_COPYIN:
+ case OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ case OACC_CLAUSE_PRESENT_OR_CREATE:
+ case OACC_CLAUSE_HOST:
+ case OACC_CLAUSE_DEVICE:
+ case OACC_CLAUSE_DEVICEPTR:
+ case OACC_CLAUSE_DEVICE_RESIDENT:
+ case OACC_CLAUSE_USE_DEVICE:
+ case OACC_CLAUSE_DELETE:
+ case OACC_CLAUSE_ASYNC:
+ case OACC_CLAUSE_GANG:
+ case OACC_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_CACHE:
+ case OACC_CLAUSE_SEQ:
+ case OACC_CLAUSE_INDEPENDENT:
+ case OACC_CLAUSE_WORKER:
+ case OACC_CLAUSE_VECTOR:
+ case OACC_CLAUSE_NUM_GANGS:
+ case OACC_CLAUSE_NUM_WORKERS:
+ case OACC_CLAUSE_VECTOR_LENGTH:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_UNTIED:
@@ -6339,6 +6374,32 @@ gimplify_adjust_omp_clauses (tree *list_p)
}
break;
+ case OACC_CLAUSE_COPY:
+ case OACC_CLAUSE_COPYOUT:
+ case OACC_CLAUSE_CREATE:
+ case OACC_CLAUSE_PRESENT:
+ case OACC_CLAUSE_PRESENT_OR_COPY:
+ case OACC_CLAUSE_PRESENT_OR_COPYIN:
+ case OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ case OACC_CLAUSE_PRESENT_OR_CREATE:
+ case OACC_CLAUSE_HOST:
+ case OACC_CLAUSE_DEVICE:
+ case OACC_CLAUSE_DEVICEPTR:
+ case OACC_CLAUSE_DEVICE_RESIDENT:
+ case OACC_CLAUSE_USE_DEVICE:
+ case OACC_CLAUSE_DELETE:
+ case OACC_CLAUSE_ASYNC:
+ case OACC_CLAUSE_GANG:
+ case OACC_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_CACHE:
+ case OACC_CLAUSE_SEQ:
+ case OACC_CLAUSE_INDEPENDENT:
+ case OACC_CLAUSE_WORKER:
+ case OACC_CLAUSE_VECTOR:
+ case OACC_CLAUSE_NUM_GANGS:
+ case OACC_CLAUSE_NUM_WORKERS:
+ case OACC_CLAUSE_VECTOR_LENGTH:
case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_COPYIN:
case OMP_CLAUSE_COPYPRIVATE:
@@ -7846,6 +7907,18 @@ 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_CACHE:
+ case OACC_WAIT:
+ case OACC_HOST_DATA:
+ case OACC_DECLARE:
+ case OACC_UPDATE:
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ ret = GS_ALL_DONE;
+ break;
+
case OMP_PARALLEL:
gimplify_omp_parallel (expr_p, pre_p);
ret = GS_ALL_DONE;
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index eb755c3..6da5977 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1543,10 +1543,12 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
/* FALLTHRU */
- case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ /* FALLTHRU. */
+
+ case OMP_CLAUSE_FIRSTPRIVATE: /* == OACC_CLAUSE_FIRSTPRIVATE. */
+ case OMP_CLAUSE_REDUCTION: /* == OACC_CLAUSE_REDUCTION. */
decl = OMP_CLAUSE_DECL (c);
do_private:
if (is_variable_sized (decl))
@@ -1583,8 +1585,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_COPYPRIVATE:
- case OMP_CLAUSE_COPYIN:
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ /* FALLTHRU. */
+
+ case OMP_CLAUSE_COPYIN: /* == OACC_CLAUSE_COPYIN. */
decl = OMP_CLAUSE_DECL (c);
by_ref = use_pointer_for_field (decl, NULL);
install_var_field (decl, by_ref, 3, ctx);
@@ -1596,7 +1600,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_FINAL:
- case OMP_CLAUSE_IF:
case OMP_CLAUSE_NUM_THREADS:
case OMP_CLAUSE_NUM_TEAMS:
case OMP_CLAUSE_THREAD_LIMIT:
@@ -1605,6 +1608,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_DEPEND:
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ /* FALLTHRU. */
+
+ case OMP_CLAUSE_IF: /* == OACC_CLAUSE_IF. */
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_OPERAND (c, 0), ctx->outer);
break;
@@ -1713,6 +1719,35 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_local (decl, ctx);
break;
+ case OACC_CLAUSE_COPY:
+ case OACC_CLAUSE_COPYOUT:
+ case OACC_CLAUSE_CREATE:
+ case OACC_CLAUSE_PRESENT:
+ case OACC_CLAUSE_PRESENT_OR_COPY:
+ case OACC_CLAUSE_PRESENT_OR_COPYIN:
+ case OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ case OACC_CLAUSE_PRESENT_OR_CREATE:
+ case OACC_CLAUSE_HOST:
+ case OACC_CLAUSE_DEVICE:
+ case OACC_CLAUSE_DEVICEPTR:
+ case OACC_CLAUSE_DEVICE_RESIDENT:
+ case OACC_CLAUSE_USE_DEVICE:
+ case OACC_CLAUSE_DELETE:
+ case OACC_CLAUSE_ASYNC:
+ case OACC_CLAUSE_GANG:
+ case OACC_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_CACHE:
+ case OACC_CLAUSE_SEQ:
+ case OACC_CLAUSE_INDEPENDENT:
+ case OACC_CLAUSE_WORKER:
+ case OACC_CLAUSE_VECTOR:
+ case OACC_CLAUSE_NUM_GANGS:
+ case OACC_CLAUSE_NUM_WORKERS:
+ case OACC_CLAUSE_VECTOR_LENGTH:
+ /* Not implemented yet. */
+ break;
+
default:
gcc_unreachable ();
}
@@ -1733,10 +1768,12 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* FALLTHRU */
case OMP_CLAUSE_PRIVATE:
- case OMP_CLAUSE_FIRSTPRIVATE:
- case OMP_CLAUSE_REDUCTION:
case OMP_CLAUSE_LINEAR:
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
+ /* FALLTHRU. */
+
+ case OMP_CLAUSE_FIRSTPRIVATE:
+ case OMP_CLAUSE_REDUCTION:
decl = OMP_CLAUSE_DECL (c);
if (is_variable_sized (decl))
install_var_local (decl, ctx);
@@ -1797,9 +1834,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
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:
@@ -1808,7 +1843,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_DIST_SCHEDULE:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
- case OMP_CLAUSE_COLLAPSE:
case OMP_CLAUSE_UNTIED:
case OMP_CLAUSE_FINAL:
case OMP_CLAUSE_MERGEABLE:
@@ -1822,6 +1856,38 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
break;
+ case OACC_CLAUSE_COPY:
+ case OACC_CLAUSE_COPYOUT:
+ case OACC_CLAUSE_CREATE:
+ case OACC_CLAUSE_PRESENT:
+ case OACC_CLAUSE_PRESENT_OR_COPY:
+ case OACC_CLAUSE_PRESENT_OR_COPYIN:
+ case OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ case OACC_CLAUSE_PRESENT_OR_CREATE:
+ case OACC_CLAUSE_HOST:
+ case OACC_CLAUSE_DEVICE:
+ case OACC_CLAUSE_DEVICEPTR:
+ case OACC_CLAUSE_DEVICE_RESIDENT:
+ case OACC_CLAUSE_USE_DEVICE:
+ case OACC_CLAUSE_DELETE:
+ case OACC_CLAUSE_ASYNC:
+ case OACC_CLAUSE_GANG:
+ case OACC_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_CACHE:
+ case OACC_CLAUSE_SEQ:
+ case OACC_CLAUSE_INDEPENDENT:
+ case OACC_CLAUSE_WORKER:
+ case OACC_CLAUSE_VECTOR:
+ case OACC_CLAUSE_NUM_GANGS:
+ case OACC_CLAUSE_NUM_WORKERS:
+ case OACC_CLAUSE_VECTOR_LENGTH:
+ case OMP_CLAUSE_COPYIN:
+ case OMP_CLAUSE_IF:
+ case OMP_CLAUSE_COLLAPSE:
+ /* Not implemented yet. */
+ break;
+
default:
gcc_unreachable ();
}
diff --git a/gcc/tree-core.h b/gcc/tree-core.h
index e2750e0..68e044f 100644
--- a/gcc/tree-core.h
+++ b/gcc/tree-core.h
@@ -216,12 +216,18 @@ enum omp_clause_code {
/* OpenMP clause: private (variable_list). */
OMP_CLAUSE_PRIVATE,
+ /* OpenACC clause: private (variable_list). */
+ OACC_CLAUSE_PRIVATE = OMP_CLAUSE_PRIVATE,
+
/* OpenMP clause: shared (variable_list). */
OMP_CLAUSE_SHARED,
/* OpenMP clause: firstprivate (variable_list). */
OMP_CLAUSE_FIRSTPRIVATE,
+ /* OpenACC clause: firstprivate (variable_list). */
+ OACC_CLAUSE_FIRSTPRIVATE = OMP_CLAUSE_FIRSTPRIVATE,
+
/* OpenMP clause: lastprivate (variable_list). */
OMP_CLAUSE_LASTPRIVATE,
@@ -234,9 +240,15 @@ enum omp_clause_code {
placeholder used in OMP_CLAUSE_REDUCTION_{INIT,MERGE}. */
OMP_CLAUSE_REDUCTION,
+ /* OpenACC clause: reduction (operator:variable_list). */
+ OACC_CLAUSE_REDUCTION = OMP_CLAUSE_REDUCTION,
+
/* OpenMP clause: copyin (variable_list). */
OMP_CLAUSE_COPYIN,
+ /* OpenACC clause: copyin (variable_list). */
+ OACC_CLAUSE_COPYIN = OMP_CLAUSE_COPYIN,
+
/* OpenMP clause: copyprivate (variable_list). */
OMP_CLAUSE_COPYPRIVATE,
@@ -261,12 +273,79 @@ enum omp_clause_code {
/* OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
+ /* OpenACC clause: copy (variable_list). */
+ OACC_CLAUSE_COPY,
+
+ /* OpenACC clause: copyout (variable_list). */
+ OACC_CLAUSE_COPYOUT,
+
+ /* OpenACC clause: create (variable_list). */
+ OACC_CLAUSE_CREATE,
+
+ /* OpenACC clause: present (variable_list). */
+ OACC_CLAUSE_PRESENT,
+
+ /* OpenACC clause: present_or_copy (variable_list). */
+ OACC_CLAUSE_PRESENT_OR_COPY,
+
+ /* OpenACC clause: present_or_copyin (variable_list). */
+ OACC_CLAUSE_PRESENT_OR_COPYIN,
+
+ /* OpenACC clause: present_or_copyout (variable_list). */
+ OACC_CLAUSE_PRESENT_OR_COPYOUT,
+
+ /* OpenACC clause: present_or_create (variable_list). */
+ OACC_CLAUSE_PRESENT_OR_CREATE,
+
+ /* OpenACC clause: host (variable_list). */
+ OACC_CLAUSE_HOST,
+
+ /* OpenACC clause: device (variable_list). */
+ OACC_CLAUSE_DEVICE,
+
+ /* OpenACC clause: deviceptr (variable_list). */
+ OACC_CLAUSE_DEVICEPTR,
+
+ /* OpenACC clause: device_resident (variable_list). */
+ OACC_CLAUSE_DEVICE_RESIDENT,
+
+ /* OpenACC clause: use_device (variable_list). */
+ OACC_CLAUSE_USE_DEVICE,
+
+ /* OpenACC clause: delete (variable_list). */
+ OACC_CLAUSE_DELETE,
+
+ /* OpenACC clause: async [(integer-expression)]. */
+ OACC_CLAUSE_ASYNC,
+
+ /* 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. */
+ OACC_CLAUSE_GANG,
+
+ /* OpenACC clause: wait [(integer-expression-list)]. */
+ OACC_CLAUSE_WAIT,
+
+ /* Internal structure to hold OpenACC wait directive's variable_list.
+ #pragma acc wait [(integer-expression-list)]. */
+ OACC_NO_CLAUSE_WAIT,
+
+ /* Internal structure to hold OpenACC cache directive's variable-list.
+ #pragma acc cache (variable-_ist). */
+ OACC_NO_CLAUSE_CACHE,
+
/* Internal clause: temporary for combined loops expansion. */
OMP_CLAUSE__LOOPTEMP_,
/* OpenMP clause: if (scalar-expression). */
OMP_CLAUSE_IF,
+ /* OpenACC clause: if (scalar-expression). */
+ OACC_CLAUSE_IF = OMP_CLAUSE_IF,
+
/* OpenMP clause: num_threads (integer-expression). */
OMP_CLAUSE_NUM_THREADS,
@@ -285,6 +364,9 @@ enum omp_clause_code {
/* OpenMP clause: collapse (constant-integer-expression). */
OMP_CLAUSE_COLLAPSE,
+ /* OpenACC clause: collapse (constant-integer-expression). */
+ OACC_CLAUSE_COLLAPSE = OMP_CLAUSE_COLLAPSE,
+
/* OpenMP clause: untied. */
OMP_CLAUSE_UNTIED,
@@ -334,7 +416,28 @@ enum omp_clause_code {
OMP_CLAUSE_TASKGROUP,
/* Internally used only clause, holding SIMD uid. */
- OMP_CLAUSE__SIMDUID_
+ OMP_CLAUSE__SIMDUID_,
+
+ /* OpenACC clause: seq. */
+ OACC_CLAUSE_SEQ,
+
+ /* OpenACC clause: independent. */
+ OACC_CLAUSE_INDEPENDENT,
+
+ /* OpenACC clause: worker [( [num:] integer-expression)]. */
+ OACC_CLAUSE_WORKER,
+
+ /* OpenACC clause: worker [( [length:] integer-expression)]. */
+ OACC_CLAUSE_VECTOR,
+
+ /* OpenACC clause: num_gangs (integer-expression). */
+ OACC_CLAUSE_NUM_GANGS,
+
+ /* OpenACC clause: num_workers (integer-expression). */
+ OACC_CLAUSE_NUM_WORKERS,
+
+ /* OpenACC clause: vector_length (integer-expression). */
+ OACC_CLAUSE_VECTOR_LENGTH
};
#undef DEFTREESTRUCT
@@ -660,7 +763,6 @@ enum annot_expr_kind {
annot_expr_ivdep_kind
};
-
/*---------------------------------------------------------------------------
Type definitions
---------------------------------------------------------------------------*/
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 320c35b..9b994e4 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -54,6 +54,248 @@ static void do_niy (pretty_printer *, const_tree);
static pretty_printer buffer;
static int initialized = 0;
+unsigned char
+dump_oacc_body (int flags, tree node, int spc,
+ unsigned char is_expr, pretty_printer* buffer)
+{
+ if (!(flags & TDF_SLIM) && OACC_BODY (node))
+ {
+ newline_and_indent (buffer, spc + 2);
+ pp_character (buffer, '{');
+ newline_and_indent (buffer, spc + 4);
+ dump_generic_node (buffer, OACC_BODY (node), spc + 4, flags, false);
+ newline_and_indent (buffer, spc + 2);
+ pp_character (buffer, '}');
+ }
+ is_expr = false;
+ return is_expr;
+}
+
+void
+dump_oacc_clause_remap (const char* name,
+ tree clause,
+ int spc,
+ int flags,
+ pretty_printer* buffer)
+{
+ pp_string (buffer, name);
+ pp_character (buffer, '(');
+ dump_generic_node (buffer, OACC_CLAUSE_DECL (clause), spc, flags, false);
+
+ pp_character (buffer, ')');
+}
+
+void
+dump_oacc_clause (pretty_printer *buffer, tree clause, int spc, int flags)
+{
+ const char *name;
+ enum omp_clause_code clause_code;
+
+ if (clause == NULL)
+ return;
+
+ clause_code = OACC_CLAUSE_CODE (clause);
+
+ switch (clause_code)
+ {
+ case OACC_CLAUSE_ASYNC:
+ pp_string (buffer, "async");
+ if (OACC_CLAUSE_DECL (clause))
+ {
+ pp_character(buffer, '(');
+ dump_generic_node (buffer, OACC_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ }
+ break;
+
+ case OACC_CLAUSE_VECTOR:
+ pp_string (buffer, "vector(");
+ dump_generic_node (buffer, OACC_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OACC_CLAUSE_GANG:
+ pp_string (buffer, "gang(");
+ dump_generic_node (buffer, OACC_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OACC_CLAUSE_WORKER:
+ pp_string (buffer, "worker(");
+ dump_generic_node (buffer, OACC_CLAUSE_DECL (clause), spc, flags, false);
+ pp_character(buffer, ')');
+ break;
+
+ case OACC_CLAUSE_INDEPENDENT:
+ name = "independent";
+ pp_string (buffer, name);
+ break;
+
+ case OACC_CLAUSE_SEQ:
+ name = "seq";
+ pp_string (buffer, name);
+ break;
+
+ case OACC_CLAUSE_COPY:
+ name = "copy";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_COPYIN:
+ name = "copyin";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_COPYOUT:
+ name = "copyout";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_CREATE:
+ name = "create";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_PRESENT:
+ name = "present";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_PRESENT_OR_COPY:
+ name = "present_or_copy";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_PRESENT_OR_COPYIN:
+ name = "present_or_copyin";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ name = "present_or_copyout";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_PRESENT_OR_CREATE:
+ name = "present_or_create";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_FIRSTPRIVATE:
+ name = "firstprivate";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_PRIVATE:
+ name = "private";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_DEVICEPTR:
+ name = "deviceptr";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_USE_DEVICE:
+ name = "use_device";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_DEVICE_RESIDENT:
+ name = "device_resident";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_HOST:
+ name = "host";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_DEVICE:
+ name = "device";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_NO_CLAUSE_CACHE:
+ name = "";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+
+ case OACC_CLAUSE_REDUCTION:
+ pp_string (buffer, "reduction(");
+ pp_string (buffer, op_symbol_code (OACC_CLAUSE_REDUCTION_CODE (clause)));
+ pp_character (buffer, ':');
+ dump_generic_node (buffer, OACC_CLAUSE_DECL (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OACC_CLAUSE_IF:
+ pp_string (buffer, "if(");
+ dump_generic_node (buffer, OACC_CLAUSE_IF_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OACC_CLAUSE_NUM_GANGS:
+ pp_string (buffer, "num_gangs(");
+ dump_generic_node (buffer, OACC_CLAUSE_NUM_GANGS_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OACC_CLAUSE_NUM_WORKERS:
+ pp_string (buffer, "num_workers(");
+ dump_generic_node (buffer, OACC_CLAUSE_NUM_WORKERS_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OACC_CLAUSE_VECTOR_LENGTH:
+ pp_string (buffer, "vector_length(");
+ dump_generic_node (buffer, OACC_CLAUSE_VECTOR_LENGTH_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OACC_CLAUSE_COLLAPSE:
+ pp_string (buffer, "collapse(");
+ dump_generic_node (buffer, OACC_CLAUSE_COLLAPSE_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ case OACC_NO_CLAUSE_WAIT:
+ pp_string (buffer, "(");
+ dump_generic_node (buffer, OACC_WAIT_EXPR (clause),
+ spc, flags, false);
+ pp_character (buffer, ')');
+ break;
+
+ default:
+ name = "unrecognized_oacc_clause";
+ dump_oacc_clause_remap (name, clause, spc, flags, buffer);
+ break;
+ }
+}
+
+void
+dump_oacc_clauses (pretty_printer *buffer, tree clause, int spc, int flags)
+{
+ if (clause == NULL)
+ return;
+
+ pp_space (buffer);
+ while (1)
+ {
+ dump_oacc_clause (buffer, clause, spc, flags);
+ clause = OACC_CLAUSE_CHAIN (clause);
+ if (clause == NULL)
+ return;
+ pp_space (buffer);
+ }
+}
+
/* Try to print something for an unknown tree code. */
static void
@@ -2358,10 +2600,50 @@ dump_generic_node (pretty_printer *buffer, tree node, int spc, int flags,
pp_string (buffer, " > ");
break;
+
case OACC_PARALLEL:
pp_string (buffer, "#pragma acc parallel");
- dump_omp_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags);
- goto dump_omp_body;
+ dump_oacc_clauses (buffer, OACC_PARALLEL_CLAUSES (node), spc, flags);
+ is_expr = dump_oacc_body(flags, node, spc, is_expr, buffer);
+ break;
+
+ case OACC_KERNELS:
+ pp_string (buffer, "#pragma acc kernels");
+ dump_oacc_clauses (buffer, OACC_KERNELS_CLAUSES(node), spc, flags);
+ is_expr = dump_oacc_body(flags, node, spc, is_expr, buffer);
+ break;
+
+ case OACC_HOST_DATA:
+ pp_string (buffer, "#pragma acc host_data");
+ dump_oacc_clauses (buffer, OACC_HOST_DATA_CLAUSES(node), spc, flags);
+ is_expr = dump_oacc_body(flags, node, spc, is_expr, buffer);
+ break;
+
+ case OACC_DATA:
+ pp_string (buffer, "#pragma acc data");
+ dump_oacc_clauses (buffer, OACC_DATA_CLAUSES(node), spc, flags);
+ is_expr = dump_oacc_body(flags, node, spc, is_expr, buffer);
+ break;
+
+ case OACC_WAIT:
+ pp_string (buffer, "#pragma acc wait");
+ dump_oacc_clauses (buffer, OACC_WAIT_CLAUSES(node), spc, flags);
+ break;
+
+ case OACC_UPDATE:
+ pp_string (buffer, "#pragma acc update");
+ dump_oacc_clauses (buffer, OACC_UPDATE_CLAUSES(node), spc, flags);
+ break;
+
+ case OACC_DECLARE:
+ pp_string (buffer, "#pragma acc declare");
+ dump_oacc_clauses (buffer, OACC_DECLARE_CLAUSES(node), spc, flags);
+ break;
+
+ case OACC_CACHE:
+ pp_string (buffer, "#pragma acc cache");
+ dump_oacc_clauses (buffer, OACC_CACHE_CLAUSES(node), spc, flags);
+ break;
case OMP_PARALLEL:
pp_string (buffer, "#pragma omp parallel");
diff --git a/gcc/tree-pretty-print.h b/gcc/tree-pretty-print.h
index 8754b0a..1014c91 100644
--- a/gcc/tree-pretty-print.h
+++ b/gcc/tree-pretty-print.h
@@ -36,6 +36,13 @@ extern void debug_generic_expr (tree);
extern void debug_generic_stmt (tree);
extern void debug_tree_chain (tree);
extern void print_generic_decl (FILE *, tree, int);
+extern unsigned char dump_oacc_body
+ (int, tree, int, unsigned char, pretty_printer*);
+extern void dump_oacc_clause_remap
+ (const char*, tree, int, int, pretty_printer*);
+extern void dump_oacc_clause (pretty_printer *, tree, int, int);
+extern void dump_oacc_clauses (pretty_printer *, tree, int, int);
+
extern void print_generic_stmt (FILE *, tree, int);
extern void print_generic_stmt_indented (FILE *, tree, int, int);
extern void print_generic_expr (FILE *, tree, int);
diff --git a/gcc/tree.c b/gcc/tree.c
index 25aa3e2..def3dc2 100644
--- a/gcc/tree.c
+++ b/gcc/tree.c
@@ -155,7 +155,8 @@ static const char * const tree_node_kind_names[] = {
"random kinds",
"lang_decl kinds",
"lang_type kinds",
- "omp clauses",
+ "acc clauses",
+ "omp clauses"
};
/* Unique id for next decl created. */
@@ -242,12 +243,12 @@ unsigned char tree_contains_struct[MAX_TREE_CODES][64];
unsigned const char omp_clause_num_ops[] =
{
0, /* OMP_CLAUSE_ERROR */
- 1, /* OMP_CLAUSE_PRIVATE */
+ 1, /* OMP_CLAUSE_PRIVATE, OACC_CLAUSE_PRIVATE */
1, /* OMP_CLAUSE_SHARED */
- 1, /* OMP_CLAUSE_FIRSTPRIVATE */
+ 1, /* OMP_CLAUSE_FIRSTPRIVATE, OACC_CLAUSE_FIRSTPRIVATE */
2, /* OMP_CLAUSE_LASTPRIVATE */
- 4, /* OMP_CLAUSE_REDUCTION */
- 1, /* OMP_CLAUSE_COPYIN */
+ 4, /* OMP_CLAUSE_REDUCTION, OACC_CLAUSE_REDUCTION */
+ 1, /* OMP_CLAUSE_COPYIN, OACC_CLAUSE_COPYIN */
1, /* OMP_CLAUSE_COPYPRIVATE */
2, /* OMP_CLAUSE_LINEAR */
2, /* OMP_CLAUSE_ALIGNED */
@@ -256,14 +257,33 @@ unsigned const char omp_clause_num_ops[] =
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
2, /* OMP_CLAUSE_MAP */
+ 3, /* OACC_CLAUSE_COPY */
+ 3, /* OACC_CLAUSE_COPYOUT */
+ 3, /* OACC_CLAUSE_CREATE */
+ 3, /* OACC_CLAUSE_PRESENT */
+ 3, /* OACC_CLAUSE_PRESENT_OR_COPY */
+ 3, /* OACC_CLAUSE_PRESENT_OR_COPYIN */
+ 3, /* OACC_CLAUSE_PRESENT_OR_COPYOUT */
+ 3, /* OACC_CLAUSE_PRESENT_OR_CREATE */
+ 1, /* OACC_CLAUSE_HOST */
+ 1, /* OACC_CLAUSE_DEVICE */
+ 1, /* OACC_CLAUSE_DEVICEPTR */
+ 1, /* OACC_CLAUSE_DEVICE_RESIDENT */
+ 1, /* OACC_CLAUSE_USE_DEVICE */
+ 3, /* OACC_CLAUSE_DELETE */
+ 1, /* OACC_CLAUSE_ASYNC */
+ 1, /* OACC_CLAUSE_GANG */
+ 1, /* OACC_CLAUSE_WAIT */
+ 1, /* OACC_NO_CLAUSE_WAIT */
+ 1, /* OACC_NO_CLAUSE_CACHE */
1, /* OMP_CLAUSE__LOOPTEMP_ */
- 1, /* OMP_CLAUSE_IF */
+ 1, /* OMP_CLAUSE_IF, OACC_CLAUSE_IF */
1, /* OMP_CLAUSE_NUM_THREADS */
1, /* OMP_CLAUSE_SCHEDULE */
0, /* OMP_CLAUSE_NOWAIT */
0, /* OMP_CLAUSE_ORDERED */
0, /* OMP_CLAUSE_DEFAULT */
- 3, /* OMP_CLAUSE_COLLAPSE */
+ 3, /* OMP_CLAUSE_COLLAPSE, OACC_CLAUSE_COLLAPSE */
0, /* OMP_CLAUSE_UNTIED */
1, /* OMP_CLAUSE_FINAL */
0, /* OMP_CLAUSE_MERGEABLE */
@@ -281,6 +301,14 @@ unsigned const char omp_clause_num_ops[] =
0, /* OMP_CLAUSE_SECTIONS */
0, /* OMP_CLAUSE_TASKGROUP */
1, /* OMP_CLAUSE__SIMDUID_ */
+
+ 0, /* OACC_CLAUSE_SEQ */
+ 0, /* OACC_CLAUSE_INDEPENDENT */
+ 1, /* OACC_CLAUSE_WORKER */
+ 1, /* OACC_CLAUSE_VECTOR */
+ 1, /* OACC_CLAUSE_NUM_GANGS */
+ 1, /* OACC_CLAUSE_NUM_WORKERS */
+ 1, /* OACC_CLAUSE_VECTOR_LENGTH */
};
const char * const omp_clause_code_name[] =
@@ -300,6 +328,25 @@ const char * const omp_clause_code_name[] =
"from",
"to",
"map",
+ "copy",
+ "copyout",
+ "create",
+ "present",
+ "present_or_copy",
+ "present_or_copyin",
+ "present_or_copyout",
+ "present_or_create",
+ "host",
+ "device",
+ "deviceptr",
+ "device_resident",
+ "use_device",
+ "delete",
+ "async",
+ "gang",
+ "wait",
+ "_wait_",
+ "_cache_",
"_looptemp_",
"if",
"num_threads",
@@ -324,7 +371,14 @@ const char * const omp_clause_code_name[] =
"parallel",
"sections",
"taskgroup",
- "_simduid_"
+ "_simduid_",
+ "seq",
+ "indepentend",
+ "worker",
+ "vector",
+ "num_gangs",
+ "num_workers",
+ "vector_length"
};
@@ -10359,7 +10413,6 @@ build_empty_stmt (location_t loc)
return t;
}
-
/* Build an OpenMP clause with code CODE. LOC is the location of the
clause. */
@@ -11022,6 +11075,29 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
switch (OMP_CLAUSE_CODE (*tp))
{
+ case OACC_CLAUSE_COPY:
+ case OACC_CLAUSE_COPYOUT:
+ case OACC_CLAUSE_CREATE:
+ case OACC_CLAUSE_PRESENT:
+ case OACC_CLAUSE_PRESENT_OR_COPY:
+ case OACC_CLAUSE_PRESENT_OR_COPYIN:
+ case OACC_CLAUSE_PRESENT_OR_COPYOUT:
+ case OACC_CLAUSE_PRESENT_OR_CREATE:
+ case OACC_CLAUSE_HOST:
+ case OACC_CLAUSE_DEVICE:
+ case OACC_CLAUSE_DEVICEPTR:
+ case OACC_CLAUSE_DEVICE_RESIDENT:
+ case OACC_CLAUSE_USE_DEVICE:
+ case OACC_CLAUSE_DELETE:
+ case OACC_NO_CLAUSE_CACHE:
+ case OACC_CLAUSE_ASYNC:
+ case OACC_CLAUSE_WORKER:
+ case OACC_CLAUSE_VECTOR:
+ case OACC_CLAUSE_NUM_GANGS:
+ case OACC_CLAUSE_NUM_WORKERS:
+ case OACC_CLAUSE_VECTOR_LENGTH:
+ case OACC_CLAUSE_WAIT:
+ case OACC_NO_CLAUSE_WAIT:
case OMP_CLAUSE_PRIVATE:
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_FIRSTPRIVATE:
@@ -11044,6 +11120,8 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
WALK_SUBTREE (OMP_CLAUSE_OPERAND (*tp, 0));
/* FALLTHRU */
+ case OACC_CLAUSE_SEQ:
+ case OACC_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_NOWAIT:
case OMP_CLAUSE_ORDERED:
case OMP_CLAUSE_DEFAULT:
diff --git a/gcc/tree.def b/gcc/tree.def
index 0edddda..562da21 100644
--- a/gcc/tree.def
+++ b/gcc/tree.def
@@ -1012,6 +1012,56 @@ DEFTREECODE (MEM_REF, "mem_ref", tcc_reference, 2)
DEFTREECODE (OACC_PARALLEL, "oacc_parallel", tcc_statement, 2)
+/* #pragma acc kernels */
+/* Operand 0: BODY
+ Operand 1: CLAUSES
+*/
+DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
+
+/* #pragma acc data */
+/* Operand 0: BODY
+ Operand 1: CLAUSES
+*/
+DEFTREECODE (OACC_DATA, "oacc_data", tcc_statement, 2)
+
+/* #pragma acc host_data */
+/* Operand 0: BODY
+ Operand 1: CLAUSES
+*/
+DEFTREECODE (OACC_HOST_DATA, "oacc_host_data", tcc_statement, 2)
+
+/* #pragma acc declare */
+/* Operand 0: CLAUSES
+*/
+DEFTREECODE (OACC_DECLARE, "oacc_declare", tcc_statement, 1)
+
+/* #pragma acc update */
+/* Operand 0: CLAUSES
+*/
+DEFTREECODE (OACC_UPDATE, "oacc_update", tcc_statement, 1)
+
+/* #pragma acc enter data */
+/* Operand 0: CLAUSES
+*/
+
+DEFTREECODE (OACC_ENTER_DATA, "oacc_enter_data", tcc_statement, 1)
+
+/* #pragma acc exit data */
+/* Operand 0: CLAUSES
+*/
+
+DEFTREECODE (OACC_EXIT_DATA, "oacc_exit_data", tcc_statement, 1)
+
+/* #pragma acc wait */
+/* Operand 0: INT_EXPR
+*/
+DEFTREECODE (OACC_WAIT, "oacc_wait", tcc_statement, 1)
+
+/* #pragma acc cache */
+/* Operand 0: LIST
+*/
+DEFTREECODE (OACC_CACHE, "oacc_cache", tcc_statement, 1)
+
/* 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. */
diff --git a/gcc/tree.h b/gcc/tree.h
index ad9901c..717a41f 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1185,11 +1185,6 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSES(NODE) \
TREE_OPERAND (TREE_RANGE_CHECK (NODE, OACC_PARALLEL, OMP_SINGLE), 1)
-#define OACC_PARALLEL_BODY(NODE) \
- TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 0)
-#define OACC_PARALLEL_CLAUSES(NODE) \
- TREE_OPERAND (OACC_PARALLEL_CHECK (NODE), 1)
-
#define OMP_PARALLEL_BODY(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 0)
#define OMP_PARALLEL_CLAUSES(NODE) TREE_OPERAND (OMP_PARALLEL_CHECK (NODE), 1)
@@ -1390,6 +1385,91 @@ extern void protected_set_expr_location (tree, location_t);
#define OMP_CLAUSE_DEFAULT_KIND(NODE) \
(OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_DEFAULT)->omp_clause.subcode.default_kind)
+/* OpenACC directives and clause accessors. */
+
+#define OACC_BODY(NODE) \
+ TREE_OPERAND (NODE, 0)
+
+#define OACC_CLAUSE_CHAIN(NODE) OMP_CLAUSE_CHAIN(NODE)
+
+#define OACC_CLAUSE(NODE) OMP_CLAUSE(NODE)
+
+#define OACC_CLAUSE_DECL(NODE) OMP_CLAUSE_DECL(NODE)
+
+#define OACC_CLAUSE_CODE(NODE) OMP_CLAUSE_CODE(NODE)
+
+#define OACC_CLAUSE_OPERAND(T, i) OMP_CLAUSE_OPERAND(T, i)
+
+#define OACC_CLAUSE_SET_CODE(NODE, CODE) OMP_CLAUSE_SET_CODE(NODE, CODE)
+
+#define OACC_CLAUSE_LOCATION(NODE) OMP_CLAUSE_LOCATION(NODE)
+
+#define OACC_PARALLEL_BODY(NODE) TREE_OPERAND (OACC_PARALLEL_CHECK(NODE), 0)
+#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)
+
+/* OpenACC clauses */
+#define OACC_CLAUSE_NUM_GANGS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_NUM_GANGS), 0)
+#define OACC_CLAUSE_NUM_WORKERS_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_NUM_WORKERS), 0)
+#define OACC_CLAUSE_VECTOR_LENGTH_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_VECTOR_LENGTH), 0)
+#define OACC_CLAUSE_VECTOR_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_VECTOR_LENGTH), 0)
+#define OACC_CLAUSE_WORKER_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_WORKER), 0)
+#define OACC_CLAUSE_GANG_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_GANG), 0)
+#define OACC_CLAUSE_COLLAPSE_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_COLLAPSE), 0)
+#define OACC_CLAUSE_IF_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_IF), 0)
+#define OACC_CLAUSE_ASYNC_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_CLAUSE_ASYNC), 0)
+#define OACC_WAIT_EXPR(NODE) \
+ OMP_CLAUSE_OPERAND ( \
+ OMP_CLAUSE_SUBCODE_CHECK (NODE, OACC_NO_CLAUSE_WAIT), 0)
+
+#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_DECLARE_CLAUSES(NODE) \
+ TREE_OPERAND (OACC_DECLARE_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 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_CLAUSE_REDUCTION_CODE(NODE) OMP_CLAUSE_REDUCTION_CODE(NODE)
+#define OACC_CLAUSE_REDUCTION_INIT(NODE) OMP_CLAUSE_REDUCTION_INIT(NODE)
+#define OACC_CLAUSE_REDUCTION_MERGE(NODE) OMP_CLAUSE_REDUCTION_MERGE(NODE)
+#define OACC_CLAUSE_REDUCTION_PLACEHOLDER(NODE) \
+ OMP_CLAUSE_REDUCTION_PLACEHOLDER(NODE)
+
/* SSA_NAME accessors. */
/* Returns the IDENTIFIER_NODE giving the SSA name a name or NULL_TREE
@@ -3585,6 +3665,11 @@ extern tree build_translation_unit_decl (tree);
extern tree build_block (tree, tree, tree, tree);
extern tree build_empty_stmt (location_t);
extern tree build_omp_clause (location_t, enum omp_clause_code);
+static inline tree
+build_oacc_clause (location_t loc, enum omp_clause_code code)
+{
+ return build_omp_clause (loc, code);
+}
extern tree build_vl_exp_stat (enum tree_code, int MEM_STAT_DECL);
#define build_vl_exp(c, n) build_vl_exp_stat (c, n MEM_STAT_INFO)
--
1.8.3.2