[PR89433] Refer to OpenACC 'routine' clauses from "omp declare target" attribute

	gcc/c-family/
	PR c/89433
	* c-attribs.c (c_common_attribute_table): Set min_len to -1 for
	"omp declare target".
	gcc/c/
	PR c/89433
	* c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
	clauses from "omp declare target" attribute.
	gcc/cp/
	PR c++/89433
	* parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
	clauses from "omp declare target" attribute.
	gcc/fortran/
	PR fortran/89433
	* f95-lang.c (gfc_attribute_table): Set min_len to -1 for "omp
	declare target".
	* trans-decl.c (add_attributes_to_decl): Refer to OpenACC
	'routine' clauses from "omp declare target" attribute.
	gcc/testsuite/
	PR testsuite/89433
	* c-c++-common/goacc/classify-routine.c: Update.
	* gfortran.dg/goacc/classify-routine.f95: Likewise.

From-SVN: r271343

12 files changed, 51 insertions, 15 deletions

diff --git a/gcc/c-family/ChangeLog b/gcc/c-family/ChangeLog
index 47c1d3d..5068776 100644
--- a/gcc/c-family/ChangeLog
+++ b/gcc/c-family/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR c/89433
+	* c-attribs.c (c_common_attribute_table): Set min_len to -1 for
+	"omp declare target".
+
 2019-05-16  Martin Sebor  <msebor@redhat.com>
 
         * c-attribs.c (handle_no_sanitize_attribute): Quote identifiers,
diff --git a/gcc/c-family/c-attribs.c b/gcc/c-family/c-attribs.c
index 12c0b9b..0320347 100644
--- a/gcc/c-family/c-attribs.c
+++ b/gcc/c-family/c-attribs.c
@@ -437,7 +437,7 @@ const struct attribute_spec c_common_attribute_table[] =
 			      handle_omp_declare_simd_attribute, NULL },
   { "simd",		      0, 1, true,  false, false, false,
 			      handle_simd_attribute, NULL },
-  { "omp declare target",     0, 0, true, false, false, false,
+  { "omp declare target",     0, -1, true, false, false, false,
 			      handle_omp_declare_target_attribute, NULL },
   { "omp declare target link", 0, 0, true, false, false, false,
 			      handle_omp_declare_target_attribute, NULL },
diff --git a/gcc/c/ChangeLog b/gcc/c/ChangeLog
index f558d28..f0cab2e 100644
--- a/gcc/c/ChangeLog
+++ b/gcc/c/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR c/89433
+	* c-parser.c (c_finish_oacc_routine): Refer to OpenACC 'routine'
+	clauses from "omp declare target" attribute.
+
 2019-05-16  Martin Sebor  <msebor@redhat.com>
 
         * c-decl.c (start_decl): Quote keywords, operators, and
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 993cfe0..3cbbb19 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -15904,7 +15904,7 @@ c_finish_oacc_routine (struct oacc_routine_data *data, tree fndecl,
   /* Add an "omp declare target" attribute.  */
   DECL_ATTRIBUTES (fndecl)
     = tree_cons (get_identifier ("omp declare target"),
-		 NULL_TREE, DECL_ATTRIBUTES (fndecl));
+		 data->clauses, DECL_ATTRIBUTES (fndecl));
 
   /* Remember that we've used this "#pragma acc routine".  */
   data->fndecl_seen = true;
diff --git a/gcc/cp/ChangeLog b/gcc/cp/ChangeLog
index 08b7d53..40622ac 100644
--- a/gcc/cp/ChangeLog
+++ b/gcc/cp/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR c++/89433
+	* parser.c (cp_finalize_oacc_routine): Refer to OpenACC 'routine'
+	clauses from "omp declare target" attribute.
+
 2019-05-16  Martin Sebor  <msebor@redhat.com>
 
         * call.c (print_z_candidate): Wrap diagnostic text in a gettext
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index e6ef5a9..15424b6 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40292,7 +40292,7 @@ cp_finalize_oacc_routine (cp_parser *parser, tree fndecl, bool is_defn)
       /* Add an "omp declare target" attribute.  */
       DECL_ATTRIBUTES (fndecl)
 	= tree_cons (get_identifier ("omp declare target"),
-		     NULL_TREE, DECL_ATTRIBUTES (fndecl));
+		     parser->oacc_routine->clauses, DECL_ATTRIBUTES (fndecl));
 
       /* Don't unset parser->oacc_routine here: we may still need it to
 	 diagnose wrong usage.  But, remember that we've used this "#pragma acc
diff --git a/gcc/fortran/ChangeLog b/gcc/fortran/ChangeLog
index b885106..f09e715 100644
--- a/gcc/fortran/ChangeLog
+++ b/gcc/fortran/ChangeLog
@@ -1,3 +1,11 @@
+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR fortran/89433
+	* f95-lang.c (gfc_attribute_table): Set min_len to -1 for "omp
+	declare target".
+	* trans-decl.c (add_attributes_to_decl): Refer to OpenACC
+	'routine' clauses from "omp declare target" attribute.
+
 2019-05-16  Martin Sebor  <msebor@redhat.com>
 
 	* gfortranspec.c (append_arg): Spell out the word "argument."
diff --git a/gcc/fortran/f95-lang.c b/gcc/fortran/f95-lang.c
index 3e3d304..6b9f490 100644
--- a/gcc/fortran/f95-lang.c
+++ b/gcc/fortran/f95-lang.c
@@ -91,7 +91,7 @@ static const struct attribute_spec gfc_attribute_table[] =
 {
   /* { name, min_len, max_len, decl_req, type_req, fn_type_req,
        affects_type_identity, handler, exclude } */
