Hi Thomas, this patch implements the OpenACC 2.7 addition of default(none|present) support for data constructs.
Apart from adjusting the front-ends for allowed clauses masks (for acc data), mostly implemented in gimplify. Tested on powerpc64le-linux/nvptx, x86_64-linux/amdgcn tests in progress (expect no surprises). Is this okay for trunk? Thanks, Chung-Lin gcc/c/ChangeLog: * c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. gcc/cp/ChangeLog: * parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. gcc/fortran/ChangeLog: * openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT. gcc/ChangeLog: * gimplify.cc (struct gimplify_omp_ctx): Add oacc_data_default_kind field. (new_omp_context): Initialize oacc_data_default_kind field. (gimplify_scan_omp_clauses): Set oacc_data_default_kind for data constructs. Set ctx->default_kind for compute constructs from ctx->oacc_data_default_kind. gcc/testsuite/ChangeLog: * c-c++-common/goacc/default-3.c: Adjust testcase. * c-c++-common/goacc/default-5.c: Adjust testcase. * gfortran.dg/goacc/default-3.f95: Adjust testcase. * gfortran.dg/goacc/default-5.f: Adjust testcase.
From 101305aee9b27c6df00d7c403e469bdf8d7f45a4 Mon Sep 17 00:00:00 2001 From: Chung-Lin Tang <clt...@codesourcery.com> Date: Tue, 6 Jun 2023 03:46:29 -0700 Subject: [PATCH 2/2] OpenACC 2.7: default clause support for data constructs This patch implements the OpenACC 2.7 addition of default(none|present) support for data constructs. Now, specifying "default(none|present)" on a data construct turns on same default clause behavior for all enclosed compute constructs (which don't already themselves have a default clause). gcc/c/ChangeLog: * c-parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. gcc/cp/ChangeLog: * parser.cc (OACC_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DEFAULT. gcc/fortran/ChangeLog: * openmp.cc (OACC_DATA_CLAUSES): Add OMP_CLAUSE_DEFAULT. gcc/ChangeLog: * gimplify.cc (struct gimplify_omp_ctx): Add oacc_data_default_kind field. (new_omp_context): Initialize oacc_data_default_kind field. (gimplify_scan_omp_clauses): Set oacc_data_default_kind for data constructs. Set ctx->default_kind for compute constructs from ctx->oacc_data_default_kind. gcc/testsuite/ChangeLog: * c-c++-common/goacc/default-3.c: Adjust testcase. * c-c++-common/goacc/default-5.c: Adjust testcase. * gfortran.dg/goacc/default-3.f95: Adjust testcase. * gfortran.dg/goacc/default-5.f: Adjust testcase. --- gcc/c/c-parser.cc | 1 + gcc/cp/parser.cc | 1 + gcc/fortran/openmp.cc | 3 ++- gcc/gimplify.cc | 20 +++++++++++++++++++ gcc/testsuite/c-c++-common/goacc/default-3.c | 15 +++++++++++++- gcc/testsuite/c-c++-common/goacc/default-5.c | 18 +++++++++++++++-- gcc/testsuite/gfortran.dg/goacc/default-3.f95 | 15 ++++++++++++++ gcc/testsuite/gfortran.dg/goacc/default-5.f | 17 ++++++++++++++-- 8 files changed, 84 insertions(+), 6 deletions(-) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index b61aef8b1a2..645d28b320d 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -18133,6 +18133,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 dd7638f1c93..4b4df29a406 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -45759,6 +45759,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 4c30548567f..b785e71f20f 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -3645,7 +3645,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 4aa6229fc74..368f5fd7ec8 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -225,6 +225,7 @@ struct gimplify_omp_ctx vec<tree> loop_iter_var; location_t location; enum omp_clause_default_kind default_kind; + enum omp_clause_default_kind oacc_data_default_kind; enum omp_region_type region_type; enum tree_code code; bool combined_loop; @@ -459,6 +460,8 @@ 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_data_default_kind = gimplify_omp_ctxp->oacc_data_default_kind; c->defaultmap[GDMK_SCALAR] = GOVD_MAP; c->defaultmap[GDMK_SCALAR_TARGET] = GOVD_MAP; c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP; @@ -12050,6 +12053,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_CLAUSE_DEFAULT: ctx->default_kind = OMP_CLAUSE_DEFAULT_KIND (c); + if (code == OACC_DATA) + ctx->oacc_data_default_kind = OMP_CLAUSE_DEFAULT_KIND (c); break; case OMP_CLAUSE_INCLUSIVE: @@ -12098,6 +12103,21 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, list_p = &OMP_CLAUSE_CHAIN (c); } + if ((code == OACC_PARALLEL + || code == OACC_KERNELS + || code == OACC_SERIAL) + && ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED + && ctx->oacc_data_default_kind != OMP_CLAUSE_DEFAULT_UNSPECIFIED) + { + ctx->default_kind = ctx->oacc_data_default_kind; + + /* Append actual default clause on compute construct. Not really needed + for omp_notice_variable to work properly, but for debug dump files. */ + c = build_omp_clause (UNKNOWN_LOCATION, OMP_CLAUSE_DEFAULT); + OMP_CLAUSE_DEFAULT_KIND (c) = ctx->oacc_data_default_kind; + *list_p = c; + } + ctx->clauses = *orig_list_p; gimplify_omp_ctxp = ctx; } diff --git a/gcc/testsuite/c-c++-common/goacc/default-3.c b/gcc/testsuite/c-c++-common/goacc/default-3.c index ac169a903e9..060629057c1 100644 --- a/gcc/testsuite/c-c++-common/goacc/default-3.c +++ b/gcc/testsuite/c-c++-common/goacc/default-3.c @@ -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,17 @@ 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) +#pragma acc kernels /* { dg-message "enclosing OpenACC .kernels. construct" } */ + { + 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) +#pragma acc parallel /* { 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" } */ + } } diff --git a/gcc/testsuite/c-c++-common/goacc/default-5.c b/gcc/testsuite/c-c++-common/goacc/default-5.c index 37e3c3555cd..e44a0dd49d7 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,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 default\\(present\\) 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 default\\(present\\) 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 98ed34200c6..13ccce30fd7 100644 --- a/gcc/testsuite/gfortran.dg/goacc/default-3.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/default-3.f95 @@ -15,4 +15,19 @@ 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) + !$acc kernels ! { dg-message "enclosing OpenACC .kernels. construct" } + 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) + !$acc parallel ! { 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 end subroutine f1 diff --git a/gcc/testsuite/gfortran.dg/goacc/default-5.f b/gcc/testsuite/gfortran.dg/goacc/default-5.f index 9dc83cbe601..4424f9e7523 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,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 default\\(present\\) 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 -- 2.27.0