2015-12-02 Tom de Vries <tom@codesourcery.com>
+ * gimplify.c (enum gimplify_omp_var_data): Add enum value
+ GOVD_MAP_FORCE.
+ (oacc_default_clause): Fix default for scalars in oacc kernels.
+ (gimplify_adjust_omp_clauses_1): Handle GOVD_MAP_FORCE.
+
+2015-12-02 Tom de Vries <tom@codesourcery.com>
+
* omp-low.c (install_var_field, scan_sharing_clauses): Add and handle
parameter base_pointers_restrict.
(omp_target_base_pointers_restrict_p): New function.
/* Flag for shared vars that are or might be stored to in the region. */
GOVD_WRITTEN = 131072,
+ /* Flag for GOVD_MAP, if it is a forced mapping. */
+ GOVD_MAP_FORCE = 262144,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
gcc_unreachable ();
case ORT_ACC_KERNELS:
- /* Everything under kernels are default 'present_or_copy'. */
+ /* Scalars are default 'copy' under kernels, non-scalars are default
+ 'present_or_copy'. */
flags |= GOVD_MAP;
+ if (!AGGREGATE_TYPE_P (TREE_TYPE (decl)))
+ flags |= GOVD_MAP_FORCE;
+
rkind = "kernels";
break;
}
else if (code == OMP_CLAUSE_MAP)
{
- OMP_CLAUSE_SET_MAP_KIND (clause,
- flags & GOVD_MAP_TO_ONLY
- ? GOMP_MAP_TO
- : GOMP_MAP_TOFROM);
+ int kind = (flags & GOVD_MAP_TO_ONLY
+ ? GOMP_MAP_TO
+ : GOMP_MAP_TOFROM);
+ if (flags & GOVD_MAP_FORCE)
+ kind |= GOMP_MAP_FLAG_FORCE;
+ OMP_CLAUSE_SET_MAP_KIND (clause, kind);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
2015-12-02 Tom de Vries <tom@codesourcery.com>
+ * c-c++-common/goacc/kernels-default-2.c: New test.
+ * c-c++-common/goacc/kernels-default.c: New test.
+
+2015-12-02 Tom de Vries <tom@codesourcery.com>
+
* c-c++-common/goacc/kernels-alias-2.c: New test.
* c-c++-common/goacc/kernels-alias-3.c: New test.
* c-c++-common/goacc/kernels-alias-4.c: New test.
--- /dev/null
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 2
+
+void
+foo (void)
+{
+ unsigned int a[N];
+
+#pragma acc kernels
+ {
+ a[0]++;
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(tofrom" 1 "gimple" } } */
--- /dev/null
+/* { dg-additional-options "-O2" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void
+foo (void)
+{
+ unsigned int i;
+#pragma acc kernels
+ {
+ i++;
+ }
+}
+
+/* { dg-final { scan-tree-dump-times "map\\(force_tofrom" 1 "gimple" } } */