* [gomp4] Initial support for OpenACC data clauses
@ 2014-01-14 15:09 Thomas Schwinge
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
` (3 more replies)
0 siblings, 4 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-01-14 15:09 UTC (permalink / raw)
To: jakub, gcc-patches
[-- Attachment #1: Type: text/plain, Size: 532 bytes --]
Hi!
Here is a patch series that adds initial support for OpenACC data
clauses. It is not yet complete, but I thought I might as well already
now strive to get this integrated upstream instead of "hoarding" the
patches locally.
Would it be a good idea to also commit to trunk the (portions of the)
patches that don't directly relate with OpenACC stuff? That way, trunk
and gomp-4_0-branch would diverge a little less? Or, would you first
like to see all of this stabilitize on gomp-4_0-branch?
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing.
2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas
@ 2014-01-14 15:10 ` thomas
2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas
2014-02-21 19:48 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
[not found] ` <538DF785.3050206@mentor.com>
2 siblings, 1 reply; 22+ messages in thread
From: thomas @ 2014-01-14 15:10 UTC (permalink / raw)
To: jakub, gcc-patches; +Cc: Thomas Schwinge
From: Thomas Schwinge <thomas@codesourcery.com>
gcc/c/
* c-parser.c (c_parser_oacc_all_clauses): New function.
(c_parser_oacc_parallel): Use it.
* c-typeck.c (c_finish_omp_clauses): Update comment. Remove
duplicated variable initialization.
---
gcc/c/c-parser.c | 59 +++++++++++++++++++++++++++++++++++++++++++++++++++-----
gcc/c/c-typeck.c | 4 ++--
2 files changed, 56 insertions(+), 7 deletions(-)
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index ce46f31..c8b80db 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -9465,7 +9465,7 @@ c_parser_pragma_pch_preprocess (c_parser *parser)
c_common_pch_pragma (parse_in, TREE_STRING_POINTER (name));
}
\f
-/* OpenMP 2.5 / 3.0 / 3.1 / 4.0 parsing routines. */
+/* OpenACC and OpenMP parsing routines. */
/* Returns name of the next clause.
If the clause is not recognized PRAGMA_OMP_CLAUSE_NONE is returned and
@@ -10767,9 +10767,58 @@ c_parser_omp_clause_uniform (c_parser *parser, tree list)
return list;
}
+/* Parse all OpenACC clauses. The set clauses allowed by the directive
+ is a bitmask in MASK. Return the list of clauses found. */
+
+static tree
+c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
+ const char *where, bool finish_p = true)
+{
+ tree clauses = NULL;
+ bool first = true;
+
+ while (c_parser_next_token_is_not (parser, CPP_PRAGMA_EOL))
+ {
+ location_t here;
+ pragma_omp_clause c_kind;
+ const char *c_name;
+ tree prev = clauses;
+
+ if (!first && c_parser_next_token_is (parser, CPP_COMMA))
+ c_parser_consume_token (parser);
+
+ here = c_parser_peek_token (parser)->location;
+ c_kind = c_parser_omp_clause_name (parser);
+
+ switch (c_kind)
+ {
+ default:
+ c_parser_error (parser, "expected clause");
+ goto saw_error;
+ }
+
+ first = false;
+
+ if (((mask >> c_kind) & 1) == 0 && !parser->error)
+ {
+ /* Remove the invalid clause(s) from the list to avoid
+ confusing the rest of the compiler. */
+ clauses = prev;
+ error_at (here, "%qs is not valid for %qs", c_name, where);
+ }
+ }
+
+ saw_error:
+ c_parser_skip_to_pragma_eol (parser);
+
+ if (finish_p)
+ return c_finish_omp_clauses (clauses);
+
+ return clauses;
+}
+
/* Parse all OpenMP clauses. The set clauses allowed by the directive
- is a bitmask in MASK. Return the list of clauses found; the result
- of clause default goes in *pdefault. */
+ is a bitmask in MASK. Return the list of clauses found. */
static tree
c_parser_omp_all_clauses (c_parser *parser, omp_clause_mask mask,
@@ -11019,8 +11068,8 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser)
{
tree stmt, clauses, block;
- clauses = c_parser_omp_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
- "#pragma acc parallel");
+ clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
+ "#pragma acc parallel");
gcc_assert (clauses == NULL);
block = c_begin_omp_parallel ();
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index 854e149..81f0c5c 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11661,7 +11661,7 @@ c_find_omp_placeholder_r (tree *tp, int *, void *data)
return NULL_TREE;
}
-/* For all elements of CLAUSES, validate them vs OpenMP constraints.
+/* For all elements of CLAUSES, validate them against their constraints.
Remove any elements from the list that are invalid. */
tree
@@ -11669,7 +11669,7 @@ c_finish_omp_clauses (tree clauses)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
bitmap_head aligned_head;
- tree c, t, *pc = &clauses;
+ tree c, t, *pc;
bool branch_seen = false;
bool copyprivate_seen = false;
tree *nowait_clause = NULL;
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 5/6] Initial support in the C front end for OpenACC data clauses.
2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas
@ 2014-01-14 15:10 ` thomas
2014-01-14 15:10 ` [gomp4 6/6] Enable initial " thomas
2014-02-12 11:17 ` [gomp4 5/6] Initial " Thomas Schwinge
0 siblings, 2 replies; 22+ messages in thread
From: thomas @ 2014-01-14 15:10 UTC (permalink / raw)
To: jakub, gcc-patches; +Cc: Thomas Schwinge
From: Thomas Schwinge <thomas@codesourcery.com>
gcc/c-family/
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY,
PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE,
PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEVICEPTR,
PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and
PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Handle these.
(c_parser_oacc_data_clause, c_parser_oacc_data_clause_deviceptr):
New functions.
(c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_COPY,
PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT,
PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DELETE,
PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and
PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE.
gcc/
* tree-core.h (omp_clause_code): Update description for
OMP_CLAUSE_MAP.
---
gcc/c-family/c-pragma.h | 12 +++-
gcc/c/c-parser.c | 171 +++++++++++++++++++++++++++++++++++++++++++++++-
gcc/tree-core.h | 6 +-
3 files changed, 184 insertions(+), 5 deletions(-)
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index 64eed11..2c8af67 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -63,18 +63,23 @@ typedef enum pragma_kind {
} pragma_kind;
-/* All clauses defined by OpenMP 2.5, 3.0, 3.1 and 4.0.
+/* All clauses defined by OpenACC 2.0, and OpenMP 2.5, 3.0, 3.1, and 4.0.
Used internally by both C and C++ parsers. */
typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_NONE = 0,
PRAGMA_OMP_CLAUSE_ALIGNED,
PRAGMA_OMP_CLAUSE_COLLAPSE,
+ PRAGMA_OMP_CLAUSE_COPY,
PRAGMA_OMP_CLAUSE_COPYIN,
+ PRAGMA_OMP_CLAUSE_COPYOUT,
PRAGMA_OMP_CLAUSE_COPYPRIVATE,
+ PRAGMA_OMP_CLAUSE_CREATE,
PRAGMA_OMP_CLAUSE_DEFAULT,
+ PRAGMA_OMP_CLAUSE_DELETE,
PRAGMA_OMP_CLAUSE_DEPEND,
PRAGMA_OMP_CLAUSE_DEVICE,
+ PRAGMA_OMP_CLAUSE_DEVICEPTR,
PRAGMA_OMP_CLAUSE_DIST_SCHEDULE,
PRAGMA_OMP_CLAUSE_FINAL,
PRAGMA_OMP_CLAUSE_FIRSTPRIVATE,
@@ -92,6 +97,11 @@ typedef enum pragma_omp_clause {
PRAGMA_OMP_CLAUSE_NUM_THREADS,
PRAGMA_OMP_CLAUSE_ORDERED,
PRAGMA_OMP_CLAUSE_PARALLEL,
+ PRAGMA_OMP_CLAUSE_PRESENT,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT,
+ PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE,
PRAGMA_OMP_CLAUSE_PRIVATE,
PRAGMA_OMP_CLAUSE_PROC_BIND,
PRAGMA_OMP_CLAUSE_REDUCTION,
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index c8b80db..48c55e6 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -9496,16 +9496,26 @@ c_parser_omp_clause_name (c_parser *parser)
case 'c':
if (!strcmp ("collapse", p))
result = PRAGMA_OMP_CLAUSE_COLLAPSE;
+ else if (!strcmp ("copy", p))
+ result = PRAGMA_OMP_CLAUSE_COPY;
else if (!strcmp ("copyin", p))
result = PRAGMA_OMP_CLAUSE_COPYIN;
+ else if (!strcmp ("copyout", p))
+ result = PRAGMA_OMP_CLAUSE_COPYOUT;
else if (!strcmp ("copyprivate", p))
result = PRAGMA_OMP_CLAUSE_COPYPRIVATE;
+ else if (!strcmp ("create", p))
+ result = PRAGMA_OMP_CLAUSE_CREATE;
break;
case 'd':
- if (!strcmp ("depend", p))
+ if (!strcmp ("delete", p))
+ result = PRAGMA_OMP_CLAUSE_DELETE;
+ else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
+ else if (!strcmp ("deviceptr", p))
+ result = PRAGMA_OMP_CLAUSE_DEVICEPTR;
else if (!strcmp ("dist_schedule", p))
result = PRAGMA_OMP_CLAUSE_DIST_SCHEDULE;
break;
@@ -9550,6 +9560,16 @@ c_parser_omp_clause_name (c_parser *parser)
case 'p':
if (!strcmp ("parallel", p))
result = PRAGMA_OMP_CLAUSE_PARALLEL;
+ else if (!strcmp ("present", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT;
+ else if (!strcmp ("present_or_copy", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY;
+ else if (!strcmp ("present_or_copyin", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN;
+ else if (!strcmp ("present_or_copyout", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT;
+ else if (!strcmp ("present_or_create", p))
+ result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE;
else if (!strcmp ("private", p))
result = PRAGMA_OMP_CLAUSE_PRIVATE;
else if (!strcmp ("proc_bind", p))
@@ -9611,7 +9631,7 @@ check_no_duplicate_clause (tree clauses, enum omp_clause_code code,
}
}
-/* OpenMP 2.5:
+/* OpenACC 2.0, OpenMP 2.5:
variable-list:
identifier
variable-list , identifier
@@ -9712,7 +9732,7 @@ c_parser_omp_variable_list (c_parser *parser,
}
/* Similarly, but expect leading and trailing parenthesis. This is a very
- common case for omp clauses. */
+ common case for OpenACC and OpenMP clauses. */
static tree
c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
@@ -9729,6 +9749,107 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
return list;
}
+/* OpenACC 2.0:
+ copy ( variable-list )
+ copyin ( variable-list )
+ copyout ( variable-list )
+ create ( variable-list )
+ delete ( variable-list )
+ present ( variable-list )
+ present_or_copy ( variable-list )
+ present_or_copyin ( variable-list )
+ present_or_copyout ( variable-list )
+ present_or_create ( variable-list ) */
+
+static tree
+c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
+ tree list)
+{
+ enum omp_clause_map_kind kind;
+ switch (c_kind)
+ {
+ default:
+ gcc_unreachable ();
+ case PRAGMA_OMP_CLAUSE_COPY:
+ kind = OMP_CLAUSE_MAP_FORCE_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ kind = OMP_CLAUSE_MAP_FORCE_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FORCE_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ kind = OMP_CLAUSE_MAP_FORCE_ALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_DELETE:
+ kind = OMP_CLAUSE_MAP_FORCE_DEALLOC;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ kind = OMP_CLAUSE_MAP_FORCE_PRESENT;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ kind = OMP_CLAUSE_MAP_TOFROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ kind = OMP_CLAUSE_MAP_TO;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ kind = OMP_CLAUSE_MAP_FROM;
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
+ kind = OMP_CLAUSE_MAP_ALLOC;
+ break;
+ }
+ tree nl, c;
+ nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
+
+ for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
+ OMP_CLAUSE_MAP_KIND (c) = kind;
+
+ return nl;
+}
+
+/* OpenACC 2.0:
+ deviceptr ( variable-list ) */
+
+static tree
+c_parser_oacc_data_clause_deviceptr (c_parser *parser, tree list)
+{
+ location_t loc = c_parser_peek_token (parser)->location;
+ tree vars, t;
+
+ /* Can't use OMP_CLAUSE_MAP here (that is, can't use the generic
+ c_parser_oacc_data_clause), as for PRAGMA_OMP_CLAUSE_DEVICEPTR,
+ variable-list must only allow for pointer variables. */
+ vars = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_ERROR, NULL);
+ for (t = vars; t && t; t = TREE_CHAIN (t))
+ {
+ tree v = TREE_PURPOSE (t);
+
+ /* FIXME diagnostics: Ideally we should keep individual
+ locations for all the variables in the var list to make the
+ following errors more precise. Perhaps
+ c_parser_omp_var_list_parens() should construct a list of
+ locations to go along with the var list. */
+
+ if (TREE_CODE (v) != VAR_DECL)
+ error_at (loc, "%qD is not a variable", v);
+ else if (TREE_TYPE (v) == error_mark_node)
+ ;
+ else if (!POINTER_TYPE_P (TREE_TYPE (v)))
+ error_at (loc, "%qD is not a pointer variable", v);
+
+ tree u = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_MAP_KIND (u) = OMP_CLAUSE_MAP_FORCE_DEVICEPTR;
+ OMP_CLAUSE_DECL (u) = v;
+ OMP_CLAUSE_CHAIN (u) = list;
+ list = u;
+ }
+
+ return list;
+}
+
/* OpenMP 3.0:
collapse ( constant-expression ) */
@@ -10792,6 +10913,50 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
switch (c_kind)
{
+ case PRAGMA_OMP_CLAUSE_COPY:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYIN:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_COPYOUT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_CREATE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "create";
+ break;
+ case PRAGMA_OMP_CLAUSE_DELETE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "delete";
+ break;
+ case PRAGMA_OMP_CLAUSE_DEVICEPTR:
+ clauses = c_parser_oacc_data_clause_deviceptr (parser, clauses);
+ c_name = "deviceptr";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copy";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyin";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_copyout";
+ break;
+ case PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "present_or_create";
+ break;
default:
c_parser_error (parser, "expected clause");
goto saw_error;
diff --git gcc/tree-core.h gcc/tree-core.h
index 0aedea3..bfe4943 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -258,7 +258,11 @@ enum omp_clause_code {
/* OpenMP clause: to (variable-list). */
OMP_CLAUSE_TO,
- /* OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
+ /* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr,
+ present, present_or_copy, present_or_copyin, present_or_copyout,
+ present_or_create} (variable-list).
+
+ OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
/* Internal clause: temporary for combined loops expansion. */
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 6/6] Enable initial support in the C front end for OpenACC data clauses.
2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas
@ 2014-01-14 15:10 ` thomas
2014-02-12 11:17 ` [gomp4 5/6] Initial " Thomas Schwinge
1 sibling, 0 replies; 22+ messages in thread
From: thomas @ 2014-01-14 15:10 UTC (permalink / raw)
To: jakub, gcc-patches; +Cc: Thomas Schwinge
From: Thomas Schwinge <thomas@codesourcery.com>
gcc/c/
* c-parser.c (OACC_PARALLEL_CLAUSE_MASK): Add
PRAGMA_OMP_CLAUSE_COPY, PRAGMA_OMP_CLAUSE_COPYIN,
PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE,
PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and
PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE.
gcc/testsuite/
* c-c++-common/goacc/data-clause-duplicate-1.c: New file.
* c-c++-common/goacc/deviceptr-1.c: New file.
libgomp/
* testsuite/libgomp.oacc-c/parallel-1.c: Extend.
---
gcc/c/c-parser.c | 14 +-
.../c-c++-common/goacc/data-clause-duplicate-1.c | 13 ++
gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 64 +++++++++
libgomp/testsuite/libgomp.oacc-c/parallel-1.c | 150 +++++++++++++++++++--
4 files changed, 228 insertions(+), 13 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 48c55e6..d6a2af0 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -11225,8 +11225,17 @@ c_parser_omp_structured_block (c_parser *parser)
LOC is the location of the #pragma token.
*/
-#define OACC_PARALLEL_CLAUSE_MASK \
- (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_NONE)
+#define OACC_PARALLEL_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
static tree
c_parser_oacc_parallel (location_t loc, c_parser *parser)
@@ -11235,7 +11244,6 @@ c_parser_oacc_parallel (location_t loc, c_parser *parser)
clauses = c_parser_oacc_all_clauses (parser, OACC_PARALLEL_CLAUSE_MASK,
"#pragma acc parallel");
- gcc_assert (clauses == NULL);
block = c_begin_omp_parallel ();
add_stmt (c_parser_omp_structured_block (parser));
diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
new file mode 100644
index 0000000..1bcf5be
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -0,0 +1,13 @@
+void
+fun (void)
+{
+ float *fp;
+#pragma acc parallel copy(fp[0:2],fp[0:2]) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel create(fp[:10]) deviceptr(fp)
+ /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */
+ /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */
+ ;
+}
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
new file mode 100644
index 0000000..0f0cf0c
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -0,0 +1,64 @@
+void
+fun1 (void)
+{
+#pragma acc parallel deviceptr(u) /* { dg-error "'u' undeclared" } */
+ ;
+#pragma acc parallel deviceptr(u[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+
+#pragma acc parallel deviceptr(fun1) /* { dg-error "'fun1' is not a variable" } */
+ ;
+#pragma acc parallel deviceptr(fun1[2:5])
+ /* { dg-error "'fun1' is not a variable" "not a variable" { target *-*-* } 11 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 11 } */
+ ;
+
+ int i;
+#pragma acc parallel deviceptr(i) /* { dg-error "'i' is not a pointer variable" } */
+ ;
+#pragma acc parallel deviceptr(i[0:4])
+ /* { dg-error "'i' is not a pointer variable" "not a pointer variable" { target *-*-* } 19 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 19 } */
+ ;
+
+ float fa[10];
+#pragma acc parallel deviceptr(fa) /* { dg-error "'fa' is not a pointer variable" } */
+ ;
+#pragma acc parallel deviceptr(fa[1:5])
+ /* { dg-error "'fa' is not a pointer variable" "not a pointer variable" { target *-*-* } 27 } */
+ /* { dg-error "expected '\\\)' before '\\\[' token" "array" { target *-*-* } 27 } */
+ ;
+
+ float *fp;
+#pragma acc parallel deviceptr(fp)
+ ;
+#pragma acc parallel deviceptr(fp[0:4]) /* { dg-error "expected '\\\)' before '\\\[' token" } */
+ ;
+}
+
+void
+fun2 (void)
+{
+ int i;
+ float *fp;
+#pragma acc parallel deviceptr(fp,u,fun2,i,fp)
+ /* { dg-error "'u' undeclared" "u undeclared" { target *-*-* } 44 } */
+ /* { dg-error "'fun2' is not a variable" "fun2 not a variable" { target *-*-* } 44 } */
+ /* { dg-error "'i' is not a pointer variable" "i not a pointer variable" { target *-*-* } 44 } */
+ /* { dg-error "'fp' appears more than once in map clauses" "fp more than once" { target *-*-* } 44 } */
+ ;
+}
+
+void
+fun3 (void)
+{
+ float *fp;
+#pragma acc parallel deviceptr(fp,fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel deviceptr(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+#pragma acc parallel copy(fp) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
+}
+
+/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
diff --git libgomp/testsuite/libgomp.oacc-c/parallel-1.c libgomp/testsuite/libgomp.oacc-c/parallel-1.c
index b40545d..ff54b9d 100644
--- libgomp/testsuite/libgomp.oacc-c/parallel-1.c
+++ libgomp/testsuite/libgomp.oacc-c/parallel-1.c
@@ -2,25 +2,155 @@
extern void abort ();
-volatile int i;
+int i;
int main(void)
{
- volatile int j;
+ int j, v;
- i = -0x42;
- j = -42;
-#pragma acc parallel
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) copyin (i, j)
{
- if (i != -0x42 || j != -42)
+ if (i != -1 || j != -2)
abort ();
- i = 42;
- j = 0x42;
- if (i != 42 || j != 0x42)
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
abort ();
+ v = 1;
}
- if (i != 42 || j != 0x42)
+ if (v != 1 || i != -1 || j != -2)
abort ();
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) copyout (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) copy (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) create (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != -1 || j != -2)
+ abort ();
+#endif
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyin (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != -1 || j != -2)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copyout (i, j)
+ {
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present_or_copy (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v) present (i, j)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+#endif
+
+#if 0
+ i = -1;
+ j = -2;
+ v = 0;
+#pragma acc parallel /* copyout */ present_or_copyout (v)
+ {
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ v = 1;
+ }
+ if (v != 1 || i != 2 || j != 1)
+ abort ();
+#endif
+
return 0;
}
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET.
2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
@ 2014-01-14 15:10 ` thomas
2014-01-14 15:10 ` [gomp4 2/6] Prepare for extending omp_clause_map_kind thomas
2014-01-28 9:44 ` [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
` (2 subsequent siblings)
3 siblings, 1 reply; 22+ messages in thread
From: thomas @ 2014-01-14 15:10 UTC (permalink / raw)
To: jakub, gcc-patches; +Cc: Thomas Schwinge
From: Thomas Schwinge <thomas@codesourcery.com>
gcc/
* gimplify.c (gimplify_call_expr, gimplify_modify_expr)
(omp_firstprivatize_variable, omp_notice_threadprivate_variable)
(omp_notice_variable, gimplify_adjust_omp_clauses)
(gimplify_omp_workshare): Treat ORT_TARGET as a flag, not as a
value.
---
gcc/gimplify.c | 14 +++++++-------
1 file changed, 7 insertions(+), 7 deletions(-)
diff --git gcc/gimplify.c gcc/gimplify.c
index e45bed2..90507c2 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -2363,7 +2363,7 @@ gimplify_call_expr (tree *expr_p, gimple_seq *pre_p, bool want_value)
during omplower pass instead. */
struct gimplify_omp_ctx *ctx;
for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
- if (ctx->region_type == ORT_TARGET)
+ if (ctx->region_type & ORT_TARGET)
break;
if (ctx == NULL)
fold_stmt (&gsi);
@@ -4534,7 +4534,7 @@ gimplify_modify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
during omplower pass instead. */
struct gimplify_omp_ctx *ctx;
for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
- if (ctx->region_type == ORT_TARGET)
+ if (ctx->region_type & ORT_TARGET)
break;
if (ctx == NULL)
fold_stmt (&gsi);
@@ -5317,7 +5317,7 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
else
return;
}
- else if (ctx->region_type == ORT_TARGET)
+ else if (ctx->region_type & ORT_TARGET)
omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
else if (ctx->region_type != ORT_WORKSHARE
&& ctx->region_type != ORT_SIMD
@@ -5499,7 +5499,7 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
struct gimplify_omp_ctx *octx;
for (octx = ctx; octx; octx = octx->outer_context)
- if (octx->region_type == ORT_TARGET)
+ if (octx->region_type & ORT_TARGET)
{
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
@@ -5560,7 +5560,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- if (ctx->region_type == ORT_TARGET)
+ if (ctx->region_type & ORT_TARGET)
{
if (n == NULL)
{
@@ -6285,7 +6285,7 @@ gimplify_adjust_omp_clauses (tree *list_p)
if (!DECL_P (decl))
break;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
- if (ctx->region_type == ORT_TARGET && !(n->value & GOVD_SEEN))
+ if ((ctx->region_type & ORT_TARGET) && !(n->value & GOVD_SEEN))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -6857,7 +6857,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
- if (ort == ORT_TARGET || ort == ORT_TARGET_DATA)
+ if ((ort & ORT_TARGET) || ort == ORT_TARGET_DATA)
{
push_gimplify_context ();
gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 2/6] Prepare for extending omp_clause_map_kind.
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
@ 2014-01-14 15:10 ` thomas
2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas
0 siblings, 1 reply; 22+ messages in thread
From: thomas @ 2014-01-14 15:10 UTC (permalink / raw)
To: jakub, gcc-patches; +Cc: Thomas Schwinge
From: Thomas Schwinge <thomas@codesourcery.com>
gcc/
* tree-core.h (omp_clause_map_kind): Make the identifiers' bit
patterns more obvious. Add comments.
* omp-low.c (lower_oacc_parallel, lower_omp_target): Test for
omp_clause_map_kind flags set instead of for values.
---
gcc/omp-low.c | 22 ++++++++++++++--------
gcc/tree-core.h | 16 +++++++++++-----
2 files changed, 25 insertions(+), 13 deletions(-)
diff --git gcc/omp-low.c gcc/omp-low.c
index eb755c3..899e970 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -8855,13 +8855,16 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree avar = create_tmp_var (TREE_TYPE (var), NULL);
mark_addressable (avar);
- if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
- && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+ enum omp_clause_map_kind map_kind
+ = OMP_CLAUSE_MAP_KIND (c);
+ if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_TO))
+ || map_kind == OMP_CLAUSE_MAP_POINTER)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
- || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+ if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
@@ -10331,13 +10334,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
tree avar = create_tmp_var (TREE_TYPE (var), NULL);
mark_addressable (avar);
- if (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_ALLOC
- && OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FROM)
+ enum omp_clause_map_kind map_kind
+ = OMP_CLAUSE_MAP_KIND (c);
+ if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_TO))
+ || map_kind == OMP_CLAUSE_MAP_POINTER)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FROM
- || OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_TOFROM)
+ if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
diff --git gcc/tree-core.h gcc/tree-core.h
index e2750e0..3602b5f 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -1112,14 +1112,20 @@ enum omp_clause_depend_kind
enum omp_clause_map_kind
{
- OMP_CLAUSE_MAP_ALLOC,
- OMP_CLAUSE_MAP_TO,
- OMP_CLAUSE_MAP_FROM,
- OMP_CLAUSE_MAP_TOFROM,
+ /* If not already present, allocate. */
+ OMP_CLAUSE_MAP_ALLOC = 0,
+ /* ..., and copy to device. */
+ OMP_CLAUSE_MAP_TO = 1 << 0,
+ /* ..., and copy from device. */
+ OMP_CLAUSE_MAP_FROM = 1 << 1,
+ /* ..., and copy to and from device. */
+ OMP_CLAUSE_MAP_TOFROM = OMP_CLAUSE_MAP_TO | OMP_CLAUSE_MAP_FROM,
+ /* Special map kinds. */
+ OMP_CLAUSE_MAP_SPECIAL = 1 << 2,
/* The following kind is an internal only map kind, used for pointer based
array sections. OMP_CLAUSE_SIZE for these is not the pointer size,
which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */
- OMP_CLAUSE_MAP_POINTER
+ OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL
};
enum omp_clause_proc_bind_kind
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 3/6] Initial support for OpenACC memory mapping semantics.
2014-01-14 15:10 ` [gomp4 2/6] Prepare for extending omp_clause_map_kind thomas
@ 2014-01-14 15:10 ` thomas
2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas
` (2 more replies)
0 siblings, 3 replies; 22+ messages in thread
From: thomas @ 2014-01-14 15:10 UTC (permalink / raw)
To: jakub, gcc-patches; +Cc: Thomas Schwinge
From: Thomas Schwinge <thomas@codesourcery.com>
gcc/
* tree-core.h (omp_clause_map_kind): Add OMP_CLAUSE_MAP_FORCE,
OMP_CLAUSE_MAP_FORCE_ALLOC, OMP_CLAUSE_MAP_FORCE_TO,
OMP_CLAUSE_MAP_FORCE_FROM, OMP_CLAUSE_MAP_FORCE_TOFROM,
OMP_CLAUSE_MAP_FORCE_PRESENT, OMP_CLAUSE_MAP_FORCE_DEALLOC, and
OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
* tree-pretty-print.c (dump_omp_clause): Handle these.
* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_FORCE.
(omp_region_type): Add ORT_TARGET_MAP_FORCE.
(omp_add_variable, omp_notice_threadprivate_variable)
(omp_notice_variable, gimplify_scan_omp_clauses)
(gimplify_adjust_omp_clauses_1): Extend accordingly.
(gimplify_oacc_parallel): Add ORT_TARGET_MAP_FORCE to ORT_TARGET
usage.
* omp-low.c (install_var_field, scan_sharing_clauses)
(lower_oacc_parallel, lower_omp_target): Extend accordingly.
---
gcc/gimplify.c | 92 ++++++++++++++++++++++++++++++++++++++++++-------
gcc/omp-low.c | 33 +++++++++++-------
gcc/tree-core.h | 19 +++++++++-
gcc/tree-pretty-print.c | 21 +++++++++++
4 files changed, 140 insertions(+), 25 deletions(-)
diff --git gcc/gimplify.c gcc/gimplify.c
index 90507c2..633784f 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -69,7 +69,13 @@ enum gimplify_omp_var_data
GOVD_PRIVATE_OUTER_REF = 1024,
GOVD_LINEAR = 2048,
GOVD_ALIGNED = 4096,
+
+ /* Flags for GOVD_MAP. */
+ /* Don't copy back. */
GOVD_MAP_TO_ONLY = 8192,
+ /* Force a specific behavior (or else, a run-time error). */
+ GOVD_MAP_FORCE = 16384,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -86,7 +92,11 @@ enum omp_region_type
ORT_UNTIED_TASK = 5,
ORT_TEAMS = 8,
ORT_TARGET_DATA = 16,
- ORT_TARGET = 32
+ ORT_TARGET = 32,
+
+ /* Flags for ORT_TARGET. */
+ /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
+ ORT_TARGET_MAP_FORCE = 64
};
/* Gimplify hashtable helper. */
@@ -5430,9 +5440,20 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
copy into or out of the context. */
if (!(flags & GOVD_LOCAL))
{
- nflags = flags & GOVD_MAP
- ? GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT
- : flags & GOVD_PRIVATE ? GOVD_PRIVATE : GOVD_FIRSTPRIVATE;
+ if (flags & GOVD_MAP)
+ {
+ nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
+#if 0
+ /* Not sure if this is actually needed; haven't found a case
+ where this would change anything; TODO. */
+ if (flags & GOVD_MAP_FORCE)
+ nflags |= OMP_CLAUSE_MAP_FORCE;
+#endif
+ }
+ else if (flags & GOVD_PRIVATE)
+ nflags = GOVD_PRIVATE;
+ else
+ nflags = GOVD_FIRSTPRIVATE;
nflags |= flags & GOVD_SEEN;
t = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (t) == INDIRECT_REF);
@@ -5501,6 +5522,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
for (octx = ctx; octx; octx = octx->outer_context)
if (octx->region_type & ORT_TARGET)
{
+ gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE));
+
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
{
@@ -5562,19 +5585,45 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
if (ctx->region_type & ORT_TARGET)
{
+ unsigned map_force;
+ if (ctx->region_type & ORT_TARGET_MAP_FORCE)
+ map_force = GOVD_MAP_FORCE;
+ else
+ map_force = 0;
if (n == NULL)
{
if (!lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
- omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT | flags);
}
else
- omp_add_variable (ctx, decl, GOVD_MAP | flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags);
}
else
- n->value |= flags;
+ {
+#if 0
+ /* The following fails for:
+
+ int l = 10;
+ float c[l];
+ #pragma acc parallel copy(c[2:4])
+ {
+ #pragma acc parallel
+ {
+ int t = sizeof c;
+ }
+ }
+
+ ..., which we currently don't have to care about (nesting
+ disabled), but eventually will have to; TODO. */
+ if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT))
+ gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force);
+#endif
+
+ n->value |= flags;
+ }
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
goto do_outer;
}
@@ -5858,6 +5907,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
goto do_add;
case OMP_CLAUSE_MAP:
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ input_location = OMP_CLAUSE_LOCATION (c);
+ /* TODO. */
+ sorry ("data clause not yet implemented");
+ remove = true;
+ break;
+ default:
+ break;
+ }
if (OMP_CLAUSE_SIZE (c)
&& gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p,
NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
@@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
{
- OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
- ? OMP_CLAUSE_MAP_TO
- : OMP_CLAUSE_MAP_TOFROM;
+ unsigned map_kind;
+ map_kind = (flags & GOVD_MAP_TO_ONLY
+ ? OMP_CLAUSE_MAP_TO
+ : OMP_CLAUSE_MAP_TOFROM);
+ if (flags & GOVD_MAP_FORCE)
+ map_kind |= OMP_CLAUSE_MAP_FORCE;
+ OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind;
+
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple g;
gimple_seq body = NULL;
+ enum omp_region_type ort =
+ (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE);
- gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p,
- ORT_TARGET);
+ gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort);
push_gimplify_context ();
diff --git gcc/omp-low.c gcc/omp-low.c
index 899e970..8c7df1b 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1064,6 +1064,8 @@ install_var_field (tree var, bool by_ref, int mask, omp_context *ctx)
|| !splay_tree_lookup (ctx->field_map, (splay_tree_key) var));
gcc_assert ((mask & 2) == 0 || !ctx->sfield_map
|| !splay_tree_lookup (ctx->sfield_map, (splay_tree_key) var));
+ gcc_assert ((mask & 3) == 3
+ || gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
type = TREE_TYPE (var);
if (mask & 4)
@@ -1611,6 +1613,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
case OMP_CLAUSE_MAP:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
@@ -1630,11 +1633,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
{
- gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL);
/* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
#pragma omp target data, there is nothing to map for
those. */
- if (gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
+ if (gimple_code (ctx->stmt) != GIMPLE_OACC_PARALLEL
+ && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
@@ -8709,8 +8712,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
- case OMP_CLAUSE_TO:
- case OMP_CLAUSE_FROM:
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -8797,8 +8798,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
- case OMP_CLAUSE_TO:
- case OMP_CLAUSE_FROM:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
if (!DECL_P (ovar))
@@ -8893,12 +8892,6 @@ lower_oacc_parallel (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_MAP:
tkind = OMP_CLAUSE_MAP_KIND (c);
break;
- case OMP_CLAUSE_TO:
- tkind = OMP_CLAUSE_MAP_TO;
- break;
- case OMP_CLAUSE_FROM:
- tkind = OMP_CLAUSE_MAP_FROM;
- break;
default:
gcc_unreachable ();
}
@@ -10179,6 +10172,22 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
+#ifdef ENABLE_CHECKING
+ /* First check what we're prepared to handle in the following. */
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case OMP_CLAUSE_MAP_ALLOC:
+ case OMP_CLAUSE_MAP_TO:
+ case OMP_CLAUSE_MAP_FROM:
+ case OMP_CLAUSE_MAP_TOFROM:
+ case OMP_CLAUSE_MAP_POINTER:
+ break;
+ default:
+ gcc_unreachable ();
+ }
+#endif
+ /* FALLTHRU */
+
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
var = OMP_CLAUSE_DECL (c);
diff --git gcc/tree-core.h gcc/tree-core.h
index 3602b5f..0aedea3 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -1125,7 +1125,24 @@ enum omp_clause_map_kind
/* The following kind is an internal only map kind, used for pointer based
array sections. OMP_CLAUSE_SIZE for these is not the pointer size,
which is implicitly POINTER_SIZE / BITS_PER_UNIT, but the bias. */
- OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL
+ OMP_CLAUSE_MAP_POINTER = OMP_CLAUSE_MAP_SPECIAL,
+ /* The following are only valid for OpenACC. */
+ /* Flag to force a specific behavior (or else, a run-time error). */
+ OMP_CLAUSE_MAP_FORCE = 1 << 3,
+ /* Allocate. */
+ OMP_CLAUSE_MAP_FORCE_ALLOC = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_ALLOC,
+ /* ..., and copy to device. */
+ OMP_CLAUSE_MAP_FORCE_TO = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TO,
+ /* ..., and copy from device. */
+ OMP_CLAUSE_MAP_FORCE_FROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_FROM,
+ /* ..., and copy to and from device. */
+ OMP_CLAUSE_MAP_FORCE_TOFROM = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_TOFROM,
+ /* Must already be present. */
+ OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL,
+ /* Deallocate a mapping, without copying from device. */
+ OMP_CLAUSE_MAP_FORCE_DEALLOC,
+ /* Is a device pointer. */
+ OMP_CLAUSE_MAP_FORCE_DEVICEPTR
};
enum omp_clause_proc_bind_kind
diff --git gcc/tree-pretty-print.c gcc/tree-pretty-print.c
index 320c35b..f75f181 100644
--- gcc/tree-pretty-print.c
+++ gcc/tree-pretty-print.c
@@ -506,6 +506,27 @@ dump_omp_clause (pretty_printer *buffer, tree clause, int spc, int flags)
case OMP_CLAUSE_MAP_TOFROM:
pp_string (buffer, "tofrom");
break;
+ case OMP_CLAUSE_MAP_FORCE_ALLOC:
+ pp_string (buffer, "force_alloc");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_TO:
+ pp_string (buffer, "force_to");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_FROM:
+ pp_string (buffer, "force_from");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_TOFROM:
+ pp_string (buffer, "force_tofrom");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ pp_string (buffer, "force_present");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ pp_string (buffer, "force_dealloc");
+ break;
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ pp_string (buffer, "force_deviceptr");
+ break;
default:
gcc_unreachable ();
}
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4] Initial support for OpenACC data clauses
2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
@ 2014-01-28 9:44 ` Thomas Schwinge
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC deviceptr clause Thomas Schwinge
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC present data clause Thomas Schwinge
3 siblings, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-01-28 9:44 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub
[-- Attachment #1: Type: text/plain, Size: 643 bytes --]
Hi!
On Tue, 14 Jan 2014 07:09:33 -0800, I wrote:
> Here is a patch series that adds initial support for OpenACC data
> clauses. It is not yet complete, but I thought I might as well already
> now strive to get this integrated upstream instead of "hoarding" the
> patches locally.
Committed to gomp-4_0-branch in r207173..8.
> Would it be a good idea to also commit to trunk the (portions of the)
> patches that don't directly relate with OpenACC stuff? That way, trunk
> and gomp-4_0-branch would diverge a little less? Or, would you first
> like to see all of this stabilitize on gomp-4_0-branch?
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 5/6] Initial support in the C front end for OpenACC data clauses.
2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas
2014-01-14 15:10 ` [gomp4 6/6] Enable initial " thomas
@ 2014-02-12 11:17 ` Thomas Schwinge
1 sibling, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-02-12 11:17 UTC (permalink / raw)
To: gcc-patches, Ilmir Usmanov; +Cc: jakub
[-- Attachment #1: Type: text/plain, Size: 6906 bytes --]
Hi!
On Tue, 14 Jan 2014 16:10:07 +0100, I wrote:
> gcc/c-family/
> * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY,
> PRAGMA_OMP_CLAUSE_COPYOUT, PRAGMA_OMP_CLAUSE_CREATE,
> PRAGMA_OMP_CLAUSE_DELETE, PRAGMA_OMP_CLAUSE_DEVICEPTR,
> PRAGMA_OMP_CLAUSE_PRESENT, PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
> PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
> PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and
> PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE.
> gcc/c/
> * c-parser.c (c_parser_omp_clause_name): Handle these.
> (c_parser_oacc_data_clause, c_parser_oacc_data_clause_deviceptr):
> New functions.
> (c_parser_oacc_all_clauses): Handle PRAGMA_OMP_CLAUSE_COPY,
> PRAGMA_OMP_CLAUSE_COPYIN, PRAGMA_OMP_CLAUSE_COPYOUT,
> PRAGMA_OMP_CLAUSE_CREATE, PRAGMA_OMP_CLAUSE_DELETE,
> PRAGMA_OMP_CLAUSE_DEVICEPTR, PRAGMA_OMP_CLAUSE_PRESENT,
> PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY,
> PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN,
> PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT, and
> PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE.
> gcc/
> * tree-core.h (omp_clause_code): Update description for
> OMP_CLAUSE_MAP.
This I committed to gomp-4_0-branch as r207177.
In
<http://news.gmane.org/find-root.php?message_id=%3C52E68AE3.9030706%40samsung.com%3E>,
Ilmir mentioned that I'm missing to handle the »short names: pcopy,
pcopyin, pcopyout and pcreate (see 2.6.5.9 - 2.6.5.12 of OpenACC 2.0«.
Unless there are any comments, I'll soon commit the following to
gomp-4_0-branch:
commit 9a1f6c075f6198c9ae3281387b875e6012e4387e
Author: Thomas Schwinge <thomas@codesourcery.com>
Date: Wed Feb 12 11:59:51 2014 +0100
OpenACC: pcopy, pcopyin, pcopyout, pcreate clauses.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin,
pcopyout, pcreate clauses.
(c_parser_oacc_data_clause): Update comment.
gcc/
* tree-core.h (omp_clause_code) <map>: Mention pcopy, pcopyin,
pcopyout, pcreate OpenACC clauses.
gcc/testsuite/
* c-c++-common/goacc/pcopy.c: New file.
* c-c++-common/goacc/pcopyin.c: Likewise.
* c-c++-common/goacc/pcopyout.c: Likewise.
* c-c++-common/goacc/pcreate.c: Likewise.
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 6e89471..f401cef 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -9671,13 +9671,17 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_PARALLEL;
else if (!strcmp ("present", p))
result = PRAGMA_OMP_CLAUSE_PRESENT;
- else if (!strcmp ("present_or_copy", p))
+ else if (!strcmp ("present_or_copy", p)
+ || !strcmp ("pcopy", p))
result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY;
- else if (!strcmp ("present_or_copyin", p))
+ else if (!strcmp ("present_or_copyin", p)
+ || !strcmp ("pcopyin", p))
result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN;
- else if (!strcmp ("present_or_copyout", p))
+ else if (!strcmp ("present_or_copyout", p)
+ || !strcmp ("pcopyout", p))
result = PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT;
- else if (!strcmp ("present_or_create", p))
+ else if (!strcmp ("present_or_create", p)
+ || !strcmp ("pcreate", p))
result = PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE;
else if (!strcmp ("private", p))
result = PRAGMA_OMP_CLAUSE_PRIVATE;
@@ -9870,9 +9874,13 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
delete ( variable-list )
present ( variable-list )
present_or_copy ( variable-list )
+ pcopy ( variable-list )
present_or_copyin ( variable-list )
+ pcopyin ( variable-list )
present_or_copyout ( variable-list )
- present_or_create ( variable-list ) */
+ pcopyout ( variable-list )
+ present_or_create ( variable-list )
+ pcreate ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
diff --git gcc/testsuite/c-c++-common/goacc/pcopy.c gcc/testsuite/c-c++-common/goacc/pcopy.c
new file mode 100644
index 0000000..fd16525
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/pcopy.c
@@ -0,0 +1,11 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+f (char *cp)
+{
+#pragma acc parallel pcopy(cp[3:5])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(tofrom:\\*\\(cp \\+ 3\\) \\\[len: 5]\\) map\\(alloc:cp \\\[pointer assign, bias: 3]\\)" 1 "original" } } */
+/* { dg-final { cleanup-tree-dump "original" } } */
diff --git gcc/testsuite/c-c++-common/goacc/pcopyin.c gcc/testsuite/c-c++-common/goacc/pcopyin.c
new file mode 100644
index 0000000..c009d24
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/pcopyin.c
@@ -0,0 +1,11 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+f (char *cp)
+{
+#pragma acc parallel pcopyin(cp[4:6])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(to:\\*\\(cp \\+ 4\\) \\\[len: 6]\\) map\\(alloc:cp \\\[pointer assign, bias: 4]\\)" 1 "original" } } */
+/* { dg-final { cleanup-tree-dump "original" } } */
diff --git gcc/testsuite/c-c++-common/goacc/pcopyout.c gcc/testsuite/c-c++-common/goacc/pcopyout.c
new file mode 100644
index 0000000..6099eff
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/pcopyout.c
@@ -0,0 +1,11 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+f (char *cp)
+{
+#pragma acc parallel pcopyout(cp[5:7])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(from:\\*\\(cp \\+ 5\\) \\\[len: 7]\\) map\\(alloc:cp \\\[pointer assign, bias: 5]\\)" 1 "original" } } */
+/* { dg-final { cleanup-tree-dump "original" } } */
diff --git gcc/testsuite/c-c++-common/goacc/pcreate.c gcc/testsuite/c-c++-common/goacc/pcreate.c
new file mode 100644
index 0000000..2f6e836
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/pcreate.c
@@ -0,0 +1,11 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+f (char *cp)
+{
+#pragma acc parallel pcreate(cp[6:8])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(alloc:\\*\\(cp \\+ 6\\) \\\[len: 8]\\) map\\(alloc:cp \\\[pointer assign, bias: 6]\\)" 1 "original" } } */
+/* { dg-final { cleanup-tree-dump "original" } } */
diff --git gcc/tree-core.h gcc/tree-core.h
index a5a95cd..2d9bf0c 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -259,8 +259,9 @@ enum omp_clause_code {
OMP_CLAUSE_TO,
/* OpenACC clauses: {copy, copyin, copyout, create, delete, deviceptr,
- present, present_or_copy, present_or_copyin, present_or_copyout,
- present_or_create} (variable-list).
+ present, present_or_copy (pcopy), present_or_copyin (pcopyin),
+ present_or_copyout (pcopyout), present_or_create (pcreate)}
+ (variable-list).
OpenMP clause: map ({alloc:,to:,from:,tofrom:,}variable-list). */
OMP_CLAUSE_MAP,
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.
2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas
2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas
@ 2014-02-21 19:48 ` Thomas Schwinge
2014-02-21 20:32 ` [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE Thomas Schwinge
[not found] ` <538DF785.3050206@mentor.com>
2 siblings, 1 reply; 22+ messages in thread
From: Thomas Schwinge @ 2014-02-21 19:48 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub
[-- Attachment #1: Type: text/plain, Size: 8797 bytes --]
Hi!
On Tue, 14 Jan 2014 16:10:05 +0100, I wrote:
> --- gcc/gimplify.c
> +++ gcc/gimplify.c
> @@ -86,7 +92,11 @@ enum omp_region_type
> ORT_UNTIED_TASK = 5,
> ORT_TEAMS = 8,
> ORT_TARGET_DATA = 16,
> - ORT_TARGET = 32
> + ORT_TARGET = 32,
> +
> + /* Flags for ORT_TARGET. */
> + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
> + ORT_TARGET_MAP_FORCE = 64
> };
Continuing on that route, I have now applied the following to
gomp-4_0-branch in r208014:
commit dee2965ae547af0bc90d618e7fa40fbf2f5292b4
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri Feb 21 19:45:12 2014 +0000
Gimplification: New flag ORT_TARGET_OFFLOAD replaces !ORT_TARGET_DATA.
gcc/
* gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a
flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA.
Update all users.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208014 138bc75d-0d04-0410-961f-82ee72b054a4
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 1ce952d..bf8ec96 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,9 @@
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+ * gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a
+ flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA.
+ Update all users.
+
* omp-low.c (gimple_code_is_oacc): Move to...
* gimple.h (is_gimple_omp_oacc_specifically): ... here. Update
users, and also use it in more places where currently we've only
diff --git gcc/gimplify.c gcc/gimplify.c
index 51a1b73..9aa9301c 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -100,10 +100,11 @@ enum omp_region_type
ORT_TASK = 4,
ORT_UNTIED_TASK = 5,
ORT_TEAMS = 8,
- ORT_TARGET_DATA = 16,
- ORT_TARGET = 32,
+ ORT_TARGET = 16,
/* Flags for ORT_TARGET. */
+ /* Prepare this region for offloading. */
+ ORT_TARGET_OFFLOAD = 32,
/* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
ORT_TARGET_MAP_FORCE = 64
};
@@ -2202,7 +2203,7 @@ gimplify_arg (tree *arg_p, gimple_seq *pre_p, location_t call_location)
return gimplify_expr (arg_p, pre_p, NULL, test, fb);
}
-/* Don't fold STMT inside ORT_TARGET, because it can break code by adding decl
+/* Don't fold inside offloading regsion: it can break code by adding decl
references that weren't in the source. We'll do it during omplower pass
instead. */
@@ -2211,7 +2212,8 @@ maybe_fold_stmt (gimple_stmt_iterator *gsi)
{
struct gimplify_omp_ctx *ctx;
for (ctx = gimplify_omp_ctxp; ctx; ctx = ctx->outer_context)
- if (ctx->region_type & ORT_TARGET)
+ if (ctx->region_type & ORT_TARGET
+ && ctx->region_type & ORT_TARGET_OFFLOAD)
return false;
return fold_stmt (gsi);
}
@@ -5388,10 +5390,12 @@ omp_firstprivatize_variable (struct gimplify_omp_ctx *ctx, tree decl)
return;
}
else if (ctx->region_type & ORT_TARGET)
- omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+ {
+ if (ctx->region_type & ORT_TARGET_OFFLOAD)
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_MAP_TO_ONLY);
+ }
else if (ctx->region_type != ORT_WORKSHARE
- && ctx->region_type != ORT_SIMD
- && ctx->region_type != ORT_TARGET_DATA)
+ && ctx->region_type != ORT_SIMD)
omp_add_variable (ctx, decl, GOVD_FIRSTPRIVATE);
ctx = ctx->outer_context;
@@ -5580,7 +5584,8 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
struct gimplify_omp_ctx *octx;
for (octx = ctx; octx; octx = octx->outer_context)
- if (octx->region_type & ORT_TARGET)
+ if ((octx->region_type & ORT_TARGET)
+ && (octx->region_type & ORT_TARGET_OFFLOAD))
{
gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE));
@@ -5643,7 +5648,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
}
n = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- if (ctx->region_type & ORT_TARGET)
+ if ((ctx->region_type & ORT_TARGET)
+ && (ctx->region_type & ORT_TARGET_OFFLOAD))
{
unsigned map_force;
if (ctx->region_type & ORT_TARGET_MAP_FORCE)
@@ -5695,7 +5701,8 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if (ctx->region_type == ORT_WORKSHARE
|| ctx->region_type == ORT_SIMD
- || ctx->region_type == ORT_TARGET_DATA)
+ || ((ctx->region_type & ORT_TARGET)
+ && !(ctx->region_type & ORT_TARGET_OFFLOAD)))
goto do_outer;
/* ??? Some compiler-generated variables (like SAVE_EXPRs) could be
@@ -5746,7 +5753,7 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
{
splay_tree_node n2;
- if ((octx->region_type & (ORT_TARGET_DATA | ORT_TARGET)) != 0)
+ if (octx->region_type & ORT_TARGET)
continue;
n2 = splay_tree_lookup (octx->variables, (splay_tree_key) decl);
if (n2 && (n2->value & GOVD_DATA_SHARE_CLASS) != GOVD_SHARED)
@@ -5899,7 +5906,7 @@ omp_check_private (struct gimplify_omp_ctx *ctx, tree decl, bool copyprivate)
|| (!copyprivate
&& lang_hooks.decls.omp_privatize_by_reference (decl)));
- if ((ctx->region_type & (ORT_TARGET | ORT_TARGET_DATA)) != 0)
+ if (ctx->region_type & ORT_TARGET)
continue;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
@@ -6456,7 +6463,9 @@ gimplify_adjust_omp_clauses (tree *list_p)
if (!DECL_P (decl))
break;
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
- if ((ctx->region_type & ORT_TARGET) && !(n->value & GOVD_SEEN))
+ if ((ctx->region_type & ORT_TARGET)
+ && (ctx->region_type & ORT_TARGET_OFFLOAD)
+ && !(n->value & GOVD_SEEN))
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -6574,8 +6583,9 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple g;
gimple_seq body = NULL;
- enum omp_region_type ort =
- (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE);
+ enum omp_region_type ort = (enum omp_region_type) (ORT_TARGET
+ | ORT_TARGET_OFFLOAD
+ | ORT_TARGET_MAP_FORCE);
gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort);
@@ -7031,11 +7041,11 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
case OMP_SINGLE:
break;
case OMP_TARGET:
+ ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
+ break;
+ case OMP_TARGET_DATA:
ort = ORT_TARGET;
break;
- case OMP_TARGET_DATA:
- ort = ORT_TARGET_DATA;
- break;
case OMP_TEAMS:
ort = ORT_TEAMS;
break;
@@ -7043,7 +7053,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
gcc_unreachable ();
}
gimplify_scan_omp_clauses (&OMP_CLAUSES (expr), pre_p, ort);
- if ((ort & ORT_TARGET) || ort == ORT_TARGET_DATA)
+ if (ort & ORT_TARGET)
{
push_gimplify_context ();
gimple g = gimplify_and_return_first (OMP_BODY (expr), &body);
@@ -7051,7 +7061,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (g);
else
pop_gimplify_context (NULL);
- if (ort == ORT_TARGET_DATA)
+ if (!(ort & ORT_TARGET_OFFLOAD))
{
gimple_seq cleanup = NULL;
tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
@@ -8697,7 +8707,9 @@ gimplify_body (tree fndecl, bool do_parms)
{
gcc_assert (gimplify_omp_ctxp == NULL);
if (lookup_attribute ("omp declare target", DECL_ATTRIBUTES (fndecl)))
- gimplify_omp_ctxp = new_omp_context (ORT_TARGET);
+ gimplify_omp_ctxp
+ = new_omp_context ((enum omp_region_type) (ORT_TARGET
+ | ORT_TARGET_OFFLOAD));
}
/* Unshare most shared trees in the body and in that of any nested functions.
diff --git gcc/omp-low.c gcc/omp-low.c
index b975dad..9fef4c1 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -10858,8 +10858,8 @@ lower_omp (gimple_seq *body, omp_context *ctx)
gimple_stmt_iterator gsi;
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
lower_omp_1 (&gsi, ctx);
- /* During gimplification, we have not always invoked fold_stmt
- (gimplify.c:maybe_fold_stmt); call it now. */
+ /* During gimplification, we haven't folded statments inside offloading
+ regions (gimplify.c:maybe_fold_stmt); do that now. */
if (target_nesting_level)
for (gsi = gsi_start (*body); !gsi_end_p (gsi); gsi_next (&gsi))
fold_stmt (&gsi);
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 3/3] OpenACC data construct support in the C front end.
2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
@ 2014-02-21 20:32 ` Thomas Schwinge
2014-03-12 13:48 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-03-20 14:39 ` [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) Thomas Schwinge
2 siblings, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-02-21 20:32 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/c-family/
* c-pragma.c (oacc_pragmas): Add "data".
* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA.
gcc/c/
* c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition.
(c_parser_oacc_data): New function.
(c_parser_omp_construct): Handle PRAGMA_OACC_DATA.
* c-tree.h (c_finish_oacc_data): New prototype.
* c-typeck.c (c_finish_oacc_data): New function.
gcc/testsuite/
* c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC
data construct.
* c-c++-common/goacc/nesting-fail-1.c: Likewise.
* c-c++-common/goacc/parallel-fail-1.c: Rename to...
* c-c++-common/goacc/clauses-fail.c: ... this new file. Extend
for OpenACC data construct.
* c-c++-common/goacc/data-1.c: New file.
libgomp/
* testsuite/libgomp.oacc-c/data-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208017 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/c-family/ChangeLog.gomp | 5 +
gcc/c-family/c-pragma.c | 1 +
gcc/c-family/c-pragma.h | 1 +
gcc/c/ChangeLog.gomp | 8 +
gcc/c/c-parser.c | 42 +++++
gcc/c/c-tree.h | 1 +
gcc/c/c-typeck.c | 19 +++
gcc/testsuite/ChangeLog.gomp | 10 ++
.../c-c++-common/goacc-gomp/nesting-fail-1.c | 92 ++++++++++-
gcc/testsuite/c-c++-common/goacc/clauses-fail.c | 9 ++
gcc/testsuite/c-c++-common/goacc/data-1.c | 6 +
gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c | 18 ++-
gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c | 6 -
libgomp/ChangeLog.gomp | 2 +
libgomp/testsuite/libgomp.oacc-c/data-1.c | 170 +++++++++++++++++++++
15 files changed, 380 insertions(+), 10 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/clauses-fail.c
create mode 100644 gcc/testsuite/c-c++-common/goacc/data-1.c
delete mode 100644 gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c/data-1.c
diff --git gcc/c-family/ChangeLog.gomp gcc/c-family/ChangeLog.gomp
index e092d53..3da377f 100644
--- gcc/c-family/ChangeLog.gomp
+++ gcc/c-family/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-pragma.c (oacc_pragmas): Add "data".
+ * c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DATA.
+
2014-01-28 Thomas Schwinge <thomas@codesourcery.com>
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_COPY,
diff --git gcc/c-family/c-pragma.c gcc/c-family/c-pragma.c
index f69486a..08374aa 100644
--- gcc/c-family/c-pragma.c
+++ gcc/c-family/c-pragma.c
@@ -1169,6 +1169,7 @@ static vec<pragma_ns_name> registered_pp_pragmas;
struct omp_pragma_def { const char *name; unsigned int id; };
static const struct omp_pragma_def oacc_pragmas[] = {
+ { "data", PRAGMA_OACC_DATA },
{ "parallel", PRAGMA_OACC_PARALLEL },
};
static const struct omp_pragma_def omp_pragmas[] = {
diff --git gcc/c-family/c-pragma.h gcc/c-family/c-pragma.h
index 1ea5b1d..d092f9f 100644
--- gcc/c-family/c-pragma.h
+++ gcc/c-family/c-pragma.h
@@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see
typedef enum pragma_kind {
PRAGMA_NONE = 0,
+ PRAGMA_OACC_DATA,
PRAGMA_OACC_PARALLEL,
PRAGMA_OMP_ATOMIC,
PRAGMA_OMP_BARRIER,
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index b199957..9b95725 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,11 @@
+2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.c (OACC_DATA_CLAUSE_MASK): New macro definition.
+ (c_parser_oacc_data): New function.
+ (c_parser_omp_construct): Handle PRAGMA_OACC_DATA.
+ * c-tree.h (c_finish_oacc_data): New prototype.
+ * c-typeck.c (c_finish_oacc_data): New function.
+
2014-02-17 Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c (c_parser_omp_clause_name): Accept pcopy, pcopyin,
diff --git gcc/c/c-parser.c gcc/c/c-parser.c
index 7850eab..4643722 100644
--- gcc/c/c-parser.c
+++ gcc/c/c-parser.c
@@ -4776,10 +4776,14 @@ c_parser_label (c_parser *parser)
openacc-construct:
parallel-construct
+ data-construct
parallel-construct:
parallel-directive structured-block
+ data-construct:
+ data-directive structured-block
+
OpenMP:
statement:
@@ -11362,6 +11366,41 @@ c_parser_omp_structured_block (c_parser *parser)
}
/* OpenACC 2.0:
+ # pragma acc data oacc-data-clause[optseq] new-line
+ structured-block
+
+ LOC is the location of the #pragma token.
+*/
+
+#define OACC_DATA_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_PRESENT_OR_CREATE) )
+
+static tree
+c_parser_oacc_data (location_t loc, c_parser *parser)
+{
+ tree stmt, clauses, block;
+
+ clauses = c_parser_oacc_all_clauses (parser, OACC_DATA_CLAUSE_MASK,
+ "#pragma acc data");
+
+ block = c_begin_omp_parallel ();
+ add_stmt (c_parser_omp_structured_block (parser));
+
+ stmt = c_finish_oacc_data (loc, clauses, block);
+
+ return stmt;
+}
+
+/* OpenACC 2.0:
# pragma acc parallel oacc-parallel-clause[optseq] new-line
structured-block
@@ -13675,6 +13714,9 @@ c_parser_omp_construct (c_parser *parser)
switch (p_kind)
{
+ case PRAGMA_OACC_DATA:
+ stmt = c_parser_oacc_data (loc, parser);
+ break;
case PRAGMA_OACC_PARALLEL:
stmt = c_parser_oacc_parallel (loc, parser);
break;
diff --git gcc/c/c-tree.h gcc/c/c-tree.h
index c174c7a..c84d3d7 100644
--- gcc/c/c-tree.h
+++ gcc/c/c-tree.h
@@ -634,6 +634,7 @@ extern tree c_finish_goto_label (location_t, tree);
extern tree c_finish_goto_ptr (location_t, tree);
extern tree c_expr_to_decl (tree, bool *, bool *);
extern tree c_finish_oacc_parallel (location_t, tree, tree);
+extern tree c_finish_oacc_data (location_t, tree, tree);
extern tree c_begin_omp_parallel (void);
extern tree c_finish_omp_parallel (location_t, tree, tree);
extern tree c_begin_omp_task (void);
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index 76d655b..8c4445b 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11122,6 +11122,25 @@ c_finish_oacc_parallel (location_t loc, tree clauses, tree block)
return add_stmt (stmt);
}
+/* Generate OACC_DATA, with CLAUSES and BLOCK as its compound
+ statement. LOC is the location of the OACC_DATA. */
+
+tree
+c_finish_oacc_data (location_t loc, tree clauses, tree block)
+{
+ tree stmt;
+
+ block = c_end_compound_stmt (loc, block, true);
+
+ stmt = make_node (OACC_DATA);
+ TREE_TYPE (stmt) = void_type_node;
+ OACC_DATA_CLAUSES (stmt) = clauses;
+ OACC_DATA_BODY (stmt) = block;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ return add_stmt (stmt);
+}
+
/* Like c_begin_compound_stmt, except force the retention of the BLOCK. */
tree
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index fbccfa3..41d73b6 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,13 @@
+2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc-gomp/nesting-fail-1.c: Extend for OpenACC
+ data construct.
+ * c-c++-common/goacc/nesting-fail-1.c: Likewise.
+ * c-c++-common/goacc/parallel-fail-1.c: Rename to...
+ * c-c++-common/goacc/clauses-fail.c: ... this new file. Extend
+ for OpenACC data construct.
+ * c-c++-common/goacc/data-1.c: New file.
+
2014-02-18 Thomas Schwinge <thomas@codesourcery.com>
* gcc.dg/goacc/parallel-sb-1.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
index 875ec66..78fb45b 100644
--- gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c
@@ -1,7 +1,7 @@
/* TODO: Some of these should either be allowed or fail with a more sensible
error message. */
void
-f1 (void)
+f_omp (void)
{
int i;
@@ -9,6 +9,8 @@ f1 (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp for
@@ -16,49 +18,68 @@ f1 (void)
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp sections
{
+ {
#pragma acc parallel /* { dg-error "may not be nested" } */
- ;
+ ;
+ }
+#pragma omp section
+ {
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
+ }
}
#pragma omp single
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp task
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp master
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp critical
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
#pragma omp ordered
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
}
/* TODO: Some of these should either be allowed or fail with a more sensible
error message. */
void
-f2 (void)
+f_acc_parallel (void)
{
#pragma acc parallel
{
@@ -119,3 +140,68 @@ f2 (void)
;
}
}
+
+/* TODO: Some of these should either be allowed or fail with a more sensible
+ error message. */
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma omp parallel /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+ int i;
+#pragma omp for /* { dg-error "may not be nested" } */
+ for (i = 0; i < 3; i++)
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp sections /* { dg-error "may not be nested" } */
+ {
+ ;
+ }
+ }
+
+#pragma acc data
+ {
+#pragma omp single /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp task /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp master /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+#pragma omp critical /* { dg-error "may not be nested" } */
+ ;
+ }
+
+#pragma acc data
+ {
+ int i;
+#pragma omp atomic write
+ i = 0; /* { dg-error "may not be nested" } */
+ }
+
+#pragma acc data
+ {
+#pragma omp ordered /* { dg-error "may not be nested" } */
+ ;
+ }
+}
diff --git gcc/testsuite/c-c++-common/goacc/clauses-fail.c gcc/testsuite/c-c++-common/goacc/clauses-fail.c
new file mode 100644
index 0000000..b0dd042
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/clauses-fail.c
@@ -0,0 +1,9 @@
+void
+f (void)
+{
+#pragma acc parallel one /* { dg-error "expected clause before 'one'" } */
+ ;
+
+#pragma acc data two /* { dg-error "expected clause before 'two'" } */
+ ;
+}
diff --git gcc/testsuite/c-c++-common/goacc/data-1.c gcc/testsuite/c-c++-common/goacc/data-1.c
new file mode 100644
index 0000000..8094575
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/data-1.c
@@ -0,0 +1,6 @@
+void
+foo (void)
+{
+#pragma acc data
+ ;
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index 6501397..24a4c11 100644
--- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -1,11 +1,27 @@
/* TODO: While the OpenACC specification does allow for certain kinds of
nesting, we don't support that yet. */
void
-f1 (void)
+f_acc_parallel (void)
{
#pragma acc parallel
{
#pragma acc parallel /* { dg-error "may not be nested" } */
;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
+ }
+}
+
+/* TODO: While the OpenACC specification does allow for certain kinds of
+ nesting, we don't support that yet. */
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma acc parallel /* { dg-error "may not be nested" } */
+ ;
+#pragma acc data /* { dg-error "may not be nested" } */
+ ;
}
}
diff --git gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
deleted file mode 100644
index efc6f14..0000000
--- gcc/testsuite/c-c++-common/goacc/parallel-fail-1.c
+++ /dev/null
@@ -1,6 +0,0 @@
-void
-foo (void)
-{
-#pragma acc parallel foo /* { dg-error "expected clause before 'foo'" } */
- foo ();
-}
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 5c15656..b90b09b 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,5 +1,7 @@
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+ * testsuite/libgomp.oacc-c/data-1.c: New file.
+
* libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start.
* libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes.
* oacc-parallel.c (GOACC_data_start, GOACC_data_end): New
diff --git libgomp/testsuite/libgomp.oacc-c/data-1.c libgomp/testsuite/libgomp.oacc-c/data-1.c
new file mode 100644
index 0000000..8f9a17a
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c/data-1.c
@@ -0,0 +1,170 @@
+/* { dg-do run } */
+
+extern void abort ();
+
+int i;
+
+int main(void)
+{
+ int j;
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data copyin (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != 2 || j != 1)
+ abort ();
+
+ i = -1;
+ j = -2;
+#pragma acc data copyout (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+
+ i = -1;
+ j = -2;
+#pragma acc data copy (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+
+ i = -1;
+ j = -2;
+#pragma acc data create (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+ i = -1;
+ j = -2;
+#pragma acc data present_or_copyin (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != 2 || j != 1)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data present_or_copyout (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+ i = -1;
+ j = -2;
+#pragma acc data present_or_copy (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data present_or_create (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data present (i, j)
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+#if 0
+ i = -1;
+ j = -2;
+#pragma acc data
+ {
+ // TODO: check that variables have been mapped.
+ if (i != -1 || j != -2)
+ abort ();
+ i = 2;
+ j = 1;
+ if (i != 2 || j != 1)
+ abort ();
+ }
+ if (i != -1 || j != -2)
+ abort ();
+#endif
+
+ return 0;
+}
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.
2014-02-21 20:32 ` [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE Thomas Schwinge
@ 2014-02-21 20:32 ` Thomas Schwinge
2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge
` (2 more replies)
0 siblings, 3 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-02-21 20:32 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/
* gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA.
(is_gimple_omp_oacc_specifically): Handle it.
* gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
* gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise.
* omp-low.c (scan_sharing_clauses, scan_omp_target)
(expand_omp_target, lower_omp_target, lower_omp_1): Likewise.
* gimple.def (GIMPLE_OMP_TARGET): Update comment.
* gimple.c (gimple_build_omp_target): Likewise.
(gimple_copy): Catch unimplemented case.
* tree-inline.c (remap_gimple_stmt): Likewise.
* tree-nested.c (convert_nonlocal_reference_stmt)
(convert_local_reference_stmt, convert_gimple_call): Likewise.
* oacc-builtins.def (BUILT_IN_GOACC_DATA_START)
(BUILT_IN_GOACC_DATA_END): New builtins.
libgomp/
* libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start.
* libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes.
* oacc-parallel.c (GOACC_data_start, GOACC_data_end): New
functions.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208016 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 15 ++++++
gcc/gimple-pretty-print.c | 3 ++
gcc/gimple.c | 4 +-
gcc/gimple.def | 1 +
gcc/gimple.h | 9 ++++
gcc/gimplify.c | 33 +++++++++---
gcc/oacc-builtins.def | 6 ++-
gcc/omp-low.c | 132 ++++++++++++++++++++++++++++++++++++----------
gcc/tree-inline.c | 1 +
gcc/tree-nested.c | 3 ++
libgomp/ChangeLog.gomp | 7 +++
libgomp/libgomp.map | 2 +
libgomp/libgomp_g.h | 3 ++
libgomp/oacc-parallel.c | 34 +++++++++++-
14 files changed, 213 insertions(+), 40 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index bd46f2e..824ec94 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,20 @@
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+ * gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DATA.
+ (is_gimple_omp_oacc_specifically): Handle it.
+ * gimple-pretty-print.c (dump_gimple_omp_target): Likewise.
+ * gimplify.c (gimplify_omp_workshare, gimplify_expr): Likewise.
+ * omp-low.c (scan_sharing_clauses, scan_omp_target)
+ (expand_omp_target, lower_omp_target, lower_omp_1): Likewise.
+ * gimple.def (GIMPLE_OMP_TARGET): Update comment.
+ * gimple.c (gimple_build_omp_target): Likewise.
+ (gimple_copy): Catch unimplemented case.
+ * tree-inline.c (remap_gimple_stmt): Likewise.
+ * tree-nested.c (convert_nonlocal_reference_stmt)
+ (convert_local_reference_stmt, convert_gimple_call): Likewise.
+ * oacc-builtins.def (BUILT_IN_GOACC_DATA_START)
+ (BUILT_IN_GOACC_DATA_END): New builtins.
+
* omp-low.c (scan_sharing_clauses): Catch unexpected occurrences
of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP.
diff --git gcc/gimple-pretty-print.c gcc/gimple-pretty-print.c
index 91a3eb2..ad9369c 100644
--- gcc/gimple-pretty-print.c
+++ gcc/gimple-pretty-print.c
@@ -1289,6 +1289,9 @@ dump_gimple_omp_target (pretty_printer *buffer, gimple gs, int spc, int flags)
case GF_OMP_TARGET_KIND_UPDATE:
kind = " update";
break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ kind = " oacc_data";
+ break;
default:
gcc_unreachable ();
}
diff --git gcc/gimple.c gcc/gimple.c
index 2a967aa..30561b1 100644
--- gcc/gimple.c
+++ gcc/gimple.c
@@ -1051,7 +1051,8 @@ gimple_build_omp_single (gimple_seq body, tree clauses)
/* Build a GIMPLE_OMP_TARGET statement.
BODY is the sequence of statements that will be executed.
- CLAUSES are any of the OMP target construct's clauses. */
+ KIND is the kind of target region.
+ CLAUSES are any of the construct's clauses. */
gimple
gimple_build_omp_target (gimple_seq body, int kind, tree clauses)
@@ -1747,6 +1748,7 @@ gimple_copy (gimple stmt)
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
copy_omp_body:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
new_seq = gimple_seq_copy (gimple_omp_body (stmt));
gimple_omp_set_body (copy, new_seq);
break;
diff --git gcc/gimple.def gcc/gimple.def
index 2b78c06..ce800bd 100644
--- gcc/gimple.def
+++ gcc/gimple.def
@@ -360,6 +360,7 @@ DEFGSCODE(GIMPLE_OMP_SECTIONS_SWITCH, "gimple_omp_sections_switch", GSS_BASE)
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
+ #pragma acc data
#pragma omp target {,data,update}
BODY is the sequence of statements inside the target construct
(NULL for target update).
diff --git gcc/gimple.h gcc/gimple.h
index 0d250ef..b4ee9fa 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -102,6 +102,7 @@ enum gf_mask {
GF_OMP_TARGET_KIND_REGION = 0 << 0,
GF_OMP_TARGET_KIND_DATA = 1 << 0,
GF_OMP_TARGET_KIND_UPDATE = 2 << 0,
+ GF_OMP_TARGET_KIND_OACC_DATA = 3 << 0,
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
a thread synchronization via some sort of barrier. The exact barrier
@@ -5684,6 +5685,14 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
{
case GIMPLE_OACC_PARALLEL:
return true;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ return true;
+ default:
+ return false;
+ }
default:
return false;
}
diff --git gcc/gimplify.c gcc/gimplify.c
index 9aa9301c..fd4305c 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -7023,9 +7023,7 @@ gimplify_omp_for (tree *expr_p, gimple_seq *pre_p)
return GS_ALL_DONE;
}
-/* Gimplify the gross structure of other OpenMP constructs.
- In particular, OMP_SECTIONS, OMP_SINGLE, OMP_TARGET, OMP_TARGET_DATA
- and OMP_TEAMS. */
+/* Gimplify the gross structure of several OpenACC or OpenMP constructs. */
static void
gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
@@ -7033,12 +7031,17 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
tree expr = *expr_p;
gimple stmt;
gimple_seq body = NULL;
- enum omp_region_type ort = ORT_WORKSHARE;
+ enum omp_region_type ort;
switch (TREE_CODE (expr))
{
+ case OACC_DATA:
+ ort = (enum omp_region_type) (ORT_TARGET
+ | ORT_TARGET_MAP_FORCE);
+ break;
case OMP_SECTIONS:
case OMP_SINGLE:
+ ort = ORT_WORKSHARE;
break;
case OMP_TARGET:
ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
@@ -7063,9 +7066,21 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
pop_gimplify_context (NULL);
if (!(ort & ORT_TARGET_OFFLOAD))
{
- gimple_seq cleanup = NULL;
- tree fn = builtin_decl_explicit (BUILT_IN_GOMP_TARGET_END_DATA);
+ enum built_in_function end_ix;
+ switch (TREE_CODE (expr))
+ {
+ case OACC_DATA:
+ end_ix = BUILT_IN_GOACC_DATA_END;
+ break;
+ case OMP_TARGET_DATA:
+ end_ix = BUILT_IN_GOMP_TARGET_END_DATA;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ tree fn = builtin_decl_explicit (end_ix);
g = gimple_build_call (fn, 0);
+ gimple_seq cleanup = NULL;
gimple_seq_add_stmt (&cleanup, g);
g = gimple_build_try (body, cleanup, GIMPLE_TRY_FINALLY);
body = NULL;
@@ -7078,6 +7093,10 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
+ case OACC_DATA:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_DATA,
+ OACC_DATA_CLAUSES (expr));
+ break;
case OMP_SECTIONS:
stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
break;
@@ -8047,7 +8066,6 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
break;
case OACC_KERNELS:
- case OACC_DATA:
case OACC_HOST_DATA:
case OACC_DECLARE:
case OACC_UPDATE:
@@ -8076,6 +8094,7 @@ gimplify_expr (tree *expr_p, gimple_seq *pre_p, gimple_seq *post_p,
ret = gimplify_omp_for (expr_p, pre_p);
break;
+ case OACC_DATA:
case OMP_SECTIONS:
case OMP_SINGLE:
case OMP_TARGET:
diff --git gcc/oacc-builtins.def gcc/oacc-builtins.def
index a75e42d..eaf3228 100644
--- gcc/oacc-builtins.def
+++ gcc/oacc-builtins.def
@@ -1,7 +1,7 @@
/* This file contains the definitions and documentation for the
OpenACC builtins used in the GNU compiler.
- Copyright (C) 2013 Free Software Foundation, Inc.
+ Copyright (C) 2013-2014 Free Software Foundation, Inc.
Contributed by Thomas Schwinge <thomas@codesourcery.com>.
@@ -29,3 +29,7 @@ along with GCC; see the file COPYING3. If not see
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel",
BT_FN_VOID_INT_OMPFN_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
+ BT_FN_VOID_INT_PTR_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
+DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
+ BT_FN_VOID, ATTR_NOTHROW_LIST)
diff --git gcc/omp-low.c gcc/omp-low.c
index bca4599..6dec687 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree c, decl;
bool scan_array_reductions = false;
+ bool offloaded;
+ switch (gimple_code (ctx->stmt))
+ {
+ case GIMPLE_OACC_PARALLEL:
+ offloaded = true;
+ break;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (ctx->stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ offloaded = true;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ offloaded = false;
+ break;
+ default:
+ gcc_unreachable ();
+ }
+ break;
+ default:
+ offloaded = false;
+ }
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
@@ -1669,11 +1693,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
{
/* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
- #pragma omp target data, there is nothing to map for
+ target regions that are not offloaded; there is nothing to map for
those. */
- if (!gimple_code_is_oacc (ctx->stmt)
- && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA
- && !POINTER_TYPE_P (TREE_TYPE (decl)))
+ if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
if (DECL_P (decl))
@@ -1698,9 +1720,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, true, 7, ctx);
else
install_var_field (decl, true, 3, ctx);
- if (gimple_code (ctx->stmt) == GIMPLE_OACC_PARALLEL
- || (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_REGION))
+ if (offloaded)
install_var_local (decl, ctx);
}
}
@@ -1824,8 +1844,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
|| (gimple_omp_target_kind (ctx->stmt)
!= GF_OMP_TARGET_KIND_UPDATE));
- if (!gimple_code_is_oacc (ctx->stmt)
- && gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
+ if (!offloaded)
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
@@ -2340,7 +2359,7 @@ scan_omp_single (gimple stmt, omp_context *outer_ctx)
layout_type (ctx->record_type);
}
-/* Scan an OpenMP target{, data, update} directive. */
+/* Scan a GIMPLE_OMP_TARGET. */
static void
scan_omp_target (gimple stmt, omp_context *outer_ctx)
@@ -2349,6 +2368,12 @@ scan_omp_target (gimple stmt, omp_context *outer_ctx)
tree name;
int kind = gimple_omp_target_kind (stmt);
+ if (kind == GF_OMP_TARGET_KIND_OACC_DATA)
+ {
+ gcc_assert (taskreg_nesting_level == 0);
+ gcc_assert (target_nesting_level == 0);
+ }
+
ctx = new_omp_context (stmt, outer_ctx);
ctx->field_map = splay_tree_new (splay_tree_compare_pointers, 0, 0);
ctx->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
@@ -8218,7 +8243,7 @@ expand_omp_atomic (struct omp_region *region)
}
-/* Expand the OpenMP target{, data, update} directive starting at REGION. */
+/* Expand the GIMPLE_OMP_TARGET starting at REGION. */
static void
expand_omp_target (struct omp_region *region)
@@ -8401,12 +8426,23 @@ expand_omp_target (struct omp_region *region)
clauses = gimple_omp_target_clauses (entry_stmt);
- if (kind == GF_OMP_TARGET_KIND_REGION)
- start_ix = BUILT_IN_GOMP_TARGET;
- else if (kind == GF_OMP_TARGET_KIND_DATA)
- start_ix = BUILT_IN_GOMP_TARGET_DATA;
- else
- start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ switch (kind)
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ start_ix = BUILT_IN_GOMP_TARGET;
+ break;
+ case GF_OMP_TARGET_KIND_DATA:
+ start_ix = BUILT_IN_GOMP_TARGET_DATA;
+ break;
+ case GF_OMP_TARGET_KIND_UPDATE:
+ start_ix = BUILT_IN_GOMP_TARGET_UPDATE;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ start_ix = BUILT_IN_GOACC_DATA_START;
+ break;
+ default:
+ gcc_unreachable ();
+ }
/* By default, the value of DEVICE is -1 (let runtime library choose)
and there is no conditional. */
@@ -8414,10 +8450,12 @@ expand_omp_target (struct omp_region *region)
device = build_int_cst (integer_type_node, -1);
c = find_omp_clause (clauses, OMP_CLAUSE_IF);
+ gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
if (c)
cond = OMP_CLAUSE_IF_EXPR (c);
c = find_omp_clause (clauses, OMP_CLAUSE_DEVICE);
+ gcc_assert (!c || kind != GF_OMP_TARGET_KIND_OACC_DATA);
if (c)
{
device = OMP_CLAUSE_DEVICE_ID (c);
@@ -8433,6 +8471,7 @@ expand_omp_target (struct omp_region *region)
(cond ? device : -2). */
if (cond)
{
+ gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
cond = gimple_boolify (cond);
basic_block cond_bb, then_bb, else_bb;
@@ -8523,7 +8562,9 @@ expand_omp_target (struct omp_region *region)
gcc_assert (g && gimple_code (g) == GIMPLE_OMP_TARGET);
gsi_remove (&gsi, true);
}
- if (kind == GF_OMP_TARGET_KIND_DATA && region->exit)
+ if ((kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
+ && region->exit)
{
gsi = gsi_last_bb (region->exit);
g = gsi_stmt (gsi);
@@ -10277,7 +10318,7 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
}
-/* Lower the OpenMP target directive in the current statement
+/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
static void
@@ -10298,7 +10339,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tgt_bind = gimple_seq_first_stmt (gimple_omp_body (stmt));
tgt_body = gimple_bind_body (tgt_bind);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
tgt_body = gimple_omp_body (stmt);
child_fn = ctx->cb.dst_fn;
@@ -10322,6 +10364,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_MAP_TOFROM:
case OMP_CLAUSE_MAP_POINTER:
break;
+ case OMP_CLAUSE_MAP_FORCE_ALLOC:
+ case OMP_CLAUSE_MAP_FORCE_TO:
+ case OMP_CLAUSE_MAP_FORCE_FROM:
+ case OMP_CLAUSE_MAP_FORCE_TOFROM:
+ case OMP_CLAUSE_MAP_FORCE_PRESENT:
+ case OMP_CLAUSE_MAP_FORCE_DEALLOC:
+ case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
+ gcc_assert (kind == GF_OMP_TARGET_KIND_OACC_DATA);
+ break;
default:
gcc_unreachable ();
}
@@ -10330,6 +10381,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ gcc_assert (kind != GF_OMP_TARGET_KIND_OACC_DATA);
var = OMP_CLAUSE_DECL (c);
if (!DECL_P (var))
{
@@ -10373,7 +10426,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
lower_omp (&tgt_body, ctx);
target_nesting_level--;
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
lower_omp (&tgt_body, ctx);
if (kind == GF_OMP_TARGET_KIND_REGION)
@@ -10400,9 +10454,25 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
DECL_NAMELESS (TREE_VEC_ELT (t, 1)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 1)) = 1;
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 1;
+ tree tkind_type;
+ int talign_shift;
+ switch (kind)
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ case GF_OMP_TARGET_KIND_DATA:
+ case GF_OMP_TARGET_KIND_UPDATE:
+ tkind_type = unsigned_char_type_node;
+ talign_shift = 3;
+ break;
+ case GF_OMP_TARGET_KIND_OACC_DATA:
+ tkind_type = short_unsigned_type_node;
+ talign_shift = 8;
+ break;
+ default:
+ gcc_unreachable ();
+ }
TREE_VEC_ELT (t, 2)
- = create_tmp_var (build_array_type_nelts (unsigned_char_type_node,
- map_cnt),
+ = create_tmp_var (build_array_type_nelts (tkind_type, map_cnt),
".omp_data_kinds");
DECL_NAMELESS (TREE_VEC_ELT (t, 2)) = 1;
TREE_ADDRESSABLE (TREE_VEC_ELT (t, 2)) = 1;
@@ -10515,7 +10585,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (TREE_CODE (s) != INTEGER_CST)
TREE_STATIC (TREE_VEC_ELT (t, 1)) = 0;
- unsigned char tkind = 0;
+ unsigned HOST_WIDE_INT tkind;
switch (OMP_CLAUSE_CODE (c))
{
case OMP_CLAUSE_MAP:
@@ -10530,14 +10600,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
gcc_unreachable ();
}
- unsigned int talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
+ gcc_assert (tkind < (HOST_WIDE_INT_C (1U) << talign_shift));
+ unsigned HOST_WIDE_INT talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
talign = ceil_log2 (talign);
- tkind |= talign << 3;
+ tkind |= talign << talign_shift;
+ gcc_assert (tkind <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
CONSTRUCTOR_APPEND_ELT (vkind, purpose,
- build_int_cst (unsigned_char_type_node,
- tkind));
+ build_int_cstu (tkind_type, tkind));
if (nc && nc != c)
c = nc;
}
@@ -10589,7 +10660,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_seq_add_seq (&new_body, tgt_body);
new_body = maybe_catch_exception (new_body);
}
- else if (kind == GF_OMP_TARGET_KIND_DATA)
+ else if (kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_OACC_DATA)
new_body = tgt_body;
if (kind != GF_OMP_TARGET_KIND_UPDATE)
{
@@ -10810,6 +10882,8 @@ lower_omp_1 (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GIMPLE_OMP_TARGET:
ctx = maybe_lookup_ctx (stmt);
gcc_assert (ctx);
+ if (gimple_omp_target_kind (stmt) == GF_OMP_TARGET_KIND_OACC_DATA)
+ gcc_assert (!ctx->cancellable);
lower_omp_target (gsi_p, ctx);
break;
case GIMPLE_OMP_TEAMS:
diff --git gcc/tree-inline.c gcc/tree-inline.c
index 99903333..61c1cc8 100644
--- gcc/tree-inline.c
+++ gcc/tree-inline.c
@@ -1397,6 +1397,7 @@ remap_gimple_stmt (gimple stmt, copy_body_data *id)
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
s1 = remap_gimple_seq (gimple_omp_body (stmt), id);
copy = gimple_build_omp_target
(s1, gimple_omp_target_kind (stmt),
diff --git gcc/tree-nested.c gcc/tree-nested.c
index 8933d02..afa7abb 100644
--- gcc/tree-nested.c
+++ gcc/tree-nested.c
@@ -1307,6 +1307,7 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
save_suppress = info->suppress_expansion;
convert_nonlocal_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op,
@@ -1769,6 +1770,7 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
break;
case GIMPLE_OMP_TARGET:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
save_suppress = info->suppress_expansion;
convert_local_omp_clauses (gimple_omp_target_clauses_ptr (stmt), wi);
walk_body (convert_local_reference_stmt, convert_local_reference_op,
@@ -2184,6 +2186,7 @@ convert_gimple_call (gimple_stmt_iterator *gsi, bool *handled_ops_p,
case GIMPLE_OMP_TASKGROUP:
case GIMPLE_OMP_ORDERED:
case GIMPLE_OMP_CRITICAL:
+ gcc_assert (!is_gimple_omp_oacc_specifically (stmt));
walk_body (convert_gimple_call, NULL, info, gimple_omp_body_ptr (stmt));
break;
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 3dffde4..5c15656 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,10 @@
+2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+
+ * libgomp.map (GOACC_2.0): Add GOACC_data_end, GOACC_data_start.
+ * libgomp_g.h (GOACC_data_start, GOACC_data_end): New prototypes.
+ * oacc-parallel.c (GOACC_data_start, GOACC_data_end): New
+ functions.
+
2014-02-20 Thomas Schwinge <thomas@codesourcery.com>
* target.c (gomp_load_plugin_for_device): Don't call dlcose if
diff --git libgomp/libgomp.map libgomp/libgomp.map
index 2b64d05..cb52e45 100644
--- libgomp/libgomp.map
+++ libgomp/libgomp.map
@@ -233,5 +233,7 @@ OACC_2.0 {
GOACC_2.0 {
global:
+ GOACC_data_end;
+ GOACC_data_start;
GOACC_parallel;
};
diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
index 7c24317..b9083a5 100644
--- libgomp/libgomp_g.h
+++ libgomp/libgomp_g.h
@@ -218,5 +218,8 @@ extern void GOMP_teams (unsigned int, unsigned int);
extern void GOACC_parallel (int, void (*) (void *), const void *,
size_t, void **, size_t *, unsigned short *);
+extern void GOACC_data_start (int, const void *,
+ size_t, void **, size_t *, unsigned short *);
+extern void GOACC_data_end (void);
#endif /* LIBGOMP_G_H */
diff --git libgomp/oacc-parallel.c libgomp/oacc-parallel.c
index bf7b74c..3ac7e39 100644
--- libgomp/oacc-parallel.c
+++ libgomp/oacc-parallel.c
@@ -1,4 +1,4 @@
-/* Copyright (C) 2013 Free Software Foundation, Inc.
+/* Copyright (C) 2013-2014 Free Software Foundation, Inc.
Contributed by Thomas Schwinge <thomas@codesourcery.com>.
@@ -23,7 +23,7 @@
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
-/* This file handles the OpenACC parallel construct. */
+/* This file handles the OpenACC data and parallel constructs. */
#include "libgomp.h"
#include "libgomp_g.h"
@@ -51,3 +51,33 @@ GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
}
GOMP_target (device, fn, openmp_target, mapnum, hostaddrs, sizes, kinds_);
}
+
+
+void
+GOACC_data_start (int device, const void *openmp_target, size_t mapnum,
+ void **hostaddrs, size_t *sizes, unsigned short *kinds)
+{
+ unsigned char kinds_[mapnum];
+ size_t i;
+
+ /* TODO. Eventually, we'll be interpreting all mapping kinds according to
+ the OpenACC semantics; for now we're re-using what is implemented for
+ OpenMP. */
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i];
+ unsigned char align = kinds[i] >> 8;
+ if (kind > 4)
+ gomp_fatal ("memory mapping kind %x for %zd is not yet supported",
+ kind, i);
+
+ kinds_[i] = kind | align << 3;
+ }
+ GOMP_target_data (device, openmp_target, mapnum, hostaddrs, sizes, kinds_);
+}
+
+void
+GOACC_data_end (void)
+{
+ GOMP_target_end_data ();
+}
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE.
2014-02-21 19:48 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
@ 2014-02-21 20:32 ` Thomas Schwinge
2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
0 siblings, 1 reply; 22+ messages in thread
From: Thomas Schwinge @ 2014-02-21 20:32 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/
* omp-low.c (scan_sharing_clauses): Catch unexpected occurrences
of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208015 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 3 +++
gcc/omp-low.c | 25 +++++++++++++++++++++++++
2 files changed, 28 insertions(+)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index bf8ec96..bd46f2e 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,8 @@
2014-02-21 Thomas Schwinge <thomas@codesourcery.com>
+ * omp-low.c (scan_sharing_clauses): Catch unexpected occurrences
+ of OMP_CLAUSE_TO, OMP_CLAUSE_FROM, OMP_CLAUSE_MAP.
+
* gimplify.c (enum omp_region_type): Make ORT_TARGET_OFFLOAD a
flag for ORT_TARGET, in its negation replacing ORT_TARGET_DATA.
Update all users.
diff --git gcc/omp-low.c gcc/omp-low.c
index 9fef4c1..bca4599 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1630,6 +1630,26 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_FROM:
gcc_assert (!is_gimple_omp_oacc_specifically (ctx->stmt));
case OMP_CLAUSE_MAP:
+ switch (OMP_CLAUSE_CODE (c))
+ {
+ case OMP_CLAUSE_TO:
+ case OMP_CLAUSE_FROM:
+ /* The to and from clauses are only ever seen with OpenMP target
+ update constructs. */
+ gcc_assert (gimple_code (ctx->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_UPDATE));
+ break;
+ case OMP_CLAUSE_MAP:
+ /* The map clause is never seen with OpenMP target update
+ constructs. */
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
+ || (gimple_omp_target_kind (ctx->stmt)
+ != GF_OMP_TARGET_KIND_UPDATE));
+ break;
+ default:
+ gcc_unreachable ();
+ }
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
decl = OMP_CLAUSE_DECL (c);
@@ -1799,6 +1819,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
break;
case OMP_CLAUSE_MAP:
+ /* The map clause is never seen with OpenMP target update
+ constructs. */
+ gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
+ || (gimple_omp_target_kind (ctx->stmt)
+ != GF_OMP_TARGET_KIND_UPDATE));
if (!gimple_code_is_oacc (ctx->stmt)
&& gimple_omp_target_kind (ctx->stmt) == GF_OMP_TARGET_KIND_DATA)
break;
--
1.8.1.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.
2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge
@ 2014-03-12 13:48 ` Thomas Schwinge
2014-03-20 14:39 ` [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) Thomas Schwinge
2 siblings, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-03-12 13:48 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub
[-- Attachment #1: Type: text/plain, Size: 4687 bytes --]
Hi!
On Fri, 21 Feb 2014 21:32:14 +0100, I wrote:
> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -1499,6 +1499,30 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
> {
> tree c, decl;
> bool scan_array_reductions = false;
> + bool offloaded;
> + switch (gimple_code (ctx->stmt))
> + {
> + case GIMPLE_OACC_PARALLEL:
> + offloaded = true;
> + break;
> + case GIMPLE_OMP_TARGET:
> + switch (gimple_omp_target_kind (ctx->stmt))
> + {
> + case GF_OMP_TARGET_KIND_REGION:
> + offloaded = true;
> + break;
> + case GF_OMP_TARGET_KIND_DATA:
> + case GF_OMP_TARGET_KIND_UPDATE:
> + case GF_OMP_TARGET_KIND_OACC_DATA:
> + offloaded = false;
> + break;
> + default:
> + gcc_unreachable ();
> + }
> + break;
> + default:
> + offloaded = false;
> + }
I now have a need for this information elsewhere; in gomp-4_0-branch
r208513 changed as follows:
commit 326592ef8fe7501f9ba7e67157d68c6c541e5601
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Wed Mar 12 13:40:07 2014 +0000
is_gimple_omp_offloaded.
gcc/
* omp-low.c (scan_sharing_clauses): Move offloaded logic into...
* gimple.h (is_gimple_omp_offloaded): ... this new static inline
function.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208513 138bc75d-0d04-0410-961f-82ee72b054a4
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 79030d6..4ee843f 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2014-03-12 Thomas Schwinge <thomas@codesourcery.com>
+
+ * omp-low.c (scan_sharing_clauses): Move offloaded logic into...
+ * gimple.h (is_gimple_omp_offloaded): ... this new static inline
+ function.
+
2014-02-28 Thomas Schwinge <thomas@codesourcery.com>
* gimple.def (GIMPLE_OACC_KERNELS): New code.
diff --git gcc/gimple.h gcc/gimple.h
index 514af32..910072d 100644
--- gcc/gimple.h
+++ gcc/gimple.h
@@ -5823,6 +5823,31 @@ is_gimple_omp_oacc_specifically (const_gimple stmt)
}
+/* Return true if OMP_* STMT is offloaded. */
+
+static inline bool
+is_gimple_omp_offloaded (const_gimple stmt)
+{
+ gcc_assert (is_gimple_omp (stmt));
+ switch (gimple_code (stmt))
+ {
+ case GIMPLE_OACC_KERNELS:
+ case GIMPLE_OACC_PARALLEL:
+ return true;
+ case GIMPLE_OMP_TARGET:
+ switch (gimple_omp_target_kind (stmt))
+ {
+ case GF_OMP_TARGET_KIND_REGION:
+ return true;
+ default:
+ return false;
+ }
+ default:
+ return false;
+ }
+}
+
+
/* Returns TRUE if statement G is a GIMPLE_NOP. */
static inline bool
diff --git gcc/omp-low.c gcc/omp-low.c
index 2f13fb4..6b676e5 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1499,31 +1499,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree c, decl;
bool scan_array_reductions = false;
- bool offloaded;
- switch (gimple_code (ctx->stmt))
- {
- case GIMPLE_OACC_KERNELS:
- case GIMPLE_OACC_PARALLEL:
- offloaded = true;
- break;
- case GIMPLE_OMP_TARGET:
- switch (gimple_omp_target_kind (ctx->stmt))
- {
- case GF_OMP_TARGET_KIND_REGION:
- offloaded = true;
- break;
- case GF_OMP_TARGET_KIND_DATA:
- case GF_OMP_TARGET_KIND_UPDATE:
- case GF_OMP_TARGET_KIND_OACC_DATA:
- offloaded = false;
- break;
- default:
- gcc_unreachable ();
- }
- break;
- default:
- offloaded = false;
- }
for (c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
@@ -1696,7 +1671,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* Ignore OMP_CLAUSE_MAP_POINTER kind for arrays in
target regions that are not offloaded; there is nothing to map for
those. */
- if (!offloaded && !POINTER_TYPE_P (TREE_TYPE (decl)))
+ if (!is_gimple_omp_offloaded (ctx->stmt)
+ && !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
if (DECL_P (decl))
@@ -1721,7 +1697,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
install_var_field (decl, true, 7, ctx);
else
install_var_field (decl, true, 3, ctx);
- if (offloaded)
+ if (is_gimple_omp_offloaded (ctx->stmt))
install_var_local (decl, ctx);
}
}
@@ -1845,7 +1821,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
gcc_assert (gimple_code (ctx->stmt) != GIMPLE_OMP_TARGET
|| (gimple_omp_target_kind (ctx->stmt)
!= GF_OMP_TARGET_KIND_UPDATE));
- if (!offloaded)
+ if (!is_gimple_omp_offloaded (ctx->stmt))
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 489 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.)
2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge
2014-03-12 13:48 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
@ 2014-03-20 14:39 ` Thomas Schwinge
2 siblings, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-03-20 14:39 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 6560 bytes --]
Hi!
Applied in r208701 to gomp-4_0-branch:
commit 22dd36a31c433dcd8bcc890d245a9e4ac6ed9c7f
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Mar 20 14:33:28 2014 +0000
Nesting of OpenACC constructs inside of OpenACC data constructs.
gcc/
* omp-low.c (check_omp_nesting_restrictions): Allow nesting of
OpenACC constructs inside of OpenACC data constructs.
gcc/testsuite/
* c-c++-common/goacc/nesting-1.c: New file.
* c-c++-common/goacc/nesting-data-1.c: Likewise.
* c-c++-common/goacc/nesting-fail-1.c: Update.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@208701 138bc75d-0d04-0410-961f-82ee72b054a4
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 1aebc4d..f43452c 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-03-20 Thomas Schwinge <thomas@codesourcery.com>
+
+ * omp-low.c (check_omp_nesting_restrictions): Allow nesting of
+ OpenACC constructs inside of OpenACC data constructs.
+
2014-03-18 Ilmir Usmanov <i.usmanov@samsung.com>
* tree.def (OACC_LOOP): New tree code.
diff --git gcc/omp-low.c gcc/omp-low.c
index f1b0fa5..23a0dda 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -2416,26 +2416,31 @@ scan_omp_teams (gimple stmt, omp_context *outer_ctx)
static bool
check_omp_nesting_restrictions (gimple stmt, omp_context *ctx)
{
- omp_context *ctx_;
-
/* TODO: While the OpenACC specification does allow for certain kinds of
- nesting, we don't support that yet. */
- /* No nesting of STMT (which is an OpenACC or OpenMP one, or a GOMP builtin)
- inside any OpenACC CTX. */
- for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
- if (is_gimple_omp (ctx_->stmt)
- && is_gimple_omp_oacc_specifically (ctx_->stmt))
- {
- error_at (gimple_location (stmt),
- "may not be nested");
- return false;
- }
- /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX. */
+ nesting, we don't support many of these yet. */
if (is_gimple_omp (stmt)
&& is_gimple_omp_oacc_specifically (stmt))
{
- for (ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
- if (is_gimple_omp (ctx_->stmt))
+ /* No nesting of OpenACC STMT inside any OpenACC or OpenMP CTX different
+ from an OpenACC data construct. */
+ for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ if (is_gimple_omp (ctx_->stmt)
+ && !(gimple_code (ctx_->stmt) == GIMPLE_OMP_TARGET
+ && (gimple_omp_target_kind (ctx_->stmt)
+ == GF_OMP_TARGET_KIND_OACC_DATA)))
+ {
+ error_at (gimple_location (stmt),
+ "may not be nested");
+ return false;
+ }
+ }
+ else
+ {
+ /* No nesting of non-OpenACC STMT (that is, an OpenMP one, or a GOMP
+ builtin) inside any OpenACC CTX. */
+ for (omp_context *ctx_ = ctx; ctx_ != NULL; ctx_ = ctx_->outer)
+ if (is_gimple_omp (ctx_->stmt)
+ && is_gimple_omp_oacc_specifically (ctx_->stmt))
{
error_at (gimple_location (stmt),
"may not be nested");
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index fd38d80..13e99d5 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,9 @@
2014-03-20 Thomas Schwinge <thomas@codesourcery.com>
+ * c-c++-common/goacc/nesting-1.c: New file.
+ * c-c++-common/goacc/nesting-data-1.c: Likewise.
+ * c-c++-common/goacc/nesting-fail-1.c: Update.
+
* c-c++-common/goacc/nesting-fail-1.c (f_acc_kernels): Replace
OpenACC parallel with kernels directive.
diff --git gcc/testsuite/c-c++-common/goacc/nesting-1.c gcc/testsuite/c-c++-common/goacc/nesting-1.c
new file mode 100644
index 0000000..3a22292
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-1.c
@@ -0,0 +1,13 @@
+void
+f_acc_data (void)
+{
+#pragma acc data
+ {
+#pragma acc parallel
+ ;
+#pragma acc kernels
+ ;
+#pragma acc data
+ ;
+ }
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-data-1.c gcc/testsuite/c-c++-common/goacc/nesting-data-1.c
new file mode 100644
index 0000000..fefe6cd
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/nesting-data-1.c
@@ -0,0 +1,61 @@
+void
+f (void)
+{
+ unsigned char c, ca[15], caa[20][30];
+
+#pragma acc data copyin(c)
+ {
+ c = 5;
+ ca[3] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+
+#pragma acc data copyin(ca[2:4])
+ {
+ c = 6;
+ ca[4] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+ }
+
+#pragma acc parallel copyout(ca[3:4])
+ {
+ c = 7;
+ ca[5] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+ }
+
+#pragma acc kernels copy(ca[4:4])
+ {
+ c = 8;
+ ca[6] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+ }
+
+#pragma acc data pcopy(ca[5:7])
+ {
+ c = 15;
+ ca[7] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+
+#pragma acc data pcopyin(caa[3:7][0:30])
+ {
+ c = 16;
+ ca[8] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+ }
+
+#pragma acc parallel pcopyout(caa[3:7][0:30])
+ {
+ c = 17;
+ ca[9] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+ }
+
+#pragma acc kernels pcopy(caa[3:7][0:30])
+ {
+ c = 18;
+ ca[10] = c;
+ caa[3][12] = ca[3] + caa[3][12];
+ }
+ }
+ }
+}
diff --git gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
index ca8921f..00dc602 100644
--- gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
+++ gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c
@@ -1,5 +1,5 @@
/* TODO: While the OpenACC specification does allow for certain kinds of
- nesting, we don't support that yet. */
+ nesting, we don't support many of these yet. */
void
f_acc_parallel (void)
{
@@ -15,7 +15,7 @@ f_acc_parallel (void)
}
/* TODO: While the OpenACC specification does allow for certain kinds of
- nesting, we don't support that yet. */
+ nesting, we don't support many of these yet. */
void
f_acc_kernels (void)
{
@@ -29,19 +29,3 @@ f_acc_kernels (void)
;
}
}
-
-/* TODO: While the OpenACC specification does allow for certain kinds of
- nesting, we don't support that yet. */
-void
-f_acc_data (void)
-{
-#pragma acc data
- {
-#pragma acc parallel /* { dg-error "may not be nested" } */
- ;
-#pragma acc kernels /* { dg-error "may not be nested" } */
- ;
-#pragma acc data /* { dg-error "may not be nested" } */
- ;
- }
-}
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* [GOMP4, COMMITTED] OpenACC deviceptr clause.
2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
2014-01-28 9:44 ` [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
@ 2014-06-05 14:00 ` Thomas Schwinge
[not found] ` <5460F49F.3040904@mentor.com>
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC present data clause Thomas Schwinge
3 siblings, 1 reply; 22+ messages in thread
From: Thomas Schwinge @ 2014-06-05 14:00 UTC (permalink / raw)
To: gcc-patches
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/c/
* c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
gcc/
* gimplify.c (gimplify_scan_omp_clauses)
(gimplify_adjust_omp_clauses): Handle
OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
* omp-low.c (scan_sharing_clauses, lower_oacc_offload)
(lower_omp_target): Likewise.
* tree-core.h (enum omp_clause_map_kind)
<OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
gcc/testsuite/
* c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
deviceptr clause is now supported.
* c-c++-common/goacc/deviceptr-1.c: Extend.
* c-c++-common/goacc/deviceptr-2.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211278 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 8 +++
gcc/c/ChangeLog.gomp | 5 ++
gcc/c/c-typeck.c | 5 +-
gcc/gimplify.c | 7 ++-
gcc/omp-low.c | 60 +++++++++++++++++++---
gcc/testsuite/ChangeLog.gomp | 5 ++
.../c-c++-common/goacc/data-clause-duplicate-1.c | 4 +-
gcc/testsuite/c-c++-common/goacc/deviceptr-1.c | 22 +++++++-
gcc/testsuite/c-c++-common/goacc/deviceptr-2.c | 23 +++++++++
gcc/tree-core.h | 3 +-
10 files changed, 127 insertions(+), 15 deletions(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 7371aa5..88f09b3 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,13 @@
2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
+ * gimplify.c (gimplify_scan_omp_clauses)
+ (gimplify_adjust_omp_clauses): Handle
+ OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
+ * omp-low.c (scan_sharing_clauses, lower_oacc_offload)
+ (lower_omp_target): Likewise.
+ * tree-core.h (enum omp_clause_map_kind)
+ <OMP_CLAUSE_MAP_FORCE_DEVICEPTR>: Update comment.
+
* gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>:
Don't block OMP_CLAUSE_MAP_FORCE_PRESENT.
diff --git gcc/c/ChangeLog.gomp gcc/c/ChangeLog.gomp
index 91978db..1e80031 100644
--- gcc/c/ChangeLog.gomp
+++ gcc/c/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-typeck.c (handle_omp_array_sections, c_finish_omp_clauses):
+ Handle OMP_CLAUSE_MAP_FORCE_DEVICEPTR.
+
2014-03-20 Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c: Update comments.
diff --git gcc/c/c-typeck.c gcc/c/c-typeck.c
index c4ba531..839cdf7 100644
--- gcc/c/c-typeck.c
+++ gcc/c/c-typeck.c
@@ -11747,6 +11747,7 @@ handle_omp_array_sections (tree c)
OMP_CLAUSE_SIZE (c) = size;
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
return false;
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
OMP_CLAUSE_MAP_KIND (c2) = OMP_CLAUSE_MAP_POINTER;
if (!c_mark_addressable (t))
@@ -12168,7 +12169,9 @@ c_finish_omp_clauses (tree clauses)
else if (!c_mark_addressable (t))
remove = true;
else if (!(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)))
&& !lang_hooks.types.omp_mappable_type (TREE_TYPE (t)))
{
error_at (OMP_CLAUSE_LOCATION (c),
diff --git gcc/gimplify.c gcc/gimplify.c
index 6eaf6fd..a1b6be6 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6015,7 +6015,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
switch (OMP_CLAUSE_MAP_KIND (c))
{
case OMP_CLAUSE_MAP_FORCE_DEALLOC:
- case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
input_location = OMP_CLAUSE_LOCATION (c);
/* TODO. */
sorry ("data clause not yet implemented");
@@ -6533,6 +6532,12 @@ gimplify_adjust_omp_clauses (tree *list_p)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != OMP_CLAUSE_MAP_POINTER)
{
+ /* For OMP_CLAUSE_MAP_FORCE_DEVICEPTR, we'll never enter here,
+ because for these, TREE_CODE (DECL_SIZE (decl)) will always be
+ INTEGER_CST. */
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
diff --git gcc/omp-low.c gcc/omp-low.c
index 3e282c0..39f0598 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1708,6 +1708,18 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& !POINTER_TYPE_P (TREE_TYPE (decl)))
break;
}
+#if 0
+ /* In target regions that are not offloaded, libgomp won't pay
+ attention to OMP_CLAUSE_MAP_FORCE_DEVICEPTR -- but I think we need
+ to handle it here anyway, in order to create a visible copy of the
+ variable. */
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ {
+ if (!is_gimple_omp_offloaded (ctx->stmt))
+ break;
+ }
+#endif
if (DECL_P (decl))
{
if (DECL_SIZE (decl)
@@ -1723,6 +1735,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
else
{
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -1738,6 +1754,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree base = get_base_address (decl);
tree nc = OMP_CLAUSE_CHAIN (c);
+ gcc_assert (nc == NULL_TREE
+ || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (nc)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR));
if (DECL_P (base)
&& nc != NULL_TREE
&& OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
@@ -1867,6 +1887,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
if (DECL_P (decl))
{
+ gcc_assert ((OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE);
if (OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
&& !COMPLETE_TYPE_P (TREE_TYPE (decl)))
@@ -1878,6 +1901,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
+ gcc_assert (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR);
+
tree decl2 = DECL_VALUE_EXPR (decl);
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
@@ -9100,6 +9126,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -9199,6 +9229,10 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree var = lookup_decl_in_outer_ctx (ovar, ctx);
tree x = build_sender_ref (ovar, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -9219,12 +9253,14 @@ lower_oacc_offload (gimple_stmt_iterator *gsi_p, omp_context *ctx)
= OMP_CLAUSE_MAP_KIND (c);
if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
&& (map_kind & OMP_CLAUSE_MAP_TO))
- || map_kind == OMP_CLAUSE_MAP_POINTER)
+ || map_kind == OMP_CLAUSE_MAP_POINTER
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
- && (map_kind & OMP_CLAUSE_MAP_FROM))
+ if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
@@ -10606,6 +10642,10 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
x = build_receiver_ref (var, true, ctx);
tree new_var = lookup_decl (var, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (var)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
@@ -10732,12 +10772,15 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree var = lookup_decl_in_outer_ctx (ovar, ctx);
tree x = build_sender_ref (ovar, ctx);
+ gcc_assert (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (OMP_CLAUSE_MAP_KIND (c)
+ != OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
+ || TREE_CODE (TREE_TYPE (ovar)) != ARRAY_TYPE);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == OMP_CLAUSE_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
tree avar
= create_tmp_var (TREE_TYPE (TREE_TYPE (x)), NULL);
mark_addressable (avar);
@@ -10747,19 +10790,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
else if (is_gimple_reg (var))
{
- gcc_assert (kind == GF_OMP_TARGET_KIND_REGION);
tree avar = create_tmp_var (TREE_TYPE (var), NULL);
mark_addressable (avar);
enum omp_clause_map_kind map_kind
= OMP_CLAUSE_MAP_KIND (c);
if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
&& (map_kind & OMP_CLAUSE_MAP_TO))
- || map_kind == OMP_CLAUSE_MAP_POINTER)
+ || map_kind == OMP_CLAUSE_MAP_POINTER
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
gimplify_assign (avar, var, &ilist);
avar = build_fold_addr_expr (avar);
gimplify_assign (x, avar, &ilist);
- if ((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
- && (map_kind & OMP_CLAUSE_MAP_FROM))
+ if (((!(map_kind & OMP_CLAUSE_MAP_SPECIAL)
+ && (map_kind & OMP_CLAUSE_MAP_FROM))
+ || map_kind == OMP_CLAUSE_MAP_FORCE_DEVICEPTR)
&& !TYPE_READONLY (TREE_TYPE (var)))
{
x = build_sender_ref (ovar, ctx);
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 4e0ee28..08ec907 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,5 +1,10 @@
2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
+ * c-c++-common/goacc/data-clause-duplicate-1.c: The OpenACC
+ deviceptr clause is now supported.
+ * c-c++-common/goacc/deviceptr-1.c: Extend.
+ * c-c++-common/goacc/deviceptr-2.c: New file.
+
* c-c++-common/goacc/data-clause-duplicate-1.c: Extend.
* c-c++-common/goacc/present-1.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 5c5ab02..7a1cf68 100644
--- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -6,9 +6,7 @@ fun (void)
;
#pragma acc kernels present_or_copyin(fp[3]) present_or_copyout(fp[7:4]) /* { dg-error "'fp' appears more than once in map clauses" } */
;
-#pragma acc data create(fp[:10]) deviceptr(fp)
- /* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */
- /* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */
+#pragma acc data create(fp[:10]) deviceptr(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
;
#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
;
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-1.c gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
index 1ac63bd..cf2d809 100644
--- gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-1.c
@@ -61,4 +61,24 @@ fun3 (void)
;
}
-/* { dg-prune-output "sorry, unimplemented: data clause not yet implemented" } */
+extern struct s s1;
+extern struct s s2[1]; /* { dg-error "array type has incomplete element type" "" { target c } } */
+
+void
+fun4 (void)
+{
+ struct s *s1_p = &s1;
+ struct s *s2_p = &s2;
+
+#pragma acc parallel deviceptr(s1) /* { dg-error "'s1' is not a pointer variable" } */
+ ;
+
+#pragma acc parallel deviceptr(s2)
+ ;
+
+#pragma acc parallel deviceptr(s1_p)
+ s1_p = 0;
+
+#pragma acc parallel deviceptr(s2_p)
+ s2_p = 0;
+}
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-2.c gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
new file mode 100644
index 0000000..ac162b4
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-2.c
@@ -0,0 +1,23 @@
+void
+fun1 (void)
+{
+ char *a = 0;
+
+#pragma acc data deviceptr(a)
+ ++a;
+
+#pragma acc data deviceptr(a)
+#pragma acc parallel
+ ++a;
+
+#pragma acc data deviceptr(a)
+#pragma acc parallel deviceptr(a)
+ ++a;
+
+#pragma acc data
+#pragma acc parallel deviceptr(a)
+ ++a;
+
+#pragma acc parallel deviceptr(a)
+ ++a;
+}
diff --git gcc/tree-core.h gcc/tree-core.h
index 8603553..8b70c5b 100644
--- gcc/tree-core.h
+++ gcc/tree-core.h
@@ -1225,7 +1225,8 @@ enum omp_clause_map_kind
OMP_CLAUSE_MAP_FORCE_PRESENT = OMP_CLAUSE_MAP_FORCE | OMP_CLAUSE_MAP_SPECIAL,
/* Deallocate a mapping, without copying from device. */
OMP_CLAUSE_MAP_FORCE_DEALLOC,
- /* Is a device pointer. */
+ /* Is a device pointer. OMP_CLAUSE_SIZE for these is unused; is implicitly
+ POINTER_SIZE / BITS_PER_UNIT. */
OMP_CLAUSE_MAP_FORCE_DEVICEPTR,
/* End marker. */
--
1.9.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* [GOMP4, COMMITTED] OpenACC present data clause.
2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
` (2 preceding siblings ...)
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC deviceptr clause Thomas Schwinge
@ 2014-06-05 14:00 ` Thomas Schwinge
3 siblings, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-06-05 14:00 UTC (permalink / raw)
To: gcc-patches
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
gcc/
* gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>:
Don't block OMP_CLAUSE_MAP_FORCE_PRESENT.
gcc/testsuite/
* c-c++-common/goacc/data-clause-duplicate-1.c: Extend.
* c-c++-common/goacc/present-1.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@211277 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 5 +++++
gcc/gimplify.c | 1 -
gcc/testsuite/ChangeLog.gomp | 5 +++++
gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c | 2 ++
gcc/testsuite/c-c++-common/goacc/present-1.c | 11 +++++++++++
5 files changed, 23 insertions(+), 1 deletion(-)
create mode 100644 gcc/testsuite/c-c++-common/goacc/present-1.c
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 011fe77..7371aa5 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gimplify.c (gimplify_scan_omp_clauses) <case OMP_CLAUSE_MAP>:
+ Don't block OMP_CLAUSE_MAP_FORCE_PRESENT.
+
2014-06-04 Thomas Schwinge <thomas@codesourcery.com>
* cgraphunit.c (ipa_passes, compile): Handle flag_openacc next to
diff --git gcc/gimplify.c gcc/gimplify.c
index e98e6e5..6eaf6fd 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -6014,7 +6014,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_MAP:
switch (OMP_CLAUSE_MAP_KIND (c))
{
- case OMP_CLAUSE_MAP_FORCE_PRESENT:
case OMP_CLAUSE_MAP_FORCE_DEALLOC:
case OMP_CLAUSE_MAP_FORCE_DEVICEPTR:
input_location = OMP_CLAUSE_LOCATION (c);
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index 78882c0..4e0ee28 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-06-05 Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-c++-common/goacc/data-clause-duplicate-1.c: Extend.
+ * c-c++-common/goacc/present-1.c: New file.
+
2014-03-20 Thomas Schwinge <thomas@codesourcery.com>
* c-c++-common/goacc-gomp/nesting-1.c: New file.
diff --git gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
index 4cb3cc2..5c5ab02 100644
--- gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
+++ gcc/testsuite/c-c++-common/goacc/data-clause-duplicate-1.c
@@ -10,4 +10,6 @@ fun (void)
/* { dg-error "'fp' appears more than once in map clauses" "" { target *-*-* } 9 } */
/* { dg-message "sorry, unimplemented: data clause not yet implemented" "" { target *-*-* } 9 } */
;
+#pragma acc data create(fp) present(fp) /* { dg-error "'fp' appears more than once in map clauses" } */
+ ;
}
diff --git gcc/testsuite/c-c++-common/goacc/present-1.c gcc/testsuite/c-c++-common/goacc/present-1.c
new file mode 100644
index 0000000..03ee592
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/present-1.c
@@ -0,0 +1,11 @@
+/* { dg-additional-options "-fdump-tree-original" } */
+
+void
+f (char *cp)
+{
+#pragma acc parallel present(cp[7:9])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "#pragma acc parallel map\\(force_present:\\*\\(cp \\+ 7\\) \\\[len: 9]\\) map\\(alloc:cp \\\[pointer assign, bias: 7]\\)" 1 "original" } } */
+/* { dg-final { cleanup-tree-dump "original" } } */
--
1.9.1
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [GOMP4, COMMITTED] OpenACC deviceptr clause.
[not found] ` <5460F49F.3040904@mentor.com>
@ 2014-11-11 21:30 ` Thomas Schwinge
0 siblings, 0 replies; 22+ messages in thread
From: Thomas Schwinge @ 2014-11-11 21:30 UTC (permalink / raw)
To: gcc-patches; +Cc: James Norris, Cesar Philippidis
[-- Attachment #1: Type: text/plain, Size: 3292 bytes --]
Hi!
On Thu, 5 Jun 2014 16:00:16 +0200, I wrote:
> --- gcc/omp-low.c
> +++ gcc/omp-low.c
> @@ -1738,6 +1754,10 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
> {
> tree base = get_base_address (decl);
> tree nc = OMP_CLAUSE_CHAIN (c);
> + gcc_assert (nc == NULL_TREE
> + || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP
> + || (OMP_CLAUSE_MAP_KIND (nc)
> + != OMP_CLAUSE_MAP_FORCE_DEVICEPTR));
> if (DECL_P (base)
> && nc != NULL_TREE
> && OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
That's a bogus assertion; removed in r217372 on gomp-4_0-branch:
commit 7ae51786d4a2aad4c82045dda780ae3e7904afa8
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Tue Nov 11 21:22:26 2014 +0000
OpenACC deviceptr clause: Remove bogus assertion.
gcc/
* omp-low.c (scan_sharing_clauses): Remove bogus assertion.
gcc/testsuite/
* c-c++-common/goacc/deviceptr-3.c: New file.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217372 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 2 ++
gcc/omp-low.c | 4 ----
gcc/testsuite/ChangeLog.gomp | 5 +++++
gcc/testsuite/c-c++-common/goacc/deviceptr-3.c | 11 +++++++++++
4 files changed, 18 insertions(+), 4 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 94a7f8c..4ea28e2 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,7 @@
2014-11-11 Thomas Schwinge <thomas@codesourcery.com>
+ * omp-low.c (scan_sharing_clauses): Remove bogus assertion.
+
* omp-low.c (delete_omp_context): Dispose of reduction_map.
* omp-low.c (maybe_lookup_reduction): Don't require an OpenACC
diff --git gcc/omp-low.c gcc/omp-low.c
index 5695ec3..1263409 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1920,10 +1920,6 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
{
tree base = get_base_address (decl);
tree nc = OMP_CLAUSE_CHAIN (c);
- gcc_assert (nc == NULL_TREE
- || OMP_CLAUSE_CODE (nc) != OMP_CLAUSE_MAP
- || (OMP_CLAUSE_MAP_KIND (nc)
- != OMP_CLAUSE_MAP_FORCE_DEVICEPTR));
if (DECL_P (base)
&& nc != NULL_TREE
&& OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
diff --git gcc/testsuite/ChangeLog.gomp gcc/testsuite/ChangeLog.gomp
index a02f58a..f8bacc3 100644
--- gcc/testsuite/ChangeLog.gomp
+++ gcc/testsuite/ChangeLog.gomp
@@ -1,3 +1,8 @@
+2014-11-07 Thomas Schwinge <thomas@codesourcery.com>
+ James Norris <jnorris@codesourcery.com>
+
+ * c-c++-common/goacc/deviceptr-3.c: New file.
+
2014-11-05 Thomas Schwinge <thomas@codesourcery.com>
* c-c++-common/goacc/update-1.c: Extend.
diff --git gcc/testsuite/c-c++-common/goacc/deviceptr-3.c gcc/testsuite/c-c++-common/goacc/deviceptr-3.c
new file mode 100644
index 0000000..bab56c3
--- /dev/null
+++ gcc/testsuite/c-c++-common/goacc/deviceptr-3.c
@@ -0,0 +1,11 @@
+float *d_a;
+
+void
+f (float *a)
+{
+#pragma acc parallel copyout (a[3:10]) deviceptr (d_a)
+ d_a[2] += 1.0;
+
+#pragma acc parallel deviceptr (d_a) copyout (a[3:10])
+ d_a[2] += 1.0;
+}
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.
[not found] ` <87egz645j4.fsf@schwinge.name>
@ 2014-11-13 12:21 ` Thomas Schwinge
2014-11-13 13:13 ` Jakub Jelinek
0 siblings, 1 reply; 22+ messages in thread
From: Thomas Schwinge @ 2014-11-13 12:21 UTC (permalink / raw)
To: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 25086 bytes --]
Hi!
On Tue, 14 Jan 2014 16:10:05 +0100, I wrote:
> --- gcc/gimplify.c
> +++ gcc/gimplify.c
> @@ -69,7 +69,13 @@ enum gimplify_omp_var_data
> + /* Force a specific behavior (or else, a run-time error). */
> + GOVD_MAP_FORCE = 16384,
> @@ -86,7 +92,11 @@ enum omp_region_type
> + /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
> + ORT_TARGET_MAP_FORCE = 64
> };
> @@ -6135,9 +6197,14 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
> OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
> else if (code == OMP_CLAUSE_MAP)
> {
> - OMP_CLAUSE_MAP_KIND (clause) = flags & GOVD_MAP_TO_ONLY
> - ? OMP_CLAUSE_MAP_TO
> - : OMP_CLAUSE_MAP_TOFROM;
> + unsigned map_kind;
> + map_kind = (flags & GOVD_MAP_TO_ONLY
> + ? OMP_CLAUSE_MAP_TO
> + : OMP_CLAUSE_MAP_TOFROM);
> + if (flags & GOVD_MAP_FORCE)
> + map_kind |= OMP_CLAUSE_MAP_FORCE;
> + OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind;
> +
> if (DECL_SIZE (decl)
> && TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
> {
> @@ -6389,9 +6456,10 @@ gimplify_oacc_parallel (tree *expr_p, gimple_seq *pre_p)
> tree expr = *expr_p;
> gimple g;
> gimple_seq body = NULL;
> + enum omp_region_type ort =
> + (enum omp_region_type) (ORT_TARGET | ORT_TARGET_MAP_FORCE);
>
> - gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p,
> - ORT_TARGET);
> + gimplify_scan_omp_clauses (&OACC_PARALLEL_CLAUSES (expr), pre_p, ort);
>
> push_gimplify_context ();
I don't remember what I have been thinking when implementing this -- per
the OpenACC specification's rules for implicitly determined data
attributes, it should be present_or_copy (that is, OpenMP's tofrom,
without "force" semantics), and firstprivate/copy for scalar variables
for the parallel/kernels constructs, respectively (which is still to be
implemented, for now not considering scalar variables different from
non-scalar ones). Committed to gomp-4_0-branch in r217482:
commit 7058203891bd6e1696763603673090f161e172b8
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Thu Nov 13 12:18:34 2014 +0000
Middle end: Don't use mapping kinds with "force" semantics for OpenACC.
..., which is the wrong thing to do. Also extend libgomp to actually
distinguish between "non-force"/"force" semantics.
gcc/
* gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS,
OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE.
(enum gimplify_omp_var_data, enum omp_region_type): Remove
GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update
all users.
include/
* gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and
_GOMP_MAP_FLAG_FORCE.
libgomp/
* target.c (gomp_map_vars_existing): Error out if "force"
semantics.
(gomp_map_vars): Actually pass kinds to gomp_map_vars_existing.
Remove FIXMEs.
* testsuite/libgomp.oacc-c-c++-common/data-already-1.c: New file.
* testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
* testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@217482 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/ChangeLog.gomp | 6 ++
gcc/gimplify.c | 65 +++-------------------
include/ChangeLog.gomp | 4 ++
include/gomp-constants.h | 3 +
libgomp/ChangeLog.gomp | 23 ++++++++
libgomp/target.c | 21 ++++---
.../libgomp.oacc-c-c++-common/data-already-1.c | 19 +++++++
.../libgomp.oacc-c-c++-common/data-already-2.c | 16 ++++++
.../libgomp.oacc-c-c++-common/data-already-3.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-4.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-5.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-6.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-7.c | 17 ++++++
.../libgomp.oacc-c-c++-common/data-already-8.c | 16 ++++++
.../libgomp.oacc-fortran/data-already-1.f | 17 ++++++
.../libgomp.oacc-fortran/data-already-2.f | 16 ++++++
.../libgomp.oacc-fortran/data-already-3.f | 15 +++++
.../libgomp.oacc-fortran/data-already-4.f | 14 +++++
.../libgomp.oacc-fortran/data-already-5.f | 14 +++++
.../libgomp.oacc-fortran/data-already-6.f | 14 +++++
.../libgomp.oacc-fortran/data-already-7.f | 14 +++++
.../libgomp.oacc-fortran/data-already-8.f | 16 ++++++
22 files changed, 311 insertions(+), 67 deletions(-)
diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 174235d..a499755 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,5 +1,11 @@
2014-11-13 Thomas Schwinge <thomas@codesourcery.com>
+ * gimplify.c (gimplify_omp_workshare) <OACC_DATA, OACC_KERNELS,
+ OACC_PARALLEL>: Don't request ORT_TARGET_MAP_FORCE.
+ (enum gimplify_omp_var_data, enum omp_region_type): Remove
+ GOVD_MAP_FORCE, and ORT_TARGET_MAP_FORCE, respectively. Update
+ all users.
+
* omp-low.c (scan_sharing_clauses) <OMP_CLAUSE_MAP>: Revert
earlier change.
diff --git gcc/gimplify.c gcc/gimplify.c
index 233ac56..2c8c666 100644
--- gcc/gimplify.c
+++ gcc/gimplify.c
@@ -94,8 +94,6 @@ enum gimplify_omp_var_data
/* Flags for GOVD_MAP. */
/* Don't copy back. */
GOVD_MAP_TO_ONLY = 8192,
- /* Force a specific behavior (or else, a run-time error). */
- GOVD_MAP_FORCE = 16384,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
@@ -116,9 +114,7 @@ enum omp_region_type
/* Flags for ORT_TARGET. */
/* Prepare this region for offloading. */
- ORT_TARGET_OFFLOAD = 32,
- /* Default to GOVD_MAP_FORCE for implicit mappings in this region. */
- ORT_TARGET_MAP_FORCE = 64
+ ORT_TARGET_OFFLOAD = 32
};
/* Gimplify hashtable helper. */
@@ -5585,15 +5581,7 @@ omp_add_variable (struct gimplify_omp_ctx *ctx, tree decl, unsigned int flags)
if (!(flags & GOVD_LOCAL))
{
if (flags & GOVD_MAP)
- {
- nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
-#if 0
- /* Not sure if this is actually needed; haven't found a case
- where this would change anything; TODO. */
- if (flags & GOVD_MAP_FORCE)
- nflags |= OMP_CLAUSE_MAP_FORCE;
-#endif
- }
+ nflags = GOVD_MAP | GOVD_MAP_TO_ONLY | GOVD_EXPLICIT;
else if (flags & GOVD_PRIVATE)
nflags = GOVD_PRIVATE;
else
@@ -5667,8 +5655,6 @@ omp_notice_threadprivate_variable (struct gimplify_omp_ctx *ctx, tree decl,
if ((octx->region_type & ORT_TARGET)
&& (octx->region_type & ORT_TARGET_OFFLOAD))
{
- gcc_assert (!(octx->region_type & ORT_TARGET_MAP_FORCE));
-
n = splay_tree_lookup (octx->variables, (splay_tree_key)decl);
if (n == NULL)
{
@@ -5731,11 +5717,6 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
if ((ctx->region_type & ORT_TARGET)
&& (ctx->region_type & ORT_TARGET_OFFLOAD))
{
- unsigned map_force;
- if (ctx->region_type & ORT_TARGET_MAP_FORCE)
- map_force = GOVD_MAP_FORCE;
- else
- map_force = 0;
ret = lang_hooks.decls.omp_disregard_value_expr (decl, true);
if (n == NULL)
{
@@ -5743,32 +5724,13 @@ omp_notice_variable (struct gimplify_omp_ctx *ctx, tree decl, bool in_code)
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
- omp_add_variable (ctx, decl, GOVD_MAP | map_force | GOVD_EXPLICIT | flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | GOVD_EXPLICIT | flags);
}
else
- omp_add_variable (ctx, decl, GOVD_MAP | map_force | flags);
+ omp_add_variable (ctx, decl, GOVD_MAP | flags);
}
else
{
-#if 0
- /* The following fails for:
-
- int l = 10;
- float c[l];
- #pragma acc parallel copy(c[2:4])
- {
- #pragma acc parallel
- {
- int t = sizeof c;
- }
- }
-
- ..., which we currently don't have to care about (nesting
- disabled), but eventually will have to; TODO. */
- if ((n->value & GOVD_MAP) && !(n->value & GOVD_EXPLICIT))
- gcc_assert ((n->value & GOVD_MAP_FORCE) == map_force);
-#endif
-
/* If nothing changed, there's nothing left to do. */
if ((n->value & flags) == flags)
return ret;
@@ -6423,13 +6385,11 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
OMP_CLAUSE_PRIVATE_OUTER_REF (clause) = 1;
else if (code == OMP_CLAUSE_MAP)
{
- unsigned map_kind;
+ enum omp_clause_map_kind map_kind;
map_kind = (flags & GOVD_MAP_TO_ONLY
? OMP_CLAUSE_MAP_TO
: OMP_CLAUSE_MAP_TOFROM);
- if (flags & GOVD_MAP_FORCE)
- map_kind |= OMP_CLAUSE_MAP_FORCE;
- OMP_CLAUSE_MAP_KIND (clause) = (enum omp_clause_map_kind) map_kind;
+ OMP_CLAUSE_MAP_KIND (clause) = map_kind;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
@@ -7258,23 +7218,16 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p)
switch (TREE_CODE (expr))
{
- case OACC_DATA:
- ort = (enum omp_region_type) (ORT_TARGET
- | ORT_TARGET_MAP_FORCE);
- break;
- case OACC_KERNELS:
- case OACC_PARALLEL:
- ort = (enum omp_region_type) (ORT_TARGET
- | ORT_TARGET_OFFLOAD
- | ORT_TARGET_MAP_FORCE);
- break;
case OMP_SECTIONS:
case OMP_SINGLE:
ort = ORT_WORKSHARE;
break;
+ case OACC_KERNELS:
+ case OACC_PARALLEL:
case OMP_TARGET:
ort = (enum omp_region_type) (ORT_TARGET | ORT_TARGET_OFFLOAD);
break;
+ case OACC_DATA:
case OMP_TARGET_DATA:
ort = ORT_TARGET;
break;
diff --git include/ChangeLog.gomp include/ChangeLog.gomp
new file mode 100644
index 0000000..9172c26
--- /dev/null
+++ include/ChangeLog.gomp
@@ -0,0 +1,4 @@
+2014-11-13 Thomas Schwinge <thomas@codesourcery.com>
+
+ * gomp-constants.h: Define _GOMP_MAP_FLAG_SPECIAL and
+ _GOMP_MAP_FLAG_FORCE.
diff --git include/gomp-constants.h include/gomp-constants.h
index e600766..15b658f 100644
--- include/gomp-constants.h
+++ include/gomp-constants.h
@@ -28,6 +28,9 @@
/* Enumerated variable mapping types used to communicate between GCC and
libgomp. These values are used for both OpenMP and OpenACC. */
+#define _GOMP_MAP_FLAG_SPECIAL (1 << 2)
+#define _GOMP_MAP_FLAG_FORCE (1 << 3)
+
#define GOMP_MAP_ALLOC 0x00
#define GOMP_MAP_ALLOC_TO 0x01
#define GOMP_MAP_ALLOC_FROM 0x02
diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp
index 0528531..254846f 100644
--- libgomp/ChangeLog.gomp
+++ libgomp/ChangeLog.gomp
@@ -1,3 +1,26 @@
+2014-11-13 Thomas Schwinge <thomas@codesourcery.com>
+
+ * target.c (gomp_map_vars_existing): Error out if "force"
+ semantics.
+ (gomp_map_vars): Actually pass kinds to gomp_map_vars_existing.
+ Remove FIXMEs.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: New file.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise.
+ * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise.
+ * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise.
+
2014-11-12 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c-c++-common/collapse-4.c: New file.
diff --git libgomp/target.c libgomp/target.c
index 052c59d..2b9f08f 100644
--- libgomp/target.c
+++ libgomp/target.c
@@ -117,9 +117,11 @@ static inline void
gomp_map_vars_existing (splay_tree_key oldn, splay_tree_key newn,
unsigned char kind)
{
- if (oldn->host_start > newn->host_start
+ if ((!(kind & _GOMP_MAP_FLAG_SPECIAL)
+ && (kind & _GOMP_MAP_FLAG_FORCE))
+ || oldn->host_start > newn->host_start
|| oldn->host_end < newn->host_end)
- gomp_fatal ("Trying to map into device [%p..%p) object when"
+ gomp_fatal ("Trying to map into device [%p..%p) object when "
"[%p..%p) is already mapped",
(void *) newn->host_start, (void *) newn->host_end,
(void *) oldn->host_start, (void *) oldn->host_end);
@@ -200,7 +202,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
if (n)
{
tgt->list[i] = n;
- gomp_map_vars_existing (n, &cur_node, kind);
+ gomp_map_vars_existing (n, &cur_node, kind & typemask);
}
else
{
@@ -323,7 +325,7 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
if (n)
{
tgt->list[i] = n;
- gomp_map_vars_existing (n, k, kind);
+ gomp_map_vars_existing (n, k, kind & typemask);
}
else
{
@@ -345,18 +347,15 @@ gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
switch (kind & typemask)
{
- case GOMP_MAP_FORCE_ALLOC:
- case GOMP_MAP_FORCE_FROM:
- /* FIXME: No special handling (see comment in
- oacc-parallel.c). */
case GOMP_MAP_ALLOC:
case GOMP_MAP_ALLOC_FROM:
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_FORCE_FROM:
break;
- case GOMP_MAP_FORCE_TO:
- case GOMP_MAP_FORCE_TOFROM:
- /* FIXME: No special handling, as above. */
case GOMP_MAP_ALLOC_TO:
case GOMP_MAP_ALLOC_TOFROM:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FORCE_TOFROM:
/* Copy from host to device memory. */
/* FIXME: Perhaps add some smarts, like if copying
several adjacent fields from host to target, use some
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
new file mode 100644
index 0000000..83c0a42
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-1.c
@@ -0,0 +1,19 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+ acc_copyin (&i, sizeof i);
+
+#pragma acc data copy (i)
+ ++i;
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "Trying to map into device .* object when .* is already mapped" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
new file mode 100644
index 0000000..137d8ce
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-2.c
@@ -0,0 +1,16 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data present_or_copy (i)
+#pragma acc data copyout (i)
+ ++i;
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "Trying to map into device .* object when .* is already mapped" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
new file mode 100644
index 0000000..b993b78
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-3.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data present_or_copy (i)
+ acc_copyin (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
new file mode 100644
index 0000000..82523f4
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-4.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+ acc_present_or_copyin (&i, sizeof i);
+ acc_copyin (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
new file mode 100644
index 0000000..4961fe5
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-5.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc enter data create (i)
+ acc_copyin (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
new file mode 100644
index 0000000..77b56a9
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-6.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+ acc_present_or_copyin (&i, sizeof i);
+#pragma acc enter data create (i)
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
new file mode 100644
index 0000000..b08417b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-7.c
@@ -0,0 +1,17 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+#include <openacc.h>
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc enter data create (i)
+ acc_create (&i, sizeof i);
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "already mapped to" } */
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
new file mode 100644
index 0000000..a50f7de
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-c-c++-common/data-already-8.c
@@ -0,0 +1,16 @@
+/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
+
+int
+main (int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data create (i)
+#pragma acc parallel copyin (i)
+ ++i;
+
+ return 0;
+}
+
+/* { dg-shouldfail "" }
+ { dg-output "Trying to map into device .* object when .* is already mapped" } */
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
new file mode 100644
index 0000000..ac220ab
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-1.f
@@ -0,0 +1,17 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+ CALL ACC_COPYIN (I)
+
+!$ACC DATA COPY (I)
+ I = 0
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "Trying to map into device .* object when .* is already mapped" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
new file mode 100644
index 0000000..2c5254b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-2.f
@@ -0,0 +1,16 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+
+ INTEGER I
+
+!$ACC DATA PRESENT_OR_COPY (I)
+!$ACC DATA COPYOUT (I)
+ I = 0
+!$ACC END DATA
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "Trying to map into device .* object when .* is already mapped" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
new file mode 100644
index 0000000..c41de28
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-3.f
@@ -0,0 +1,15 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+!$ACC DATA PRESENT_OR_COPY (I)
+ CALL ACC_COPYIN (I)
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
new file mode 100644
index 0000000..f54bf58
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-4.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+ CALL ACC_PRESENT_OR_COPYIN (I)
+ CALL ACC_COPYIN (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
new file mode 100644
index 0000000..9a3e94f
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-5.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+!$ACC ENTER DATA CREATE (I)
+ CALL ACC_COPYIN (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
new file mode 100644
index 0000000..eaf5d98
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-6.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+ CALL ACC_PRESENT_OR_COPYIN (I)
+!$ACC ENTER DATA CREATE (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
new file mode 100644
index 0000000..d96bf0b
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-7.f
@@ -0,0 +1,14 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+ INCLUDE "openacc_lib.h"
+
+ INTEGER I
+
+!$ACC ENTER DATA CREATE (I)
+ CALL ACC_CREATE (I)
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "already mapped to" }
diff --git libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
new file mode 100644
index 0000000..16da048
--- /dev/null
+++ libgomp/testsuite/libgomp.oacc-fortran/data-already-8.f
@@ -0,0 +1,16 @@
+! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } }
+
+ IMPLICIT NONE
+
+ INTEGER I
+
+!$ACC DATA CREATE (I)
+!$ACC PARALLEL COPYIN (I)
+ I = 0
+!$ACC END PARALLEL
+!$ACC END DATA
+
+ END
+
+! { dg-shouldfail "" }
+! { dg-output "Trying to map into device .* object when .* is already mapped" }
Grüße,
Thomas
[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.
2014-11-13 12:21 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
@ 2014-11-13 13:13 ` Jakub Jelinek
2014-11-13 13:39 ` gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) Thomas Schwinge
0 siblings, 1 reply; 22+ messages in thread
From: Jakub Jelinek @ 2014-11-13 13:13 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote:
> --- include/gomp-constants.h
> +++ include/gomp-constants.h
> @@ -28,6 +28,9 @@
> /* Enumerated variable mapping types used to communicate between GCC and
> libgomp. These values are used for both OpenMP and OpenACC. */
>
> +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2)
> +#define _GOMP_MAP_FLAG_FORCE (1 << 3)
I'm worried about reserved namespace issues if you use _ followed by
capital letter. Can't it be just GOMP_MAP_FLAG_* ?
Jakub
^ permalink raw reply [flat|nested] 22+ messages in thread
* gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.)
2014-11-13 13:13 ` Jakub Jelinek
@ 2014-11-13 13:39 ` Thomas Schwinge
2014-11-13 13:57 ` Jakub Jelinek
0 siblings, 1 reply; 22+ messages in thread
From: Thomas Schwinge @ 2014-11-13 13:39 UTC (permalink / raw)
To: Jakub Jelinek; +Cc: gcc-patches
[-- Attachment #1: Type: text/plain, Size: 1186 bytes --]
Hi Jakub!
On Thu, 13 Nov 2014 14:10:10 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote:
> > --- include/gomp-constants.h
> > +++ include/gomp-constants.h
> > @@ -28,6 +28,9 @@
> > /* Enumerated variable mapping types used to communicate between GCC and
> > libgomp. These values are used for both OpenMP and OpenACC. */
> >
> > +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2)
> > +#define _GOMP_MAP_FLAG_FORCE (1 << 3)
>
> I'm worried about reserved namespace issues if you use _ followed by
> capital letter.
Please remind me what those are reserved for?
> Can't it be just GOMP_MAP_FLAG_* ?
My worry is the other way round: gomp-constants.h is also #included from
<openacc.h> (to grab some of its constants), and using plain GOMP_* would
pollute the user's namespace? (I'm working on a patch to clean that up,
and also use gomp-constants.h more often, also for OpenMP code.) (Such a
shared (GCC/libgomp) header files had been discussed before, and now
introduced in
<http://news.gmane.org/find-root.php?message_id=%3C20140923191931.2177e60f%40octopus%3E>.)
Grüße,
Thomas
[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]
^ permalink raw reply [flat|nested] 22+ messages in thread
* Re: gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.)
2014-11-13 13:39 ` gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) Thomas Schwinge
@ 2014-11-13 13:57 ` Jakub Jelinek
0 siblings, 0 replies; 22+ messages in thread
From: Jakub Jelinek @ 2014-11-13 13:57 UTC (permalink / raw)
To: Thomas Schwinge; +Cc: gcc-patches
On Thu, Nov 13, 2014 at 02:38:06PM +0100, Thomas Schwinge wrote:
> Hi Jakub!
>
> On Thu, 13 Nov 2014 14:10:10 +0100, Jakub Jelinek <jakub@redhat.com> wrote:
> > On Thu, Nov 13, 2014 at 01:19:55PM +0100, Thomas Schwinge wrote:
> > > --- include/gomp-constants.h
> > > +++ include/gomp-constants.h
> > > @@ -28,6 +28,9 @@
> > > /* Enumerated variable mapping types used to communicate between GCC and
> > > libgomp. These values are used for both OpenMP and OpenACC. */
> > >
> > > +#define _GOMP_MAP_FLAG_SPECIAL (1 << 2)
> > > +#define _GOMP_MAP_FLAG_FORCE (1 << 3)
> >
> > I'm worried about reserved namespace issues if you use _ followed by
> > capital letter.
>
> Please remind me what those are reserved for?
See e.g.
http://www.gnu.org/software/libc/manual/html_node/Reserved-Names.html
http://pubs.opengroup.org/onlinepubs/007904975/functions/xsh_chap02_02.html
and remember that if you use gomp-constants.h in the compiler, it can be
built by the system compiler, which can be a very different implementation.
> > Can't it be just GOMP_MAP_FLAG_* ?
>
> My worry is the other way round: gomp-constants.h is also #included from
> <openacc.h> (to grab some of its constants), and using plain GOMP_* would
> pollute the user's namespace? (I'm working on a patch to clean that up,
> and also use gomp-constants.h more often, also for OpenMP code.) (Such a
> shared (GCC/libgomp) header files had been discussed before, and now
> introduced in
> <http://news.gmane.org/find-root.php?message_id=%3C20140923191931.2177e60f%40octopus%3E>.)
I think including gomp-constants.h in openacc.h, if that is a publicly
installed header, is a bad idea, you'll pollute namespace of that header.
Just duplicate the values in there under the right standard required names,
and you want, either add a testcase or some static assertions (e.g.
of the kind extern char typedef1[condition ? 1 : -1]; in some macros)
to verify that the openacc.h constants match the gomp-constants.h where
required.
Jakub
^ permalink raw reply [flat|nested] 22+ messages in thread
end of thread, other threads:[~2014-11-13 13:54 UTC | newest]
Thread overview: 22+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2014-01-14 15:09 [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
2014-01-14 15:10 ` [gomp4 1/6] During gimplification, allow additional flags next to ORT_TARGET thomas
2014-01-14 15:10 ` [gomp4 2/6] Prepare for extending omp_clause_map_kind thomas
2014-01-14 15:10 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics thomas
2014-01-14 15:10 ` [gomp4 4/6] C front end infrastructure for OpenACC clauses parsing thomas
2014-01-14 15:10 ` [gomp4 5/6] Initial support in the C front end for OpenACC data clauses thomas
2014-01-14 15:10 ` [gomp4 6/6] Enable initial " thomas
2014-02-12 11:17 ` [gomp4 5/6] Initial " Thomas Schwinge
2014-02-21 19:48 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
2014-02-21 20:32 ` [gomp4 1/3] Clarify to/from/map clauses usage in context of GF_OMP_TARGET_KIND_UPDATE Thomas Schwinge
2014-02-21 20:32 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-02-21 20:32 ` [gomp4 3/3] OpenACC data construct support in the C front end Thomas Schwinge
2014-03-12 13:48 ` [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA Thomas Schwinge
2014-03-20 14:39 ` [gomp4] Nesting of OpenACC constructs inside of OpenACC data constructs (was: [gomp4 2/3] OpenACC data construct implementation in terms of GF_OMP_TARGET_KIND_OACC_DATA.) Thomas Schwinge
[not found] ` <538DF785.3050206@mentor.com>
[not found] ` <87egz645j4.fsf@schwinge.name>
2014-11-13 12:21 ` [gomp4 3/6] Initial support for OpenACC memory mapping semantics Thomas Schwinge
2014-11-13 13:13 ` Jakub Jelinek
2014-11-13 13:39 ` gomp-constants.h (was: [gomp4 3/6] Initial support for OpenACC memory mapping semantics.) Thomas Schwinge
2014-11-13 13:57 ` Jakub Jelinek
2014-01-28 9:44 ` [gomp4] Initial support for OpenACC data clauses Thomas Schwinge
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC deviceptr clause Thomas Schwinge
[not found] ` <5460F49F.3040904@mentor.com>
2014-11-11 21:30 ` Thomas Schwinge
2014-06-05 14:00 ` [GOMP4, COMMITTED] OpenACC present data clause Thomas Schwinge
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).