OpenACC: Separate enter/exit data ABIs
Move the OpenACC enter and exit data directives from using a single builtin to having one each. For most purposes it was easy to tell which was which, from the clauses given, but it's overhead we can easily avoid, and there may be future uses where that isn't possible. gcc/ * omp-builtins.def (BUILT_IN_GOACC_ENTER_EXIT_DATA): Split into... (BUILT_IN_GOACC_ENTER_DATA, BUILT_IN_GOACC_EXIT_DATA): ... these. * gimple.h (enum gf_mask): Split 'GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA' into 'GF_OMP_TARGET_KIND_OACC_ENTER_DATA' and 'GF_OMP_TARGET_KIND_OACC_EXIT_DATA'. (is_gimple_omp_oacc): Update. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * gimplify.c (gimplify_omp_target_update): Likewise. * omp-expand.c (expand_omp_target, build_omp_regions_1) (omp_make_gimple_edges): Likewise. * omp-low.c (check_omp_nesting_restrictions, lower_omp_target): Likewise. gcc/testsuite/ * c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust patterns. * c-c++-common/goacc/finalize-1.c: Likewise. * c-c++-common/goacc/mdc-1.c: Likewise. * c-c++-common/goacc/nesting-fail-1.c: Likewise. * c-c++-common/goacc/struct-enter-exit-data-1.c: Likewise. * gfortran.dg/goacc/attach-descriptor.f90: Likewise. * gfortran.dg/goacc/finalize-1.f: Likewise. * gfortran.dg/goacc/mapping-tests-3.f90: Likewise. libgomp/ * libgomp.map (GOACC_2.0.2): New symbol version. * libgomp_g.h (GOACC_enter_data, GOACC_exit_data) New prototypes. * oacc-mem.c (GOACC_enter_data, GOACC_exit_data) New functions. Co-Authored-By: Thomas Schwinge <thomas@codesourcery.com>
This commit is contained in:
Andrew Stubbs
Thomas Schwinge
parent
7999363961
commit
7aefef3136
|
|
@ -1697,8 +1697,11 @@ dump_gimple_omp_target (pretty_printer *buffer, const gomp_target *gs,
|
|||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
kind = " oacc_update";
|
||||
break;
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
kind = " oacc_enter_exit_data";
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
kind = " oacc_enter_data";
|
||||
break;
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
kind = " oacc_exit_data";
|
||||
break;
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
kind = " oacc_declare";
|
||||
|
|
|
|||
18
gcc/gimple.h
18
gcc/gimple.h
|
|
@ -161,7 +161,7 @@ enum gf_mask {
|
|||
GF_OMP_FOR_KIND_SIMD = 5,
|
||||
GF_OMP_FOR_COMBINED = 1 << 3,
|
||||
GF_OMP_FOR_COMBINED_INTO = 1 << 4,
|
||||
GF_OMP_TARGET_KIND_MASK = (1 << 4) - 1,
|
||||
GF_OMP_TARGET_KIND_MASK = (1 << 5) - 1,
|
||||
GF_OMP_TARGET_KIND_REGION = 0,
|
||||
GF_OMP_TARGET_KIND_DATA = 1,
|
||||
GF_OMP_TARGET_KIND_UPDATE = 2,
|
||||
|
|
@ -172,18 +172,19 @@ enum gf_mask {
|
|||
GF_OMP_TARGET_KIND_OACC_SERIAL = 7,
|
||||
GF_OMP_TARGET_KIND_OACC_DATA = 8,
|
||||
GF_OMP_TARGET_KIND_OACC_UPDATE = 9,
|
||||
GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA = 10,
|
||||
GF_OMP_TARGET_KIND_OACC_DECLARE = 11,
|
||||
GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12,
|
||||
GF_OMP_TARGET_KIND_OACC_ENTER_DATA = 10,
|
||||
GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 11,
|
||||
GF_OMP_TARGET_KIND_OACC_DECLARE = 12,
|
||||
GF_OMP_TARGET_KIND_OACC_HOST_DATA = 13,
|
||||
/* A 'GF_OMP_TARGET_KIND_OACC_PARALLEL' representing an OpenACC 'kernels'
|
||||
decomposed part, parallelized. */
|
||||
GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED = 13,
|
||||
GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED = 14,
|
||||
/* A 'GF_OMP_TARGET_KIND_OACC_PARALLEL' representing an OpenACC 'kernels'
|
||||
decomposed part, "gang-single". */
|
||||
GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE = 14,
|
||||
GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE = 15,
|
||||
/* A 'GF_OMP_TARGET_KIND_OACC_DATA' representing an OpenACC 'kernels'
|
||||
decomposed parts' 'data' construct. */
|
||||
GF_OMP_TARGET_KIND_OACC_DATA_KERNELS = 15,
|
||||
GF_OMP_TARGET_KIND_OACC_DATA_KERNELS = 16,
|
||||
GF_OMP_TEAMS_HOST = 1 << 0,
|
||||
|
||||
/* True on an GIMPLE_OMP_RETURN statement if the return does not require
|
||||
|
|
@ -6525,7 +6526,8 @@ is_gimple_omp_oacc (const gimple *stmt)
|
|||
case GF_OMP_TARGET_KIND_OACC_SERIAL:
|
||||
case GF_OMP_TARGET_KIND_OACC_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
|
||||
|
|
|
|||
|
|
@ -13383,8 +13383,11 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
|
|||
switch (TREE_CODE (expr))
|
||||
{
|
||||
case OACC_ENTER_DATA:
|
||||
kind = GF_OMP_TARGET_KIND_OACC_ENTER_DATA;
|
||||
ort = ORT_ACC;
|
||||
break;
|
||||
case OACC_EXIT_DATA:
|
||||
kind = GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA;
|
||||
kind = GF_OMP_TARGET_KIND_OACC_EXIT_DATA;
|
||||
ort = ORT_ACC;
|
||||
break;
|
||||
case OACC_UPDATE:
|
||||
|
|
|
|||
|
|
@ -35,7 +35,10 @@ DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_START, "GOACC_data_start",
|
|||
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR, ATTR_NOTHROW_LIST)
|
||||
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DATA_END, "GOACC_data_end",
|
||||
BT_FN_VOID, ATTR_NOTHROW_LIST)
|
||||
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_EXIT_DATA, "GOACC_enter_exit_data",
|
||||
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_ENTER_DATA, "GOACC_enter_data",
|
||||
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
|
||||
ATTR_NOTHROW_LIST)
|
||||
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_EXIT_DATA, "GOACC_exit_data",
|
||||
BT_FN_VOID_INT_SIZE_PTR_PTR_PTR_INT_INT_VAR,
|
||||
ATTR_NOTHROW_LIST)
|
||||
DEF_GOACC_BUILTIN (BUILT_IN_GOACC_PARALLEL, "GOACC_parallel_keyed",
|
||||
|
|
|
|||
|
|
@ -9290,7 +9290,8 @@ expand_omp_target (struct omp_region *region)
|
|||
case GF_OMP_TARGET_KIND_OACC_KERNELS:
|
||||
case GF_OMP_TARGET_KIND_OACC_SERIAL:
|
||||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
|
||||
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
|
||||
|
|
@ -9574,8 +9575,11 @@ expand_omp_target (struct omp_region *region)
|
|||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
start_ix = BUILT_IN_GOACC_UPDATE;
|
||||
break;
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
start_ix = BUILT_IN_GOACC_ENTER_EXIT_DATA;
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
start_ix = BUILT_IN_GOACC_ENTER_DATA;
|
||||
break;
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
start_ix = BUILT_IN_GOACC_EXIT_DATA;
|
||||
break;
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
start_ix = BUILT_IN_GOACC_DECLARE;
|
||||
|
|
@ -9773,7 +9777,8 @@ expand_omp_target (struct omp_region *region)
|
|||
oacc_set_fn_attrib (child_fn, clauses, &args);
|
||||
tagging = true;
|
||||
/* FALLTHRU */
|
||||
case BUILT_IN_GOACC_ENTER_EXIT_DATA:
|
||||
case BUILT_IN_GOACC_ENTER_DATA:
|
||||
case BUILT_IN_GOACC_EXIT_DATA:
|
||||
case BUILT_IN_GOACC_UPDATE:
|
||||
{
|
||||
tree t_async = NULL_TREE;
|
||||
|
|
@ -10042,7 +10047,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent,
|
|||
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
|
||||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
/* ..., other than for those stand-alone directives... */
|
||||
region = NULL;
|
||||
|
|
@ -10299,7 +10305,8 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region,
|
|||
case GF_OMP_TARGET_KIND_OACC_HOST_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DATA_KERNELS:
|
||||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
cur_region = cur_region->outer;
|
||||
break;
|
||||
|
|
|
|||
|
|
@ -3646,8 +3646,10 @@ check_omp_nesting_restrictions (gimple *stmt, omp_context *ctx)
|
|||
case GF_OMP_TARGET_KIND_OACC_SERIAL: stmt_name = "serial"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_DATA: stmt_name = "data"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_UPDATE: stmt_name = "update"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
stmt_name = "enter/exit data"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
stmt_name = "enter data"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
stmt_name = "exit data"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE: stmt_name = "declare"; break;
|
||||
case GF_OMP_TARGET_KIND_OACC_HOST_DATA: stmt_name = "host_data";
|
||||
break;
|
||||
|
|
@ -12194,7 +12196,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
|
|||
case GF_OMP_TARGET_KIND_OACC_KERNELS:
|
||||
case GF_OMP_TARGET_KIND_OACC_SERIAL:
|
||||
case GF_OMP_TARGET_KIND_OACC_UPDATE:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_ENTER_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_EXIT_DATA:
|
||||
case GF_OMP_TARGET_KIND_OACC_DECLARE:
|
||||
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED:
|
||||
case GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE:
|
||||
|
|
|
|||
|
|
@ -162,8 +162,8 @@ f_omp (void)
|
|||
#pragma acc data /* { dg-error "OpenACC .data. construct inside of OpenMP .target. region" } */
|
||||
;
|
||||
#pragma acc update host(i) /* { dg-error "OpenACC .update. construct inside of OpenMP .target. region" } */
|
||||
#pragma acc enter data copyin(i) /* { dg-error "OpenACC .enter/exit data. construct inside of OpenMP .target. region" } */
|
||||
#pragma acc exit data delete(i) /* { dg-error "OpenACC .enter/exit data. construct inside of OpenMP .target. region" } */
|
||||
#pragma acc enter data copyin(i) /* { dg-error "OpenACC .enter data. construct inside of OpenMP .target. region" } */
|
||||
#pragma acc exit data delete(i) /* { dg-error "OpenACC .exit data. construct inside of OpenMP .target. region" } */
|
||||
#pragma acc loop /* { dg-error "loop directive must be associated with an OpenACC compute region" } */
|
||||
for (i = 0; i < 2; ++i)
|
||||
;
|
||||
|
|
|
|||
|
|
@ -13,25 +13,25 @@ void f ()
|
|||
{
|
||||
#pragma acc exit data delete (del_r)
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data finalize delete (del_f)
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data finalize delete (del_f_p[2:5])
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data copyout (cpo_r)
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data copyout (cpo_f) finalize
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
|
||||
|
||||
#pragma acc exit data copyout (cpo_f_p[4:10]) finalize
|
||||
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
|
||||
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
|
||||
}
|
||||
|
|
|
|||
|
|
@ -43,15 +43,15 @@ t1 ()
|
|||
}
|
||||
}
|
||||
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:s .len: 32.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.to:a .len: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data map.release:a .len: 8.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:a .bias: 0.." 1 "omplower" } } */
|
||||
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_exit_data finalize map.force_detach:s.a .bias: 0.." 1 "omplower" } } */
|
||||
|
|
|
|||
|
|
@ -14,8 +14,8 @@ f_acc_parallel (void)
|
|||
#pragma acc data /* { dg-error ".data. construct inside of .parallel. region" } */
|
||||
;
|
||||
#pragma acc update host(i) /* { dg-error ".update. construct inside of .parallel. region" } */
|
||||
#pragma acc enter data copyin(i) /* { dg-error ".enter/exit data. construct inside of .parallel. region" } */
|
||||
#pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .parallel. region" } */
|
||||
#pragma acc enter data copyin(i) /* { dg-error ".enter data. construct inside of .parallel. region" } */
|
||||
#pragma acc exit data delete(i) /* { dg-error ".exit data. construct inside of .parallel. region" } */
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -33,8 +33,8 @@ f_acc_kernels (void)
|
|||
#pragma acc data /* { dg-error ".data. construct inside of .kernels. region" } */
|
||||
;
|
||||
#pragma acc update host(i) /* { dg-error ".update. construct inside of .kernels. region" } */
|
||||
#pragma acc enter data copyin(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */
|
||||
#pragma acc exit data delete(i) /* { dg-error ".enter/exit data. construct inside of .kernels. region" } */
|
||||
#pragma acc enter data copyin(i) /* { dg-error ".enter data. construct inside of .kernels. region" } */
|
||||
#pragma acc exit data delete(i) /* { dg-error ".exit data. construct inside of .kernels. region" } */
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -20,8 +20,8 @@ test (int *b, int *c, int *e)
|
|||
struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
|
||||
|
||||
#pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
|
||||
/* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
|
||||
/* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
|
||||
|
||||
#pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
|
||||
/* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
|
||||
/* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
|
||||
}
|
||||
|
|
|
|||
|
|
@ -12,11 +12,11 @@ program att
|
|||
|
||||
!$acc enter data attach(myvar%arr2, myptr)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
|
||||
|
||||
!$acc exit data detach(myvar%arr2, myptr)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
|
||||
|
||||
! Test valid usage and processing of the finalize clause.
|
||||
!$acc exit data detach(myvar%arr2, myptr) finalize
|
||||
|
|
@ -24,6 +24,6 @@ program att
|
|||
! For array-descriptor detaches, we no longer generate a "release" mapping
|
||||
! for the pointed-to data for gimplify.c to turn into "delete". Make sure
|
||||
! the mapping still isn't there.
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
|
||||
|
||||
end program att
|
||||
|
|
|
|||
|
|
@ -13,25 +13,25 @@
|
|||
|
||||
!$ACC EXIT DATA DELETE (del_r)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(release:del_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA FINALIZE DELETE (del_f)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_r)
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
|
||||
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
|
||||
END SUBROUTINE f
|
||||
|
|
|
|||
|
|
@ -11,5 +11,5 @@ subroutine foo
|
|||
type(two) x
|
||||
|
||||
!$acc enter data copyin(x%A)
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_enter_exit_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
|
||||
! { dg-final { scan-tree-dump-times "omp target oacc_enter_data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } }
|
||||
end
|
||||
|
|
|
|||
|
|
@ -541,6 +541,12 @@ GOACC_2.0.1 {
|
|||
GOACC_parallel_keyed;
|
||||
} GOACC_2.0;
|
||||
|
||||
GOACC_2.0.2 {
|
||||
global:
|
||||
GOACC_enter_data;
|
||||
GOACC_exit_data;
|
||||
} GOACC_2.0.1;
|
||||
|
||||
GOMP_PLUGIN_1.0 {
|
||||
global:
|
||||
GOMP_PLUGIN_malloc;
|
||||
|
|
|
|||
|
|
@ -370,6 +370,10 @@ extern void GOACC_wait (int, int, ...);
|
|||
|
||||
extern void GOACC_enter_exit_data (int, size_t, void **, size_t *,
|
||||
unsigned short *, int, int, ...);
|
||||
extern void GOACC_enter_data (int, size_t, void **, size_t *,
|
||||
unsigned short *, int, int, ...);
|
||||
extern void GOACC_exit_data (int, size_t, void **, size_t *,
|
||||
unsigned short *, int, int, ...);
|
||||
extern void GOACC_declare (int, size_t, void **, size_t *, unsigned short *);
|
||||
|
||||
/* oacc-parallel.c */
|
||||
|
|
|
|||
|
|
@ -1422,6 +1422,8 @@ goacc_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs,
|
|||
}
|
||||
}
|
||||
|
||||
/* Legacy entry point (GCC 11 and earlier). */
|
||||
|
||||
void
|
||||
GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
|
||||
size_t *sizes, unsigned short *kinds, int async,
|
||||
|
|
@ -1468,6 +1470,30 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
|
|||
va_end (ap);
|
||||
}
|
||||
|
||||
void
|
||||
GOACC_enter_data (int flags_m, size_t mapnum, void **hostaddrs,
|
||||
size_t *sizes, unsigned short *kinds, int async,
|
||||
int num_waits, ...)
|
||||
{
|
||||
va_list ap;
|
||||
va_start (ap, num_waits);
|
||||
goacc_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds,
|
||||
true, async, num_waits, &ap);
|
||||
va_end (ap);
|
||||
}
|
||||
|
||||
void
|
||||
GOACC_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
|
||||
size_t *sizes, unsigned short *kinds, int async,
|
||||
int num_waits, ...)
|
||||
{
|
||||
va_list ap;
|
||||
va_start (ap, num_waits);
|
||||
goacc_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds,
|
||||
false, async, num_waits, &ap);
|
||||
va_end (ap);
|
||||
}
|
||||
|
||||
void
|
||||
GOACC_declare (int flags_m, size_t mapnum,
|
||||
void **hostaddrs, size_t *sizes, unsigned short *kinds)
|
||||
|
|
|
|||
Loading…
Reference in New Issue