* [PATCH] OpenACC: Separate enter/exit data APIs @ 2020-07-29 14:05 Andrew Stubbs 2020-07-30 11:10 ` Andrew Stubbs 0 siblings, 1 reply; 4+ messages in thread From: Andrew Stubbs @ 2020-07-29 14:05 UTC (permalink / raw) To: gcc-patches [-- Attachment #1: Type: text/plain, Size: 614 bytes --] This patch does not implement anything new, but simply separates OpenACC 'enter data' and 'exit data' into two libgomp API functions. The original API name is kept for backward compatibility, but no longer referenced by the compiler. The previous implementation assumed that it would always be possible to infer which kind of pragma it was dealing with from the context, but there are a few exceptions, and I want to add one more: zero-length arrays. By cleaning this up I will be free to add the new feature without the reference counting getting broken. OK for mainline (and backport to OG10)? Andrew [-- Attachment #2: 200729-separate-enter-exit.patch --] [-- Type: text/x-patch, Size: 14630 bytes --] OpenACC: Separate enter/exit data APIs 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 directives given, but there are some exceptions. In particular, zero-length array copies are indistiguishable, but we still want reference counting to work. gcc/ChangeLog: * gimple-pretty-print.c (dump_gimple_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. * gimple.h (enum gf_mask): Likewise. (is_gimple_omp_oacc): Likewise. * gimplify.c (gimplify_omp_target_update): Likewise. * omp-builtins.def (BUILT_IN_GOACC_ENTER_EXIT_DATA): Delete. (BUILT_IN_GOACC_ENTER_DATA): Add new. (BUILT_IN_GOACC_EXIT_DATA): Add new. * omp-expand.c (expand_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. (build_omp_regions_1): Likewise. (omp_make_gimple_edges): Likewise. * omp-low.c (check_omp_nesting_restrictions): Likewise. (lower_omp_target): Likewise. libgomp/ChangeLog: * libgomp.map: Add GOACC_enter_data and GOACC_exit_data. * libgomp_g.h (GOACC_enter_exit_data): Delete. (GOACC_enter_data): New prototype. (GOACC_exit_data) New prototype.: * oacc-mem.c (GOACC_enter_exit_data): Move most of the content ... (GOACC_enter_exit_data_internal): ... here. (GOACC_enter_data): New function. (GOACC_exit_data) New function.: * oacc-parallel.c (GOACC_declare): Replace GOACC_enter_exit_data with GOACC_enter_data and GOACC_exit_data. diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index e05b770138e..a4dfc493588 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1694,8 +1694,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"; diff --git a/gcc/gimple.h b/gcc/gimple.h index d64c47a7d0d..681000d2ae4 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -180,9 +180,10 @@ 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_ENTER_DATA = 10, GF_OMP_TARGET_KIND_OACC_DECLARE = 11, GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12, + GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 13, GF_OMP_TEAMS_GRID_PHONY = 1 << 0, GF_OMP_TEAMS_HOST = 1 << 1, @@ -6587,7 +6588,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: return true; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 15dfee903ab..6948c1ccdbb 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -12950,8 +12950,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: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index f461d60e52b..ab45eecb752 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -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", diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index ee354b7ef9c..0e21ea90d90 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -8787,7 +8787,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: data_region = false; break; @@ -9052,8 +9053,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; @@ -9251,7 +9255,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; @@ -9822,7 +9827,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_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: /* ..., other than for those stand-alone directives... */ region = NULL; @@ -10077,7 +10083,8 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_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: cur_region = cur_region->outer; break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index da6c275f4a0..df9d108c8a7 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3414,8 +3414,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; @@ -11402,7 +11404,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: data_region = false; break; diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index c808e810702..3965f036c43 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -517,6 +517,8 @@ GOACC_2.0 { global: GOACC_data_end; GOACC_data_start; + GOACC_enter_data; + GOACC_exit_data; GOACC_enter_exit_data; GOACC_parallel; GOACC_update; diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index 59e3697bfd8..51df36bc8db 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -363,8 +363,10 @@ extern void GOACC_wait (int, int, ...); /* oacc-mem.c */ -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, ...); /* oacc-parallel.c */ diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 65757ab2ffc..45162d24786 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1317,56 +1317,21 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_mutex_unlock (&acc_dev->lock); } -void -GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds, int async, - int num_waits, ...) +static void +GOACC_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, + int async, bool data_enter, int num_waits, ...) { int flags = GOACC_FLAGS_UNMARSHAL (flags_m); struct goacc_thread *thr; struct gomp_device_descr *acc_dev; - bool data_enter = false; - size_t i; goacc_lazy_initialize (); thr = goacc_thread (); acc_dev = thr->dev; - /* Determine if this is an "acc enter data". */ - for (i = 0; i < mapnum; ++i) - { - unsigned char kind = kinds[i] & 0xff; - - if (kind == GOMP_MAP_POINTER - || kind == GOMP_MAP_TO_PSET - || kind == GOMP_MAP_STRUCT) - continue; - - if (kind == GOMP_MAP_FORCE_ALLOC - || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_ATTACH - || kind == GOMP_MAP_FORCE_TO - || kind == GOMP_MAP_TO - || kind == GOMP_MAP_ALLOC) - { - data_enter = true; - break; - } - - if (kind == GOMP_MAP_RELEASE - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_DETACH - || kind == GOMP_MAP_FORCE_DETACH - || kind == GOMP_MAP_FROM - || kind == GOMP_MAP_FORCE_FROM) - break; - - gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", - kind); - } - bool profiling_p = GOACC_PROFILING_DISPATCH_P (true); acc_prof_info prof_info; @@ -1458,3 +1423,71 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, thr->api_info = NULL; } } + +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; + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, true, num_waits, 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; + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, false, num_waits, ap); +} + +/* This function is not used. It is provided for backwards compatibility + with older user-binaries only. */ + +void +GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + bool data_enter = false; + + /* Determine if this is an "acc enter data". */ + for (int i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER + || kind == GOMP_MAP_TO_PSET + || kind == GOMP_MAP_STRUCT) + continue; + + if (kind == GOMP_MAP_FORCE_ALLOC + || kind == GOMP_MAP_FORCE_PRESENT + || kind == GOMP_MAP_ATTACH + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) + { + data_enter = true; + break; + } + + if (kind == GOMP_MAP_RELEASE + || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH + || kind == GOMP_MAP_FROM + || kind == GOMP_MAP_FORCE_FROM) + break; + + gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", + kind); + } + + va_list ap; + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, data_enter, num_waits, ap); +} diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index c7e46e35bd6..bca31b51427 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -745,12 +745,15 @@ GOACC_declare (int flags_m, size_t mapnum, switch (kind) { case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TO: - case GOMP_MAP_POINTER: + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); + break; + + case GOMP_MAP_FORCE_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], &kinds[i], GOMP_ASYNC_SYNC, 0); break; @@ -759,19 +762,19 @@ GOACC_declare (int flags_m, size_t mapnum, case GOMP_MAP_ALLOC: if (!acc_is_present (hostaddrs[i], sizes[i])) - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_TO: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FROM: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FORCE_PRESENT: ^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenACC: Separate enter/exit data APIs 2020-07-29 14:05 [PATCH] OpenACC: Separate enter/exit data APIs Andrew Stubbs @ 2020-07-30 11:10 ` Andrew Stubbs 2020-09-25 15:22 ` Andrew Stubbs 0 siblings, 1 reply; 4+ messages in thread From: Andrew Stubbs @ 2020-07-30 11:10 UTC (permalink / raw) To: gcc-patches [-- Attachment #1: Type: text/plain, Size: 806 bytes --] On 29/07/2020 15:05, Andrew Stubbs wrote: > This patch does not implement anything new, but simply separates OpenACC > 'enter data' and 'exit data' into two libgomp API functions. The > original API name is kept for backward compatibility, but no longer > referenced by the compiler. > > The previous implementation assumed that it would always be possible to > infer which kind of pragma it was dealing with from the context, but > there are a few exceptions, and I want to add one more: zero-length arrays. > > By cleaning this up I will be free to add the new feature without the > reference counting getting broken. Here's an updated patch that handles the variadic functions properly. Thanks to Julian for pointing out my misunderstanding. OK for mainline (and backport to OG10)? Andrew [-- Attachment #2: 200730-separate-enter-exit.patch --] [-- Type: text/x-patch, Size: 15122 bytes --] OpenACC: Separate enter/exit data APIs 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 directives given, but there are some exceptions. In particular, zero-length array copies are indistiguishable, but we still want reference counting to work. gcc/ChangeLog: * gimple-pretty-print.c (dump_gimple_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. * gimple.h (enum gf_mask): Likewise. (is_gimple_omp_oacc): Likewise. * gimplify.c (gimplify_omp_target_update): Likewise. * omp-builtins.def (BUILT_IN_GOACC_ENTER_EXIT_DATA): Delete. (BUILT_IN_GOACC_ENTER_DATA): Add new. (BUILT_IN_GOACC_EXIT_DATA): Add new. * omp-expand.c (expand_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. (build_omp_regions_1): Likewise. (omp_make_gimple_edges): Likewise. * omp-low.c (check_omp_nesting_restrictions): Likewise. (lower_omp_target): Likewise. libgomp/ChangeLog: * libgomp.map: Add GOACC_enter_data and GOACC_exit_data. * libgomp_g.h (GOACC_enter_exit_data): Delete. (GOACC_enter_data): New prototype. (GOACC_exit_data) New prototype.: * oacc-mem.c (GOACC_enter_exit_data): Move most of the content ... (GOACC_enter_exit_data_internal): ... here. (GOACC_enter_data): New function. (GOACC_exit_data) New function.: * oacc-parallel.c (GOACC_declare): Replace GOACC_enter_exit_data with GOACC_enter_data and GOACC_exit_data. diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index e05b770138e..a4dfc493588 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1694,8 +1694,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"; diff --git a/gcc/gimple.h b/gcc/gimple.h index d64c47a7d0d..681000d2ae4 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -180,9 +180,10 @@ 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_ENTER_DATA = 10, GF_OMP_TARGET_KIND_OACC_DECLARE = 11, GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12, + GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 13, GF_OMP_TEAMS_GRID_PHONY = 1 << 0, GF_OMP_TEAMS_HOST = 1 << 1, @@ -6587,7 +6588,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: return true; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 15dfee903ab..6948c1ccdbb 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -12950,8 +12950,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: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index f461d60e52b..ab45eecb752 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -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", diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index ee354b7ef9c..0e21ea90d90 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -8787,7 +8787,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: data_region = false; break; @@ -9052,8 +9053,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; @@ -9251,7 +9255,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; @@ -9822,7 +9827,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_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: /* ..., other than for those stand-alone directives... */ region = NULL; @@ -10077,7 +10083,8 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_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: cur_region = cur_region->outer; break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index da6c275f4a0..df9d108c8a7 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3414,8 +3414,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; @@ -11402,7 +11404,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: data_region = false; break; diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index c808e810702..3965f036c43 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -517,6 +517,8 @@ GOACC_2.0 { global: GOACC_data_end; GOACC_data_start; + GOACC_enter_data; + GOACC_exit_data; GOACC_enter_exit_data; GOACC_parallel; GOACC_update; diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index 59e3697bfd8..51df36bc8db 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -363,8 +363,10 @@ extern void GOACC_wait (int, int, ...); /* oacc-mem.c */ -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, ...); /* oacc-parallel.c */ diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 65757ab2ffc..a46c6a02626 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1317,56 +1317,22 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_mutex_unlock (&acc_dev->lock); } -void -GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds, int async, - int num_waits, ...) +static void +GOACC_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, + int async, bool data_enter, int num_waits, + va_list *ap) { int flags = GOACC_FLAGS_UNMARSHAL (flags_m); struct goacc_thread *thr; struct gomp_device_descr *acc_dev; - bool data_enter = false; - size_t i; goacc_lazy_initialize (); thr = goacc_thread (); acc_dev = thr->dev; - /* Determine if this is an "acc enter data". */ - for (i = 0; i < mapnum; ++i) - { - unsigned char kind = kinds[i] & 0xff; - - if (kind == GOMP_MAP_POINTER - || kind == GOMP_MAP_TO_PSET - || kind == GOMP_MAP_STRUCT) - continue; - - if (kind == GOMP_MAP_FORCE_ALLOC - || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_ATTACH - || kind == GOMP_MAP_FORCE_TO - || kind == GOMP_MAP_TO - || kind == GOMP_MAP_ALLOC) - { - data_enter = true; - break; - } - - if (kind == GOMP_MAP_RELEASE - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_DETACH - || kind == GOMP_MAP_FORCE_DETACH - || kind == GOMP_MAP_FROM - || kind == GOMP_MAP_FORCE_FROM) - break; - - gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", - kind); - } - bool profiling_p = GOACC_PROFILING_DISPATCH_P (true); acc_prof_info prof_info; @@ -1430,13 +1396,7 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, } if (num_waits) - { - va_list ap; - - va_start (ap, num_waits); - goacc_wait (async, num_waits, &ap); - va_end (ap); - } + goacc_wait (async, num_waits, ap); goacc_aq aq = get_goacc_asyncqueue (async); @@ -1458,3 +1418,77 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, thr->api_info = NULL; } } + +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, + async, true, 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, + async, false, num_waits, &ap); + va_end (ap); +} + +/* This function is not used. It is provided for backwards compatibility + with older user-binaries only. */ + +void +GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + bool data_enter = false; + + /* Determine if this is an "acc enter data". */ + for (int i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER + || kind == GOMP_MAP_TO_PSET + || kind == GOMP_MAP_STRUCT) + continue; + + if (kind == GOMP_MAP_FORCE_ALLOC + || kind == GOMP_MAP_FORCE_PRESENT + || kind == GOMP_MAP_ATTACH + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) + { + data_enter = true; + break; + } + + if (kind == GOMP_MAP_RELEASE + || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH + || kind == GOMP_MAP_FROM + || kind == GOMP_MAP_FORCE_FROM) + break; + + gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", + kind); + } + + va_list ap; + va_start (ap, num_waits); + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, data_enter, num_waits, &ap); + va_end (ap); +} diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index c7e46e35bd6..bca31b51427 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -745,12 +745,15 @@ GOACC_declare (int flags_m, size_t mapnum, switch (kind) { case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TO: - case GOMP_MAP_POINTER: + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); + break; + + case GOMP_MAP_FORCE_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], &kinds[i], GOMP_ASYNC_SYNC, 0); break; @@ -759,19 +762,19 @@ GOACC_declare (int flags_m, size_t mapnum, case GOMP_MAP_ALLOC: if (!acc_is_present (hostaddrs[i], sizes[i])) - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_TO: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FROM: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FORCE_PRESENT: ^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenACC: Separate enter/exit data APIs 2020-07-30 11:10 ` Andrew Stubbs @ 2020-09-25 15:22 ` Andrew Stubbs 2021-06-10 11:51 ` Thomas Schwinge 0 siblings, 1 reply; 4+ messages in thread From: Andrew Stubbs @ 2020-09-25 15:22 UTC (permalink / raw) To: gcc-patches [-- Attachment #1: Type: text/plain, Size: 814 bytes --] On 30/07/2020 12:10, Andrew Stubbs wrote: > On 29/07/2020 15:05, Andrew Stubbs wrote: >> This patch does not implement anything new, but simply separates >> OpenACC 'enter data' and 'exit data' into two libgomp API functions. >> The original API name is kept for backward compatibility, but no >> longer referenced by the compiler. >> >> The previous implementation assumed that it would always be possible >> to infer which kind of pragma it was dealing with from the context, >> but there are a few exceptions, and I want to add one more: >> zero-length arrays. >> >> By cleaning this up I will be free to add the new feature without the >> reference counting getting broken. This update fixes a new conflict and updates the patterns in a number of testcases that were affected. OK to commit? Andrew [-- Attachment #2: 200925-separate-enter-exit.patch --] [-- Type: text/x-patch, Size: 29024 bytes --] OpenACC: Separate enter/exit data APIs 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 directives given, but there are some exceptions. In particular, zero-length array copies are indistiguishable, but we still want reference counting to work. gcc/ChangeLog: * gimple-pretty-print.c (dump_gimple_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. * gimple.h (enum gf_mask): Likewise. (is_gimple_omp_oacc): Likewise. * gimplify.c (gimplify_omp_target_update): Likewise. * omp-builtins.def (BUILT_IN_GOACC_ENTER_EXIT_DATA): Delete. (BUILT_IN_GOACC_ENTER_DATA): Add new. (BUILT_IN_GOACC_EXIT_DATA): Add new. * omp-expand.c (expand_omp_target): Replace GF_OMP_TARGET_KIND_OACC_ENTER_EXIT_DATA with GF_OMP_TARGET_KIND_OACC_ENTER_DATA and GF_OMP_TARGET_KIND_OACC_EXIT_DATA. (build_omp_regions_1): Likewise. (omp_make_gimple_edges): Likewise. * omp-low.c (check_omp_nesting_restrictions): Likewise. (lower_omp_target): Likewise. gcc/testsuite/ChangeLog: * c-c++-common/goacc-gomp/nesting-fail-1.c: Adjust patterns. * c-c++-common/goacc/finalize-1.c: Adjust patterns. * c-c++-common/goacc/mdc-1.c: Adjust patterns. * c-c++-common/goacc/nesting-fail-1.c: Adjust patterns. * c-c++-common/goacc/struct-enter-exit-data-1.c: Adjust patterns. libgomp/ChangeLog: * libgomp.map: Add GOACC_enter_data and GOACC_exit_data. * libgomp_g.h (GOACC_enter_exit_data): Delete. (GOACC_enter_data): New prototype. (GOACC_exit_data) New prototype.: * oacc-mem.c (GOACC_enter_exit_data): Move most of the content ... (GOACC_enter_exit_data_internal): ... here. (GOACC_enter_data): New function. (GOACC_exit_data) New function.: * oacc-parallel.c (GOACC_declare): Replace GOACC_enter_exit_data with GOACC_enter_data and GOACC_exit_data. * testsuite/libgomp.oacc-c-c++-common/lib-26.c: Delete file. * testsuite/libgomp.oacc-c-c++-common/lib-36.c: Delete file. * testsuite/libgomp.oacc-c-c++-common/lib-40.c: Delete file. diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index a01bf901657..26978ec1ab5 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -1691,8 +1691,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"; diff --git a/gcc/gimple.h b/gcc/gimple.h index 6cc7e66059d..3f17b1c0739 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -171,9 +171,10 @@ 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_ENTER_DATA = 10, GF_OMP_TARGET_KIND_OACC_DECLARE = 11, GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12, + GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 13, GF_OMP_TEAMS_HOST = 1 << 0, /* True on an GIMPLE_OMP_RETURN statement if the return does not require @@ -6482,7 +6483,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: return true; diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 2dea03cce3d..8fcba8b5b18 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -12976,8 +12976,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: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index f461d60e52b..ab45eecb752 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -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", diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 8f1286e3176..70880b78ffc 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -8917,7 +8917,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: data_region = false; break; @@ -9182,8 +9183,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; @@ -9381,7 +9385,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; @@ -9654,7 +9659,8 @@ build_omp_regions_1 (basic_block bb, struct omp_region *parent, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_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: /* ..., other than for those stand-alone directives... */ region = NULL; @@ -9908,7 +9914,8 @@ omp_make_gimple_edges (basic_block bb, struct omp_region **region, case GF_OMP_TARGET_KIND_ENTER_DATA: case GF_OMP_TARGET_KIND_EXIT_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: cur_region = cur_region->outer; break; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 3d2a9d77c1c..12346689f36 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -3387,8 +3387,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; @@ -11344,7 +11346,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: data_region = false; break; diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 1a3324200e2..ddbd247342f 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -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) ; diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c index 3d64b2e7cb3..54bf1b76a1b 100644 --- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -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" } } */ } diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c index 337c1f7cc77..d961c77e3bb 100644 --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -43,14 +43,14 @@ 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.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_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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 93a911120d4..5cfb327f4b6 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -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" } */ } } diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c index df405e448b2..9e5d3f2c9d2 100644 --- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -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 } } */ } diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index c808e810702..3965f036c43 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -517,6 +517,8 @@ GOACC_2.0 { global: GOACC_data_end; GOACC_data_start; + GOACC_enter_data; + GOACC_exit_data; GOACC_enter_exit_data; GOACC_parallel; GOACC_update; diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index 59e3697bfd8..51df36bc8db 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -363,8 +363,10 @@ extern void GOACC_wait (int, int, ...); /* oacc-mem.c */ -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, ...); /* oacc-parallel.c */ diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 65757ab2ffc..a46c6a02626 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1317,56 +1317,22 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_mutex_unlock (&acc_dev->lock); } -void -GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds, int async, - int num_waits, ...) +static void +GOACC_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, + int async, bool data_enter, int num_waits, + va_list *ap) { int flags = GOACC_FLAGS_UNMARSHAL (flags_m); struct goacc_thread *thr; struct gomp_device_descr *acc_dev; - bool data_enter = false; - size_t i; goacc_lazy_initialize (); thr = goacc_thread (); acc_dev = thr->dev; - /* Determine if this is an "acc enter data". */ - for (i = 0; i < mapnum; ++i) - { - unsigned char kind = kinds[i] & 0xff; - - if (kind == GOMP_MAP_POINTER - || kind == GOMP_MAP_TO_PSET - || kind == GOMP_MAP_STRUCT) - continue; - - if (kind == GOMP_MAP_FORCE_ALLOC - || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_ATTACH - || kind == GOMP_MAP_FORCE_TO - || kind == GOMP_MAP_TO - || kind == GOMP_MAP_ALLOC) - { - data_enter = true; - break; - } - - if (kind == GOMP_MAP_RELEASE - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_DETACH - || kind == GOMP_MAP_FORCE_DETACH - || kind == GOMP_MAP_FROM - || kind == GOMP_MAP_FORCE_FROM) - break; - - gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", - kind); - } - bool profiling_p = GOACC_PROFILING_DISPATCH_P (true); acc_prof_info prof_info; @@ -1430,13 +1396,7 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, } if (num_waits) - { - va_list ap; - - va_start (ap, num_waits); - goacc_wait (async, num_waits, &ap); - va_end (ap); - } + goacc_wait (async, num_waits, ap); goacc_aq aq = get_goacc_asyncqueue (async); @@ -1458,3 +1418,77 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, thr->api_info = NULL; } } + +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, + async, true, 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, + async, false, num_waits, &ap); + va_end (ap); +} + +/* This function is not used. It is provided for backwards compatibility + with older user-binaries only. */ + +void +GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + bool data_enter = false; + + /* Determine if this is an "acc enter data". */ + for (int i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER + || kind == GOMP_MAP_TO_PSET + || kind == GOMP_MAP_STRUCT) + continue; + + if (kind == GOMP_MAP_FORCE_ALLOC + || kind == GOMP_MAP_FORCE_PRESENT + || kind == GOMP_MAP_ATTACH + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) + { + data_enter = true; + break; + } + + if (kind == GOMP_MAP_RELEASE + || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH + || kind == GOMP_MAP_FROM + || kind == GOMP_MAP_FORCE_FROM) + break; + + gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", + kind); + } + + va_list ap; + va_start (ap, num_waits); + GOACC_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + async, data_enter, num_waits, &ap); + va_end (ap); +} diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c index c7e46e35bd6..bca31b51427 100644 --- a/libgomp/oacc-parallel.c +++ b/libgomp/oacc-parallel.c @@ -745,12 +745,15 @@ GOACC_declare (int flags_m, size_t mapnum, switch (kind) { case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_FROM: case GOMP_MAP_FORCE_TO: - case GOMP_MAP_POINTER: + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); + break; + + case GOMP_MAP_FORCE_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], &kinds[i], GOMP_ASYNC_SYNC, 0); break; @@ -759,19 +762,19 @@ GOACC_declare (int flags_m, size_t mapnum, case GOMP_MAP_ALLOC: if (!acc_is_present (hostaddrs[i], sizes[i])) - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_TO: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FROM: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], GOMP_ASYNC_SYNC, 0); break; case GOMP_MAP_FORCE_PRESENT: diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c deleted file mode 100644 index 8e1a911abd2..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-do run } */ - -#include <stdio.h> -#include <stdlib.h> -#include <openacc.h> - -int -main (int argc, char **argv) -{ - const int N = 256; - unsigned char *h; - void *d; - - h = (unsigned char *) malloc (N); - - fprintf (stderr, "CheCKpOInT\n"); - d = acc_create (h, 0); - if (!d) - abort (); - - acc_delete (h, N); - - free (h); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+0\\\] is a bad range" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c deleted file mode 100644 index 8ff61c76491..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c +++ /dev/null @@ -1,30 +0,0 @@ -/* { dg-do run } */ - -#include <stdio.h> -#include <stdlib.h> -#include <openacc.h> - -int -main (int argc, char **argv) -{ - const int N = 256; - unsigned char *h; - void *d; - - h = (unsigned char *) malloc (N); - - fprintf (stderr, "CheCKpOInT\n"); - d = acc_present_or_create (h, 0); - if (!d) - abort (); - - acc_delete (h, N); - - free (h); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+0\\\] is a bad range" } */ -/* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c deleted file mode 100644 index f4ab6d8fc3e..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c +++ /dev/null @@ -1,46 +0,0 @@ -/* { dg-do run } */ - -#include <stdio.h> -#include <stdlib.h> -#include <string.h> -#include <unistd.h> -#include <openacc.h> - -int -main (int argc, char **argv) -{ - const int N = 256; - int i; - unsigned char *h; - void *d; - - h = (unsigned char *) malloc (N); - - for (i = 0; i < N; i++) - { - h[i] = i; - } - - fprintf (stderr, "CheCKpOInT\n"); - d = acc_present_or_copyin (h, 0); - if (!d) - abort (); - - memset (&h[0], 0, N); - - acc_copyout (h, N); - - for (i = 0; i < N; i++) - { - if (h[i] != i) - abort (); - } - - free (h); - - return 0; -} - -/* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "\\\[\[0-9a-fA-FxX\]+,\\\+0\\\] is a bad range" } */ -/* { dg-shouldfail "" } */ ^ permalink raw reply [flat|nested] 4+ messages in thread
* Re: [PATCH] OpenACC: Separate enter/exit data APIs 2020-09-25 15:22 ` Andrew Stubbs @ 2021-06-10 11:51 ` Thomas Schwinge 0 siblings, 0 replies; 4+ messages in thread From: Thomas Schwinge @ 2021-06-10 11:51 UTC (permalink / raw) To: Andrew Stubbs, gcc-patches [-- Attachment #1: Type: text/plain, Size: 5691 bytes --] Hi! On 2020-09-25T16:22:47+0100, Andrew Stubbs <ams@codesourcery.com> wrote: > On 30/07/2020 12:10, Andrew Stubbs wrote: >> On 29/07/2020 15:05, Andrew Stubbs wrote: >>> This patch does not implement anything new, but simply separates >>> OpenACC 'enter data' and 'exit data' into two libgomp API functions. >>> The original API name is kept for backward compatibility, but no >>> longer referenced by the compiler. (ABI, not API.) >>> The previous implementation assumed that it would always be possible >>> to infer which kind of pragma it was dealing with from the context, >>> but there are a few exceptions, and I want to add one more: >>> zero-length arrays. >>> >>> By cleaning this up I will be free to add the new feature without the >>> reference counting getting broken. ACK, that's a good cleanup already for its own sake. > This update [...] updates the patterns in a number of > testcases that were affected. ... just that you missed all the Fortran testcases, had me do the "dirty Fortran work", tsk... ;-P A few comments, hope that's useful: > --- a/gcc/gimple.h > +++ b/gcc/gimple.h > @@ -171,9 +171,10 @@ 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_ENTER_DATA = 10, > GF_OMP_TARGET_KIND_OACC_DECLARE = 11, > GF_OMP_TARGET_KIND_OACC_HOST_DATA = 12, > + GF_OMP_TARGET_KIND_OACC_EXIT_DATA = 13, > GF_OMP_TEAMS_HOST = 1 << 0, We shall re-number, so that 'GF_OMP_TARGET_KIND_OACC_ENTER_DATA' and 'GF_OMP_TARGET_KIND_OACC_EXIT_DATA' stay in proximity. > --- a/libgomp/libgomp.map > +++ b/libgomp/libgomp.map > @@ -517,6 +517,8 @@ GOACC_2.0 { > global: > GOACC_data_end; > GOACC_data_start; > + GOACC_enter_data; > + GOACC_exit_data; > GOACC_enter_exit_data; > GOACC_parallel; > GOACC_update; As we've been told, new symbols need a new symbol version if 'GOACC_2.0' (here) has already been part of a GCC release. > --- a/libgomp/libgomp_g.h > +++ b/libgomp/libgomp_g.h > @@ -363,8 +363,10 @@ extern void GOACC_wait (int, int, ...); > > /* oacc-mem.c */ > > -extern void GOACC_enter_exit_data (int, size_t, void **, size_t *, > - unsigned short *, int, int, ...); Don't remove: "This file contains prototypes of functions in the external ABI" -- and 'GOACC_enter_exit_data' remains part of that. > +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, ...); > --- a/libgomp/oacc-mem.c > +++ b/libgomp/oacc-mem.c > +static void > +GOACC_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs, > + size_t *sizes, unsigned short *kinds, > + int async, bool data_enter, int num_waits, > + va_list *ap) The order 'async', 'data_enter', 'num_waits' seemed a bit awkward to me, so I changed that to 'data_enter', 'async', 'num_waits'. > --- a/libgomp/oacc-parallel.c > +++ b/libgomp/oacc-parallel.c > @@ -745,12 +745,15 @@ GOACC_declare (int flags_m, size_t mapnum, > switch (kind) > { > case GOMP_MAP_FORCE_ALLOC: > - case GOMP_MAP_FORCE_FROM: > case GOMP_MAP_FORCE_TO: > - case GOMP_MAP_POINTER: I've separately pushed "Clean up 'GOMP_MAP_POINTER' handling in 'libgomp/oacc-parallel.c:GOACC_declare'". > + GOACC_enter_data (flags_m, 1, &hostaddrs[i], &sizes[i], > + &kinds[i], GOMP_ASYNC_SYNC, 0); > + break; > + > + case GOMP_MAP_FORCE_FROM: > case GOMP_MAP_RELEASE: > case GOMP_MAP_DELETE: > - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], > + GOACC_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], > &kinds[i], GOMP_ASYNC_SYNC, 0); > break; > [...] And, after separately pushed "Move 'libgomp/oacc-parallel.c:GOACC_declare' into 'libgomp/oacc-mem.c'", we can then directly call the internal 'goacc_enter_exit_data_internal' here, instead of the external 'GOACC_enter_data'/'GOACC_exit_data', which is benefitial for certain reasons. > --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-26.c > +++ /dev/null > --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-36.c > +++ /dev/null > --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-40.c > +++ /dev/null No reason to remove these as part of this patch. Fixed these items, plus some minor additional clean-up. I've then first pushed "Extract 'goacc_enter_exit_data_internal' from 'libgomp/oacc-mem.c:GOACC_enter_exit_data'" to master branch in commit 7999363961dc6feeb0976cc6d85ea91a120d0e1d, see attached. Doing that as a separate step makes sure that we don't break anything for 'GOACC_enter_exit_data'. Then I've pushed "OpenACC: Separate enter/exit data ABIs" to master branch in commit 7aefef31365b9c3d32a0edb6ea0d3b8864d7e91a, see attached. This now doesn't anymore touch the old 'GOACC_enter_exit_data' ABI, so cannot possibly break anything there. Grüße Thomas ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstrasse 201, 80634 München Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Frank Thürauf [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #2: 0001-Extract-goacc_enter_exit_data_internal-from-libgomp-.patch --] [-- Type: text/x-diff, Size: 5593 bytes --] From 7999363961dc6feeb0976cc6d85ea91a120d0e1d Mon Sep 17 00:00:00 2001 From: Thomas Schwinge <thomas@codesourcery.com> Date: Tue, 8 Jun 2021 17:39:25 +0200 Subject: [PATCH 1/2] Extract 'goacc_enter_exit_data_internal' from 'libgomp/oacc-mem.c:GOACC_enter_exit_data' libgomp/ * oacc-mem.c (goacc_enter_exit_data_internal): New function, extracted from... (GOACC_enter_exit_data): ... here. (GOACC_declare): Use it. Co-Authored-By: Andrew Stubbs <ams@codesourcery.com> --- libgomp/oacc-mem.c | 130 ++++++++++++++++++++++----------------------- 1 file changed, 65 insertions(+), 65 deletions(-) diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 056600aca52..f6173b91fdd 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -1320,56 +1320,22 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, size_t mapnum, gomp_mutex_unlock (&acc_dev->lock); } -void -GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, - size_t *sizes, unsigned short *kinds, int async, - int num_waits, ...) +static void +goacc_enter_exit_data_internal (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, + bool data_enter, int async, int num_waits, + va_list *ap) { int flags = GOACC_FLAGS_UNMARSHAL (flags_m); struct goacc_thread *thr; struct gomp_device_descr *acc_dev; - bool data_enter = false; - size_t i; goacc_lazy_initialize (); thr = goacc_thread (); acc_dev = thr->dev; - /* Determine if this is an "acc enter data". */ - for (i = 0; i < mapnum; ++i) - { - unsigned char kind = kinds[i] & 0xff; - - if (kind == GOMP_MAP_POINTER - || kind == GOMP_MAP_TO_PSET - || kind == GOMP_MAP_STRUCT) - continue; - - if (kind == GOMP_MAP_FORCE_ALLOC - || kind == GOMP_MAP_FORCE_PRESENT - || kind == GOMP_MAP_ATTACH - || kind == GOMP_MAP_FORCE_TO - || kind == GOMP_MAP_TO - || kind == GOMP_MAP_ALLOC) - { - data_enter = true; - break; - } - - if (kind == GOMP_MAP_RELEASE - || kind == GOMP_MAP_DELETE - || kind == GOMP_MAP_DETACH - || kind == GOMP_MAP_FORCE_DETACH - || kind == GOMP_MAP_FROM - || kind == GOMP_MAP_FORCE_FROM) - break; - - gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", - kind); - } - bool profiling_p = GOACC_PROFILING_DISPATCH_P (true); acc_prof_info prof_info; @@ -1433,13 +1399,7 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, } if (num_waits) - { - va_list ap; - - va_start (ap, num_waits); - goacc_wait (async, num_waits, &ap); - va_end (ap); - } + goacc_wait (async, num_waits, ap); goacc_aq aq = get_goacc_asyncqueue (async); @@ -1462,6 +1422,52 @@ GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, } } +void +GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs, + size_t *sizes, unsigned short *kinds, int async, + int num_waits, ...) +{ + /* Determine if this is an OpenACC "enter data". */ + bool data_enter = false; + for (size_t i = 0; i < mapnum; ++i) + { + unsigned char kind = kinds[i] & 0xff; + + if (kind == GOMP_MAP_POINTER + || kind == GOMP_MAP_TO_PSET + || kind == GOMP_MAP_STRUCT) + continue; + + if (kind == GOMP_MAP_FORCE_ALLOC + || kind == GOMP_MAP_FORCE_PRESENT + || kind == GOMP_MAP_ATTACH + || kind == GOMP_MAP_FORCE_TO + || kind == GOMP_MAP_TO + || kind == GOMP_MAP_ALLOC) + { + data_enter = true; + break; + } + + if (kind == GOMP_MAP_RELEASE + || kind == GOMP_MAP_DELETE + || kind == GOMP_MAP_DETACH + || kind == GOMP_MAP_FORCE_DETACH + || kind == GOMP_MAP_FROM + || kind == GOMP_MAP_FORCE_FROM) + break; + + gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x", + kind); + } + + va_list ap; + va_start (ap, num_waits); + goacc_enter_exit_data_internal (flags_m, mapnum, hostaddrs, sizes, kinds, + data_enter, async, num_waits, &ap); + va_end (ap); +} + void GOACC_declare (int flags_m, size_t mapnum, void **hostaddrs, size_t *sizes, unsigned short *kinds) @@ -1475,34 +1481,28 @@ GOACC_declare (int flags_m, size_t mapnum, switch (kind) { + case GOMP_MAP_ALLOC: + if (acc_is_present (hostaddrs[i], sizes[i])) + continue; + /* FALLTHRU */ case GOMP_MAP_FORCE_ALLOC: - case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_TO: case GOMP_MAP_FORCE_TO: + goacc_enter_exit_data_internal (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], true, GOMP_ASYNC_SYNC, 0, NULL); + break; + + case GOMP_MAP_FROM: + case GOMP_MAP_FORCE_FROM: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); + goacc_enter_exit_data_internal (flags_m, 1, &hostaddrs[i], &sizes[i], + &kinds[i], false, GOMP_ASYNC_SYNC, 0, NULL); break; case GOMP_MAP_FORCE_DEVICEPTR: break; - case GOMP_MAP_ALLOC: - if (!acc_is_present (hostaddrs[i], sizes[i])) - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); - break; - - case GOMP_MAP_TO: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); - break; - - case GOMP_MAP_FROM: - GOACC_enter_exit_data (flags_m, 1, &hostaddrs[i], &sizes[i], - &kinds[i], GOMP_ASYNC_SYNC, 0); - break; - case GOMP_MAP_FORCE_PRESENT: if (!acc_is_present (hostaddrs[i], sizes[i])) gomp_fatal ("[%p,%ld] is not mapped", hostaddrs[i], -- 2.30.2 [-- Warning: decoded text below may be mangled, UTF-8 assumed --] [-- Attachment #3: 0002-OpenACC-Separate-enter-exit-data-ABIs.patch --] [-- Type: text/x-diff, Size: 31983 bytes --] From 7aefef31365b9c3d32a0edb6ea0d3b8864d7e91a Mon Sep 17 00:00:00 2001 From: Andrew Stubbs <ams@codesourcery.com> Date: Fri, 25 Sep 2020 16:22:47 +0100 Subject: [PATCH 2/2] 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> --- gcc/gimple-pretty-print.c | 7 +++-- gcc/gimple.h | 18 +++++++------ gcc/gimplify.c | 5 +++- gcc/omp-builtins.def | 5 +++- gcc/omp-expand.c | 19 +++++++++----- gcc/omp-low.c | 9 ++++--- .../c-c++-common/goacc-gomp/nesting-fail-1.c | 4 +-- gcc/testsuite/c-c++-common/goacc/finalize-1.c | 12 ++++----- gcc/testsuite/c-c++-common/goacc/mdc-1.c | 18 ++++++------- .../c-c++-common/goacc/nesting-fail-1.c | 8 +++--- .../goacc/struct-enter-exit-data-1.c | 4 +-- .../gfortran.dg/goacc/attach-descriptor.f90 | 6 ++--- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 12 ++++----- .../gfortran.dg/goacc/mapping-tests-3.f90 | 2 +- libgomp/libgomp.map | 6 +++++ libgomp/libgomp_g.h | 4 +++ libgomp/oacc-mem.c | 26 +++++++++++++++++++ 17 files changed, 111 insertions(+), 54 deletions(-) diff --git a/gcc/gimple-pretty-print.c b/gcc/gimple-pretty-print.c index c9c0a66ee30..8be40416dd2 100644 --- a/gcc/gimple-pretty-print.c +++ b/gcc/gimple-pretty-print.c @@ -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"; diff --git a/gcc/gimple.h b/gcc/gimple.h index 91b92b4a4d1..e7dc2a45a13 100644 --- a/gcc/gimple.h +++ b/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: diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 39f5b973d18..c96d611b115 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -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: diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def index cfbf1e67b8e..97964f866ec 100644 --- a/gcc/omp-builtins.def +++ b/gcc/omp-builtins.def @@ -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", diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c index 0f843bad79a..f8b1558b1d1 100644 --- a/gcc/omp-expand.c +++ b/gcc/omp-expand.c @@ -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; diff --git a/gcc/omp-low.c b/gcc/omp-low.c index 2d5cdf671eb..2cc2a186080 100644 --- a/gcc/omp-low.c +++ b/gcc/omp-low.c @@ -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: diff --git a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c index 1a3324200e2..ddbd247342f 100644 --- a/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc-gomp/nesting-fail-1.c @@ -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) ; diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c index 3d64b2e7cb3..54bf1b76a1b 100644 --- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c +++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c @@ -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" } } */ } diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c index c9ab7c24074..c2b8dc6c880 100644 --- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c +++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c @@ -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" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c index 93a911120d4..5cfb327f4b6 100644 --- a/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c +++ b/gcc/testsuite/c-c++-common/goacc/nesting-fail-1.c @@ -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" } */ } } diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c index df405e448b2..9e5d3f2c9d2 100644 --- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c +++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c @@ -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 } } */ } diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 index 373bdcb2114..8c2ee4a5cca 100644 --- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 @@ -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 diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index a7788580819..b706b385aeb 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -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 diff --git a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 index 890ca781967..662104f8aaa 100644 --- a/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/mapping-tests-3.f90 @@ -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 diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 4ad190a52af..8ea27b5565f 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -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; diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h index b66b6978202..f890a204fb8 100644 --- a/libgomp/libgomp_g.h +++ b/libgomp/libgomp_g.h @@ -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 */ diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index f6173b91fdd..123fe154089 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.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) -- 2.30.2 ^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2021-06-10 11:51 UTC | newest] Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed) -- links below jump to the message on this page -- 2020-07-29 14:05 [PATCH] OpenACC: Separate enter/exit data APIs Andrew Stubbs 2020-07-30 11:10 ` Andrew Stubbs 2020-09-25 15:22 ` Andrew Stubbs 2021-06-10 11:51 ` Thomas Schwinge
This is a public inbox, see mirroring instructions for how to clone and mirror all data and code used for this inbox; as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).