-  { "omp declare target", 0, 0, true,  false, false, false,
+  { "omp declare target", 0, -1, true,  false, false, false,
     gfc_handle_omp_declare_target_attribute, NULL },
   { "omp declare target link", 0, 0, true,  false, false, false,
     gfc_handle_omp_declare_target_attribute, NULL },
diff --git a/gcc/fortran/trans-decl.c b/gcc/fortran/trans-decl.c
index 07d1c33..8420870 100644
--- a/gcc/fortran/trans-decl.c
+++ b/gcc/fortran/trans-decl.c
@@ -1400,12 +1400,7 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
 	list = chainon (list, attr);
       }
 
-  if (sym_attr.omp_declare_target_link)
-    list = tree_cons (get_identifier ("omp declare target link"),
-		      NULL_TREE, list);
-  else if (sym_attr.omp_declare_target)
-    list = tree_cons (get_identifier ("omp declare target"),
-		      NULL_TREE, list);
+  tree clauses = NULL_TREE;
 
   if (sym_attr.oacc_routine_lop != OACC_ROUTINE_LOP_NONE)
     {
@@ -1430,11 +1425,20 @@ add_attributes_to_decl (symbol_attribute sym_attr, tree list)
 	  gcc_unreachable ();
 	}
       tree c = build_omp_clause (UNKNOWN_LOCATION, code);
+      OMP_CLAUSE_CHAIN (c) = clauses;
+      clauses = c;
 
-      tree dims = oacc_build_routine_dims (c);
+      tree dims = oacc_build_routine_dims (clauses);
       list = oacc_replace_fn_attrib_attr (list, dims);
     }
 
+  if (sym_attr.omp_declare_target_link)
+    list = tree_cons (get_identifier ("omp declare target link"),
+		      NULL_TREE, list);
+  else if (sym_attr.omp_declare_target)
+    list = tree_cons (get_identifier ("omp declare target"),
+		      clauses, list);
+
   return list;
 }
 
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index e1aa257..0f44c6d 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-05-17  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR testsuite/89433
+	* c-c++-common/goacc/classify-routine.c: Update.
+	* gfortran.dg/goacc/classify-routine.f95: Likewise.
+
 2019-05-16  Martin Sebor  <msebor@redhat.com>
 
         * c-c++-common/Wbool-operation-1.c: Adjust text of expected diagnostics.
diff --git a/gcc/testsuite/c-c++-common/goacc/classify-routine.c b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
index a723d2c..0b9ba6e 100644
--- a/gcc/testsuite/c-c++-common/goacc/classify-routine.c
+++ b/gcc/testsuite/c-c++-common/goacc/classify-routine.c
@@ -22,10 +22,10 @@ void ROUTINE ()
 }
 
 /* Check the offloaded function's attributes.
-   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
+   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } } */
 
 /* Check the offloaded function's classification and compute dimensions (will
    always be 1 x 1 x 1 for non-offloading compilation).
    { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
    { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
-   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */
+   { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\), oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95 b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
index e435f5d..401d527 100644
--- a/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/classify-routine.f95
@@ -21,10 +21,10 @@ subroutine ROUTINE
 end subroutine ROUTINE
 
 ! Check the offloaded function's attributes.
-! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "ompexp" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 0, 1 0\\), omp declare target \\(worker\\)\\)\\)" 1 "ompexp" } }
 
 ! Check the offloaded function's classification and compute dimensions (will
 ! always be 1 x 1 x 1 for non-offloading compilation).
 ! { dg-final { scan-tree-dump-times "(?n)Function is OpenACC routine level 1" 1 "oaccdevlow" } }
 ! { dg-final { scan-tree-dump-times "(?n)Compute dimensions \\\[1, 1, 1\\\]" 1 "oaccdevlow" } }
-! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target, oacc function \\(0 1, 1 0, 1 0\\)\\)\\)" 1 "oaccdevlow" } }
+! { dg-final { scan-tree-dump-times "(?n)__attribute__\\(\\(oacc function \\(0 1, 1 1, 1 1\\), omp declare target \\(worker\\)\\)\\)" 1 "oaccdevlow" } }