From: Thomas Schwinge <thomas@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>, Julian Brown <julian@codesourcery.com>
Cc: "Joseph S. Myers" <joseph@codesourcery.com>,
Nathan Sidwell <Nathan_Sidwell@mentor.com>,
Jakub Jelinek <jakub@redhat.com>,
James Norris <jnorris@codesourcery.com>
Subject: Re: [Bulk] [OpenACC 0/7] host_data construct
Date: Wed, 23 Dec 2015 11:02:00 -0000 [thread overview]
Message-ID: <87twn9h1hg.fsf@kepler.schwinge.homeip.net> (raw)
In-Reply-To: <20151026183422.GW478@tucnak.redhat.com>
[-- Attachment #1: Type: text/plain, Size: 13600 bytes --]
Hi!
On Mon, 26 Oct 2015 19:34:22 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> Your use_device sounds very similar to use_device_ptr clause in OpenMP,
> which is allowed on #pragma omp target data construct and is implemented
> quite a bit differently from this; it is unclear if the OpenACC standard
> requires this kind of implementation, or you just chose to implement it this
> way. In particular, the GOMP_target_data call puts the variables mentioned
> in the use_device_ptr clauses into the mapping structures (similarly how
> map clause appears) and the corresponding vars are privatized within the
> target data region (which is a host region, basically a fancy { } braces),
> where the private variables contain the offloading device's pointers.
ACK. As the OpenACC use_device clause implementation now completely
matches the OpenMP use_device_ptr clause implementation, there is no use
anymore for a separate OMP_CLAUSE_USE_DEVICE, so in r231926 I cleaned
that up, as obvious:
commit 9d5fd7c608fef6e7a9efbfc940545d49452c4e01
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Dec 23 11:01:18 2015 +0000
Merge OMP_CLAUSE_USE_DEVICE into OMP_CLAUSE_USE_DEVICE_PTR
gcc/c/
* c-parser.c (c_parser_oacc_clause_use_device): Merge function
into...
(c_parser_omp_clause_use_device_ptr): ... this function. Adjust
all users.
gcc/
* tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE
into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@231926 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog | 5 +++++
gcc/c/ChangeLog | 7 +++++++
gcc/c/c-parser.c | 16 +++++-----------
gcc/c/c-typeck.c | 1 -
gcc/cp/parser.c | 2 +-
gcc/cp/pt.c | 2 --
gcc/cp/semantics.c | 1 -
gcc/fortran/trans-openmp.c | 2 +-
gcc/gimplify.c | 2 --
gcc/omp-low.c | 11 ++---------
gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 | 2 +-
gcc/tree-core.h | 6 ++----
gcc/tree-nested.c | 2 --
gcc/tree-pretty-print.c | 3 ---
gcc/tree.c | 3 ---
15 files changed, 24 insertions(+), 41 deletions(-)
diff --git gcc/ChangeLog gcc/ChangeLog
index aa28d10..d67b9c6 100644
--- gcc/ChangeLog
+++ gcc/ChangeLog
@@ -1,3 +1,8 @@
+2015-12-23 Thomas Schwinge <thomas@codesourcery.com>
+
+ * tree-core.h (enum omp_clause_code): Merge OMP_CLAUSE_USE_DEVICE
+ into OMP_CLAUSE_USE_DEVICE_PTR. Adjust all users.
+
2015-12-23 David Sherwood <david.sherwood@arm.com>
* config/arm/iterators.md (VMAXMINFNM): New int iterator.
diff --git gcc/c/ChangeLog gcc/c/ChangeLog
index f99f426..7b275d8 100644
--- gcc/c/ChangeLog
+++ gcc/c/ChangeLog
@@ -1,3 +1,10 @@
+2015-12-23 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.c (c_parser_oacc_clause_use_device): Merge function
+ into...
+ (c_parser_omp_clause_use_device_ptr): ... this function. Adjust
+ all users.
+
2015-12-22 Marek Polacek <polacek@redhat.com>
PR c/69002
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 353e3da..8e754d0 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -11395,7 +11395,10 @@ c_parser_omp_clause_defaultmap (c_parser *parser, tree list)
return list;
}
-/* OpenMP 4.5:
+/* OpenACC 2.0:
+ use_device ( variable-list )
+
+ OpenMP 4.5:
use_device_ptr ( variable-list ) */
static tree
@@ -11730,15 +11733,6 @@ c_parser_oacc_clause_tile (c_parser *parser, tree list)
return c;
}
-/* OpenACC 2.0:
- use_device ( variable-list ) */
-
-static tree
-c_parser_oacc_clause_use_device (c_parser *parser, tree list)
-{
- return c_parser_omp_var_list_parens (parser, OMP_CLAUSE_USE_DEVICE, list);
-}
-
/* OpenACC:
wait ( int-expr-list ) */
@@ -13058,7 +13052,7 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
c_name = "tile";
break;
case PRAGMA_OACC_CLAUSE_USE_DEVICE:
- clauses = c_parser_oacc_clause_use_device (parser, clauses);
+ clauses = c_parser_omp_clause_use_device_ptr (parser, clauses);
c_name = "use_device";
break;
case PRAGMA_OACC_CLAUSE_VECTOR:
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index 928fcd5..7406bd4 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -13125,7 +13125,6 @@ c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
bitmap_set_bit (&map_head, DECL_UID (t));
goto check_dup_generic;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
t = OMP_CLAUSE_DECL (c);
diff --git gcc/cp/parser.c gcc/cp/parser.c
index 842dded..4829a77 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -32097,7 +32097,7 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
c_name = "tile";
break;
case PRAGMA_OACC_CLAUSE_USE_DEVICE:
- clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE,
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE_PTR,
clauses);
c_name = "use_device";
break;
diff --git gcc/cp/pt.c gcc/cp/pt.c
index dab15bd..4555b32 100644
--- gcc/cp/pt.c
+++ gcc/cp/pt.c
@@ -14425,7 +14425,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
case OMP_CLAUSE_MAP:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
OMP_CLAUSE_DECL (nc)
@@ -14552,7 +14551,6 @@ tsubst_omp_clauses (tree clauses, bool declare_simd, bool allow_fields,
case OMP_CLAUSE_COPYPRIVATE:
case OMP_CLAUSE_LINEAR:
case OMP_CLAUSE_REDUCTION:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
/* tsubst_expr on SCOPE_REF results in returning
diff --git gcc/cp/semantics.c gcc/cp/semantics.c
index ab9989a..37bf050 100644
--- gcc/cp/semantics.c
+++ gcc/cp/semantics.c
@@ -6886,7 +6886,6 @@ finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
}
break;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
field_ok = allow_fields;
diff --git gcc/fortran/trans-openmp.c gcc/fortran/trans-openmp.c
index 227964c..70a7722 100644
--- gcc/fortran/trans-openmp.c
+++ gcc/fortran/trans-openmp.c
@@ -1771,7 +1771,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
clause_code = OMP_CLAUSE_UNIFORM;
goto add_clause;
case OMP_LIST_USE_DEVICE:
- clause_code = OMP_CLAUSE_USE_DEVICE;
+ clause_code = OMP_CLAUSE_USE_DEVICE_PTR;
goto add_clause;
case OMP_LIST_DEVICE_RESIDENT:
clause_code = OMP_CLAUSE_DEVICE_RESIDENT;
diff --git gcc/gimplify.c gcc/gimplify.c
index 62b0e64..bc90401 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7139,7 +7139,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
goto do_notice;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
flags = GOVD_FIRSTPRIVATE | GOVD_EXPLICIT;
goto do_add;
@@ -8051,7 +8050,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_DEVICE_RESIDENT:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_NUM_GANGS:
case OMP_CLAUSE_NUM_WORKERS:
diff --git gcc/omp-low.c gcc/omp-low.c
index 676b1df..a0c3e1c 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1957,7 +1957,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
install_var_local (decl, ctx);
break;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
decl = OMP_CLAUSE_DECL (c);
if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
@@ -2314,7 +2313,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx,
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_NOGROUP:
case OMP_CLAUSE_DEFAULTMAP:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE__CILK_FOR_COUNT_:
case OMP_CLAUSE_ASYNC:
@@ -15288,7 +15286,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
break;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
@@ -15674,14 +15671,12 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
build_int_cstu (tkind_type, tkind));
break;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
ovar = OMP_CLAUSE_DECL (c);
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
- || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
tkind = GOMP_MAP_USE_DEVICE_PTR;
else
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
@@ -15884,12 +15879,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_build_assign (new_var, x));
}
break;
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
var = OMP_CLAUSE_DECL (c);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
- || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR)
x = build_sender_ref (var, ctx);
else
x = build_receiver_ref (var, false, ctx);
diff --git gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95 gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
index 7a5eea6..23aba8c 100644
--- gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
+++ gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
@@ -8,4 +8,4 @@ program test
!$acc host_data use_device(i)
!$acc end host_data
end program test
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device\\(i\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(i\\)" 1 "original" } }
diff --git gcc/tree-core.h gcc/tree-core.h
index 9cc64d9..5371378 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -302,7 +302,8 @@ enum omp_clause_code {
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
- /* OpenMP clause: use_device_ptr (variable-list). */
+ /* OpenACC clause: use_device (variable_list).
+ OpenMP clause: use_device_ptr (variable-list). */
OMP_CLAUSE_USE_DEVICE_PTR,
/* OpenMP clause: is_device_ptr (variable-list). */
@@ -315,9 +316,6 @@ enum omp_clause_code {
/* 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
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 3a9479a..8211a12 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1072,7 +1072,6 @@ convert_nonlocal_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
do_decl_clause:
@@ -1744,7 +1743,6 @@ convert_local_omp_clauses (tree *pclauses, struct walk_stmt_info *wi)
case OMP_CLAUSE_SHARED:
case OMP_CLAUSE_TO_DECLARE:
case OMP_CLAUSE_LINK:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_USE_DEVICE_PTR:
case OMP_CLAUSE_IS_DEVICE_PTR:
do_decl_clause:
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index caec760..c4a180e 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -327,9 +327,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
case OMP_CLAUSE_DEVICE_RESIDENT:
name = "device_resident";
goto print_remap;
- case OMP_CLAUSE_USE_DEVICE:
- name = "use_device";
- goto print_remap;
case OMP_CLAUSE_TO_DECLARE:
name = "to";
goto print_remap;
diff --git gcc/tree.c gcc/tree.c
index 2190cae..d837609 100644
--- gcc/tree.c
+++ gcc/tree.c
@@ -282,7 +282,6 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_IS_DEVICE_PTR */
2, /* OMP_CLAUSE__CACHE_ */
1, /* OMP_CLAUSE_DEVICE_RESIDENT */
- 1, /* OMP_CLAUSE_USE_DEVICE */
2, /* OMP_CLAUSE_GANG */
1, /* OMP_CLAUSE_ASYNC */
1, /* OMP_CLAUSE_WAIT */
@@ -354,7 +353,6 @@ const char * const omp_clause_code_name[] =
"is_device_ptr",
"_cache_",
"device_resident",
- "use_device",
"gang",
"async",
"wait",
@@ -11612,7 +11610,6 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
/* FALLTHRU */
case OMP_CLAUSE_DEVICE_RESIDENT:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_ASYNC:
case OMP_CLAUSE_WAIT:
case OMP_CLAUSE_WORKER:
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
prev parent reply other threads:[~2015-12-23 11:02 UTC|newest]
Thread overview: 33+ messages / expand[flat|nested] mbox.gz Atom feed top
2015-10-22 19:14 James Norris
2015-10-22 19:15 ` [OpenACC 2/7] host_data construct (C FE) James Norris
2015-10-22 19:15 ` [OpenACC 1/7] host_data construct (C/C++ common) James Norris
2015-10-22 19:16 ` [OpenACC 3/7] host_data construct (C front-end) James Norris
2015-10-22 19:18 ` [OpenACC 4/7] host_data construct (middle end) James Norris
2015-10-22 19:19 ` [OpenACC 5/7] host_data construct (gcc tests) James Norris
2015-10-22 19:20 ` [OpenACC 6/7] host_data construct James Norris
2015-10-22 19:22 ` [OpenACC 7/7] host_data construct (runtime tests) James Norris
2015-10-22 20:42 ` [OpenACC 0/7] host_data construct Joseph Myers
2015-10-22 20:53 ` James Norris
2015-10-23 16:01 ` [Bulk] " James Norris
2015-10-26 18:36 ` Jakub Jelinek
2015-10-27 15:57 ` Cesar Philippidis
2015-11-02 18:33 ` Julian Brown
2015-11-02 19:29 ` Jakub Jelinek
2015-11-12 11:16 ` Julian Brown
2015-11-18 12:48 ` Julian Brown
2015-11-19 13:13 ` Jakub Jelinek
2015-11-19 14:29 ` Julian Brown
2015-11-19 15:57 ` Jakub Jelinek
2015-11-30 19:34 ` Julian Brown
2015-12-01 8:30 ` Jakub Jelinek
2015-12-02 15:27 ` Tom de Vries
2015-12-02 15:59 ` Thomas Schwinge
2015-12-02 19:16 ` Cesar Philippidis
2015-12-02 19:28 ` Steve Kargl
2015-12-02 19:35 ` Jakub Jelinek
2015-12-02 19:54 ` Cesar Philippidis
2015-12-02 22:14 ` [gomp4] " Thomas Schwinge
2016-04-08 13:41 ` Fortran OpenACC host_data construct ICE (was: [gomp4] Re: [OpenACC 0/7] host_data construct) Thomas Schwinge
2016-02-02 13:57 ` [OpenACC 0/7] host_data construct Thomas Schwinge
2015-11-13 15:31 ` [Bulk] " Jakub Jelinek
2015-12-23 11:02 ` Thomas Schwinge [this message]
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=87twn9h1hg.fsf@kepler.schwinge.homeip.net \
--to=thomas@codesourcery.com \
--cc=Nathan_Sidwell@mentor.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@redhat.com \
--cc=jnorris@codesourcery.com \
--cc=joseph@codesourcery.com \
--cc=julian@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).