From: Thomas Schwinge <thomas@codesourcery.com>
To: Julian Brown <julian@codesourcery.com>,
GCC Patches <gcc-patches@gcc.gnu.org>,
Cesar Philippidis <cesar@codesourcery.com>,
James Norris <jnorris@codesourcery.com>
Cc: "Joseph S. Myers" <joseph@codesourcery.com>,
Nathan Sidwell <Nathan_Sidwell@mentor.com>,
Jakub Jelinek <jakub@redhat.com>
Subject: [gomp4] Re: [OpenACC 0/7] host_data construct
Date: Wed, 02 Dec 2015 22:14:00 -0000 [thread overview]
Message-ID: <87r3j4lcrd.fsf@kepler.schwinge.homeip.net> (raw)
In-Reply-To: <877fkwn8p6.fsf@kepler.schwinge.homeip.net>
[-- Attachment #1: Type: text/plain, Size: 45078 bytes --]
Hi!
On Wed, 2 Dec 2015 16:58:45 +0100, I wrote:
> Cesar and Jim copied, for help with Fortran and generally testsuite
> things.
>
> On Mon, 30 Nov 2015 19:30:34 +0000, Julian Brown <julian@codesourcery.com> wrote:
> > [patch]
>
> First, thanks!
Aside from a number of formatting/re-ordering changes, the front end
changes were basically still the same, but otherwise (middle end,
libgomp) the patch as committed to trunk in r231118 was quite (totally?)
;-) different from the code we had on gomp-4_0-branch, so I had to spend
some time on merging, cleaning things up.
> What about the test cases present on gomp-4_0-branch,
> gcc/testsuite/c-c++-common/goacc/host_data-1.c,
> gcc/testsuite/c-c++-common/goacc/host_data-2.c,
> gcc/testsuite/c-c++-common/goacc/host_data-3.c, and
> gcc/testsuite/c-c++-common/goacc/host_data-4.c, [...]
In the merge, I had to move two use_device usages from
c-c++-common/goacc/host_data-1.c (was accepted) to
c-c++-common/goacc/host_data-2.c (now rejected); I hope that's correct.
> Your submission/commit didn't have any execution tests for OpenACC
> host_data in Fortran. On gomp-4_0-branch, there is
> libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 at least.
..., but this one now FAILs (ICE) as follows:
[...]/source-gcc/libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90:11:0: internal compiler error: in scan_omp_target, at omp-low.c:3218
0xa33e80 scan_omp_target
[...]/source-gcc/gcc/omp-low.c:3218
0xa33e80 scan_omp_1_stmt
[...]/source-gcc/gcc/omp-low.c:3980
0x8e4e7e walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:555
0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:51
0x8e4f62 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:583
0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:51
0x8e4ff2 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:619
0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:51
0xa02479 scan_omp
[...]/source-gcc/gcc/omp-low.c:4024
0xa32ea5 scan_omp_target
[...]/source-gcc/gcc/omp-low.c:3204
0xa32ea5 scan_omp_1_stmt
[...]/source-gcc/gcc/omp-low.c:3980
0x8e4e7e walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:555
0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:51
0x8e4ff2 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:619
0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:51
0x8e4f62 walk_gimple_stmt(gimple_stmt_iterator*, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:583
0x8e50b8 walk_gimple_seq_mod(gimple**, tree_node* (*)(gimple_stmt_iterator*, bool*, walk_stmt_info*), tree_node* (*)(tree_node**, int*, void*), walk_stmt_info*)
[...]/source-gcc/gcc/gimple-walk.c:51
0xa02479 scan_omp
[...]/source-gcc/gcc/omp-low.c:4024
0xa3f35a execute_lower_omp
[...]/source-gcc/gcc/omp-low.c:16735
0xa3f35a execute
[...]/source-gcc/gcc/omp-low.c:16782
Maybe that's due to the gcc/gimplify.c:gimplify_scan_omp_clauses issue
mentioned in
<http://news.gmane.org/find-root.php?message_id=%3C877fkwn8p6.fsf%40kepler.schwinge.homeip.net%3E>,
or maybe something else? (XFAILed for now.)
(For avoidance of doubt, the merge does not include my "Some OpenACC
host_data cleanup" commit, trunk r231184, which will get merged into
gomp-4_0-branch later.) So, merging trunk r231118 into gomp-4_0-branch,
I effectively applied the following patch, in r231207. Please verify.
For instance, do we need to re-instantiate any of the testsuite code that
we've lost here, or is all of that actually not supported?
commit 15723d76ae42dfe3f7201e0e3c6cbd9f4fc480b2
Merge: e08db3c 571b348
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Dec 2 21:52:25 2015 +0000
svn merge -r 231117:231118 svn+ssh://gcc.gnu.org/svn/gcc/trunk
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@231207 138bc75d-0d04-0410-961f-82ee72b054a4
gcc/ChangeLog | 29 ++++
gcc/c-family/ChangeLog | 8 +
gcc/c/ChangeLog | 14 ++
gcc/c/c-parser.c | 12 +-
gcc/c/c-typeck.c | 2 +-
gcc/cp/ChangeLog | 14 ++
gcc/cp/parser.c | 62 ++++----
gcc/cp/semantics.c | 6 +-
gcc/gimple-pretty-print.c | 3 +
gcc/gimple.h | 2 +
gcc/gimplify.c | 177 +++++----------------
gcc/omp-builtins.def | 4 +-
gcc/omp-low.c | 25 ++-
gcc/testsuite/c-c++-common/goacc/host_data-1.c | 4 +-
gcc/testsuite/c-c++-common/goacc/host_data-2.c | 10 ++
gcc/testsuite/gfortran.dg/goacc/coarray.f95 | 2 -
gcc/tree-nested.c | 2 +
libgomp/ChangeLog | 12 ++
libgomp/libgomp.map | 2 +-
libgomp/oacc-mem.c | 32 ----
libgomp/oacc-parallel.c | 40 +++++
.../libgomp.oacc-c-c++-common/host_data-1.c | 39 +----
.../libgomp.oacc-c-c++-common/host_data-2.c | 57 +++----
.../libgomp.oacc-c-c++-common/host_data-3.c | 29 ++++
.../libgomp.oacc-c-c++-common/host_data-4.c | 29 ++++
.../libgomp.oacc-c-c++-common/host_data-5.c | 38 +++++
.../libgomp.oacc-c-c++-common/host_data-6.c | 31 ++++
.../testsuite/libgomp.oacc-fortran/host_data-1.f90 | 5 +-
28 files changed, 394 insertions(+), 296 deletions(-)
[diff --git gcc/ChangeLog gcc/ChangeLog]
[diff --git gcc/c-family/ChangeLog gcc/c-family/ChangeLog]
[diff --git gcc/c/ChangeLog gcc/c/ChangeLog]
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 7191665..0251b80 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -10315,10 +10315,10 @@ c_parser_omp_clause_name (c_parser *parser, bool consume_token = true)
result = PRAGMA_OMP_CLAUSE_UNIFORM;
else if (!strcmp ("untied", p))
result = PRAGMA_OMP_CLAUSE_UNTIED;
- else if (!strcmp ("use_device", p))
- result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("use_device", p))
+ result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
break;
case 'v':
if (!strcmp ("vector", p))
@@ -13113,6 +13113,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_omp_clause_reduction (parser, clauses);
c_name = "reduction";
break;
+ case PRAGMA_OACC_CLAUSE_USE_DEVICE:
+ clauses = c_parser_oacc_clause_use_device (parser, clauses);
+ c_name = "use_device";
+ break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = c_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
clauses);
@@ -13122,10 +13126,6 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_clause_tile (parser, clauses);
c_name = "tile";
break;
- case PRAGMA_OACC_CLAUSE_USE_DEVICE:
- clauses = c_parser_oacc_clause_use_device (parser, clauses);
- c_name = "use_device";
- break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
clauses = c_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index c40f6da..4659814 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -13168,6 +13168,7 @@ c_finish_omp_clauses (tree clauses, bool is_oacc, 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);
@@ -13230,7 +13231,6 @@ c_finish_omp_clauses (tree clauses, bool is_oacc, bool is_omp, bool declare_simd
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
[diff --git gcc/cp/ChangeLog gcc/cp/ChangeLog]
diff --git gcc/cp/parser.c gcc/cp/parser.c
index ac3f45c..d32aa91 100644
--- gcc/cp/parser.c
+++ gcc/cp/parser.c
@@ -29242,10 +29242,10 @@ cp_parser_omp_clause_name (cp_parser *parser, bool consume_token = true)
result = PRAGMA_OMP_CLAUSE_UNIFORM;
else if (!strcmp ("untied", p))
result = PRAGMA_OMP_CLAUSE_UNTIED;
- else if (!strcmp ("use_device", p))
- result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
else if (!strcmp ("use_device_ptr", p))
result = PRAGMA_OMP_CLAUSE_USE_DEVICE_PTR;
+ else if (!strcmp ("use_device", p))
+ result = PRAGMA_OACC_CLAUSE_USE_DEVICE;
break;
case 'v':
if (!strcmp ("vector", p))
@@ -31752,6 +31752,11 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_omp_clause_reduction (parser, clauses);
c_name = "reduction";
break;
+ case PRAGMA_OACC_CLAUSE_USE_DEVICE:
+ clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE,
+ clauses);
+ c_name = "use_device";
+ break;
case PRAGMA_OACC_CLAUSE_SEQ:
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_SEQ,
clauses, here);
@@ -31761,11 +31766,6 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_clause_tile (parser, here, clauses);
c_name = "tile";
break;
- case PRAGMA_OACC_CLAUSE_USE_DEVICE:
- clauses = cp_parser_omp_var_list (parser, OMP_CLAUSE_USE_DEVICE,
- clauses);
- c_name = "use_device";
- break;
case PRAGMA_OACC_CLAUSE_VECTOR:
c_name = "vector";
clauses = cp_parser_oacc_shape_clause (parser, OMP_CLAUSE_VECTOR,
@@ -34671,6 +34671,30 @@ cp_parser_oacc_data (cp_parser *parser, cp_token *pragma_tok)
return stmt;
}
+#define OACC_HOST_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
+
+/* OpenACC 2.0:
+ # pragma acc host_data <clauses> new-line
+ structured-block */
+
+static tree
+cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok)
+{
+ tree stmt, clauses, block;
+ unsigned int save;
+
+ clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK,
+ "#pragma acc host_data", pragma_tok);
+
+ block = begin_omp_parallel ();
+ save = cp_parser_begin_omp_structured_block (parser);
+ cp_parser_statement (parser, NULL_TREE, false, NULL);
+ cp_parser_end_omp_structured_block (parser, save);
+ stmt = finish_oacc_host_data (clauses, block);
+ return stmt;
+}
+
/* OpenACC 2.0:
# pragma acc declare oacc-data-clause[optseq] new-line
*/
@@ -34823,30 +34847,6 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token *pragma_tok)
return NULL_TREE;
}
-#define OACC_HOST_DATA_CLAUSE_MASK \
- ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_USE_DEVICE) )
-
-/* OpenACC 2.0:
- # pragma acc host_data <clauses> new-line
- structured-block */
-
-static tree
-cp_parser_oacc_host_data (cp_parser *parser, cp_token *pragma_tok)
-{
- tree stmt, clauses, block;
- unsigned int save;
-
- clauses = cp_parser_oacc_all_clauses (parser, OACC_HOST_DATA_CLAUSE_MASK,
- "#pragma acc host_data", pragma_tok);
-
- block = begin_omp_parallel ();
- save = cp_parser_begin_omp_structured_block (parser);
- cp_parser_statement (parser, NULL_TREE, false, NULL);
- cp_parser_end_omp_structured_block (parser, save);
- stmt = finish_oacc_host_data (clauses, block);
- return stmt;
-}
-
/* OpenACC 2.0:
# pragma acc enter data oacc-enter-data-clause[optseq] new-line
diff --git gcc/cp/semantics.c gcc/cp/semantics.c
index 0d7e23d..a9a6671 100644
--- gcc/cp/semantics.c
+++ gcc/cp/semantics.c
@@ -6911,6 +6911,7 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
}
break;
+ case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_IS_DEVICE_PTR:
case OMP_CLAUSE_USE_DEVICE_PTR:
field_ok = allow_fields;
@@ -6948,7 +6949,6 @@ finish_omp_clauses (tree clauses, bool is_oacc, bool allow_fields,
case OMP_CLAUSE_SIMD:
case OMP_CLAUSE_DEFAULTMAP:
case OMP_CLAUSE__CILK_FOR_COUNT_:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_INDEPENDENT:
case OMP_CLAUSE_SEQ:
@@ -7483,9 +7483,9 @@ tree
finish_oacc_host_data (tree clauses, tree block)
{
tree stmt;
-
+
block = finish_omp_structured_block (block);
-
+
stmt = make_node (OACC_HOST_DATA);
TREE_TYPE (stmt) = void_type_node;
OACC_HOST_DATA_CLAUSES (stmt) = clauses;
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index 6c4e42c..c0f7c20 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1356,6 +1356,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gomp_target *gs,
case GF_OMP_TARGET_KIND_OACC_DECLARE:
kind = " oacc_declare";
break;
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+ kind = " oacc_host_data";
+ break;
default:
gcc_unreachable ();
}
diff --git gcc/gimple.h gcc/gimple.h
index 4c90bd7..7aaf785 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -171,6 +171,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_OACC_UPDATE = 8,
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9,
GF_OMP_TARGET_KIND_OACC_DECLARE = 10,
+ GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -6006,6 +6007,7 @@ is_gimple_omp_oacc (const gimple *stmt)
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
return true;
default:
return false;
diff --git gcc/gimplify.c gcc/gimplify.c
index 3bb3bfe..b00de81 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -90,10 +90,8 @@ enum gimplify_omp_var_data
/* Flag for shared vars that are or might be stored to in the region. */
GOVD_WRITTEN = 131072,
- GOVD_USE_DEVICE = 1 << 18,
-
/* OpenACC deviceptr clause. */
- GOVD_USE_DEVPTR = 1 << 19,
+ GOVD_USE_DEVPTR = 1 << 18,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
@@ -122,18 +120,16 @@ enum omp_region_type
ORT_TARGET = 0x20,
ORT_COMBINED_TARGET = 0x21,
- ORT_HOST_DATA = 0x40,
-
/* OpenACC variants. */
- ORT_ACC = 0x80, /* A generic OpenACC region. */
+ ORT_ACC = 0x40, /* A generic OpenACC region. */
ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
- ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x100, /* Kernels construct. */
- ORT_ACC_HOST = ORT_ACC | ORT_HOST_DATA,
+ ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 0x80, /* Kernels construct. */
+ ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 0x80, /* Host data. */
/* Dummy OpenMP region, used to disable expansion of
DECL_VALUE_EXPRs in taskloop pre body. */
- ORT_NONE = 0x200
+ ORT_NONE = 0x100
};
/* Gimplify hashtable helper. */
@@ -6126,8 +6122,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
for (; octx; octx = octx->outer_context)
{
- if (octx->region_type & ORT_HOST_DATA)
- continue;
if (!(octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)))
break;
splay_tree_node n2
@@ -6135,6 +6129,9 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
(splay_tree_key) decl);
if (n2)
{
+ if (octx->region_type == ORT_ACC_HOST_DATA)
+ error ("variable %qE declared in enclosing "
+ "%<host_data%> region", DECL_NAME (decl));
nflags |= GOVD_MAP;
goto found_outer;
}
@@ -6436,6 +6433,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_HOST_DATA:
ctx->target_firstprivatize_array_bases = true;
default:
break;
@@ -6571,10 +6569,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|| outer_ctx->region_type == ORT_ACC_DATA))
redvec.safe_push (OMP_CLAUSE_DECL (c));
goto do_add_decl;
- case OMP_CLAUSE_USE_DEVICE:
- flags = GOVD_USE_DEVICE | GOVD_EXPLICIT;
- check_non_private = "use_device";
- goto do_add;
case OMP_CLAUSE_LINEAR:
if (gimplify_expr (&OMP_CLAUSE_LINEAR_STEP (c), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
@@ -6709,6 +6703,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
+ case OACC_HOST_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
@@ -6721,6 +6716,22 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (remove)
break;
+ if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
+ {
+ struct gimplify_omp_ctx *octx;
+ for (octx = outer_ctx; octx; octx = octx->outer_context)
+ {
+ if (octx->region_type != ORT_ACC_HOST_DATA)
+ break;
+ splay_tree_node n2
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) decl);
+ if (n2)
+ error_at (OMP_CLAUSE_LOCATION (c), "variable %qE "
+ "declared in enclosing %<host_data%> region",
+ DECL_NAME (decl));
+ }
+ }
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_P (decl) ? DECL_SIZE_UNIT (decl)
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
@@ -7120,6 +7131,7 @@ 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;
@@ -7639,7 +7651,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
code = OMP_CLAUSE_FIRSTPRIVATE;
else if (flags & GOVD_LASTPRIVATE)
code = OMP_CLAUSE_LASTPRIVATE;
- else if (flags & (GOVD_ALIGNED | GOVD_USE_DEVICE))
+ else if (flags & GOVD_ALIGNED)
return 0;
else
gcc_unreachable ();
@@ -8244,126 +8256,6 @@ gimplify_oacc_declare (tree *expr_p, gimple_seq *pre_p)
*expr_p = NULL_TREE;
}
-static tree
-gimplify_oacc_host_data_1 (tree *tp, int *walk_subtrees, void *data ATTRIBUTE_UNUSED)
-{
- splay_tree_node n = NULL;
- location_t loc = EXPR_LOCATION (*tp);
-
- switch (TREE_CODE (*tp))
- {
- case ADDR_EXPR:
- {
- tree decl = TREE_OPERAND (*tp, 0);
-
- switch (TREE_CODE (decl))
- {
- case ARRAY_REF:
- case ARRAY_RANGE_REF:
- case COMPONENT_REF:
- case VIEW_CONVERT_EXPR:
- case REALPART_EXPR:
- case IMAGPART_EXPR:
- if (TREE_CODE (TREE_OPERAND (decl, 0)) == VAR_DECL)
- n = splay_tree_lookup (gimplify_omp_ctxp->variables,
- (splay_tree_key) TREE_OPERAND (decl, 0));
- break;
-
- case VAR_DECL:
- n = splay_tree_lookup (gimplify_omp_ctxp->variables,
- (splay_tree_key) decl);
- break;
-
- default:
- ;
- }
-
- if (n != NULL && (n->value & GOVD_USE_DEVICE) != 0)
- {
- tree t = builtin_decl_explicit (BUILT_IN_GOACC_DEVICEPTR);
- *tp = build_call_expr_loc (loc, t, 1, *tp);
- }
-
- *walk_subtrees = 0;
- }
- break;
-
- case VAR_DECL:
- {
- tree decl = *tp;
-
- n = splay_tree_lookup (gimplify_omp_ctxp->variables,
- (splay_tree_key) decl);
-
- if (n != NULL && (n->value & GOVD_USE_DEVICE) != 0)
- {
- if (!POINTER_TYPE_P (TREE_TYPE (decl)))
- return decl;
-
- tree t = builtin_decl_explicit (BUILT_IN_GOACC_DEVICEPTR);
- *tp = build_call_expr_loc (loc, t, 1, *tp);
- *walk_subtrees = 0;
- }
- }
- break;
-
- case OACC_PARALLEL:
- case OACC_KERNELS:
- case OACC_LOOP:
- *walk_subtrees = 0;
- break;
-
- default:
- ;
- }
-
- return NULL_TREE;
-}
-
-static enum gimplify_status
-gimplify_oacc_host_data (tree *expr_p, gimple_seq *pre_p)
-{
- tree expr = *expr_p, orig_body;
- gimple_seq body = NULL;
-
- gimplify_scan_omp_clauses (&OACC_HOST_DATA_CLAUSES (expr), pre_p,
- ORT_ACC_HOST, OACC_HOST_DATA);
-
- orig_body = OACC_HOST_DATA_BODY (expr);
-
- /* Perform a pre-pass over the host_data region's body, inserting calls to
- GOACC_deviceptr where appropriate. */
-
- tree ret = walk_tree_without_duplicates (&orig_body,
- &gimplify_oacc_host_data_1, 0);
-
- if (ret)
- {
- error_at (EXPR_LOCATION (expr),
- "undefined use of variable %qE in host_data region",
- DECL_NAME (ret));
- gimplify_adjust_omp_clauses (pre_p, body, &OACC_HOST_DATA_CLAUSES (expr),
- OACC_HOST_DATA);
- return GS_ERROR;
- }
-
- push_gimplify_context ();
-
- gimple *g = gimplify_and_return_first (orig_body, &body);
-
- if (gimple_code (g) == GIMPLE_BIND)
- pop_gimplify_context (g);
- else
- pop_gimplify_context (NULL);
-
- gimplify_adjust_omp_clauses (pre_p, body, &OACC_HOST_DATA_CLAUSES (expr),
- OACC_HOST_DATA);
-
- gimplify_seq_add_stmt (pre_p, g);
-
- return GS_ALL_DONE;
-}
-
/* Gimplify the contents of an OMP_PARALLEL statement. This involves
gimplification of the body, as well as scanning the body for used
variables. We need to do this scan now, because variable-sized
@@ -9648,6 +9540,9 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
case OMP_TEAMS:
ort = OMP_TEAMS_COMBINED (expr) ? ORT_COMBINED_TEAMS : ORT_TEAMS;
break;
+ case OACC_HOST_DATA:
+ ort = ORT_ACC_HOST_DATA;
+ break;
default:
gcc_unreachable ();
}
@@ -9673,6 +9568,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
case OACC_DATA:
+ case OACC_HOST_DATA:
end_ix = BUILT_IN_GOACC_DATA_END;
break;
case OMP_TARGET_DATA:
@@ -9705,6 +9601,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_KERNELS,
OMP_CLAUSES (expr));
break;
+ case OACC_HOST_DATA:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_HOST_DATA,
+ OMP_CLAUSES (expr));
+ break;
case OACC_PARALLEL:
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
OMP_CLAUSES (expr));
@@ -10814,15 +10714,12 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = GS_ALL_DONE;
break;
- case OACC_HOST_DATA:
- ret = gimplify_oacc_host_data (expr_p, pre_p);
- break;
-
case OACC_DECLARE:
gimplify_oacc_declare (expr_p, pre_p);
ret = GS_ALL_DONE;
break;
+ case OACC_HOST_DATA:
case OACC_DATA:
case OACC_KERNELS:
case OACC_PARALLEL:
diff --git gcc/omp-builtins.def gcc/omp-builtins.def
index 63e5e6e..35f5014 100644
--- gcc/omp-builtins.def
+++ gcc/omp-builtins.def
@@ -47,8 +47,8 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_UPDATE, "GOACC_update",
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait",
BT_FN_VOID_INT_INT_VAR,
ATTR_NOTHROW_LIST)
-DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DEVICEPTR, "GOACC_deviceptr",
- BT_FN_PTR_PTR, ATTR_CONST_NOTHROW_LEAF_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_HOST_DATA, "GOACC_host_data",
+ BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
BT_FN_INT_INT, ATTR_CONST_NOTHROW_LEAF_LIST)
diff --git gcc/omp-low.c gcc/omp-low.c
index a1e7a14..88e41b8 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -2071,6 +2071,7 @@ scan_sharing_clauses (tree clauses, omp_context *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)
@@ -2274,7 +2275,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE__CACHE_:
sorry ("Clause not supported yet");
break;
@@ -2430,6 +2430,7 @@ 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:
@@ -2448,7 +2449,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_DEVICE_RESIDENT:
- case OMP_CLAUSE_USE_DEVICE:
case OMP_CLAUSE__CACHE_:
sorry ("Clause not supported yet");
break;
@@ -3763,6 +3763,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
stmt_name = "enter/exit data"; break;
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
+ break;
default: gcc_unreachable ();
}
switch (gimple_omp_target_kind (ctx->stmt))
@@ -3774,6 +3776,8 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
case GF_OMP_TARGET_KIND_OACC_KERNELS:
ctx_stmt_name = "kernels"; break;
case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+ ctx_stmt_name = "host_data"; break;
default: gcc_unreachable ();
}
@@ -12730,6 +12734,7 @@ expand_omp_target (struct omp_region *region)
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA:
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
data_region = true;
break;
default:
@@ -12980,6 +12985,9 @@ expand_omp_target (struct omp_region *region)
case GF_OMP_TARGET_KIND_OACC_DECLARE:
start_ix = BUILT_IN_GOACC_DECLARE;
break;
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
+ start_ix = BUILT_IN_GOACC_HOST_DATA;
+ break;
default:
gcc_unreachable ();
}
@@ -13104,6 +13112,7 @@ expand_omp_target (struct omp_region *region)
case BUILT_IN_GOACC_DATA_START:
case BUILT_IN_GOACC_DECLARE:
case BUILT_IN_GOMP_TARGET_DATA:
+ case BUILT_IN_GOACC_HOST_DATA:
break;
case BUILT_IN_GOMP_TARGET:
case BUILT_IN_GOMP_TARGET_UPDATE:
@@ -13445,6 +13454,7 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_DATA:
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
if (is_gimple_omp_oacc (stmt))
region->kind = gimple_omp_target_kind (stmt);
break;
@@ -15277,6 +15287,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_DATA:
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
data_region = true;
break;
default:
@@ -15485,6 +15496,7 @@ 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);
@@ -15870,12 +15882,14 @@ 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)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
tkind = GOMP_MAP_USE_DEVICE_PTR;
else
tkind = GOMP_MAP_FIRSTPRIVATE_INT;
@@ -16078,10 +16092,12 @@ 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)
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE_PTR
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_USE_DEVICE)
x = build_sender_ref (var, ctx);
else
x = build_receiver_ref (var, false, ctx);
@@ -17076,6 +17092,7 @@ make_gimple_omp_edges (basic_block bb, struct omp_region **region,
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_DATA:
+ case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
break;
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_ENTER_DATA:
diff --git gcc/testsuite/c-c++-common/goacc/host_data-1.c gcc/testsuite/c-c++-common/goacc/host_data-1.c
index 521c854..a8922df 100644
--- gcc/testsuite/c-c++-common/goacc/host_data-1.c
+++ gcc/testsuite/c-c++-common/goacc/host_data-1.c
@@ -1,13 +1,11 @@
/* Test valid use of host_data directive. */
/* { dg-do compile } */
-int v0;
int v1[3][3];
void
f (void)
{
- int v2 = 3;
-#pragma acc host_data use_device(v2, v0, v1)
+#pragma acc host_data use_device(v1)
;
}
diff --git gcc/testsuite/c-c++-common/goacc/host_data-2.c gcc/testsuite/c-c++-common/goacc/host_data-2.c
index e5213a0..1dd5be7 100644
--- gcc/testsuite/c-c++-common/goacc/host_data-2.c
+++ gcc/testsuite/c-c++-common/goacc/host_data-2.c
@@ -10,4 +10,14 @@ f (void)
int v2 = 3;
#pragma acc host_data copy(v2) /* { dg-error "not valid for" } */
;
+
+#pragma acc host_data use_device(v2)
+ ;
+ /* { dg-error ".use_device. variable is neither a pointer nor an array" "" { target c } 14 } */
+ /* { dg-error ".use_device. variable is neither a pointer, nor an arraynor reference to pointer or array" "" { target c++ } 14 } */
+
+#pragma acc host_data use_device(v0)
+ ;
+ /* { dg-error ".use_device. variable is neither a pointer nor an array" "" { target c } 19 } */
+ /* { dg-error ".use_device. variable is neither a pointer, nor an arraynor reference to pointer or array" "" { target c++ } 19 } */
}
diff --git gcc/testsuite/gfortran.dg/goacc/coarray.f95 gcc/testsuite/gfortran.dg/goacc/coarray.f95
index 0ca14e2..d2f10d5 100644
--- gcc/testsuite/gfortran.dg/goacc/coarray.f95
+++ gcc/testsuite/gfortran.dg/goacc/coarray.f95
@@ -5,8 +5,6 @@
! { dg-xfail-if "<http://gcc.gnu.org/PR63861>" { *-*-* } }
! { dg-excess-errors "TODO" }
-! TODO: These cases must fail
-
module test
contains
subroutine oacc1(a)
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 8b5aba2..da19e8d 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1072,6 +1072,7 @@ 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:
@@ -1743,6 +1744,7 @@ 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 libgomp/ChangeLog libgomp/ChangeLog]
diff --git libgomp/libgomp.map libgomp/libgomp.map
index cceb92d..a42142f 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -394,11 +394,11 @@ GOACC_2.0.1 {
global:
GOACC_declare;
GOACC_parallel_keyed;
+ GOACC_host_data;
} GOACC_2.0;
GOACC_2.0.GOMP_4_BRANCH {
global:
- GOACC_deviceptr;
GOMP_set_offload_targets;
} GOACC_2.0.1;
diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c
index 11edcce..588782b 100644
--- libgomp/oacc-mem.c
+++ libgomp/oacc-mem.c
@@ -203,38 +203,6 @@ acc_deviceptr (void *h)
return d;
}
-/* This function is used as a helper in generated code to implement pointer
- lookup in host_data regions. Unlike acc_deviceptr, it returns its argument
- unchanged on a shared-memory system (e.g. the host). */
-
-void *
-GOACC_deviceptr (void *h)
-{
- splay_tree_key n;
- void *d;
- void *offset;
-
- goacc_lazy_initialize ();
-
- struct goacc_thread *thr = goacc_thread ();
-
- if ((thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) == 0)
- {
- n = lookup_host (thr->dev, h, 1);
-
- if (!n)
- return NULL;
-
- offset = h - n->host_start;
-
- d = n->tgt->tgt_start + n->tgt_offset + offset;
-
- return d;
- }
- else
- return h;
-}
-
/* Return the host pointer that corresponds to device data D. Or NULL
if no mapping. */
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index d66e343..e60a61b 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -555,6 +555,46 @@ GOACC_wait (int async, int num_waits, ...)
goacc_thread ()->dev->openacc.async_wait_all_async_func (acc_async_noval);
}
+void
+GOACC_host_data (int device, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+ bool host_fallback = device == GOMP_DEVICE_HOST_FALLBACK;
+ struct target_mem_desc *tgt;
+
+#ifdef HAVE_INTTYPES_H
+ gomp_debug (0, "%s: mapnum=%"PRIu64", hostaddrs=%p, size=%p, kinds=%p\n",
+ __FUNCTION__, (uint64_t) mapnum, hostaddrs, sizes, kinds);
+#else
+ gomp_debug (0, "%s: mapnum=%lu, hostaddrs=%p, sizes=%p, kinds=%p\n",
+ __FUNCTION__, (unsigned long) mapnum, hostaddrs, sizes, kinds);
+#endif
+
+ goacc_lazy_initialize ();
+
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+
+ /* Host fallback or 'do nothing'. */
+ if ((acc_dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ || host_fallback)
+ {
+ tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true,
+ GOMP_MAP_VARS_OPENACC);
+ tgt->prev = thr->mapped_data;
+ thr->mapped_data = tgt;
+
+ return;
+ }
+
+ gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
+ tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
+ GOMP_MAP_VARS_OPENACC);
+ gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
+ tgt->prev = thr->mapped_data;
+ thr->mapped_data = tgt;
+}
+
int
GOACC_get_num_threads (int gang, int worker, int vector)
{
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
index 15ccb27..51745ba 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c
@@ -1,7 +1,6 @@
/* { dg-do run { target openacc_nvidia_accel_selected } } */
/* { dg-additional-options "-lcuda -lcublas -lcudart" } */
-#include <stdio.h>
#include <stdlib.h>
#include <openacc.h>
#include <cuda.h>
@@ -30,35 +29,13 @@ saxpy_target (int n, float a, float *x, float *y)
int
main(int argc, char **argv)
{
- const int N = 8;
+#define N 8
int i;
- float *x_ref, *y_ref;
- float *x, *y;
+ float x_ref[N], y_ref[N];
+ float x[N], y[N];
cublasHandle_t h;
float a = 2.0;
- x_ref = (float*) malloc (N * sizeof(float));
- y_ref = (float*) malloc (N * sizeof(float));
-
- x = (float*) malloc (N * sizeof(float));
- y = (float*) malloc (N * sizeof(float));
-
-#pragma acc data copyin (x[0:N]) copy (y[0:N])
- {
- float *xp, *yp;
-#pragma acc host_data use_device (x, y)
- {
-#pragma acc parallel pcopy (xp, yp) present (x, y)
- {
- xp = x;
- yp = y;
- }
- }
-
- if (xp != acc_deviceptr (x) || yp != acc_deviceptr (y))
- abort ();
- }
-
for (i = 0; i < N; i++)
{
x[i] = x_ref[i] = 4.0 + i;
@@ -106,13 +83,11 @@ main(int argc, char **argv)
for (i = 0; i < N; i++)
y[i] = 3.0;
-#pragma acc data copyin (x[0:N]) copyin (a, N) copy (y[0:N])
+ /* There's no need to use host_data here. */
+#pragma acc data copyin (x[0:N]) copyin (a) copy (y[0:N])
{
-#pragma acc host_data use_device (x, y)
- {
-#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a, N)
- saxpy_target (N, a, x, y);
- }
+#pragma acc parallel present (x[0:N]) pcopy (y[0:N]) present (a)
+ saxpy_target (N, a, x, y);
}
for (i = 0; i < N; i++)
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
index 511ec64..614f143 100644
--- libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-2.c
@@ -1,50 +1,31 @@
-/* { dg-do run } */
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
#include <stdlib.h>
+#include <openacc.h>
-struct by_lightning {
- int a;
- int b;
- int c;
-};
+char *global_in_host;
-int main (int argc, char* argv[])
+void foo (char *in)
{
- int x;
- void *q = NULL, *r = NULL, *p = NULL, *s = NULL, *t = NULL;
- long u;
- struct by_lightning on_the_head = {1, 2, 3};
- int arr[10], *f = NULL;
- _Complex float cf;
- #pragma acc enter data copyin (x, arr, on_the_head, cf)
- #pragma acc host_data use_device (x, arr, on_the_head, cf)
+ if (!acc_is_present (global_in_host, sizeof (*global_in_host))
+ || in != acc_deviceptr (global_in_host))
+ abort ();
+}
+
+int
+main (int argc, char **argv)
+{
+ char mydata[1024];
+
+ global_in_host = mydata;
+
+#pragma acc data copyin(mydata)
{
- q = &x;
+#pragma acc host_data use_device (mydata)
{
- f = &arr[5];
- r = f;
- s = &__real__ cf;
- t = &on_the_head.c;
- u = (long) &__imag__ cf;
- #pragma acc parallel copyout(p) present (x, arr, on_the_head, cf)
- {
- /* This will not (and must not) call GOACC_deviceptr, but '&x' will be
- the address on the device (if appropriate) regardless. */
- p = &x;
- }
+ foo (mydata);
}
}
- #pragma acc exit data delete (x)
-
-#if ACC_MEM_SHARED
- if (q != &x || f != &arr[5] || r != f || s != &(__real__ cf)
- || t != &on_the_head.c || u != (long) &(__imag__ cf) || p != &x)
- abort ();
-#else
- if (q == &x || f == &arr[5] || r != f || s == &(__real__ cf)
- || t == &on_the_head.c || u == (long) &(__imag__ cf) || p == &x)
- abort ();
-#endif
return 0;
}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c
new file mode 100644
index 0000000..7d9b5f7
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-3.c
@@ -0,0 +1,29 @@
+/* { dg-do compile } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+ int x[N];
+
+#pragma acc data copyin (x[0:N])
+ {
+ int *xp;
+#pragma acc host_data use_device (x)
+ {
+ /* This use of the present clause is undefined behaviour for OpenACC. */
+#pragma acc parallel present (x) copyout (xp) /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */
+ {
+ xp = x;
+ }
+ }
+
+ if (xp != acc_deviceptr (x))
+ abort ();
+ }
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c
new file mode 100644
index 0000000..0ab5a35
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-4.c
@@ -0,0 +1,29 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+ int x[N], *xp2;
+
+#pragma acc data copyin (x[0:N])
+ {
+ int *xp;
+#pragma acc host_data use_device (x)
+ {
+#pragma acc data
+ {
+ xp = x;
+ }
+ xp2 = x;
+ }
+
+ if (xp != acc_deviceptr (x) || xp2 != xp)
+ abort ();
+ }
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c
new file mode 100644
index 0000000..a3737a7
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-5.c
@@ -0,0 +1,38 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+ int x[N], y[N], *yp;
+
+ yp = y + 1;
+
+#pragma acc data copyin (x[0:N])
+ {
+ int *xp, *yp2;
+#pragma acc host_data use_device (x)
+ {
+#pragma acc data copyin (y)
+ {
+#pragma acc host_data use_device (yp)
+ {
+ xp = x;
+ yp2 = yp;
+ }
+
+ if (yp2 != acc_deviceptr (yp))
+ abort ();
+ }
+ }
+
+ if (xp != acc_deviceptr (x))
+ abort ();
+
+ }
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
new file mode 100644
index 0000000..a841488
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-6.c
@@ -0,0 +1,31 @@
+/* { dg-do compile } */
+
+#include <openacc.h>
+#include <stdlib.h>
+
+#define N 1024
+
+int main (int argc, char* argv[])
+{
+ int x[N];
+
+#pragma acc data copyin (x[0:N])
+ {
+ int *xp;
+#pragma acc host_data use_device (x)
+ {
+ /* Here 'x' being implicitly firstprivate for the parallel region
+ conflicts with it being declared as use_device in the enclosing
+ host_data region. */
+#pragma acc parallel copyout (xp)
+ {
+ xp = x; /* { dg-error "variable 'x' declared in enclosing 'host_data' region" } */
+ }
+ }
+
+ if (xp != acc_deviceptr (x))
+ abort ();
+ }
+
+ return 0;
+}
diff --git libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90 libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90
index a219eaf..9bb79c3 100644
--- libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90
+++ libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90
@@ -1,6 +1,9 @@
-! { dg-do run } */
+! { dg-do run }
! { dg-additional-options "-cpp" }
+! { dg-xfail-if "TODO" { *-*-* } }
+! { dg-excess-errors "TODO" }
+
program test
implicit none
Grüße
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
next prev parent reply other threads:[~2015-12-02 22:14 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 ` Thomas Schwinge [this message]
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
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=87r3j4lcrd.fsf@kepler.schwinge.homeip.net \
--to=thomas@codesourcery.com \
--cc=Nathan_Sidwell@mentor.com \
--cc=cesar@codesourcery.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).