procyberian/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

OpenACC 2.5 default (present) clause

Browse source 
gcc/c/
	* c-parser.c (c_parser_omp_clause_default): Handle
	"OMP_CLAUSE_DEFAULT_PRESENT".
	gcc/cp/
	* parser.c (cp_parser_omp_clause_default): Handle
	"OMP_CLAUSE_DEFAULT_PRESENT".
	gcc/fortran/
	* gfortran.h (enum gfc_omp_default_sharing): Add
	"OMP_DEFAULT_PRESENT".
	* dump-parse-tree.c (show_omp_clauses): Handle it.
	* openmp.c (gfc_match_omp_clauses): Likewise.
	* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
	gcc/
	* tree-core.h (enum omp_clause_default_kind): Add
	"OMP_CLAUSE_DEFAULT_PRESENT".
	* tree-pretty-print.c (dump_omp_clause): Handle it.
	* gimplify.c (enum gimplify_omp_var_data): Add
	"GOVD_MAP_FORCE_PRESENT".
	(gimplify_adjust_omp_clauses_1): Map it to
	"GOMP_MAP_FORCE_PRESENT".
	(oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
	gcc/testsuite/
	* c-c++-common/goacc/default-1.c: Update.
	* c-c++-common/goacc/default-2.c: Likewise.
	* c-c++-common/goacc/default-4.c: Likewise.
	* gfortran.dg/goacc/default-1.f95: Likewise.
	* gfortran.dg/goacc/default-4.f: Likewise.
	* c-c++-common/goacc/default-5.c: New file.
	* gfortran.dg/goacc/default-5.f: Likewise.
	libgomp/
	* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
	* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
	* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.

From-SVN: r248280
This commit is contained in:
Thomas Schwinge 2017-05-19 15:32:48 +02:00 • committed by Thomas Schwinge
parent 0d0afa9faf
commit 7fd549d24f
27 changed files with 346 additions and 52 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

9
gcc/ChangeLog
View file

@ -1,5 +1,14 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
* tree-core.h (enum omp_clause_default_kind): Add
"OMP_CLAUSE_DEFAULT_PRESENT".
* tree-pretty-print.c (dump_omp_clause): Handle it.
* gimplify.c (enum gimplify_omp_var_data): Add
"GOVD_MAP_FORCE_PRESENT".
(gimplify_adjust_omp_clauses_1): Map it to
"GOMP_MAP_FORCE_PRESENT".
(oacc_default_clause): Handle "OMP_CLAUSE_DEFAULT_PRESENT".
* gimplify.c (oacc_default_clause): Clarify.
2017-05-19 Nathan Sidwell <nathan@acm.org>

5
gcc/c/ChangeLog
View file

@ -1,3 +1,8 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
* c-parser.c (c_parser_omp_clause_default): Handle
"OMP_CLAUSE_DEFAULT_PRESENT".
2017-05-18 Bernd Edlinger <bernd.edlinger@hotmail.de>
* config-lang.in (gtfiles): Add c-family/c-format.c.

14
gcc/c/c-parser.c
View file

@ -11057,10 +11057,10 @@ c_parser_omp_clause_copyprivate (c_parser *parser, tree list)
}
/* OpenMP 2.5:
default ( shared | none )
default ( none | shared )
OpenACC 2.0:
default (none) */
OpenACC:
default ( none | present ) */
static tree
c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
@ -11083,6 +11083,12 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
kind = OMP_CLAUSE_DEFAULT_NONE;
break;
case 'p':
if (strcmp ("present", p) != 0 || !is_oacc)
goto invalid_kind;
kind = OMP_CLAUSE_DEFAULT_PRESENT;
break;
case 's':
if (strcmp ("shared", p) != 0 || is_oacc)
goto invalid_kind;
@ -11099,7 +11105,7 @@ c_parser_omp_clause_default (c_parser *parser, tree list, bool is_oacc)
{
invalid_kind:
if (is_oacc)
c_parser_error (parser, "expected %<none%>");
c_parser_error (parser, "expected %<none%> or %<present%>");
else
c_parser_error (parser, "expected %<none%> or %<shared%>");
}

3
gcc/cp/ChangeLog
View file

@ -1,5 +1,8 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
* parser.c (cp_parser_omp_clause_default): Handle
"OMP_CLAUSE_DEFAULT_PRESENT".
* parser.c (cp_parser_omp_clause_default): Avoid printing more
than one syntax error message.

