+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * omp-low.c (lower_omp_target): Support GOMP_MAP_NO_ALLOC.
+ * tree-pretty-print.c (dump_omp_clause): Likewise.
+
2019-12-18 Eric Botcazou <ebotcazou@adacore.com>
* ira.c (ira): Use simple LRA algorithm when not optimizing.
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-pragma.h (pragma_omp_clause): Add
+ PRAGMA_OACC_CLAUSE_NO_CREATE.
+
2019-12-17 Martin Sebor <msebor@redhat.com>
PR c++/61339
PRAGMA_OACC_CLAUSE_GANG,
PRAGMA_OACC_CLAUSE_HOST,
PRAGMA_OACC_CLAUSE_INDEPENDENT,
+ PRAGMA_OACC_CLAUSE_NO_CREATE,
PRAGMA_OACC_CLAUSE_NUM_GANGS,
PRAGMA_OACC_CLAUSE_NUM_WORKERS,
PRAGMA_OACC_CLAUSE_PRESENT,
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * c-parser.c (c_parser_omp_clause_name): Support no_create.
+ (c_parser_oacc_data_clause): Likewise.
+ (c_parser_oacc_all_clauses): Likewise.
+ (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
+ (OACC_PARALLEL_CLAUSE_MASK, OACC_SERIAL_CLAUSE_MASK): Add
+ PRAGMA_OACC_CLAUSE_NO_CREATE.
+ * c-typeck.c (handle_omp_array_sections): Support
+ GOMP_MAP_NO_ALLOC.
+
2019-12-09 David Malcolm <dmalcolm@redhat.com>
* c-objc-common.c (range_label_for_type_mismatch::get_text):
result = PRAGMA_OMP_CLAUSE_MERGEABLE;
break;
case 'n':
- if (!strcmp ("nogroup", p))
+ if (!strcmp ("no_create", p))
+ result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
copyout ( variable-list )
create ( variable-list )
delete ( variable-list )
- present ( variable-list ) */
+ present ( variable-list )
+
+ OpenACC 2.6:
+ no_create ( variable-list ) */
static tree
c_parser_oacc_data_clause (c_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "no_create";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
clauses = c_parser_oacc_single_int_clause (parser,
OMP_CLAUSE_NUM_GANGS,
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT))
static tree
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * parser.c (cp_parser_omp_clause_name): Support no_create.
+ (cp_parser_oacc_data_clause): Likewise.
+ (cp_parser_oacc_all_clauses): Likewise.
+ (OACC_DATA_CLAUSE_MASK, OACC_KERNELS_CLAUSE_MASK)
+ (OACC_PARALLEL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_NO_CREATE.
+ * semantics.c (handle_omp_array_sections): Support no_create.
+
2019-12-18 Paolo Carlini <paolo.carlini@oracle.com>
* typeck.c (cxx_sizeof_or_alignof_type): Add location_t parameter
result = PRAGMA_OMP_CLAUSE_MERGEABLE;
break;
case 'n':
- if (!strcmp ("nogroup", p))
+ if (!strcmp ("no_create", p))
+ result = PRAGMA_OACC_CLAUSE_NO_CREATE;
+ else if (!strcmp ("nogroup", p))
result = PRAGMA_OMP_CLAUSE_NOGROUP;
else if (!strcmp ("nontemporal", p))
result = PRAGMA_OMP_CLAUSE_NONTEMPORAL;
copyout ( variable-list )
create ( variable-list )
delete ( variable-list )
- present ( variable-list ) */
+ present ( variable-list )
+
+ OpenACC 2.6:
+ no_create ( variable-list ) */
static tree
cp_parser_oacc_data_clause (cp_parser *parser, pragma_omp_clause c_kind,
case PRAGMA_OACC_CLAUSE_LINK:
kind = GOMP_MAP_LINK;
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ kind = GOMP_MAP_IF_PRESENT;
+ break;
case PRAGMA_OACC_CLAUSE_PRESENT:
kind = GOMP_MAP_FORCE_PRESENT;
break;
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "link";
break;
+ case PRAGMA_OACC_CLAUSE_NO_CREATE:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "no_create";
+ break;
case PRAGMA_OACC_CLAUSE_NUM_GANGS:
code = OMP_CLAUSE_NUM_GANGS;
c_name = "num_gangs";
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
static tree
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (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_FIRSTPRIVATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NO_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_GANGS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_NUM_WORKERS) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
| (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) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FIRSTPRIVATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) \
switch (OMP_CLAUSE_MAP_KIND (c))
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * gfortran.h (gfc_omp_map_op): Add OMP_MAP_NO_ALLOC.
+ * openmp.c (omp_mask2): Add OMP_CLAUSE_NO_CREATE.
+ (gfc_match_omp_clauses): Support no_create.
+ (OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES)
+ (OACC_DATA_CLAUSES): Add OMP_CLAUSE_NO_CREATE.
+ * trans-openmp.c (gfc_trans_omp_clauses_1): Support
+ OMP_MAP_NO_ALLOC.
+
2019-12-18 Harald Anlauf <anlauf@gmx.de>
PR fortran/70853
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_IF_PRESENT,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
OMP_CLAUSE_COPY,
OMP_CLAUSE_COPYOUT,
OMP_CLAUSE_CREATE,
+ OMP_CLAUSE_NO_CREATE,
OMP_CLAUSE_PRESENT,
OMP_CLAUSE_DEVICEPTR,
OMP_CLAUSE_GANG,
}
break;
case 'n':
+ if ((mask & OMP_CLAUSE_NO_CREATE)
+ && gfc_match ("no_create ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_IF_PRESENT, true))
+ continue;
if ((mask & OMP_CLAUSE_NOGROUP)
&& !c->nogroup
&& gfc_match ("nogroup") == MATCH_YES)
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
- | OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_CREATE | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT \
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_WAIT)
#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_PRESENT)
+ | OMP_CLAUSE_NO_CREATE | OMP_CLAUSE_PRESENT)
#define OACC_LOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
| OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
#define OMP_TASKLOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_LASTPRIVATE | OMP_CLAUSE_SHARED | OMP_CLAUSE_IF \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_UNTIED | OMP_CLAUSE_FINAL \
| OMP_CLAUSE_MERGEABLE | OMP_CLAUSE_PRIORITY | OMP_CLAUSE_GRAINSIZE \
| OMP_CLAUSE_NUM_TASKS | OMP_CLAUSE_COLLAPSE | OMP_CLAUSE_NOGROUP)
#define OMP_TARGET_CLAUSES \
| OMP_CLAUSE_FROM | OMP_CLAUSE_DEPEND | OMP_CLAUSE_NOWAIT)
#define OMP_TEAMS_CLAUSES \
(omp_mask (OMP_CLAUSE_NUM_TEAMS) | OMP_CLAUSE_THREAD_LIMIT \
- | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
| OMP_CLAUSE_SHARED | OMP_CLAUSE_REDUCTION)
#define OMP_DISTRIBUTE_CLAUSES \
(omp_mask (OMP_CLAUSE_PRIVATE) | OMP_CLAUSE_FIRSTPRIVATE \
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_IF_PRESENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
+ break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
break;
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
switch (tkind)
{
case GOMP_MAP_ALLOC:
+ case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_TOFROM:
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * gfortran.dg/goacc/common-block-1.f90: Add no_create-clause tests.
+ * gfortran.dg/goacc/common-block-1.f90: Likewise.
+ * gfortran.dg/goacc/data-clauses.f95: Likewise.
+ * gfortran.dg/goacc/data-tree.f95: Likewise.
+ * gfortran.dg/goacc/kernels-tree.f95: Likewise.
+ * gfortran.dg/goacc/parallel-tree.f95: Likewise.
+
2019-12-18 Paolo Carlini <paolo.carlini@oracle.com>
* g++.dg/diagnostic/alignof2.C: New.
!$acc data pcopyout(/blockA/, /blockB/, e, v)
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v)
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v)
!$acc end parallel
!$acc data pcopyout(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end data
+ !$acc data no_create(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
+ !$acc end data
+
!$acc parallel private(/blockA/, /blockB/, e, v, a) ! { dg-error "Symbol .a. present on multiple clauses" }
!$acc end parallel
!$acc end data
+ !$acc parallel no_create (tip) ! { dg-error "POINTER" }
+ !$acc end parallel
+ !$acc parallel no_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc end parallel
+ !$acc parallel deviceptr (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copy (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyin (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+ !$acc parallel copyout (i) no_create (i) ! { dg-error "multiple clauses" }
+ !$acc end parallel
+
+ !$acc parallel no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end parallel
+ !$acc kernels no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end kernels
+ !$acc data no_create (i, c, r, ia, ca, ra, asa, rp, ti, vi, aa)
+ !$acc end data
+
+
!$acc parallel present (tip) ! { dg-error "POINTER" }
!$acc end parallel
!$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
logical :: l = .true.
!$acc data if(l) copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end data
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
-
+! { dg-final { scan-tree-dump-times "map\\(no_alloc:n\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
!$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u)
!$acc end kernels
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
-
+! { dg-final { scan-tree-dump-times "map\\(no_alloc:n\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
!$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) &
!$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) &
+ !$acc no_create(n) &
!$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) &
!$acc deviceptr(u), private(v), firstprivate(w)
!$acc end parallel
! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(alloc:m\\)" 1 "original" } }
-
+! { dg-final { scan-tree-dump-times "map\\(no_alloc:n\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(force_present:o\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(tofrom:p\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "map\\(to:r\\)" 1 "original" } }
case GOMP_MAP_POINTER:
pp_string (pp, "alloc");
break;
+ case GOMP_MAP_IF_PRESENT:
+ pp_string (pp, "no_alloc");
+ break;
case GOMP_MAP_TO:
case GOMP_MAP_TO_PSET:
pp_string (pp, "to");
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * gomp-constants.h (gomp_map_kind): Support GOMP_MAP_NO_ALLOC.
+
2019-11-16 Tim Ruehsen <tim.ruehsen@gmx.de>
* demangle.h (struct demangle_component): Add member
GOMP_MAP_DEVICE_RESIDENT = (GOMP_MAP_FLAG_SPECIAL_1 | 1),
/* OpenACC link. */
GOMP_MAP_LINK = (GOMP_MAP_FLAG_SPECIAL_1 | 2),
+ /* Use device data if present, fall back to host address otherwise. */
+ GOMP_MAP_IF_PRESENT = (GOMP_MAP_FLAG_SPECIAL_1 | 3),
/* Do not map, copy bits for firstprivate instead. */
GOMP_MAP_FIRSTPRIVATE = (GOMP_MAP_FLAG_SPECIAL | 0),
/* Similarly, but store the value in the pointer rather than
+2019-12-19 Julian Brown <julian@codesourcery.com>
+ Maciej W. Rozycki <macro@codesourcery.com>
+ Tobias Burnus <tobias@codesourcery.com>
+ Thomas Schwinge <thomas@codesourcery.com>
+
+ * target.c (gomp_map_vars_async): Support GOMP_MAP_NO_ALLOC.
+ * testsuite/libgomp.oacc-c-c++-common/no_create-1.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/no_create-2.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/no_create-3.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/no_create-4.c: New test.
+ * testsuite/libgomp.oacc-c-c++-common/no_create-5.c: New test.
+ * testsuite/libgomp.oacc-fortran/no_create-1.f90: New test.
+ * testsuite/libgomp.oacc-fortran/no_create-2.f90: New test.
+ * testsuite/libgomp.oacc-fortran/no_create-3.F90: New test.
+
2019-12-18 Thomas Schwinge <thomas@codesourcery.com>
* oacc-mem.c (goacc_enter_data): Refactor, so that it can be
{
tgt->list[i].key = NULL;
+ if ((kind & typemask) == GOMP_MAP_IF_PRESENT)
+ {
+ /* Not present, hence, skip entry - including its MAP_POINTER,
+ when existing. */
+ tgt->list[i].offset = 0;
+ if (i + 1 < mapnum
+ && ((typemask & get_kind (short_mapkind, kinds, i + 1))
+ == GOMP_MAP_POINTER))
+ {
+ ++i;
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = 0;
+ }
+ continue;
+ }
size_t align = (size_t) 1 << (kind >> rshift);
not_found_cnt++;
if (tgt_align < align)
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
+ case GOMP_MAP_IF_PRESENT:
+ /* Not present - otherwise handled above. Skip over its
+ MAP_POINTER as well. */
+ if (i + 1 < mapnum
+ && ((typemask & get_kind (short_mapkind, kinds, i + 1))
+ == GOMP_MAP_POINTER))
+ ++i;
+ continue;
default:
break;
}
--- /dev/null
+/* Test 'no_create' clause on compute construct, with data present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+ free (arr);
+
+ return 0;
+}
--- /dev/null
+/* Test 'no_create' clause on compute construct, with data not present on the
+ device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc parallel no_create(var, arr[0:N]) copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+
+ free (arr);
+
+ return 0;
+}
--- /dev/null
+#include <float.h> /* For FLT_EPSILON. */
+#include <math.h> /* For fabs. */
+#include <stdlib.h> /* For abort. */
+
+
+int main()
+{
+#define N 100
+ float b[N];
+ float c[N];
+
+#pragma acc enter data create(b)
+
+#pragma acc parallel loop no_create(b) no_create(c)
+ for (int i = 0; i < N; ++i)
+ b[i] = i;
+
+#pragma acc exit data copyout(b)
+
+ for (int i = 0; i < N; ++i)
+ if (fabs (b[i] - i) > 10.0*FLT_EPSILON)
+ abort ();
+
+ return 0;
+}
--- /dev/null
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+ acc_copyin (&var, sizeof (var));
+ acc_copyin (arr, N * sizeof (*arr));
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr)
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+
+ if (acc_hostptr (devptr[0]) != (void *) &var)
+ __builtin_abort ();
+ if (acc_hostptr (devptr[1]) != (void *) &arr[2])
+ __builtin_abort ();
+
+#if ACC_MEM_SHARED
+ if (devptr[0] != &var)
+ __builtin_abort ();
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+#else
+ if (devptr[0] == &var)
+ __builtin_abort ();
+ if (devptr[1] == &arr[2])
+ __builtin_abort ();
+#endif
+ }
+
+ acc_delete (&var, sizeof (var));
+ acc_delete (arr, N * sizeof (*arr));
+
+ free (arr);
+
+ return 0;
+}
--- /dev/null
+/* Test 'no_create' clause on 'data' construct and nested compute construct,
+ with data not present on the device. */
+
+#include <stdlib.h>
+#include <stdio.h>
+#include <openacc.h>
+
+#define N 128
+
+int
+main (int argc, char *argv[])
+{
+ int var;
+ int *arr = (int *) malloc (N * sizeof (*arr));
+ int *devptr[2];
+
+#pragma acc data no_create(var, arr[0:N])
+ {
+ devptr[0] = (int *) acc_deviceptr (&var);
+ devptr[1] = (int *) acc_deviceptr (&arr[2]);
+
+#if ACC_MEM_SHARED
+ if (devptr[0] == NULL)
+ __builtin_abort ();
+ if (devptr[1] == NULL)
+ __builtin_abort ();
+#else
+ if (devptr[0] != NULL)
+ __builtin_abort ();
+ if (devptr[1] != NULL)
+ __builtin_abort ();
+#endif
+
+#pragma acc parallel copyout(devptr) // TODO implicit 'copy(var)' -- huh?!
+ {
+ devptr[0] = &var;
+ devptr[1] = &arr[2];
+ }
+
+ if (devptr[0] != &var)
+ __builtin_abort (); // { dg-xfail-run-if "TODO" { *-*-* } { "-DACC_MEM_SHARED=0" } }
+ if (devptr[1] != &arr[2])
+ __builtin_abort ();
+ }
+
+ free (arr);
+
+ return 0;
+}
--- /dev/null
+! { dg-do run }
+
+! Test no_create clause with data construct when data is present/not present.
+
+program no_create
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myvar, myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ myvar = 77
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .neqv. shared_memory) stop 10
+ if (acc_is_present (myarr) .neqv. shared_memory) stop 11
+ !$acc end data
+
+ !$acc enter data copyin (myvar, myarr)
+ !$acc data no_create (myvar, myarr)
+ if (acc_is_present (myvar) .eqv. .false.) stop 20
+ if (acc_is_present (myarr) .eqv. .false.) stop 21
+ !$acc end data
+ !$acc exit data copyout (myvar, myarr)
+
+ if (myvar .ne. 77) stop 30
+ do i = 1, n
+ if (myarr(i) .ne. 0) stop 31
+ end do
+end program no_create
--- /dev/null
+! { dg-do run }
+
+! Test no_create clause with data/parallel constructs.
+
+program no_create
+ use openacc
+ implicit none
+ logical :: shared_memory
+ integer, parameter :: n = 512
+ integer :: myvar, myarr(n)
+ integer i
+
+ shared_memory = .false.
+ !$acc kernels copyin (shared_memory)
+ shared_memory = .true.
+ !$acc end kernels
+
+ myvar = 55
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ call do_on_target(myvar, n, myarr)
+
+ if (shared_memory) then
+ if (myvar .ne. 44) stop 10
+ else
+ if (myvar .ne. 33) stop 11
+ end if
+ do i = 1, n
+ if (shared_memory) then
+ if (myarr(i) .ne. i * 2) stop 20
+ else
+ if (myarr(i) .ne. i) stop 21
+ end if
+ end do
+
+ myvar = 55
+ do i = 1, n
+ myarr(i) = 0
+ end do
+
+ !$acc enter data copyin(myvar, myarr)
+ call do_on_target(myvar, n, myarr)
+ !$acc exit data copyout(myvar, myarr)
+
+ if (myvar .ne. 44) stop 30
+ do i = 1, n
+ if (myarr(i) .ne. i * 2) stop 31
+ end do
+end program no_create
+
+subroutine do_on_target (var, n, arr)
+ use openacc
+ implicit none
+ integer :: var, n, arr(n)
+ integer :: i
+
+!$acc data no_create (var, arr)
+
+if (acc_is_present(var)) then
+ ! The no_create clause is meant for partially shared-memory machines. This
+ ! test is written to work on non-shared-memory machines, though this is not
+ ! necessarily a useful way to use the no_create clause in practice.
+
+ !$acc parallel !no_create (var)
+ var = 44
+ !$acc end parallel
+else
+ var = 33
+end if
+if (acc_is_present(arr)) then
+ ! The no_create clause is meant for partially shared-memory machines. This
+ ! test is written to work on non-shared-memory machines, though this is not
+ ! necessarily a useful way to use the no_create clause in practice.
+
+ !$acc parallel loop !no_create (arr)
+ do i = 1, n
+ arr(i) = i * 2
+ end do
+ !$acc end parallel loop
+else
+ do i = 1, n
+ arr(i) = i
+ end do
+end if
+
+!$acc end data
+
+end subroutine do_on_target
--- /dev/null
+! { dg-do run }
+
+program main
+ use iso_c_binding, only: c_sizeof
+ use openacc, only: acc_is_present
+ implicit none
+ integer i
+ integer, parameter :: n = 100
+ real*4 b(n), c(n)
+ real :: d(n), e(n)
+ common /BLOCK/ d, e
+
+ !$acc enter data create(b) create(d)
+
+ if (.not. acc_is_present(b, c_sizeof(b))) stop 1
+ if (.not. acc_is_present(d, c_sizeof(d))) stop 2
+#if !ACC_MEM_SHARED
+ if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 3
+ if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(d))) stop 4
+#endif
+
+ !$acc parallel loop no_create(b) no_create(c) no_create(/BLOCK/)
+ do i = 1, n
+ b(i) = i
+ d(i) = -i
+ end do
+ !$acc end parallel loop
+
+ if (.not. acc_is_present(b, c_sizeof(b))) stop 5
+ if (.not. acc_is_present(d, c_sizeof(d))) stop 6
+#if !ACC_MEM_SHARED
+ if (acc_is_present(c, 1) .or. acc_is_present(c, c_sizeof(c))) stop 7
+ if (acc_is_present(e, 1) .or. acc_is_present(e, c_sizeof(e))) stop 8
+#endif
+
+ !$acc exit data copyout(b) copyout(d)
+ if (any(abs(b - [(real(i), i = 1, n)]) > 10*epsilon(b))) stop 9
+ if (any(abs(d - [(real(-i), i = 1, n)]) > 10*epsilon(d))) stop 10
+end program main
projects / gcc.git / commitdiff