@@ -18196,6 +18196,7 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
@@ -45860,6 +45860,7 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token *pragma_tok)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEFAULT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
@@ -3802,7 +3802,8 @@ error:
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
- | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH)
+ | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH \
+ | OMP_CLAUSE_DEFAULT)
#define OACC_LOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
| OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
@@ -225,6 +225,7 @@ struct gimplify_omp_ctx
vec<tree> loop_iter_var;
location_t location;
enum omp_clause_default_kind default_kind;
+ struct gimplify_omp_ctx *oacc_default_clause_ctx;
enum omp_region_type region_type;
enum tree_code code;
bool combined_loop;
@@ -459,6 +460,10 @@ new_omp_context (enum omp_region_type region_type)
c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
else
c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+ if (gimplify_omp_ctxp)
+ c->oacc_default_clause_ctx = gimplify_omp_ctxp->oacc_default_clause_ctx;
+ else
+ c->oacc_default_clause_ctx = c;
c->defaultmap[GDMK_SCALAR] = GOVD_MAP;
c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP;
c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
@@ -7699,6 +7704,25 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
return flags;
}
+/* Return string name for types of OpenACC constructs from ORT_* values. */
+
+static const char *
+oacc_region_type_name (enum omp_region_type region_type)
+{
+ switch (region_type)
+ {
+ case ORT_ACC_DATA:
+ return "data";
+ case ORT_ACC_PARALLEL:
+ return "parallel";
+ case ORT_ACC_KERNELS:
+ return "kernels";
+ case ORT_ACC_SERIAL:
+ return "serial";
+ default:
+ gcc_unreachable ();
+ }
+}
/* Determine outer default flags for DECL mentioned in an OACC region
but not declared in an enclosing clause. */
@@ -7706,7 +7730,6 @@ omp_default_clause (struct gimplify_omp_ctx *ctx, tree decl,
static unsigned
oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
{
- const char *rkind;
bool on_device = false;
bool is_private = false;
bool declared = is_oacc_declared (decl);
@@ -7735,17 +7758,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
flags |= GOVD_MAP_TO_ONLY;
}
+ /* Use the enclosing construct with a default clause to set the current
+ default kind. */
+ enum omp_clause_default_kind default_kind
+ = ctx->oacc_default_clause_ctx->default_kind;
+
switch (ctx->region_type)
{
case ORT_ACC_KERNELS:
- rkind = "kernels";
-
if (is_private)
flags |= GOVD_FIRSTPRIVATE;
else if (AGGREGATE_TYPE_P (type))
{
/* Aggregates default to 'present_or_copy', or 'present'. */
- if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+ if (default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
flags |= GOVD_MAP;
else
flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7758,8 +7784,6 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
case ORT_ACC_PARALLEL:
case ORT_ACC_SERIAL:
- rkind = ctx->region_type == ORT_ACC_PARALLEL ? "parallel" : "serial";
-
if (is_private)
flags |= GOVD_FIRSTPRIVATE;
else if (on_device || declared)
@@ -7767,7 +7791,7 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
else if (AGGREGATE_TYPE_P (type))
{
/* Aggregates default to 'present_or_copy', or 'present'. */
- if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
+ if (default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
flags |= GOVD_MAP;
else
flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
@@ -7785,16 +7809,20 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
if (DECL_ARTIFICIAL (decl))
; /* We can get compiler-generated decls, and should not complain
about them. */
- else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_NONE)
+ else if (default_kind == OMP_CLAUSE_DEFAULT_NONE)
{
error ("%qE not specified in enclosing OpenACC %qs construct",
- DECL_NAME (lang_hooks.decls.omp_report_decl (decl)), rkind);
- inform (ctx->location, "enclosing OpenACC %qs construct", rkind);
- }
- else if (ctx->default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
+ DECL_NAME (lang_hooks.decls.omp_report_decl (decl)),
+ oacc_region_type_name (ctx->region_type));
+ inform (ctx->oacc_default_clause_ctx->location,
+ "enclosing OpenACC %qs construct",
+ oacc_region_type_name
+ (ctx->oacc_default_clause_ctx->region_type));
+ }
+ else if (default_kind == OMP_CLAUSE_DEFAULT_PRESENT)
; /* Handled above. */
else
- gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
+ gcc_checking_assert (default_kind == OMP_CLAUSE_DEFAULT_SHARED);
return flags;
}
@@ -12124,6 +12152,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_CLAUSE_DEFAULT:
ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c);
+ if (flag_openacc)
+ /* For OpenACC, set current context to the 'active' one for
+ default-clause lookup. */
+ ctx->oacc_default_clause_ctx = ctx;
break;
case OMP_CLAUSE_INCLUSIVE:
@@ -4,7 +4,7 @@ void f1 ()
{
int f1_a = 2;
float f1_b[2];
-
+
#pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */
{
f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
@@ -15,4 +15,49 @@ void f1 ()
f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
= f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
}
+
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc kernels
+ {
+ f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" } */
+ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" } */
+ }
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc parallel
+ {
+ f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+ }
+
+#pragma acc data default (none)
+#pragma acc parallel default (none) /* { dg-message "enclosing OpenACC .parallel. construct" } */
+ {
+ f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+ }
+
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc data
+#pragma acc data
+#pragma acc parallel
+ {
+ f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+ }
+#pragma acc data
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc data
+#pragma acc parallel
+ {
+ f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+ }
+#pragma acc data
+#pragma acc data
+#pragma acc data default (none) /* { dg-message "enclosing OpenACC .data. construct" } */
+#pragma acc parallel
+ {
+ f1_b[0] /* { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" } */
+ = f1_a; /* { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" } */
+ }
}
@@ -4,8 +4,8 @@
void f1 ()
{
- int f1_a = 2;
- float f1_b[2];
+ int f1_a = 2, f1_c = 3;
+ float f1_b[2], f1_d[2];
#pragma acc kernels default (present)
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } } */
@@ -17,4 +17,18 @@ void f1 ()
{
f1_b[0] = f1_a;
}
+
+ /* { dg-final { scan-tree-dump-times "omp target oacc_data default\\(present\\)" 2 "gimple" } } */
+#pragma acc data default (present)
+#pragma acc kernels
+ /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } } */
+ {
+ f1_d[0] = f1_c;
+ }
+#pragma acc data default (present)
+#pragma acc parallel
+ /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } } */
+ {
+ f1_d[0] = f1_c;
+ }
}
@@ -15,4 +15,65 @@ subroutine f1
= f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
!$acc end parallel
+
+ !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc kernels
+ f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } }
+ = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .kernels. construct" }
+ ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .kernels. construct" "" { xfail *-*-* } .-1 }
+ !$acc end kernels
+ !$acc end data
+
+ !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc parallel
+ f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+ = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+ ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+ !$acc end parallel
+ !$acc end data
+
+ !$acc data default (none)
+ !$acc parallel default (none) ! { dg-message "enclosing OpenACC .parallel. construct" }
+ f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+ = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+ ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+ !$acc end parallel
+ !$acc end data
+
+ !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc data
+ !$acc data
+ !$acc parallel
+ f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+ = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+ ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+ !$acc end parallel
+ !$acc end data
+ !$acc end data
+ !$acc end data
+
+ !$acc data
+ !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc data
+ !$acc parallel
+ f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+ = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+ ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+ !$acc end parallel
+ !$acc end data
+ !$acc end data
+ !$acc end data
+
+ !$acc data
+ !$acc data
+ !$acc data default (none) ! { dg-message "enclosing OpenACC .data. construct" }
+ !$acc parallel
+ f1_b(1) & ! { dg-error ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } }
+ = f1_a; ! { dg-error ".f1_a. not specified in enclosing OpenACC .parallel. construct" }
+ ! { dg-bogus ".f1_b. not specified in enclosing OpenACC .parallel. construct" "" { xfail *-*-* } .-1 }
+ !$acc end parallel
+ !$acc end data
+ !$acc end data
+ !$acc end data
+
end subroutine f1
@@ -4,8 +4,8 @@
SUBROUTINE F1
IMPLICIT NONE
- INTEGER :: F1_A = 2
- REAL, DIMENSION (2) :: F1_B
+ INTEGER :: F1_A = 2, F1_C = 3
+ REAL, DIMENSION (2) :: F1_B, F1_D
!$ACC KERNELS DEFAULT (PRESENT)
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) map\\(force_tofrom:f1_a" 1 "gimple" } }
@@ -15,4 +15,17 @@
! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } }
F1_B(1) = F1_A;
!$ACC END PARALLEL
+
+!$ACC DATA DEFAULT (PRESENT)
+!$ACC KERNELS
+! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(force_present:f1_d \[^\\)\]+\\) map\\(force_tofrom:f1_c" 1 "gimple" } }
+ F1_D(1) = F1_C;
+!$ACC END KERNELS
+!$ACC END DATA
+!$ACC DATA DEFAULT (PRESENT)
+!$ACC PARALLEL DEFAULT (PRESENT)
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_d \[^\\)\]+\\) firstprivate\\(f1_c\\)" 1 "gimple" } }
+ F1_D(1) = F1_C;
+!$ACC END PARALLEL
+!$ACC END DATA
END SUBROUTINE F1