14
gcc/cp/parser.c
View file

@ -31456,10 +31456,10 @@ cp_parser_omp_clause_collapse (cp_parser *parser, tree list, location_t location
}
/* OpenMP 2.5:
default ( shared | none )
default ( none | shared )
OpenACC 2.0
default (none) */
OpenACC:
default ( none | present ) */
static tree
cp_parser_omp_clause_default (cp_parser *parser, tree list,
@ -31483,6 +31483,12 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
kind = OMP_CLAUSE_DEFAULT_NONE;
break;
case 'p':
if (strcmp ("present", p) != 0 || !is_oacc)
goto invalid_kind;
kind = OMP_CLAUSE_DEFAULT_PRESENT;
break;
case 's':
if (strcmp ("shared", p) != 0 || is_oacc)
goto invalid_kind;
@ -31499,7 +31505,7 @@ cp_parser_omp_clause_default (cp_parser *parser, tree list,
{
invalid_kind:
if (is_oacc)
cp_parser_error (parser, "expected %<none%>");
cp_parser_error (parser, "expected %<none%> or %<present%>");
else
cp_parser_error (parser, "expected %<none%> or %<shared%>");
}

8
gcc/fortran/ChangeLog
View file

@ -1,3 +1,11 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
* gfortran.h (enum gfc_omp_default_sharing): Add
"OMP_DEFAULT_PRESENT".
* dump-parse-tree.c (show_omp_clauses): Handle it.
* openmp.c (gfc_match_omp_clauses): Likewise.
* trans-openmp.c (gfc_trans_omp_clauses): Likewise.
2017-05-18 Fritz Reese <fritzoreese@gmail.com>
PR fortran/79968

1
gcc/fortran/dump-parse-tree.c
View file

@ -1283,6 +1283,7 @@ show_omp_clauses (gfc_omp_clauses *omp_clauses)
case OMP_DEFAULT_PRIVATE: type = "PRIVATE"; break;
case OMP_DEFAULT_SHARED: type = "SHARED"; break;
case OMP_DEFAULT_FIRSTPRIVATE: type = "FIRSTPRIVATE"; break;
case OMP_DEFAULT_PRESENT: type = "PRESENT"; break;
default:
gcc_unreachable ();
}

3
gcc/fortran/gfortran.h
View file

@ -1241,7 +1241,8 @@ enum gfc_omp_default_sharing
OMP_DEFAULT_NONE,
OMP_DEFAULT_PRIVATE,
OMP_DEFAULT_SHARED,
OMP_DEFAULT_FIRSTPRIVATE
OMP_DEFAULT_FIRSTPRIVATE,
OMP_DEFAULT_PRESENT
};
enum gfc_omp_proc_bind_kind

20
gcc/fortran/openmp.c
View file

