OpenMP, C++: Add template support for the has_device_addr clause.

This patch adds support for list items in the has_device_addr clause which type
is given by C++ template parameters.

gcc/cp/ChangeLog:

	* pt.cc (tsubst_omp_clauses): Added OMP_CLAUSE_HAS_DEVICE_ADDR.
	* semantics.cc (finish_omp_clauses): Added template decl processing.

libgomp/ChangeLog:

	* testsuite/libgomp.c++/target-has-device-addr-7.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-8.C: New test.
	* testsuite/libgomp.c++/target-has-device-addr-9.C: New test.

gcc/cp/pt.cc

@@ -17722,6 +17722,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	case OMP_CLAUSE_USE_DEVICE_PTR:
 	case OMP_CLAUSE_USE_DEVICE_ADDR:
 	case OMP_CLAUSE_IS_DEVICE_PTR:
+	case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	case OMP_CLAUSE_INCLUSIVE:
 	case OMP_CLAUSE_EXCLUSIVE:
 	  OMP_CLAUSE_DECL (nc)
@@ -17867,6 +17868,7 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort,
 	  case OMP_CLAUSE_USE_DEVICE_PTR:
 	  case OMP_CLAUSE_USE_DEVICE_ADDR:
 	  case OMP_CLAUSE_IS_DEVICE_PTR:
+	  case OMP_CLAUSE_HAS_DEVICE_ADDR:
 	  case OMP_CLAUSE_INCLUSIVE:
 	  case OMP_CLAUSE_EXCLUSIVE:
 	  case OMP_CLAUSE_ALLOCATE:

gcc/cp/semantics.cc

@@ -8575,14 +8575,20 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      else
 		{
 		  t = OMP_CLAUSE_DECL (c);
+		  while (TREE_CODE (t) == TREE_LIST)
+		    t = TREE_CHAIN (t);
 		  while (TREE_CODE (t) == INDIRECT_REF
 			 || TREE_CODE (t) == ARRAY_REF)
 		    t = TREE_OPERAND (t, 0);
 		}
 	    }
-	  bitmap_set_bit (&is_on_device_head, DECL_UID (t));
 	  if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
-	    cxx_mark_addressable (t);
+	    {
+	      bitmap_set_bit (&is_on_device_head, DECL_UID (t));
+	      if (!processing_template_decl
+		  && !cxx_mark_addressable (t))
+		remove = true;
+	    }
 	  goto check_dup_generic_t;
 
 	case OMP_CLAUSE_USE_DEVICE_ADDR:

libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C

@@ -0,0 +1,36 @@
+/* Testing 'has_device_addr' clause on the target construct with template.  */
+
+template <typename T>
+void
+foo (T x)
+{
+  x = 24;
+  #pragma omp target data map(x) use_device_addr(x)
+    #pragma omp target has_device_addr(x)
+      x = 42;
+
+  if (x != 42)
+    __builtin_abort ();
+}
+
+template <typename T>
+void
+bar (T (&x)[])
+{
+  x[0] = 24;
+  #pragma omp target data map(x[:2]) use_device_addr(x)
+    #pragma omp target has_device_addr(x[:2])
+      x[0] = 42;
+
+  if (x[0] != 42)
+    __builtin_abort ();
+}
+
+int
+main ()
+{
+  int a[] = { 24, 42};
+  foo <int> (42);
+  bar <int> (a);
+  return 0;
+}

libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C

@@ -0,0 +1,47 @@
+/* Testing 'has_device_addr' clause on the target construct with template.  */
+
+#include <omp.h>
+
+template <typename T>
+void
+foo (T (&x)[])
+{
+  #pragma omp target has_device_addr(x)
+    for (int i = 0; i < 15; i++)
+      x[i] = 2 * i;
+
+  #pragma omp target has_device_addr(x[15:15])
+    for (int i = 15; i < 30; i++)
+      x[i] = 3 * i;
+}
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc (30*sizeof(int), 0);
+
+  #pragma omp target is_device_ptr(dp)
+    for (int i = 0; i < 30; i++)
+      dp[i] = i;
+
+  int (&x)[30] = *static_cast<int(*)[30]>(static_cast<void*>(dp));
+
+  foo <int> (x);
+
+  int y[30];
+  for (int i = 0; i < 30; ++i)
+    y[i] = 0;
+  int h = omp_get_initial_device ();
+  int t = omp_get_default_device ();
+  omp_target_memcpy (&y, dp, 30 * sizeof(int), 0, 0, h, t);
+  for (int i = 0; i < 15; ++i)
+    if (y[i] != 2 * i)
+      __builtin_abort ();
+  for (int i = 15; i < 30; ++i)
+    if (y[i] != 3 * i)
+      __builtin_abort ();
+
+  omp_target_free (dp, 0);
+
+  return 0;
+}

libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C

@@ -0,0 +1,30 @@
+/* Testing 'has_device_addr' clause on the target construct with template.  */
+
+#include <omp.h>
+
+template <typename T>
+void
+foo (T (&x))
+{
+  #pragma omp target has_device_addr(x)
+    x = 24;
+}
+
+int
+main ()
+{
+  int *dp = (int*)omp_target_alloc (sizeof(int), 0);
+  int &x = *dp;
+
+  foo <int> (x);
+
+  int y = 42;
+  int h = omp_get_initial_device ();
+  int t = omp_get_default_device ();
+  omp_target_memcpy (&y, dp, sizeof(int), 0, 0, h, t);
+  if (y != 24)
+    __builtin_abort ();
+
+  omp_target_free (dp, 0);
+  return 0;
+}