diff options
| -rw-r--r-- | gcc/c/c-parser.cc | 1 | ||||
| -rw-r--r-- | gcc/cp/parser.cc | 1 | ||||
| -rw-r--r-- | gcc/fortran/openmp.cc | 3 | ||||
| -rw-r--r-- | gcc/gimplify.cc | 64 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-3.c | 59 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-4.c | 42 | ||||
| -rw-r--r-- | gcc/testsuite/c-c++-common/goacc/default-5.c | 19 | ||||
| -rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 77 | ||||
| -rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/default-4.f | 36 | ||||
| -rw-r--r-- | gcc/testsuite/gfortran.dg/goacc/default-5.f | 19 |
10 files changed, 298 insertions, 23 deletions
diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index cabb18d..33fe7b1 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18284,6 +18284,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) \ diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index 7f64670..774706a 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45854,6 +45854,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) \ diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index 234d896..bee3015 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -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 \ diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 320920e..7549436 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -7699,6 +7699,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 +7725,23 @@ 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; + struct gimplify_omp_ctx *ctx_default = ctx; + /* If no 'default' clause appears on this compute construct... */ + if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_SHARED) + { + /* ..., see if one appears on a lexically containing 'data' + construct. */ + while ((ctx_default = ctx_default->outer_context)) + { + if (ctx_default->region_type == ORT_ACC_DATA + && ctx_default->default_kind != OMP_CLAUSE_DEFAULT_SHARED) + break; + } + /* If not, reset. */ + if (!ctx_default) + ctx_default = ctx; + } + bool on_device = false; bool is_private = false; bool declared = is_oacc_declared (decl); @@ -7738,14 +7773,12 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags) 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 (ctx_default->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) flags |= GOVD_MAP; else flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; @@ -7758,8 +7791,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 +7798,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 (ctx_default->default_kind != OMP_CLAUSE_DEFAULT_PRESENT) flags |= GOVD_MAP; else flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT; @@ -7785,16 +7816,23 @@ 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 (ctx_default->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)); + if (ctx_default != ctx) + inform (ctx->location, "enclosing OpenACC %qs construct and", + oacc_region_type_name (ctx->region_type)); + inform (ctx_default->location, + "enclosing OpenACC %qs construct with %qs clause", + oacc_region_type_name (ctx_default->region_type), + "default(none)"); + } + else if (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_PRESENT) ; /* Handled above. */ else - gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED); + gcc_checking_assert (ctx_default->default_kind == OMP_CLAUSE_DEFAULT_SHARED); return flags; } diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c index ac169a9..73dbc90 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-3.c +++ b/gcc/testsuite/c-c++-common/goacc/default-3.c @@ -4,13 +4,66 @@ void f1 () { int f1_a = 2; float f1_b[2]; - -#pragma acc kernels default (none) /* { dg-message "enclosing OpenACC .kernels. construct" } */ + +#pragma acc kernels default (none) /* { dg-note "enclosing OpenACC 'kernels' construct with 'default\\\(none\\\)' clause" } */ + { + 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 parallel default (none) /* { dg-note "enclosing OpenACC 'parallel' construct with 'default\\\(none\\\)' clause" } */ + { + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc kernels /* { dg-note "enclosing OpenACC 'kernels' construct and" } */ { 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 parallel default (none) /* { dg-message "enclosing OpenACC .parallel. construct" } */ +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + 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-note "enclosing OpenACC 'parallel' construct with 'default\\\(none\\\)' clause" } */ + { + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc data +#pragma acc data +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc data +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ + { + 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) +#pragma acc data default (none) /* { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } */ +#pragma acc parallel /* { dg-note "enclosing OpenACC 'parallel' construct and" } */ { 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" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/default-4.c b/gcc/testsuite/c-c++-common/goacc/default-4.c index 867175d..e12cb86 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-4.c +++ b/gcc/testsuite/c-c++-common/goacc/default-4.c @@ -44,6 +44,27 @@ void f2 () } } +void f2_ () +{ + int f2__a = 2; + float f2__b[2]; + +#pragma acc data default (none) copyin (f2__a) copyout (f2__b) + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f2__b \[^\\)\]+\\) map\\(to:f2__a \[^\\)\]+\\) default\\(none\\)" 1 "gimple" } } */ + { +#pragma acc kernels + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } */ + { + f2__b[0] = f2__a; + } +#pragma acc parallel + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } */ + { + f2__b[0] = f2__a; + } + } +} + void f3 () { int f3_a = 2; @@ -64,3 +85,24 @@ void f3 () } } } + +void f3_ () +{ + int f3__a = 2; + float f3__b[2]; + +#pragma acc data default (present) copyin (f3__a) copyout (f3__b) + /* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:f3__b \[^\\)\]+\\) map\\(to:f3__a \[^\\)\]+\\) default\\(present\\)" 1 "gimple" } } */ + { +#pragma acc kernels + /* { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } */ + { + f3__b[0] = f3__a; + } +#pragma acc parallel + /* { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } */ + { + f3__b[0] = f3__a; + } + } +} diff --git a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c index 37e3c355..59ac1d7 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-5.c +++ b/gcc/testsuite/c-c++-common/goacc/default-5.c @@ -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,19 @@ 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 (none) +#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; + } } diff --git a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 b/gcc/testsuite/gfortran.dg/goacc/default-3.f95 index 98ed342..c1edf4c 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95 @@ -5,14 +5,87 @@ subroutine f1 integer :: f1_a = 2 real, dimension (2) :: f1_b - !$acc kernels default (none) ! { dg-message "enclosing OpenACC .kernels. construct" } + !$acc kernels default (none) ! { dg-note "enclosing OpenACC .kernels. construct with 'default\\\(none\\\)' clause" } 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 parallel default (none) ! { dg-message "enclosing OpenACC .parallel. construct" } + !$acc parallel default (none) ! { dg-note "enclosing OpenACC .parallel. construct with 'default\\\(none\\\)' clause" } 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 data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc kernels ! { dg-note "enclosing OpenACC 'kernels' construct and" } + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + 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-note "enclosing OpenACC .parallel. construct with 'default\\\(none\\\)' clause" } + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc data + !$acc data + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc data + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + 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-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + 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) + !$acc data default (none) ! { dg-note "enclosing OpenACC 'data' construct with 'default\\\(none\\\)' clause" } + !$acc parallel ! { dg-note "enclosing OpenACC 'parallel' construct and" } + 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 diff --git a/gcc/testsuite/gfortran.dg/goacc/default-4.f b/gcc/testsuite/gfortran.dg/goacc/default-4.f index 30f411f..4e89b68 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-4.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-4.f @@ -38,6 +38,24 @@ !$ACC END DATA END SUBROUTINE F2 + SUBROUTINE F2_ + IMPLICIT NONE + INTEGER :: F2__A = 2 + REAL, DIMENSION (2) :: F2__B + +!$ACC DATA DEFAULT (NONE) COPYIN (F2__A) COPYOUT (F2__B) +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f2__a \[^\\)\]+\\) map\\(from:f2__b \[^\\)\]+\\) default\\(none\\)" 1 "gimple" } } +!$ACC KERNELS +! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } + F2__B(1) = F2__A; +!$ACC END KERNELS +!$ACC PARALLEL +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f2__b \[^\\)\]+\\) map\\(tofrom:f2__a" 1 "gimple" } } + F2__B(1) = F2__A; +!$ACC END PARALLEL +!$ACC END DATA + END SUBROUTINE F2_ + SUBROUTINE F3 IMPLICIT NONE INTEGER :: F3_A = 2 @@ -55,3 +73,21 @@ !$ACC END PARALLEL !$ACC END DATA END SUBROUTINE F3 + + SUBROUTINE F3_ + IMPLICIT NONE + INTEGER :: F3__A = 2 + REAL, DIMENSION (2) :: F3__B + +!$ACC DATA DEFAULT (PRESENT) COPYIN (F3__A) COPYOUT (F3__B) +! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(to:f3__a \[^\\)\]+\\) map\\(from:f3__b \[^\\)\]+\\) default\\(present\\)" 1 "gimple" } } +!$ACC KERNELS +! { dg-final { scan-tree-dump-times "omp target oacc_kernels map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } + F3__B(1) = F3__A; +!$ACC END KERNELS +!$ACC PARALLEL +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(tofrom:f3__b \[^\\)\]+\\) map\\(tofrom:f3__a" 1 "gimple" } } + F3__B(1) = F3__A; +!$ACC END PARALLEL +!$ACC END DATA + END SUBROUTINE F3_ diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f index 9dc83cb..2cb07a8 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-5.f +++ b/gcc/testsuite/gfortran.dg/goacc/default-5.f @@ -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,19 @@ ! { 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 (NONE) +!$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 +!$ACC END DATA END SUBROUTINE F1 |