+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * tree-core.h (omp_clause_code): Add OMP_CLAUSE_NOHOST.
+ * tree.c (omp_clause_num_ops, omp_clause_code_name, walk_tree_1):
+ Update for these.
+ * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_NOHOST.
+ * gimplify.c (gimplify_scan_omp_clauses)
+ (gimplify_adjust_omp_clauses): Handle OMP_CLAUSE_NOHOST.
+ * tree-nested.c (convert_nonlocal_omp_clauses)
+ (convert_local_omp_clauses): Likewise.
+ * omp-low.c (scan_sharing_clauses): Likewise.
+ * omp-offload.c (maybe_discard_oacc_function): New function.
+ (execute_oacc_device_lower) [!ACCEL_COMPILER]: Handle OpenACC
+ nohost clauses.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
--- /dev/null
+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_NOHST.
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
PRAGMA_OACC_CLAUSE_NO_CREATE,
+ PRAGMA_OACC_CLAUSE_NOHOST,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * c-parser.c (c_parser_omp_clause_name): Handle "nohost".
+ (c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST.
+ (c_parser_oacc_routine): Update.
+ * c-typeck.c (c_finish_omp_clauses): Handle OMP_CLAUSE_NOHOST.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
result = PRAGMA_OMP_CLAUSE_NOTINBRANCH;
else if (!strcmp ("nowait", p))
result = PRAGMA_OMP_CLAUSE_NOWAIT;
+ else if (!strcmp ("nohost", p))
+ result = PRAGMA_OACC_CLAUSE_NOHOST;
else if (!strcmp ("num_gangs", p))
result = PRAGMA_OACC_CLAUSE_NUM_GANGS;
else if (!strcmp ("num_tasks", p))
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "no_create";
break;
+ case PRAGMA_OACC_CLAUSE_NOHOST:
+ clauses = c_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST,
+ clauses);
+ c_name = "nohost";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) )
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST) )
/* Parse an OpenACC routine directive. For named directives, we apply
immediately to the named function. For unnamed ones we then parse
case OMP_CLAUSE_GANG:
case OMP_CLAUSE_WORKER:
case OMP_CLAUSE_VECTOR:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * parser.c (cp_parser_omp_clause_name): Handle "nohost".
+ (cp_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_NOHOST,
+ (cp_parser_oacc_routine): Update.
+ * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_NOHOST.
+ * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_NOHOST.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
-
* semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
parameter, adjust recursive call site, add cases for allowing
pointer based multi-dimensional arrays for OpenACC.
result = PRAGMA_OACC_CLAUSE_NO_CREATE;
else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
+ else if (!strcmp ("nohost", p))
+ result = PRAGMA_OACC_CLAUSE_NOHOST;
else if (!strcmp ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
else if (!strcmp ("notinbranch", p))
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "no_create";
break;
+ case PRAGMA_OACC_CLAUSE_NOHOST:
+ clauses = cp_parser_oacc_simple_clause (here, OMP_CLAUSE_NOHOST,
+ clauses);
+ c_name = "nohost";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_GANG) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WORKER) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_VECTOR) \
- | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ))
-
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_SEQ) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NOHOST) )
/* Parse the OpenACC routine pragma. This has an optional '( name )'
component, which must resolve to a declared namespace-scope
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
case OMP_CLAUSE_SEQ:
case OMP_CLAUSE_IF_PRESENT:
case OMP_CLAUSE_FINALIZE:
+ case OMP_CLAUSE_NOHOST:
break;
case OMP_CLAUSE_MERGEABLE:
+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * gfortran.h (gfc_omp_clauses): Add nohost members.
+ * openmp.c (omp_mask2): Add OMP_CLAUSE_NOHOST.
+ (gfc_match_omp_clauses): Handle OMP_CLAUSE_NOHOST.
+ (gfc_match_oacc_routine): Set oacc_function_nohost when appropriate.
+ * gfortran.h (symbol_attribute): Add oacc_function_nohost member.
+ * trans-openmp.c (gfc_add_omp_offload_attributes): Use it to decide
+ whether to generate an OMP_CLAUSE_NOHOST clause.
+ (gfc_trans_omp_clauses_1): Unreachable code to generate an
+ OMP_CLAUSE_NOHOST clause.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
/* OpenACC 'routine' directive's level of parallelism. */
ENUM_BITFIELD (oacc_routine_lop) oacc_routine_lop:3;
+ unsigned oacc_function_nohost:1;
/* Attributes set by compiler extensions (!GCC$ ATTRIBUTES). */
unsigned ext_attr:EXT_ATTR_NUM;
gfc_expr_list *tile_list;
unsigned async:1, gang:1, worker:1, vector:1, seq:1, independent:1;
unsigned par_auto:1, gang_static:1;
- unsigned if_present:1, finalize:1;
+ unsigned if_present:1, finalize:1, nohost:1;
locus loc;
}
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_ATTACH,
+ OMP_CLAUSE_NOHOST,
/* This must come last. */
OMP_MASK2_LAST
};
&c->lists[OMP_LIST_NONTEMPORAL],
true) == MATCH_YES)
continue;
+ if ((mask & OMP_CLAUSE_NOHOST)
+ && !c->nohost
+ && gfc_match ("nohost") == MATCH_YES)
+ {
+ c->nohost = true;
+ continue;
+ }
if ((mask & OMP_CLAUSE_NOTINBRANCH)
&& !c->notinbranch
&& !c->inbranch
omp_mask (OMP_CLAUSE_ASYNC)
#define OACC_ROUTINE_CLAUSES \
(omp_mask (OMP_CLAUSE_GANG) | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR \
- | OMP_CLAUSE_SEQ)
+ | OMP_CLAUSE_SEQ \
+ | OMP_CLAUSE_NOHOST)
static match
&old_loc))
goto cleanup;
gfc_current_ns->proc_name->attr.oacc_routine_lop = lop;
+ gfc_current_ns->proc_name->attr.oacc_function_nohost
+ = c ? c->nohost : false;
}
else
/* Something has gone wrong, possibly a syntax error. */
|| sym_attr.oacc_declare_copyin
|| sym_attr.oacc_declare_deviceptr
|| sym_attr.oacc_declare_device_resident)
- list = tree_cons (get_identifier ("omp declare target"),
- clauses, list);
+ {
+ list = tree_cons (get_identifier ("omp declare target"),
+ clauses, list);
+ tree c = clauses;
+ if (sym_attr.oacc_function_nohost)
+ c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_NOHOST);
+ list = tree_cons (get_identifier ("omp declare target"), c, list);
+ }
return list;
}
}
}
-
static inline tree
gfc_trans_add_clause (tree node, tree tail)
{
OMP_CLAUSE_GANG_STATIC_EXPR (c) = arg;
}
}
+ if (clauses->nohost)
+ {
+ c = build_omp_clause (where.lb->location, OMP_CLAUSE_NOHOST);
+ omp_clauses = gfc_trans_add_clause (c, omp_clauses);
+ //TODO
+ gcc_unreachable();
+ }
return nreverse (omp_clauses);
}
}
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
case OMP_CLAUSE_EXCLUSIVE:
break;
+ case OMP_CLAUSE_NOHOST:
default:
gcc_unreachable ();
}
install_var_local (decl, ctx);
break;
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE__CACHE_:
default:
gcc_unreachable ();
case OMP_CLAUSE__CONDTEMP_:
break;
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE__CACHE_:
default:
gcc_unreachable ();
gsi_replace_with_seq (&gsi, seq, true);
}
+/* Determine whether DECL should be discarded in this offload
+ compilation. */
+
+static bool
+maybe_discard_oacc_function (tree decl)
+{
+ tree attr = lookup_attribute ("omp declare target", DECL_ATTRIBUTES (decl));
+
+ if (!attr)
+ return false;
+
+ enum omp_clause_code kind = OMP_CLAUSE_NOHOST;
+
+ if (omp_find_clause (TREE_VALUE (attr), kind))
+ return true;
+
+ return false;
+}
+
/* Main entry point for oacc transformations which run on the device
compiler after LTO, so we know what the target device is at this
point (including the host fallback). */
static unsigned int
execute_oacc_device_lower ()
{
- tree attrs = oacc_get_fn_attrib (current_function_decl);
-
- if (!attrs)
+ tree attr = oacc_get_fn_attrib (current_function_decl);
+ if (!attr)
/* Not an offloaded function. */
return 0;
+ if (maybe_discard_oacc_function (current_function_decl))
+ {
+ if (dump_file)
+ fprintf (dump_file, "Discarding function\n");
+ TREE_ASM_WRITTEN (current_function_decl) = 1;
+ return TODO_discard_function;
+ }
+
/* Parse the default dim argument exactly once. */
if ((const void *)flag_openacc_dims != &flag_openacc_dims)
{
bool is_oacc_parallel_kernels_gang_single
= (lookup_attribute ("oacc parallel_kernels_gang_single",
DECL_ATTRIBUTES (current_function_decl)) != NULL);
- int fn_level = oacc_fn_attrib_level (attrs);
+ int fn_level = oacc_fn_attrib_level (attr);
bool is_oacc_routine = (fn_level >= 0);
gcc_checking_assert (is_oacc_parallel
+ is_oacc_kernels
if (is_oacc_kernels && !is_oacc_kernels_parallelized)
{
oacc_set_fn_attrib (current_function_decl, NULL, NULL);
- attrs = oacc_get_fn_attrib (current_function_decl);
+ attr = oacc_get_fn_attrib (current_function_decl);
}
/* Discover, partition and process the loops. */
oacc_loop *loops = oacc_loop_discovery ();
+ fn_level = oacc_fn_attrib_level (attr);
unsigned outer_mask = 0;
if (is_oacc_routine)
}
int dims[GOMP_DIM_MAX];
- oacc_validate_dims (current_function_decl, attrs, dims, fn_level, used_mask);
+ oacc_validate_dims (current_function_decl, attr, dims, fn_level, used_mask);
if (dump_file)
{
+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * c-c++-common/goacc/routine-1.c: Adjust test.
+ * c-c++-common/goacc/routine-2.c: Likewise.
+ * c-c++-common/goacc/routine-nohost-1.c: New test.
+ * g++.dg/goacc/routine-2.C: Adjust test.
+ * gfortran.dg/goacc/pr72741.f90: New test.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
+/* Test valid use of clauses with routine. */
#pragma acc routine gang
void gang (void)
{
}
+#pragma acc routine nohost
+void nohost (void)
+{
+}
+
int main ()
{
#pragma acc kernels num_gangs (32) num_workers (32) vector_length (32)
worker ();
vector ();
seq ();
+ nohost ();
}
#pragma acc parallel num_gangs (32) num_workers (32) vector_length (32)
worker ();
vector ();
seq ();
+ nohost ();
}
return 0;
/* Test invalid use of the OpenACC 'routine' directive. */
+#pragma acc routine gang worker /* { dg-error "conflicting level" } */
+void gang (void)
+{
+}
+
+#pragma acc routine worker vector /* { dg-error "conflicting level" } */
+void worker (void)
+{
+}
+
+#pragma acc routine vector seq /* { dg-error "conflicting level" } */
+void vector (void)
+{
+}
+
+#pragma acc routine seq gang /* { dg-error "conflicting level" } */
+void seq (void)
+{
+}
+
#pragma acc routine (nothing) gang /* { dg-error "not been declared" } */
--- /dev/null
+/* Test the nohost clause for OpenACC routine directive. Exercising different
+ variants for declaring routines. */
+
+/* { dg-additional-options "-fdump-tree-oaccdevlow" } */
+
+#pragma acc routine nohost
+int THREE(void)
+{
+ return 3;
+}
+
+#pragma acc routine nohost
+extern void NOTHING(void);
+
+void NOTHING(void)
+{
+}
+
+extern float ADD(float, float);
+
+#pragma acc routine (ADD) nohost
+
+float ADD(float x, float y)
+{
+ return x + y;
+}
+
+/* { dg-final { scan-tree-dump-times "Discarding function" 3 "oaccdevlow" } } */
template <typename T>
extern T one_d();
-#pragma acc routine (one_d) /* { dg-error "names a set of overloads" } */
+#pragma acc routine (one_d) nohost /* { dg-error "names a set of overloads" } */
-template <typename T>
-T
-one()
-{
- return 1;
-}
-#pragma acc routine (one) /* { dg-error "names a set of overloads" } */
int incr (int);
float incr (float);
--- /dev/null
+SUBROUTINE v_1
+ !$ACC ROUTINE VECTOR WORKER ! { dg-error "Multiple loop axes" }
+END SUBROUTINE v_1
+
+SUBROUTINE sub_1
+ IMPLICIT NONE
+ EXTERNAL :: g_1
+ !$ACC ROUTINE (g_1) GANG WORKER ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE (ABORT) SEQ VECTOR ! { dg-error "Multiple loop axes" "" { xfail *-*-* } }
+! { dg-bogus "invalid function name abort" "" { xfail *-*-* } .-1 }
+
+ CALL v_1
+ CALL g_1
+ CALL ABORT
+END SUBROUTINE sub_1
+
+MODULE m_w_1
+ IMPLICIT NONE
+ EXTERNAL :: w_1
+ !$ACC ROUTINE (w_1) WORKER SEQ ! { dg-error "Multiple loop axes" }
+ !$ACC ROUTINE (ABORT) VECTOR GANG ! { dg-error "Multiple loop axes" "" { xfail *-*-* } }
+! { dg-bogus "invalid function name abort" "" { xfail *-*-* } .-1 }
+
+CONTAINS
+ SUBROUTINE sub_2
+ CALL v_1
+ CALL w_1
+ CALL ABORT
+ END SUBROUTINE sub_2
+END MODULE m_w_1
/* OpenACC clause: vector_length (integer-expression). */
OMP_CLAUSE_VECTOR_LENGTH,
+ /* OpenACC clause: nohost. */
+ OMP_CLAUSE_NOHOST,
+
/* OpenACC clause: tile ( size-expr-list ). */
OMP_CLAUSE_TILE,
case OMP_CLAUSE__SCANTEMP_:
break;
+ /* OpenACC nohost clause is not yet handled here. */
+ case OMP_CLAUSE_NOHOST:
/* The following clause belongs to the OpenACC cache directive, which
is discarded during gimplification. */
case OMP_CLAUSE__CACHE_:
case OMP_CLAUSE__SCANTEMP_:
break;
+ /* OpenACC nohost clauses is not yet handled here. */
+ case OMP_CLAUSE_NOHOST:
/* The following clause belongs to the OpenACC cache directive, which
is discarded during gimplification. */
case OMP_CLAUSE__CACHE_:
spc, flags, false);
pp_right_paren (pp);
break;
+ case OMP_CLAUSE_NOHOST:
+ pp_string (pp, "nohost");
+ break;
case OMP_CLAUSE_IF_PRESENT:
pp_string (pp, "if_present");
1, /* OMP_CLAUSE_NUM_GANGS */
1, /* OMP_CLAUSE_NUM_WORKERS */
1, /* OMP_CLAUSE_VECTOR_LENGTH */
+ 0, /* OMP_CLAUSE_NOHOST */
3, /* OMP_CLAUSE_TILE */
0, /* OMP_CLAUSE_IF_PRESENT */
0, /* OMP_CLAUSE_FINALIZE */
"num_gangs",
"num_workers",
"vector_length",
+ "nohost",
"tile",
"if_present",
"finalize",
case OMP_CLAUSE_BIND:
case OMP_CLAUSE_AUTO:
case OMP_CLAUSE_SEQ:
+ case OMP_CLAUSE_NOHOST:
case OMP_CLAUSE_TILE:
case OMP_CLAUSE__SIMT_:
case OMP_CLAUSE_IF_PRESENT:
+2018-10-02 Thomas Schwinge <thomas@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/routine-3.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/routine-nohost-1.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/routine-bind-nohost-1.c:
+ Update test.
+ * testsuite/libgomp.oacc-fortran/routine-6.f90: Likewise.
+
2020-04-19 Chung-Lin Tang <cltang@codesourcery.com>
PR other/76739
--- /dev/null
+/* At -O0, we do get the expected "undefined reference to `foo'" link-time
+ error message (but the check needs to be done differently; compare to
+ routine-nohost-1.c), but for -O2 we don't; presumably because the function
+ gets inlined.
+ { dg-xfail-if "TODO" { *-*-* } { "-O0" } { "" } } */
+
+#include <stdlib.h>
+
+#pragma acc routine nohost
+int
+foo (int n)
+{
+ if (n == 0 || n == 1)
+ return 1;
+
+ return n * n;
+}
+
+int
+main()
+{
+ int a, n = 10;
+
+#pragma acc parallel copy (a, n)
+ {
+ a = foo (n);
+ }
+
+ if (a != n * n)
+ abort ();
+
+ return 0;
+}
--- /dev/null
+/* { dg-do link } */
+
+extern int three (void);
+
+#pragma acc routine (three) nohost
+__attribute__((noinline))
+int three(void)
+{
+ return 3;
+}
+
+int main(void)
+{
+ return (three() == 3) ? 0 : 1;
+}
+
+/* Expecting link to fail; "undefined reference to `three'" (or similar).
+ { dg-excess-errors "" } */
--- /dev/null
+! { dg-do run }
+! { dg-xfail-if "TODO" { *-*-* } }
+
+program main
+ integer :: a, n
+
+ n = 10
+
+ !$acc parallel copy (a, n)
+ a = foo (n)
+ !$acc end parallel
+
+ if (a .ne. n * n) call abort
+
+contains
+
+function foo (n) result (rc)
+ !$acc routine nohost
+
+ integer, intent (in) :: n
+ integer :: rc
+
+ rc = n * n
+
+end function
+
+end program main
+
projects / thirdparty / gcc.git / commitdiff
