+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ gcc/
+ * gimple.h (gf_mask): Add GF_OMP_TARGET_KIND_OACC_SERIAL
+ enumeration constant.
+ (is_gimple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
+ (is_gimple_omp_offloaded): Likewise.
+ * gimplify.c (omp_region_type): Add ORT_ACC_SERIAL enumeration
+ constant. Adjust the value of ORT_NONE accordingly.
+ (is_gimple_stmt): Handle OACC_SERIAL.
+ (oacc_default_clause): Handle ORT_ACC_SERIAL.
+ (gomp_needs_data_present): Likewise.
+ (gimplify_adjust_omp_clauses): Likewise.
+ (gimplify_omp_workshare): Handle OACC_SERIAL.
+ (gimplify_expr): Likewise.
+ * omp-expand.c (expand_omp_target):
+ Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
+ (build_omp_regions_1, omp_make_gimple_edges): Likewise.
+ * omp-low.c (is_oacc_parallel): Rename function to...
+ (is_oacc_parallel_or_serial): ... this.
+ Handle GF_OMP_TARGET_KIND_OACC_SERIAL.
+ (scan_sharing_clauses): Adjust accordingly.
+ (scan_omp_for): Likewise.
+ (lower_oacc_head_mark): Likewise.
+ (convert_from_firstprivate_int): Likewise.
+ (lower_omp_target): Likewise.
+ (check_omp_nesting_restrictions): Handle
+ GF_OMP_TARGET_KIND_OACC_SERIAL.
+ (lower_oacc_reductions): Likewise.
+ (lower_omp_target): Likewise.
+ * tree.def (OACC_SERIAL): New tree code.
+ * tree-pretty-print.c (dump_generic_node): Handle OACC_SERIAL.
+
+ * doc/generic.texi (OpenACC): Document OACC_SERIAL.
+
2019-11-12 Jakub Jelinek <jakub@redhat.com>
PR target/92449
+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+
+ gcc/c-family/
+ * c-pragma.h (pragma_kind): Add PRAGMA_OACC_SERIAL enumeration
+ constant.
+ * c-pragma.c (oacc_pragmas): Add "serial" entry.
+
2019-11-08 Richard Sandiford <richard.sandiford@arm.com>
* c-common.h (gnu_vector_type_p): New function.
{ "loop", PRAGMA_OACC_LOOP },
{ "parallel", PRAGMA_OACC_PARALLEL },
{ "routine", PRAGMA_OACC_ROUTINE },
+ { "serial", PRAGMA_OACC_SERIAL },
{ "update", PRAGMA_OACC_UPDATE },
{ "wait", PRAGMA_OACC_WAIT }
};
PRAGMA_OACC_LOOP,
PRAGMA_OACC_PARALLEL,
PRAGMA_OACC_ROUTINE,
+ PRAGMA_OACC_SERIAL,
PRAGMA_OACC_UPDATE,
PRAGMA_OACC_WAIT,
+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+
+ gcc/c/
+ * c-parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
+ (c_parser_oacc_kernels_parallel): Rename function to...
+ (c_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL.
+ (c_parser_omp_construct): Update accordingly.
+
+
2019-11-11 Jakub Jelinek <jakub@redhat.com>
* c-parser.c (c_parser_translation_unit): Diagnose declare target
# pragma acc parallel oacc-parallel-clause[optseq] new-line
structured-block
+ OpenACC 2.6:
+
+ # pragma acc serial oacc-serial-clause[optseq] new-line
+ structured-block
+
LOC is the location of the #pragma token.
*/
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+#define OACC_SERIAL_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
static tree
-c_parser_oacc_kernels_parallel (location_t loc, c_parser *parser,
- enum pragma_kind p_kind, char *p_name,
- bool *if_p)
+c_parser_oacc_compute (location_t loc, c_parser *parser,
+ enum pragma_kind p_kind, char *p_name, bool *if_p)
{
omp_clause_mask mask;
enum tree_code code;
mask = OACC_PARALLEL_CLAUSE_MASK;
code = OACC_PARALLEL;
break;
+ case PRAGMA_OACC_SERIAL:
+ strcat (p_name, " serial");
+ mask = OACC_SERIAL_CLAUSE_MASK;
+ code = OACC_SERIAL;
+ break;
default:
gcc_unreachable ();
}
break;
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
+ case PRAGMA_OACC_SERIAL:
strcpy (p_name, "#pragma acc");
- stmt = c_parser_oacc_kernels_parallel (loc, parser, p_kind, p_name,
- if_p);
+ stmt = c_parser_oacc_compute (loc, parser, p_kind, p_name, if_p);
break;
case PRAGMA_OACC_LOOP:
strcpy (p_name, "#pragma acc");
+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+
+ gcc/cp/
+ * constexpr.c (potential_constant_expression_1): Handle
+ OACC_SERIAL.
+ * parser.c (OACC_SERIAL_CLAUSE_MASK): New macro.
+ (cp_parser_oacc_kernels_parallel): Rename function to...
+ (cp_parser_oacc_compute): ... this. Handle PRAGMA_OACC_SERIAL.
+ (cp_parser_omp_construct): Update accordingly.
+ (cp_parser_pragma): Handle PRAGMA_OACC_SERIAL. Fix alphabetic
+ order.
+ * pt.c (tsubst_expr): Handle OACC_SERIAL.
+
2019-11-11 Jason Merrill <jason@redhat.com>
Implement P1946R0, Allow defaulting comparisons by value.
case OMP_DEPOBJ:
case OACC_PARALLEL:
case OACC_KERNELS:
+ case OACC_SERIAL:
case OACC_DATA:
case OACC_HOST_DATA:
case OACC_LOOP:
# pragma acc parallel oacc-parallel-clause[optseq] new-line
structured-block
+
+ OpenACC 2.6:
+
+ # pragma acc serial oacc-serial-clause[optseq] new-line
*/
#define OACC_KERNELS_CLAUSE_MASK \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR_LENGTH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+#define OACC_SERIAL_CLAUSE_MASK \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_REDUCTION) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
+
static tree
-cp_parser_oacc_kernels_parallel (cp_parser *parser, cp_token *pragma_tok,
- char *p_name, bool *if_p)
+cp_parser_oacc_compute (cp_parser *parser, cp_token *pragma_tok,
+ char *p_name, bool *if_p)
{
omp_clause_mask mask;
enum tree_code code;
mask = OACC_PARALLEL_CLAUSE_MASK;
code = OACC_PARALLEL;
break;
+ case PRAGMA_OACC_SERIAL:
+ strcat (p_name, " serial");
+ mask = OACC_SERIAL_CLAUSE_MASK;
+ code = OACC_SERIAL;
+ break;
default:
gcc_unreachable ();
}
break;
case PRAGMA_OACC_KERNELS:
case PRAGMA_OACC_PARALLEL:
+ case PRAGMA_OACC_SERIAL:
strcpy (p_name, "#pragma acc");
- stmt = cp_parser_oacc_kernels_parallel (parser, pragma_tok, p_name,
- if_p);
+ stmt = cp_parser_oacc_compute (parser, pragma_tok, p_name, if_p);
break;
case PRAGMA_OACC_LOOP:
strcpy (p_name, "#pragma acc");
case PRAGMA_OACC_DATA:
case PRAGMA_OACC_HOST_DATA:
case PRAGMA_OACC_KERNELS:
- case PRAGMA_OACC_PARALLEL:
case PRAGMA_OACC_LOOP:
+ case PRAGMA_OACC_PARALLEL:
+ case PRAGMA_OACC_SERIAL:
case PRAGMA_OMP_ATOMIC:
case PRAGMA_OMP_CRITICAL:
case PRAGMA_OMP_DISTRIBUTE:
case OACC_KERNELS:
case OACC_PARALLEL:
+ case OACC_SERIAL:
tmp = tsubst_omp_clauses (OMP_CLAUSES (t), C_ORT_ACC, args, complain,
in_decl);
stmt = begin_omp_parallel ();
@tindex OACC_KERNELS
@tindex OACC_LOOP
@tindex OACC_PARALLEL
+@tindex OACC_SERIAL
@tindex OACC_UPDATE
All the statements starting with @code{OACC_} represent directives and
Represents @code{#pragma acc parallel [clause1 @dots{} clauseN]}.
+@item OACC_SERIAL
+
+Represents @code{#pragma acc serial [clause1 @dots{} clauseN]}.
+
@item OACC_UPDATE
Represents @code{#pragma acc update [clause1 @dots{} clauseN]}.
+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+
+ gcc/fortran/
+ * gfortran.h (gfc_statement): Add ST_OACC_SERIAL_LOOP,
+ ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL and ST_OACC_END_SERIAL
+ enumeration constants.
+ (gfc_exec_op): Add EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL
+ enumeration constants.
+ * match.h (gfc_match_oacc_serial): New prototype.
+ (gfc_match_oacc_serial_loop): Likewise.
+ * dump-parse-tree.c (show_omp_node, show_code_node): Handle
+ EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
+ * match.c (match_exit_cycle): Handle EXEC_OACC_SERIAL_LOOP.
+ * openmp.c (OACC_SERIAL_CLAUSES): New macro.
+ (gfc_match_oacc_serial_loop): New function.
+ (gfc_match_oacc_serial): Likewise.
+ (oacc_is_loop): Handle EXEC_OACC_SERIAL_LOOP.
+ (resolve_omp_clauses): Handle EXEC_OACC_SERIAL.
+ (oacc_code_to_statement): Handle EXEC_OACC_SERIAL and
+ EXEC_OACC_SERIAL_LOOP.
+ (gfc_resolve_oacc_directive): Likewise.
+ * parse.c (decode_oacc_directive) <'s'>: Add case for "serial"
+ and "serial loop".
+ (next_statement): Handle ST_OACC_SERIAL_LOOP and ST_OACC_SERIAL.
+ (gfc_ascii_statement): Likewise. Handle ST_OACC_END_SERIAL_LOOP
+ and ST_OACC_END_SERIAL.
+ (parse_oacc_structured_block): Handle ST_OACC_SERIAL.
+ (parse_oacc_loop): Handle ST_OACC_SERIAL_LOOP and
+ ST_OACC_END_SERIAL_LOOP.
+ (parse_executable): Handle ST_OACC_SERIAL_LOOP and
+ ST_OACC_SERIAL.
+ (is_oacc): Handle EXEC_OACC_SERIAL_LOOP and EXEC_OACC_SERIAL.
+ * resolve.c (gfc_resolve_blocks, gfc_resolve_code): Likewise.
+ * st.c (gfc_free_statement): Likewise.
+ * trans-openmp.c (gfc_trans_oacc_construct): Handle
+ EXEC_OACC_SERIAL.
+ (gfc_trans_oacc_combined_directive): Handle
+ EXEC_OACC_SERIAL_LOOP.
+ (gfc_trans_oacc_directive): Handle EXEC_OACC_SERIAL_LOOP and
+ EXEC_OACC_SERIAL.
+ * trans.c (trans_code): Likewise.
+
2019-11-11 Janne Blomqvist <jb@gcc.gnu.org>
PR fortran/91828
case EXEC_OACC_PARALLEL: name = "PARALLEL"; is_oacc = true; break;
case EXEC_OACC_KERNELS_LOOP: name = "KERNELS LOOP"; is_oacc = true; break;
case EXEC_OACC_KERNELS: name = "KERNELS"; is_oacc = true; break;
+ case EXEC_OACC_SERIAL_LOOP: name = "SERIAL LOOP"; is_oacc = true; break;
+ case EXEC_OACC_SERIAL: name = "SERIAL"; is_oacc = true; break;
case EXEC_OACC_DATA: name = "DATA"; is_oacc = true; break;
case EXEC_OACC_HOST_DATA: name = "HOST_DATA"; is_oacc = true; break;
case EXEC_OACC_LOOP: name = "LOOP"; is_oacc = true; break;
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
ST_OACC_END_DATA, ST_OACC_HOST_DATA, ST_OACC_END_HOST_DATA, ST_OACC_LOOP,
ST_OACC_END_LOOP, ST_OACC_DECLARE, ST_OACC_UPDATE, ST_OACC_WAIT,
ST_OACC_CACHE, ST_OACC_KERNELS_LOOP, ST_OACC_END_KERNELS_LOOP,
- ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE,
+ ST_OACC_SERIAL_LOOP, ST_OACC_END_SERIAL_LOOP, ST_OACC_SERIAL,
+ ST_OACC_END_SERIAL, ST_OACC_ENTER_DATA, ST_OACC_EXIT_DATA, ST_OACC_ROUTINE,
ST_OACC_ATOMIC, ST_OACC_END_ATOMIC,
ST_OMP_ATOMIC, ST_OMP_BARRIER, ST_OMP_CRITICAL, ST_OMP_END_ATOMIC,
ST_OMP_END_CRITICAL, ST_OMP_END_DO, ST_OMP_END_MASTER, ST_OMP_END_ORDERED,
EXEC_BACKSPACE, EXEC_ENDFILE, EXEC_INQUIRE, EXEC_REWIND, EXEC_FLUSH,
EXEC_FORM_TEAM, EXEC_CHANGE_TEAM, EXEC_END_TEAM, EXEC_SYNC_TEAM,
EXEC_LOCK, EXEC_UNLOCK, EXEC_EVENT_POST, EXEC_EVENT_WAIT, EXEC_FAIL_IMAGE,
- EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_ROUTINE,
- EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_DATA, EXEC_OACC_HOST_DATA,
- EXEC_OACC_LOOP, EXEC_OACC_UPDATE, EXEC_OACC_WAIT, EXEC_OACC_CACHE,
- EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA, EXEC_OACC_ATOMIC,
- EXEC_OACC_DECLARE,
+ EXEC_OACC_KERNELS_LOOP, EXEC_OACC_PARALLEL_LOOP, EXEC_OACC_SERIAL_LOOP,
+ EXEC_OACC_ROUTINE, EXEC_OACC_PARALLEL, EXEC_OACC_KERNELS, EXEC_OACC_SERIAL,
+ EXEC_OACC_DATA, EXEC_OACC_HOST_DATA, EXEC_OACC_LOOP, EXEC_OACC_UPDATE,
+ EXEC_OACC_WAIT, EXEC_OACC_CACHE, EXEC_OACC_ENTER_DATA, EXEC_OACC_EXIT_DATA,
+ EXEC_OACC_ATOMIC, EXEC_OACC_DECLARE,
EXEC_OMP_CRITICAL, EXEC_OMP_DO, EXEC_OMP_FLUSH, EXEC_OMP_MASTER,
EXEC_OMP_ORDERED, EXEC_OMP_PARALLEL, EXEC_OMP_PARALLEL_DO,
EXEC_OMP_PARALLEL_SECTIONS, EXEC_OMP_PARALLEL_WORKSHARE,
&& o != NULL
&& o->state == COMP_OMP_STRUCTURED_BLOCK
&& (o->head->op == EXEC_OACC_LOOP
- || o->head->op == EXEC_OACC_PARALLEL_LOOP))
+ || o->head->op == EXEC_OACC_PARALLEL_LOOP
+ || o->head->op == EXEC_OACC_SERIAL_LOOP))
{
int collapse = 1;
gcc_assert (o->head->next != NULL
match gfc_match_oacc_kernels_loop (void);
match gfc_match_oacc_parallel (void);
match gfc_match_oacc_parallel_loop (void);
+match gfc_match_oacc_serial (void);
+match gfc_match_oacc_serial_loop (void);
match gfc_match_oacc_enter_data (void);
match gfc_match_oacc_exit_data (void);
match gfc_match_oacc_routine (void);
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
| OMP_CLAUSE_WAIT)
+#define OACC_SERIAL_CLAUSES \
+ (omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
+ | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
+ | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
+ | OMP_CLAUSE_WAIT)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
(OACC_LOOP_CLAUSES | OACC_PARALLEL_CLAUSES)
#define OACC_KERNELS_LOOP_CLAUSES \
(OACC_LOOP_CLAUSES | OACC_KERNELS_CLAUSES)
+#define OACC_SERIAL_LOOP_CLAUSES \
+ (OACC_LOOP_CLAUSES | OACC_SERIAL_CLAUSES)
#define OACC_HOST_DATA_CLAUSES omp_mask (OMP_CLAUSE_USE_DEVICE)
#define OACC_DECLARE_CLAUSES \
(omp_mask (OMP_CLAUSE_COPY) | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
}
+match
+gfc_match_oacc_serial_loop (void)
+{
+ return match_acc (EXEC_OACC_SERIAL_LOOP, OACC_SERIAL_LOOP_CLAUSES);
+}
+
+
+match
+gfc_match_oacc_serial (void)
+{
+ return match_acc (EXEC_OACC_SERIAL, OACC_SERIAL_CLAUSES);
+}
+
+
match
gfc_match_oacc_data (void)
{
{
return code->op == EXEC_OACC_PARALLEL_LOOP
|| code->op == EXEC_OACC_KERNELS_LOOP
+ || code->op == EXEC_OACC_SERIAL_LOOP
|| code->op == EXEC_OACC_LOOP;
}
n->sym->name, name, &n->where);
}
if (code
- && (oacc_is_loop (code) || code->op == EXEC_OACC_PARALLEL))
+ && (oacc_is_loop (code)
+ || code->op == EXEC_OACC_PARALLEL
+ || code->op == EXEC_OACC_SERIAL))
check_array_not_assumed (n->sym, n->where, name);
else if (n->sym->as && n->sym->as->type == AS_ASSUMED_SIZE)
gfc_error ("Assumed size array %qs in %s clause at %L",
return ST_OACC_PARALLEL;
case EXEC_OACC_KERNELS:
return ST_OACC_KERNELS;
+ case EXEC_OACC_SERIAL:
+ return ST_OACC_SERIAL;
case EXEC_OACC_DATA:
return ST_OACC_DATA;
case EXEC_OACC_HOST_DATA:
return ST_OACC_PARALLEL_LOOP;
case EXEC_OACC_KERNELS_LOOP:
return ST_OACC_KERNELS_LOOP;
+ case EXEC_OACC_SERIAL_LOOP:
+ return ST_OACC_SERIAL_LOOP;
case EXEC_OACC_LOOP:
return ST_OACC_LOOP;
case EXEC_OACC_ATOMIC:
{
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_UPDATE:
break;
case EXEC_OACC_PARALLEL_LOOP:
case EXEC_OACC_KERNELS_LOOP:
+ case EXEC_OACC_SERIAL_LOOP:
case EXEC_OACC_LOOP:
resolve_oacc_loop (code);
break;
matcha ("end parallel loop", gfc_match_omp_eos_error,
ST_OACC_END_PARALLEL_LOOP);
matcha ("end parallel", gfc_match_omp_eos_error, ST_OACC_END_PARALLEL);
+ matcha ("end serial loop", gfc_match_omp_eos_error,
+ ST_OACC_END_SERIAL_LOOP);
+ matcha ("end serial", gfc_match_omp_eos_error, ST_OACC_END_SERIAL);
matcha ("enter data", gfc_match_oacc_enter_data, ST_OACC_ENTER_DATA);
matcha ("exit data", gfc_match_oacc_exit_data, ST_OACC_EXIT_DATA);
break;
case 'r':
match ("routine", gfc_match_oacc_routine, ST_OACC_ROUTINE);
break;
+ case 's':
+ matcha ("serial loop", gfc_match_oacc_serial_loop, ST_OACC_SERIAL_LOOP);
+ matcha ("serial", gfc_match_oacc_serial, ST_OACC_SERIAL);
+ break;
case 'u':
matcha ("update", gfc_match_oacc_update, ST_OACC_UPDATE);
break;
case ST_CRITICAL: \
case ST_OACC_PARALLEL_LOOP: case ST_OACC_PARALLEL: case ST_OACC_KERNELS: \
case ST_OACC_DATA: case ST_OACC_HOST_DATA: case ST_OACC_LOOP: \
- case ST_OACC_KERNELS_LOOP: case ST_OACC_ATOMIC
+ case ST_OACC_KERNELS_LOOP: case ST_OACC_SERIAL_LOOP: case ST_OACC_SERIAL: \
+ case ST_OACC_ATOMIC
/* Declaration statements */
case ST_OACC_END_KERNELS_LOOP:
p = "!$ACC END KERNELS LOOP";
break;
+ case ST_OACC_SERIAL_LOOP:
+ p = "!$ACC SERIAL LOOP";
+ break;
+ case ST_OACC_END_SERIAL_LOOP:
+ p = "!$ACC END SERIAL LOOP";
+ break;
+ case ST_OACC_SERIAL:
+ p = "!$ACC SERIAL";
+ break;
+ case ST_OACC_END_SERIAL:
+ p = "!$ACC END SERIAL";
+ break;
case ST_OACC_DATA:
p = "!$ACC DATA";
break;
case ST_OACC_KERNELS:
acc_end_st = ST_OACC_END_KERNELS;
break;
+ case ST_OACC_SERIAL:
+ acc_end_st = ST_OACC_END_SERIAL;
+ break;
case ST_OACC_DATA:
acc_end_st = ST_OACC_END_DATA;
break;
pop_state ();
}
-/* Parse the statements of OpenACC loop/parallel loop/kernels loop. */
+/* Parse the statements of OpenACC 'loop', or combined compute 'loop'. */
static gfc_statement
parse_oacc_loop (gfc_statement acc_st)
gfc_warning (0, "Redundant !$ACC END LOOP at %C");
if ((acc_st == ST_OACC_PARALLEL_LOOP && st == ST_OACC_END_PARALLEL_LOOP) ||
(acc_st == ST_OACC_KERNELS_LOOP && st == ST_OACC_END_KERNELS_LOOP) ||
+ (acc_st == ST_OACC_SERIAL_LOOP && st == ST_OACC_END_SERIAL_LOOP) ||
(acc_st == ST_OACC_LOOP && st == ST_OACC_END_LOOP))
{
gcc_assert (new_st.op == EXEC_NOP);
case ST_OACC_PARALLEL_LOOP:
case ST_OACC_KERNELS_LOOP:
+ case ST_OACC_SERIAL_LOOP:
case ST_OACC_LOOP:
st = parse_oacc_loop (st);
if (st == ST_IMPLIED_ENDDO)
case ST_OACC_PARALLEL:
case ST_OACC_KERNELS:
+ case ST_OACC_SERIAL:
case ST_OACC_DATA:
case ST_OACC_HOST_DATA:
parse_oacc_structured_block (st);
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL_LOOP:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
case EXEC_OACC_LOOP:
return stmt;
}
-/* Trans OpenACC directives. */
-/* parallel, kernels, data and host_data. */
+/* Translate OpenACC 'parallel', 'kernels', 'serial', 'data', 'host_data'
+ construct. */
+
static tree
gfc_trans_oacc_construct (gfc_code *code)
{
case EXEC_OACC_KERNELS:
construct_code = OACC_KERNELS;
break;
+ case EXEC_OACC_SERIAL:
+ construct_code = OACC_SERIAL;
+ break;
case EXEC_OACC_DATA:
construct_code = OACC_DATA;
break;
return gfc_finish_block (&block);
}
-/* parallel loop and kernels loop. */
+/* Translate combined OpenACC 'parallel loop', 'kernels loop', 'serial loop'
+ construct. */
+
static tree
gfc_trans_oacc_combined_directive (gfc_code *code)
{
case EXEC_OACC_KERNELS_LOOP:
construct_code = OACC_KERNELS;
break;
+ case EXEC_OACC_SERIAL_LOOP:
+ construct_code = OACC_SERIAL;
+ break;
default:
gcc_unreachable ();
}
{
case EXEC_OACC_PARALLEL_LOOP:
case EXEC_OACC_KERNELS_LOOP:
+ case EXEC_OACC_SERIAL_LOOP:
return gfc_trans_oacc_combined_directive (code);
case EXEC_OACC_PARALLEL:
case EXEC_OACC_KERNELS:
+ case EXEC_OACC_SERIAL:
case EXEC_OACC_DATA:
case EXEC_OACC_HOST_DATA:
return gfc_trans_oacc_construct (code);
case EXEC_OACC_KERNELS_LOOP:
case EXEC_OACC_PARALLEL:
case EXEC_OACC_PARALLEL_LOOP:
+ case EXEC_OACC_SERIAL:
+ case EXEC_OACC_SERIAL_LOOP:
case EXEC_OACC_ENTER_DATA:
case EXEC_OACC_EXIT_DATA:
case EXEC_OACC_ATOMIC:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
kind = " oacc_parallel";
break;
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
+ kind = " oacc_serial";
+ break;
case GF_OMP_TARGET_KIND_OACC_DATA:
kind = " oacc_data";
break;
DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT)
/* GIMPLE_OMP_TARGET <BODY, CLAUSES, CHILD_FN> represents
- #pragma acc {kernels,parallel,data,enter data,exit data,update}
+ #pragma acc {kernels,parallel,serial,data,enter data,exit data,update}
#pragma omp target {,data,update}
BODY is the sequence of statements inside the construct
(NULL for some variants).
GF_OMP_TARGET_KIND_EXIT_DATA = 4,
GF_OMP_TARGET_KIND_OACC_PARALLEL = 5,
GF_OMP_TARGET_KIND_OACC_KERNELS = 6,
- GF_OMP_TARGET_KIND_OACC_DATA = 7,
- GF_OMP_TARGET_KIND_OACC_UPDATE = 8,
- GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 9,
- GF_OMP_TARGET_KIND_OACC_DECLARE = 10,
- GF_OMP_TARGET_KIND_OACC_HOST_DATA = 11,
+ GF_OMP_TARGET_KIND_OACC_SERIAL = 7,
+ GF_OMP_TARGET_KIND_OACC_DATA = 8,
+ GF_OMP_TARGET_KIND_OACC_UPDATE = 9,
+ GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 10,
+ GF_OMP_TARGET_KIND_OACC_DECLARE = 11,
+ GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12,
GF_OMP_TEAMS_GRID_PHONY = 1 << 0,
GF_OMP_TEAMS_HOST = 1 << 1,
{
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_REGION:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
return true;
default:
return false;
ORT_ACC_DATA = ORT_ACC | ORT_TARGET_DATA, /* Data construct. */
ORT_ACC_PARALLEL = ORT_ACC | ORT_TARGET, /* Parallel construct */
ORT_ACC_KERNELS = ORT_ACC | ORT_TARGET | 2, /* Kernels construct. */
+ ORT_ACC_SERIAL = ORT_ACC | ORT_TARGET | 4, /* Serial construct. */
ORT_ACC_HOST_DATA = ORT_ACC | ORT_TARGET_DATA | 2, /* Host data. */
/* Dummy OpenMP region, used to disable expansion of
case STATEMENT_LIST:
case OACC_PARALLEL:
case OACC_KERNELS:
+ case OACC_SERIAL:
case OACC_DATA:
case OACC_HOST_DATA:
case OACC_DECLARE:
break;
case ORT_ACC_PARALLEL:
- rkind = "parallel";
+ case ORT_ACC_SERIAL:
+ rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial";
if (is_private)
flags |= GOVD_FIRSTPRIVATE;
break;
}
decl = OMP_CLAUSE_DECL (c);
- /* Data clauses associated with acc parallel reductions must be
+ /* Data clauses associated with reductions must be
compatible with present_or_copy. Warn and adjust the clause
if that is not the case. */
- if (ctx->region_type == ORT_ACC_PARALLEL)
+ if (ctx->region_type == ORT_ACC_PARALLEL
+ || ctx->region_type == ORT_ACC_SERIAL)
{
tree t = DECL_P (decl) ? decl : TREE_OPERAND (decl, 0);
n = NULL;
decl = OMP_CLAUSE_DECL (c);
/* OpenACC reductions need a present_or_copy data clause.
Add one if necessary. Emit error when the reduction is private. */
- if (ctx->region_type == ORT_ACC_PARALLEL)
+ if (ctx->region_type == ORT_ACC_PARALLEL
+ || ctx->region_type == ORT_ACC_SERIAL)
{
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))
case OACC_PARALLEL:
ort = ORT_ACC_PARALLEL;
break;
+ case OACC_SERIAL:
+ ort = ORT_ACC_SERIAL;
+ break;
case OACC_DATA:
ort = ORT_ACC_DATA;
break;
stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_PARALLEL,
OMP_CLAUSES (expr));
break;
+ case OACC_SERIAL:
+ stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_OACC_SERIAL,
+ OMP_CLAUSES (expr));
+ break;
case OMP_SECTIONS:
stmt = gimple_build_omp_sections (body, OMP_CLAUSES (expr));
break;
case OACC_DATA:
case OACC_KERNELS:
case OACC_PARALLEL:
+ case OACC_SERIAL:
case OMP_SECTIONS:
case OMP_SINGLE:
case OMP_TARGET:
&& code != EH_ELSE_EXPR
&& code != OACC_PARALLEL
&& code != OACC_KERNELS
+ && code != OACC_SERIAL
&& code != OACC_DATA
&& code != OACC_HOST_DATA
&& code != OACC_DECLARE
gimple *stmt;
edge e;
bool offloaded, data_region;
+ int target_kind;
entry_stmt = as_a <gomp_target *> (last_stmt (region->entry));
+ target_kind = gimple_omp_target_kind (entry_stmt);
new_bb = region->entry;
offloaded = is_gimple_omp_offloaded (entry_stmt);
- switch (gimple_omp_target_kind (entry_stmt))
+ switch (target_kind)
{
case GF_OMP_TARGET_KIND_REGION:
case GF_OMP_TARGET_KIND_UPDATE:
case GF_OMP_TARGET_KIND_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
entry_bb = region->entry;
exit_bb = region->exit;
- if (gimple_omp_target_kind (entry_stmt) == GF_OMP_TARGET_KIND_OACC_KERNELS)
+ switch (target_kind)
{
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
mark_loops_in_oacc_kernels_region (region->entry, region->exit);
- /* Further down, both OpenACC kernels and OpenACC parallel constructs
- will be mappted to BUILT_IN_GOACC_PARALLEL, and to distinguish the
- two, there is an "oacc kernels" attribute set for OpenACC kernels. */
+ /* Further down, all OpenACC compute constructs will be mapped to
+ BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there
+ is an "oacc kernels" attribute set for OpenACC kernels. */
DECL_ATTRIBUTES (child_fn)
= tree_cons (get_identifier ("oacc kernels"),
NULL_TREE, DECL_ATTRIBUTES (child_fn));
+ break;
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
+ /* Further down, all OpenACC compute constructs will be mapped to
+ BUILT_IN_GOACC_PARALLEL, and to distinguish between them, there
+ is an "oacc serial" attribute set for OpenACC serial. */
+ DECL_ATTRIBUTES (child_fn)
+ = tree_cons (get_identifier ("oacc serial"),
+ NULL_TREE, DECL_ATTRIBUTES (child_fn));
+ break;
+ default:
+ break;
}
if (offloaded)
start_ix = BUILT_IN_GOMP_TARGET_ENTER_EXIT_DATA;
flags_i |= GOMP_TARGET_FLAG_EXIT_DATA;
break;
- case GF_OMP_TARGET_KIND_OACC_KERNELS:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
+ case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
start_ix = BUILT_IN_GOACC_PARALLEL;
break;
case GF_OMP_TARGET_KIND_OACC_DATA:
args.quick_push (get_target_arguments (&gsi, entry_stmt));
break;
case BUILT_IN_GOACC_PARALLEL:
- oacc_set_fn_attrib (child_fn, clauses, &args);
+ if (lookup_attribute ("oacc serial", DECL_ATTRIBUTES (child_fn)) != NULL)
+ {
+ tree dims = NULL_TREE;
+ unsigned int ix;
+
+ /* For serial constructs we set all dimensions to 1. */
+ for (ix = GOMP_DIM_MAX; ix--;)
+ dims = tree_cons (NULL_TREE, integer_one_node, dims);
+ oacc_replace_fn_attrib (child_fn, dims);
+ }
+ else
+ oacc_set_fn_attrib (child_fn, clauses, &args);
tagging = true;
/* FALLTHRU */
case BUILT_IN_GOACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
break;
case GF_OMP_TARGET_KIND_DATA:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_DATA:
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
break;
*handled_ops_p = false; \
break;
-/* Return true if CTX corresponds to an oacc parallel region. */
+/* Return true if CTX corresponds to an OpenACC 'parallel' or 'serial'
+ region. */
static bool
-is_oacc_parallel (omp_context *ctx)
+is_oacc_parallel_or_serial (omp_context *ctx)
{
enum gimple_code outer_type = gimple_code (ctx->stmt);
return ((outer_type == GIMPLE_OMP_TARGET)
- && (gimple_omp_target_kind (ctx->stmt)
- == GF_OMP_TARGET_KIND_OACC_PARALLEL));
+ && ((gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ || (gimple_omp_target_kind (ctx->stmt)
+ == GF_OMP_TARGET_KIND_OACC_SERIAL)));
}
/* Return true if CTX corresponds to an oacc kernels region. */
goto do_private;
case OMP_CLAUSE_REDUCTION:
- if (is_oacc_parallel (ctx) || is_oacc_kernels (ctx))
+ if (is_oacc_parallel_or_serial (ctx) || is_oacc_kernels (ctx))
ctx->local_reduction_clauses
= tree_cons (NULL, c, ctx->local_reduction_clauses);
/* FALLTHRU */
{
omp_context *tgt = enclosing_target_ctx (outer_ctx);
- if (!tgt || is_oacc_parallel (tgt))
+ if (!tgt || is_oacc_parallel_or_serial (tgt))
for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
{
char const *check = NULL;
if (check && OMP_CLAUSE_OPERAND (c, 0))
error_at (gimple_location (stmt),
"argument not permitted on %qs clause in"
- " OpenACC %<parallel%>", check);
+ " OpenACC %<parallel%> or %<serial%>", check);
}
if (tgt && is_oacc_kernels (tgt))
{
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
ok = true;
break;
stmt_name = "target exit data"; break;
case GF_OMP_TARGET_KIND_OACC_PARALLEL: stmt_name = "parallel"; break;
case GF_OMP_TARGET_KIND_OACC_KERNELS: stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break;
case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break;
case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
ctx_stmt_name = "parallel"; break;
case GF_OMP_TARGET_KIND_OACC_KERNELS:
ctx_stmt_name = "kernels"; break;
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
+ ctx_stmt_name = "serial"; break;
case GF_OMP_TARGET_KIND_OACC_DATA: ctx_stmt_name = "data"; break;
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
ctx_stmt_name = "host_data"; break;
break;
case GIMPLE_OMP_TARGET:
- if (gimple_omp_target_kind (probe->stmt)
- != GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ if ((gimple_omp_target_kind (probe->stmt)
+ != GF_OMP_TARGET_KIND_OACC_PARALLEL)
+ && (gimple_omp_target_kind (probe->stmt)
+ != GF_OMP_TARGET_KIND_OACC_SERIAL))
goto do_lookup;
cls = gimple_omp_target_clauses (probe->stmt);
/* In a parallel region, loops are implicitly INDEPENDENT. */
omp_context *tgt = enclosing_target_ctx (ctx);
- if (!tgt || is_oacc_parallel (tgt))
+ if (!tgt || is_oacc_parallel_or_serial (tgt))
tag |= OLF_INDEPENDENT;
if (tag & OLF_TILE)
case GF_OMP_TARGET_KIND_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_PARALLEL:
case GF_OMP_TARGET_KIND_OACC_KERNELS:
+ case GF_OMP_TARGET_KIND_OACC_SERIAL:
case GF_OMP_TARGET_KIND_OACC_UPDATE:
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
case GF_OMP_TARGET_KIND_OACC_DECLARE:
if (!maybe_lookup_field (var, ctx))
continue;
- /* Don't remap oacc parallel reduction variables, because the
+ /* Don't remap compute constructs' reduction variables, because the
intermediate result must be local to each gang. */
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_oacc_parallel (ctx))
+ if (is_oacc_parallel_or_serial (ctx))
goto oacc_firstprivate;
map_cnt++;
var = OMP_CLAUSE_DECL (c);
break;
case OMP_CLAUSE_FIRSTPRIVATE:
- if (is_oacc_parallel (ctx))
+ if (is_oacc_parallel_or_serial (ctx))
goto oacc_firstprivate_map;
ovar = OMP_CLAUSE_DECL (c);
if (omp_is_reference (ovar))
gimple_seq fork_seq = NULL;
gimple_seq join_seq = NULL;
- if (is_oacc_parallel (ctx))
+ if (is_oacc_parallel_or_serial (ctx))
{
/* If there are reductions on the offloaded region itself, treat
them as a dummy GANG loop. */
+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ gcc/testsuite/
+ * c-c++-common/goacc/parallel-dims.c: New test.
+ * gfortran.dg/goacc/parallel-dims.f90: New test.
+
2019-11-12 Jakub Jelinek <jakub@redhat.com>
PR tree-optimization/92452
-/* Invalid use of OpenACC parallelism dimensions clauses: num_gangs,
- num_workers, vector_length. */
+/* Invalid use of OpenACC parallelism dimensions clauses: 'num_gangs',
+ 'num_workers', 'vector_length'. */
+
+/* See also '../../gfortran.dg/goacc/parallel-dims-2.f90'. */
void f(int i, float f)
{
vector_length(&f) /* { dg-error "'vector_length' expression must be integral" } */ \
num_gangs( /* { dg-error "expected (primary-|)expression before end of line" "TODO" { xfail c } } */
;
+
+
+ /* The 'serial' construct doesn't allow these at all. */
+
+#pragma acc serial num_gangs (1) /* { dg-error "'num_gangs' is not valid for '#pragma acc serial'" } */
+ ;
+#pragma acc serial num_workers (1) /* { dg-error "'num_workers' is not valid for '#pragma acc serial'" } */
+ ;
+#pragma acc serial vector_length (1) /* { dg-error "'vector_length' is not valid for '#pragma acc serial'" } */
+ ;
}
--- /dev/null
+! Invalid use of OpenACC parallelism dimensions clauses: 'num_gangs',
+! 'num_workers', 'vector_length'.
+
+! See also '../../c-c++-common/goacc/parallel-dims-2.c'.
+
+subroutine f()
+ !TODO 'kernels', 'parallel' testing per '../../c-c++-common/goacc/parallel-dims-2.c'.
+ !TODO This should incorporate some of the testing done in 'sie.f95'.
+
+
+ ! The 'serial' construct doesn't allow these at all.
+
+!$acc serial num_gangs (1) ! { dg-error "Failed to match clause at" }
+!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" }
+
+!$acc serial num_workers (1) ! { dg-error "Failed to match clause at" }
+!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" }
+
+!$acc serial vector_length (1) ! { dg-error "Failed to match clause at" }
+!$acc end serial ! { dg-error "Unexpected !.ACC END SERIAL statement" }
+
+end subroutine f
pp_string (pp, "#pragma acc kernels");
goto dump_omp_clauses_body;
+ case OACC_SERIAL:
+ pp_string (pp, "#pragma acc serial");
+ goto dump_omp_clauses_body;
+
case OACC_DATA:
pp_string (pp, "#pragma acc data");
dump_omp_clauses (pp, OACC_DATA_CLAUSES (node), spc, flags);
DEFTREECODE (OACC_KERNELS, "oacc_kernels", tcc_statement, 2)
+/* OpenACC - #pragma acc serial [clause1 ... clauseN]
+ Operand 0: OMP_BODY: Code to be executed sequentially.
+ Operand 1: OMP_CLAUSES: List of clauses. */
+
+DEFTREECODE (OACC_SERIAL, "oacc_serial", tcc_statement, 2)
+
/* OpenACC - #pragma acc data [clause1 ... clauseN]
Operand 0: OACC_DATA_BODY: Data construct body.
Operand 1: OACC_DATA_CLAUSES: List of clauses. */
treatment if OMP_CLAUSE_SIZE is zero. */
#define OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION(NODE) \
TREE_PROTECTED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
-/* Nonzero if this map clause is for an ACC parallel reduction variable. */
+/* Nonzero if this map clause is for an OpenACC compute construct's reduction
+ variable. */
#define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+2019-11-12 Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Frederik Harwath <frederik@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ libgomp/
+ * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: New test.
+ * testsuite/libgomp.oacc-fortran/parallel-dims-aux.c: New test.
+ * testsuite/libgomp.oacc-fortran/parallel-dims.f89: New test.
+
2019-11-11 Tobias Burnus <tobias@codesourcery.com>
Kwok Cheung Yeung <kcy@codesourcery.com>
/* OpenACC parallelism dimensions clauses: num_gangs, num_workers,
vector_length. */
+/* See also '../libgomp.oacc-fortran/parallel-dims.f90'. */
+
#include <limits.h>
#include <openacc.h>
#include <gomp-constants.h>
{
acc_init (acc_device_default);
+ /* OpenACC parallel construct. */
+
/* Non-positive value. */
/* GR, WS, VS. */
}
+ /* OpenACC kernels construct. */
+
/* We can't test parallelized OpenACC kernels constructs in this way: use of
the acc_gang, acc_worker, acc_vector functions will make the construct
unparallelizable. */
}
+ /* OpenACC serial construct. */
+
+ /* GR, WS, VS. */
+ {
+ int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
+ gangs_min = workers_min = vectors_min = INT_MAX;
+ gangs_max = workers_max = vectors_max = INT_MIN;
+#pragma acc serial /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
+ reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ {
+ for (int i = 100; i > -100; i--)
+ {
+ gangs_min = gangs_max = acc_gang ();
+ workers_min = workers_max = acc_worker ();
+ vectors_min = vectors_max = acc_vector ();
+ }
+ }
+ if (gangs_min != 0 || gangs_max != 1 - 1
+ || workers_min != 0 || workers_max != 1 - 1
+ || vectors_min != 0 || vectors_max != 1 - 1)
+ __builtin_abort ();
+ }
+
+ /* Composition of GP, WP, VP. */
+ {
+ int vectors_actual = 1; /* Implicit 'vector_length (1)' clause. */
+ int gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max;
+ gangs_min = workers_min = vectors_min = INT_MAX;
+ gangs_max = workers_max = vectors_max = INT_MIN;
+#pragma acc serial copy (vectors_actual) /* { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } } */ \
+ copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max)
+ {
+ if (acc_on_device (acc_device_nvidia))
+ {
+ /* The GCC nvptx back end enforces vector_length (32). */
+ /* It's unclear if that's actually permissible here;
+ <https://github.com/OpenACC/openacc-spec/issues/238> "OpenACC
+ 'serial' construct might not actually be serial". */
+ vectors_actual = 32;
+ }
+#pragma acc loop gang reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ for (int i = 100; i > -100; i--)
+#pragma acc loop worker reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ for (int j = 100; j > -100; j--)
+#pragma acc loop vector reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ for (int k = 100 * vectors_actual; k > -100 * vectors_actual; k--)
+ {
+ gangs_min = gangs_max = acc_gang ();
+ workers_min = workers_max = acc_worker ();
+ vectors_min = vectors_max = acc_vector ();
+ }
+ }
+ if (acc_get_device_type () == acc_device_nvidia)
+ {
+ if (vectors_actual != 32)
+ __builtin_abort ();
+ }
+ else
+ if (vectors_actual != 1)
+ __builtin_abort ();
+ if (gangs_min != 0 || gangs_max != 1 - 1
+ || workers_min != 0 || workers_max != 1 - 1
+ || vectors_min != 0 || vectors_max != vectors_actual - 1)
+ __builtin_abort ();
+ }
+
+
return 0;
}
--- /dev/null
+/* OpenACC parallelism dimensions clauses: num_gangs, num_workers,
+ vector_length. */
+
+/* Copied from '../libgomp.oacc-c-c++-common/parallel-dims.c'. */
+
+/* Used by 'parallel-dims.f90'. */
+
+#include <limits.h>
+#include <openacc.h>
+#include <gomp-constants.h>
+
+/* TODO: "(int) acc_device_*" casts because of the C++ acc_on_device wrapper
+ not behaving as expected for -O0. */
+#pragma acc routine seq
+/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_gang ()
+{
+ if (acc_on_device ((int) acc_device_host))
+ return 0;
+ else if (acc_on_device ((int) acc_device_nvidia))
+ return __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
+ else
+ __builtin_abort ();
+}
+
+#pragma acc routine seq
+/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_worker ()
+{
+ if (acc_on_device ((int) acc_device_host))
+ return 0;
+ else if (acc_on_device ((int) acc_device_nvidia))
+ return __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
+ else
+ __builtin_abort ();
+}
+
+#pragma acc routine seq
+/* static */ unsigned int __attribute__ ((optimize ("O2"))) acc_vector ()
+{
+ if (acc_on_device ((int) acc_device_host))
+ return 0;
+ else if (acc_on_device ((int) acc_device_nvidia))
+ return __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
+ else
+ __builtin_abort ();
+}
--- /dev/null
+! OpenACC parallelism dimensions clauses: num_gangs, num_workers,
+! vector_length.
+
+! { dg-additional-sources parallel-dims-aux.c }
+! { dg-do run }
+! { dg-prune-output "command-line option '-fintrinsic-modules-path=.*' is valid for Fortran but not for C" }
+
+! See also '../libgomp.oacc-c-c++-common/parallel-dims.c'.
+
+module acc_routines
+ implicit none (type, external)
+
+ interface
+ integer function acc_gang() bind(C)
+ !$acc routine seq
+ end function acc_gang
+
+ integer function acc_worker() bind(C)
+ !$acc routine seq
+ end function acc_worker
+
+ integer function acc_vector() bind(C)
+ !$acc routine seq
+ end function acc_vector
+ end interface
+end module acc_routines
+
+program main
+ use iso_c_binding
+ use openacc
+ use acc_routines
+ implicit none (type, external)
+
+ integer :: gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max
+ integer :: vectors_actual
+ integer :: i, j, k
+
+ call acc_init (acc_device_default)
+
+ ! OpenACC parallel construct.
+
+ !TODO
+
+
+ ! OpenACC kernels construct.
+
+ !TODO
+
+
+ ! OpenACC serial construct.
+
+ ! GR, WS, VS.
+
+ gangs_min = huge(gangs_min) ! INT_MAX
+ workers_min = huge(workers_min) ! INT_MAX
+ vectors_min = huge(vectors_min) ! INT_MAX
+ gangs_max = -huge(gangs_max) - 1 ! INT_MIN
+ workers_max = -huge(gangs_max) - 1 ! INT_MIN
+ vectors_max = -huge(gangs_max) - 1 ! INT_MIN
+ !$acc serial &
+ !$acc reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
+ do i = 100, -99, -1
+ gangs_min = acc_gang ();
+ gangs_max = acc_gang ();
+ workers_min = acc_worker ();
+ workers_max = acc_worker ();
+ vectors_min = acc_vector ();
+ vectors_max = acc_vector ();
+ end do
+ !$acc end serial
+ if (gangs_min /= 0 .or. gangs_max /= 1 - 1 &
+ .or. workers_min /= 0 .or. workers_max /= 1 - 1 &
+ .or. vectors_min /= 0 .or. vectors_max /= 1 - 1) &
+ stop 1
+
+ ! Composition of GP, WP, VP.
+
+ vectors_actual = 1 ! Implicit 'vector_length (1)' clause.
+ gangs_min = huge(gangs_min) ! INT_MAX
+ workers_min = huge(workers_min) ! INT_MAX
+ vectors_min = huge(vectors_min) ! INT_MAX
+ gangs_max = -huge(gangs_max) - 1 ! INT_MIN
+ workers_max = -huge(gangs_max) - 1 ! INT_MIN
+ vectors_max = -huge(gangs_max) - 1 ! INT_MIN
+ !$acc serial copy (vectors_actual) &
+ !$acc copy (gangs_min, gangs_max, workers_min, workers_max, vectors_min, vectors_max) ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } }
+ if (acc_on_device (acc_device_nvidia)) then
+ ! The GCC nvptx back end enforces vector_length (32).
+ ! It's unclear if that's actually permissible here;
+ ! <https://github.com/OpenACC/openacc-spec/issues/238> "OpenACC 'serial'
+ ! construct might not actually be serial".
+ vectors_actual = 32
+ end if
+ !$acc loop gang reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ do i = 100, -99, -1
+ !$acc loop worker reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ do j = 100, -99, -1
+ !$acc loop vector reduction (min: gangs_min, workers_min, vectors_min) reduction (max: gangs_max, workers_max, vectors_max)
+ do k = 100 * vectors_actual, -99 * vectors_actual, -1
+ gangs_min = acc_gang ();
+ gangs_max = acc_gang ();
+ workers_min = acc_worker ();
+ workers_max = acc_worker ();
+ vectors_min = acc_vector ();
+ vectors_max = acc_vector ();
+ end do
+ end do
+ end do
+ !$acc end serial
+ if (acc_get_device_type () .eq. acc_device_nvidia) then
+ if (vectors_actual /= 32) stop 2
+ else
+ if (vectors_actual /= 1) stop 3
+ end if
+ if (gangs_min /= 0 .or. gangs_max /= 1 - 1 &
+ .or. workers_min /= 0 .or. workers_max /= 1 - 1 &
+ .or. vectors_min /= 0 .or. vectors_max /= vectors_actual - 1) &
+ stop 4
+
+end program main
projects / gcc.git / commitdiff