Hi,
The patch for adding the has_device_addr clause on the target construct was
recently committed (bbb7f8604e1dfc08f44354cfd93d2287f2fdd489).
Additionally, this patch adds support for list items in the has_device_addr
clause which type is given by C++ template parameters.
Marcel
-----------------
Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634
München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas
Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht
München, HRB 106955
OpenMP, C++: Add template support for the has_device_addr clause.
gcc/cp/ChangeLog:
* pt.cc (tsubst_omp_clauses): Add OMP_CLAUSE_HAS_DEVICE_ADDR.
* semantics.cc (finish_omp_clauses): Handle PARM_DECL and
NON_LVALUE_EXPR.
gcc/ChangeLog:
* gimplify.cc (gimplify_scan_omp_clauses): Handle NON_LVALUE_EXPR.
(gimplify_adjust_omp_clauses): Likewise.
* omp-low.cc (scan_sharing_clauses): Likewise.
(lower_omp_target): Likewise.
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.
diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc
index 6dda660..86446d7 100644
--- a/gcc/cp/pt.cc
+++ b/gcc/cp/pt.cc
@@ -17652,6 +17652,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)
@@ -17797,6 +17798,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:
diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc
index 0cb17a6..452ecfd 100644
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -8534,11 +8534,14 @@ finish_omp_clauses (tree clauses, enum
c_omp_region_type ort)
{
if (handle_omp_array_sections (c, ort))
remove = true;
+ else if (TREE_CODE (TREE_CHAIN (t)) == PARM_DECL)
+ t = TREE_CHAIN (t);
else
{
t = OMP_CLAUSE_DECL (c);
while (TREE_CODE (t) == INDIRECT_REF
- || TREE_CODE (t) == ARRAY_REF)
+ || TREE_CODE (t) == ARRAY_REF
+ || TREE_CODE (t) == NON_LVALUE_EXPR)
t = TREE_OPERAND (t, 0);
}
}
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index f570daa..b1bb5be 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -10285,7 +10285,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
case OMP_CLAUSE_HAS_DEVICE_ADDR:
decl = OMP_CLAUSE_DECL (c);
while (TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == ARRAY_REF)
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
flags = GOVD_EXPLICIT;
goto do_add_decl;
@@ -11443,7 +11444,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
case OMP_CLAUSE_HAS_DEVICE_ADDR:
decl = OMP_CLAUSE_DECL (c);
while (TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == ARRAY_REF)
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
remove = n == NULL || !(n->value & GOVD_SEEN);
diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 77176ef..30cc9b6 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -1384,7 +1384,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
- if (TREE_CODE (decl) == INDIRECT_REF)
+ if (TREE_CODE (decl) == INDIRECT_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
install_var_field (decl, true, 3, ctx);
}
@@ -1747,7 +1748,8 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (decl) == INDIRECT_REF
- || TREE_CODE (decl) == ARRAY_REF)
+ || TREE_CODE (decl) == ARRAY_REF
+ || TREE_CODE (decl) == NON_LVALUE_EXPR)
decl = TREE_OPERAND (decl, 0);
}
@@ -12847,7 +12849,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (var) == INDIRECT_REF
- || TREE_CODE (var) == ARRAY_REF)
+ || TREE_CODE (var) == ARRAY_REF
+ || TREE_CODE (var) == NON_LVALUE_EXPR)
var = TREE_OPERAND (var, 0);
}
map_cnt++;
@@ -13337,7 +13340,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (ovar) == INDIRECT_REF
- || TREE_CODE (ovar) == ARRAY_REF)
+ || TREE_CODE (ovar) == ARRAY_REF
+ || TREE_CODE (ovar) == NON_LVALUE_EXPR)
ovar = TREE_OPERAND (ovar, 0);
}
var = lookup_decl_in_outer_ctx (ovar, ctx);
@@ -13607,7 +13611,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR)
{
while (TREE_CODE (var) == INDIRECT_REF
- || TREE_CODE (var) == ARRAY_REF)
+ || TREE_CODE (var) == ARRAY_REF
+ || TREE_CODE (var) == NON_LVALUE_EXPR)
var = TREE_OPERAND (var, 0);
}
x = build_receiver_ref (var, false, ctx);
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C
b/libgomp/testsuite/libgomp.c++/target-has-device-addr-7.C
new file mode 100644
index 0000000..2c4571b
--- /dev/null
+++ b/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;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C
b/libgomp/testsuite/libgomp.c++/target-has-device-addr-8.C
new file mode 100644
index 0000000..2adfd30
--- /dev/null
+++ b/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;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C
b/libgomp/testsuite/libgomp.c++/target-has-device-addr-9.C
new file mode 100644
index 0000000..0c34cab
--- /dev/null
+++ b/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;
+}
