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)
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:
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:
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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;
+}