@ -1080,13 +1080,19 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
if (gfc_match ("default ( none )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_NONE;
else if (openacc)
/* c->default_sharing = OMP_DEFAULT_UNKNOWN */;
else if (gfc_match ("default ( shared )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_SHARED;
else if (gfc_match ("default ( private )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_PRIVATE;
else if (gfc_match ("default ( firstprivate )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
{
if (gfc_match ("default ( present )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_PRESENT;
}
else
{
if (gfc_match ("default ( firstprivate )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_FIRSTPRIVATE;
else if (gfc_match ("default ( private )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_PRIVATE;
else if (gfc_match ("default ( shared )") == MATCH_YES)
c->default_sharing = OMP_DEFAULT_SHARED;
}
if (c->default_sharing != OMP_DEFAULT_UNKNOWN)
continue;
}

3
gcc/fortran/trans-openmp.c
View file

@ -2564,6 +2564,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_DEFAULT_FIRSTPRIVATE:
OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_FIRSTPRIVATE;
break;
case OMP_DEFAULT_PRESENT:
OMP_CLAUSE_DEFAULT_KIND (c) = OMP_CLAUSE_DEFAULT_PRESENT;
break;
default:
gcc_unreachable ();
}

52
gcc/gimplify.c
View file

@ -99,6 +99,9 @@ enum gimplify_omp_var_data
/* Flag for GOVD_MAP, if it is a forced mapping. */
GOVD_MAP_FORCE = 262144,
/* Flag for GOVD_MAP: must be present already. */
GOVD_MAP_FORCE_PRESENT = 524288,
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@ -6956,8 +6959,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
rkind = "kernels";
if (AGGREGATE_TYPE_P (type))
/* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
{
/* Aggregates default to 'present_or_copy', or 'present'. */
if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
flags |= GOVD_MAP;
else
flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
}
else
/* Scalars default to 'copy'. */
flags |= GOVD_MAP | GOVD_MAP_FORCE;
@ -6970,8 +6978,13 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
if (on_device || declared)
flags |= GOVD_MAP;
else if (AGGREGATE_TYPE_P (type))
/* Aggregates default to 'present_or_copy'. */
flags |= GOVD_MAP;
{
/* Aggregates default to 'present_or_copy', or 'present'. */
if (ctx->default_kind != OMP_CLAUSE_DEFAULT_PRESENT)
flags |= GOVD_MAP;
else
flags |= GOVD_MAP | GOVD_MAP_FORCE_PRESENT;
}
else
/* Scalars default to 'firstprivate'. */
flags |= GOVD_FIRSTPRIVATE;
@ -6991,6 +7004,8 @@ oacc_default_clause (struct gimplify_omp_ctx *ctx, tree decl, unsigned flags)
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)
; /* Handled above. */
else
gcc_checking_assert (ctx->default_kind == OMP_CLAUSE_DEFAULT_SHARED);
@ -8708,11 +8723,30 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
}
else if (code == OMP_CLAUSE_MAP)
{
int kind = (flags & GOVD_MAP_TO_ONLY
? GOMP_MAP_TO
: GOMP_MAP_TOFROM);
if (flags & GOVD_MAP_FORCE)
kind |= GOMP_MAP_FLAG_FORCE;
int kind;
/* Not all combinations of these GOVD_MAP flags are actually valid. */
switch (flags & (GOVD_MAP_TO_ONLY
| GOVD_MAP_FORCE
| GOVD_MAP_FORCE_PRESENT))
{
case 0:
kind = GOMP_MAP_TOFROM;
break;
case GOVD_MAP_FORCE:
kind = GOMP_MAP_TOFROM | GOMP_MAP_FLAG_FORCE;
break;
case GOVD_MAP_TO_ONLY:
kind = GOMP_MAP_TO;
break;
case GOVD_MAP_TO_ONLY | GOVD_MAP_FORCE:
kind = GOMP_MAP_TO | GOMP_MAP_FLAG_FORCE;
break;
case GOVD_MAP_FORCE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
default:
gcc_unreachable ();
}
OMP_CLAUSE_SET_MAP_KIND (clause, kind);
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)

8
gcc/testsuite/ChangeLog
View file

@ -1,5 +1,13 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
* c-c++-common/goacc/default-1.c: Update.
* c-c++-common/goacc/default-2.c: Likewise.
* c-c++-common/goacc/default-4.c: Likewise.
* gfortran.dg/goacc/default-1.f95: Likewise.
* gfortran.dg/goacc/default-4.f: Likewise.
* c-c++-common/goacc/default-5.c: New file.
* gfortran.dg/goacc/default-5.f: Likewise.
* c-c++-common/goacc/default-1.c: New file.
* c-c++-common/goacc/default-2.c: Likewise.
* c-c++-common/goacc/data-default-1.c: Remove file, including its

5
gcc/testsuite/c-c++-common/goacc/default-1.c
View file

@ -6,4 +6,9 @@ void f1 ()
;
#pragma acc parallel default (none)
;
#pragma acc kernels default (present)
;
#pragma acc parallel default (present)
;
}

28
gcc/testsuite/c-c++-common/goacc/default-2.c
View file

@ -7,39 +7,39 @@ void f1 ()
#pragma acc parallel default /* { dg-error "expected .\\(. before end of line" } */
;
#pragma acc kernels default ( /* { dg-error "expected .none. before end of line" } */
#pragma acc kernels default ( /* { dg-error "expected .none. or .present. before end of line" } */
;
#pragma acc parallel default ( /* { dg-error "expected .none. before end of line" } */
#pragma acc parallel default ( /* { dg-error "expected .none. or .present. before end of line" } */
;
#pragma acc kernels default (, /* { dg-error "expected .none. before .,. token" } */
#pragma acc kernels default (, /* { dg-error "expected .none. or .present. before .,. token" } */
;
#pragma acc parallel default (, /* { dg-error "expected .none. before .,. token" } */
#pragma acc parallel default (, /* { dg-error "expected .none. or .present. before .,. token" } */
;
#pragma acc kernels default () /* { dg-error "expected .none. before .\\). token" } */
#pragma acc kernels default () /* { dg-error "expected .none. or .present. before .\\). token" } */
;
#pragma acc parallel default () /* { dg-error "expected .none. before .\\). token" } */
#pragma acc parallel default () /* { dg-error "expected .none. or .present. before .\\). token" } */
;
#pragma acc kernels default (,) /* { dg-error "expected .none. before .,. token" } */
#pragma acc kernels default (,) /* { dg-error "expected .none. or .present. before .,. token" } */
;
#pragma acc parallel default (,) /* { dg-error "expected .none. before .,. token" } */
#pragma acc parallel default (,) /* { dg-error "expected .none. or .present. before .,. token" } */
;
#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */
#pragma acc kernels default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */
;
#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. before .firstprivate." } */
#pragma acc parallel default (firstprivate) /* { dg-error "expected .none. or .present. before .firstprivate." } */
;
#pragma acc kernels default (private) /* { dg-error "expected .none. before .private." } */
#pragma acc kernels default (private) /* { dg-error "expected .none. or .present. before .private." } */
;
#pragma acc parallel default (private) /* { dg-error "expected .none. before .private." } */
#pragma acc parallel default (private) /* { dg-error "expected .none. or .present. before .private." } */
;
#pragma acc kernels default (shared) /* { dg-error "expected .none. before .shared." } */
#pragma acc kernels default (shared) /* { dg-error "expected .none. or .present. before .shared." } */
;
#pragma acc parallel default (shared) /* { dg-error "expected .none. before .shared." } */
#pragma acc parallel default (shared) /* { dg-error "expected .none. or .present. before .shared." } */
;
#pragma acc kernels default (none /* { dg-error "expected .\\). before end of line" } */

21
gcc/testsuite/c-c++-common/goacc/default-4.c
View file

@ -43,3 +43,24 @@ void f2 ()
}
}
}
void f3 ()
{
int f3_a = 2;
float f3_b[2];
#pragma acc data copyin (f3_a) copyout (f3_b)
/* { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_from:f3_b \[^\\)\]+\\) map\\(force_to:f3_a" 1 "gimple" } } */
{
#pragma acc kernels default (present)
/* { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
{
f3_b[0] = f3_a;
}
#pragma acc parallel default (present)
/* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } } */
{
f3_b[0] = f3_a;
}
}
}

20
gcc/testsuite/c-c++-common/goacc/default-5.c Normal file
View file

@ -0,0 +1,20 @@
/* OpenACC default (present) clause. */
/* { dg-additional-options "-fdump-tree-gimple" } */
void f1 ()
{
int f1_a = 2;
float f1_b[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" } } */
{
f1_b[0] = f1_a;
}
#pragma acc parallel default (present)
/* { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(force_present:f1_b \[^\\)\]+\\) firstprivate\\(f1_a\\)" 1 "gimple" } } */
{
f1_b[0] = f1_a;
}
}

5
gcc/testsuite/gfortran.dg/goacc/default-1.f95
View file

@ -7,4 +7,9 @@ subroutine f1
!$acc end kernels
!$acc parallel default (none)
!$acc end parallel
!$acc kernels default (present)
!$acc end kernels
!$acc parallel default (present)
!$acc end parallel
end subroutine f1

18
gcc/testsuite/gfortran.dg/goacc/default-4.f
View file

@ -37,3 +37,21 @@
!$ACC END PARALLEL
!$ACC END DATA
END SUBROUTINE F2
SUBROUTINE F3
IMPLICIT NONE
INTEGER :: F3_A = 2
REAL, DIMENSION (2) :: F3_B
!$ACC DATA COPYIN (F3_A) COPYOUT (F3_B)
! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(force_to:f3_a \[^\\)\]+\\) map\\(force_from:f3_b" 1 "gimple" } }
!$ACC KERNELS DEFAULT (PRESENT)
! { dg-final { scan-tree-dump-times "omp target oacc_kernels default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
F3_B(1) = F3_A;
!$ACC END KERNELS
!$ACC PARALLEL DEFAULT (PRESENT)
! { dg-final { scan-tree-dump-times "omp target oacc_parallel default\\(present\\) map\\(tofrom:f3_b \[^\\)\]+\\) map\\(tofrom:f3_a" 1 "gimple" } }
F3_B(1) = F3_A;
!$ACC END PARALLEL
!$ACC END DATA
END SUBROUTINE F3

18
gcc/testsuite/gfortran.dg/goacc/default-5.f Normal file
View file

@ -0,0 +1,18 @@
! OpenACC default (present) clause.
! { dg-additional-options "-fdump-tree-gimple" }
SUBROUTINE F1
IMPLICIT NONE
INTEGER :: F1_A = 2
REAL, DIMENSION (2) :: F1_B
!$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" } }
F1_B(1) = F1_A;
!$ACC END KERNELS
!$ACC PARALLEL DEFAULT (PRESENT)
! { 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
END SUBROUTINE F1

3
gcc/tree-core.h
View file

@ -357,7 +357,7 @@ enum omp_clause_code {
/* OpenMP clause: ordered [(constant-integer-expression)]. */
OMP_CLAUSE_ORDERED,
/* OpenMP clause: default. */
/* OpenACC/OpenMP clause: default. */
OMP_CLAUSE_DEFAULT,
/* OpenACC/OpenMP clause: collapse (constant-integer-expression). */
@ -499,6 +499,7 @@ enum omp_clause_default_kind {
OMP_CLAUSE_DEFAULT_NONE,
OMP_CLAUSE_DEFAULT_PRIVATE,
OMP_CLAUSE_DEFAULT_FIRSTPRIVATE,
OMP_CLAUSE_DEFAULT_PRESENT,
OMP_CLAUSE_DEFAULT_LAST
};

3
gcc/tree-pretty-print.c
View file

@ -502,6 +502,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
case OMP_CLAUSE_DEFAULT_FIRSTPRIVATE:
pp_string (pp, "firstprivate");
break;
case OMP_CLAUSE_DEFAULT_PRESENT:
pp_string (pp, "present");
break;
default:
gcc_unreachable ();
}

6
libgomp/ChangeLog
View file

@ -1,5 +1,11 @@
2017-05-19 Thomas Schwinge <thomas@codesourcery.com>
* testsuite/libgomp.oacc-c++/template-reduction.C: Update.
* testsuite/libgomp.oacc-c-c++-common/nested-2.c: Update.
* testsuite/libgomp.oacc-fortran/data-4-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/default-1.f90: Likewise.
* testsuite/libgomp.oacc-fortran/non-scalar-data.f90: Likewise.
* plugin/plugin-hsa.c (DLSYM_FN, init_hsa_runtime_functions):
Debug output for failure.

25
libgomp/testsuite/libgomp.oacc-c++/template-reduction.C
View file

@ -32,6 +32,28 @@ sum ()
return s;
}
// Check template with default (present)
template<typename T> T
sum_default_present ()
{
T s = 0;
T array[n];
for (int i = 0; i < n; i++)
array[i] = i+1;
#pragma acc enter data copyin (array)
#pragma acc parallel loop num_gangs (10) gang reduction (+:s) default (present)
for (int i = 0; i < n; i++)
s += array[i];
#pragma acc exit data delete (array)
return s;
}
// Check present and async
template<typename T> T
@ -86,6 +108,9 @@ main()
if (sum<int> () != result)
__builtin_abort ();
if (sum_default_present<int> () != result)
__builtin_abort ();
#pragma acc enter data copyin (a)
if (async_sum (a) != result)
__builtin_abort ();

31
libgomp/testsuite/libgomp.oacc-c-c++-common/nested-2.c
View file

@ -137,5 +137,36 @@ main (int argc, char *argv[])
abort ();
}
#pragma acc enter data create (a)
#pragma acc parallel default (present)
{
for (int j = 0; j < N; ++j)
a[j] = j + 1;
}
#pragma acc update host (a)
for (i = 0; i < N; ++i)
{
if (a[i] != i + 1)
abort ();
}
#pragma acc kernels default (present)
{
for (int j = 0; j < N; ++j)
a[j] = j + 2;
}
#pragma acc exit data copyout (a)
for (i = 0; i < N; ++i)
{
if (a[i] != i + 2)
abort ();
}
return 0;
}

21
libgomp/testsuite/libgomp.oacc-fortran/data-4-2.f90
View file

@ -1,4 +1,5 @@
! Copy of data-4.f90 with self exchanged with host for !acc update.
! Copy of data-4.f90 with self exchanged with host for !acc update, and with
! default (present) clauses added.
! { dg-do run }
@ -19,7 +20,7 @@ program asyncwait
!$acc enter data copyin (a(1:N)) copyin (b(1:N)) copyin (N) async
!$acc parallel async wait
!$acc parallel default (present) async wait
!$acc loop
do i = 1, N
b(i) = a(i)
@ -39,7 +40,7 @@ program asyncwait
!$acc update device (a(1:N), b(1:N)) async (1)
!$acc parallel async (1) wait (1)
!$acc parallel default (present) async (1) wait (1)
!$acc loop
do i = 1, N
b(i) = a(i)
@ -62,19 +63,19 @@ program asyncwait
!$acc enter data copyin (c(1:N), d(1:N)) async (1)
!$acc update device (a(1:N), b(1:N)) async (1)
!$acc parallel async (1)
!$acc parallel default (present) async (1)
do i = 1, N
b(i) = (a(i) * a(i) * a(i)) / a(i)
end do
!$acc end parallel
!$acc parallel async (1)
!$acc parallel default (present) async (1)
do i = 1, N
c(i) = (a(i) * 4) / a(i)
end do
!$acc end parallel
!$acc parallel async (1)
!$acc parallel default (present) async (1)
do i = 1, N
d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
end do
@ -100,25 +101,25 @@ program asyncwait
!$acc enter data copyin (e(1:N)) async (1)
!$acc update device (a(1:N), b(1:N), c(1:N), d(1:N)) async (1)
!$acc parallel async (1)
!$acc parallel default (present) async (1)
do i = 1, N
b(i) = (a(i) * a(i) * a(i)) / a(i)
end do
!$acc end parallel
!$acc parallel async (1)
!$acc parallel default (present) async (1)
do i = 1, N
c(i) = (a(i) * 4) / a(i)
end do
!$acc end parallel
!$acc parallel async (1)
!$acc parallel default (present) async (1)
do i = 1, N
d(i) = ((a(i) * a(i) + a(i)) / a(i)) - a(i)
end do
!$acc end parallel
!$acc parallel wait (1) async (1)
!$acc parallel default (present) wait (1) async (1)
do i = 1, N
e(i) = a(i) + b(i) + c(i) + d(i)
end do

10
libgomp/testsuite/libgomp.oacc-fortran/default-1.f90
View file

@ -51,4 +51,14 @@ program main
if (a .ne. 7.0) call abort
! The default (present) clause doesn't affect scalar variables; these will
! still get an implicit copy clause added.
!$acc kernels default (present)
c = a
a = 1.0
a = a + c
!$acc end kernels
if (a .ne. 8.0) call abort
end program main

44
libgomp/testsuite/libgomp.oacc-fortran/non-scalar-data.f90
View file

@ -1,5 +1,6 @@
! Ensure that a non-scalar dummy arguments which are implicitly used inside
! offloaded regions are properly mapped using present_or_copy.
! offloaded regions are properly mapped using present_or_copy, or (default)
! present.
! { dg-do run }
@ -12,6 +13,7 @@ program main
n = size
!$acc data copy(array)
call kernels(array, n)
!$acc update host(array)
@ -20,12 +22,29 @@ program main
if (array(i) .ne. i) call abort
end do
call kernels_default_present(array, n)
!$acc update host(array)
do i = 1, n
if (array(i) .ne. i+1) call abort
end do
call parallel(array, n)
!$acc end data
!$acc update host(array)
do i = 1, n
if (array(i) .ne. i+i) call abort
end do
call parallel_default_present(array, n)
!$acc end data
do i = 1, n
if (array(i) .ne. i+i+1) call abort
end do
end program main
subroutine kernels (array, n)
@ -39,6 +58,16 @@ subroutine kernels (array, n)
!$acc end kernels
end subroutine kernels
subroutine kernels_default_present (array, n)
integer, dimension (n) :: array
integer :: n, i
!$acc kernels default(present)
do i = 1, n
array(i) = i+1
end do
!$acc end kernels
end subroutine kernels_default_present
subroutine parallel (array, n)
integer, dimension (n) :: array
@ -50,3 +79,14 @@ subroutine parallel (array, n)
end do
!$acc end parallel
end subroutine parallel
subroutine parallel_default_present (array, n)
integer, dimension (n) :: array
integer :: n, i
!$acc parallel default(present)
do i = 1, n
array(i) = i+i+1
end do
!$acc end parallel
end subroutine parallel_default_present