diff options
| author | Ian Lance Taylor <iant@golang.org> | 2023-06-21 11:04:04 -0700 |
|---|---|---|
| committer | Ian Lance Taylor <iant@golang.org> | 2023-06-21 11:04:04 -0700 |
| commit | 97e31a0a2a2d2273687fcdb4e5416aab1a2186e1 (patch) | |
| tree | d5c1cae4de436a0fe54a5f0a2a197d309f3d654c /libgomp | |
| parent | 6612f4f8cb9b0d5af18ec69ad04e56debc3e6ced (diff) | |
| parent | 577223aebc7acdd31e62b33c1682fe54a622ae27 (diff) | |
| download | gcc-97e31a0a2a2d2273687fcdb4e5416aab1a2186e1.zip gcc-97e31a0a2a2d2273687fcdb4e5416aab1a2186e1.tar.gz gcc-97e31a0a2a2d2273687fcdb4e5416aab1a2186e1.tar.bz2 | |
Merge from trunk revision 577223aebc7acdd31e62b33c1682fe54a622ae27.
Diffstat (limited to 'libgomp')
81 files changed, 4983 insertions, 644 deletions
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index c50e591..002d802 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,385 @@ +2023-06-19 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c/target-51.c: Fix DejaGnu directive syntax + error. + +2023-06-19 Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.c/target-51.c: Accept more error msg variants + as expected dg-output. + +2023-06-19 Tobias Burnus <tobias@codesourcery.com> + + PR middle-end/110270 + * target.c (gomp_map_vars_internal): Copy host value instead of NULL + for GOMP_MAP_ZERO_LEN_ARRAY_SECTION if not mapped. + * libgomp.texi (OpenMP 5.2 Impl.): Mark as 'Y'. + * testsuite/libgomp.c/target-19.c: Update expected value. + * testsuite/libgomp.c++/target-18.C: Likewise. + * testsuite/libgomp.c++/target-19.C: Likewise. + * testsuite/libgomp.c-c++-common/requires-unified-addr-2.c: New test. + * testsuite/libgomp.c-c++-common/target-implicit-map-3.c: New test. + * testsuite/libgomp.c-c++-common/target-implicit-map-4.c: New test. + +2023-06-16 Tobias Burnus <tobias@codesourcery.com> + + * target.c (resolve_device): Call gomp_get_num_devices early to ensure + gomp_init_targets_once was called before using default-device-var. + * testsuite/libgomp.c/target-55.c: New test. + * testsuite/libgomp.c/target-55a.c: New test. + +2023-06-15 Tobias Burnus <tobias@codesourcery.com> + + * env.c (gomp_def_allocator_envvar): New var. + (parse_allocator): Handle OpenMP 5.1 syntax. + (cleanup_env): New. + (omp_display_env): Output gomp_def_allocator_envvar + for an allocator with traits. + * libgomp.texi (OMP_ALLOCATOR, OMP_AFFINITY_FORMAT, + OMP_DISPLAY_AFFINITY): New. + * testsuite/libgomp.c/allocator-1.c: New test. + * testsuite/libgomp.c/allocator-2.c: New test. + * testsuite/libgomp.c/allocator-3.c: New test. + * testsuite/libgomp.c/allocator-4.c: New test. + * testsuite/libgomp.c/allocator-5.c: New test. + * testsuite/libgomp.c/allocator-6.c: New test. + +2023-06-14 Thomas Schwinge <thomas@codesourcery.com> + + * target.c (resolve_device): Align a + 'OMP_TARGET_OFFLOAD=mandatory' diagnostic with others. + * testsuite/libgomp.c/target-51.c: Adjust. + +2023-06-14 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.fortran/fortran.exp (lang_link_flags): Don't + set. + * testsuite/libgomp.oacc-fortran/fortran.exp (lang_link_flags): + Likewise. + * testsuite/libgomp.c/simd-math-1.c: Remove + '-foffload-options=-lm'. + * testsuite/libgomp.fortran/fortran-torture_execute_math.f90: + Likewise. + * testsuite/libgomp.oacc-fortran/fortran-torture_execute_math.f90: + Likewise. + +2023-06-14 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.fortran/fortran-torture_execute_math.f90: New. + * testsuite/libgomp.oacc-fortran/fortran-torture_execute_math.f90: + Likewise. + +2023-06-14 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c/target-51.c: Fix typo. + +2023-06-14 Tobias Burnus <tobias@codesourcery.com> + + * env.c (gomp_default_icv_values): Init default_device_var to + an nonconforming value - INT_MIN. + (initialize_env): After env-var parsing, set default_device_var to + device 0 unless OMP_TARGET_OFFLOAD=mandatory. + (omp_display_env): If default_device_var is INT_MIN, call + gomp_init_targets_once. + * icv-device.c (omp_get_default_device): Likewise. + * libgomp.texi (OMP_DEFAULT_DEVICE): Update init description. + (OpenMP 5.2 Impl. Status): Mark OMP_TARGET_OFFLOAD=mandatory as 'Y'. + * target.c (resolve_device): Improve error message device-num < 0 + with 'mandatory' and no no-host devices available. + (gomp_target_init): Set default-device-var if INT_MIN. + * testsuite/libgomp.c/target-48.c: New test. + * testsuite/libgomp.c/target-49.c: New test. + * testsuite/libgomp.c/target-50.c: New test. + * testsuite/libgomp.c/target-50a.c: New test. + * testsuite/libgomp.c/target-51.c: New test. + * testsuite/libgomp.c/target-52.c: New test. + * testsuite/libgomp.c/target-53.c: New test. + * testsuite/libgomp.c/target-54.c: New test. + +2023-06-13 Tobias Burnus <tobias@codesourcery.com> + + PR libgomp/109837 + * testsuite/libgomp.c-c++-common/requires-unified-addr-1.c: New test. + * testsuite/libgomp.fortran/requires-unified-addr-1.f90: New test. + +2023-06-12 Tobias Burnus <tobias@codesourcery.com> + + * target.c (gomp_to_device_kind_p, gomp_map_vars_internal): Replace + GOMP_MAP_PRESENT_{FROM,TO,TOFROM,ACLLOC} by GOMP_MAP_FORCE_PRESENT. + (gomp_map_vars_internal, gomp_update): Likewise; unify and improve + error message. + * testsuite/libgomp.c-c++-common/target-present-2.c: Update for + changed error message. + * testsuite/libgomp.fortran/target-present-1.f90: Likewise. + * testsuite/libgomp.fortran/target-present-2.f90: Likewise. + * testsuite/libgomp.oacc-c-c++-common/present-1.c: Likewise. + * testsuite/libgomp.c-c++-common/target-present-1.c: Likewise and + extend testcase to check that data is copied when needed. + * testsuite/libgomp.c-c++-common/target-present-3.c: Likewise. + * testsuite/libgomp.fortran/target-present-3.f90: Likewise. + +2023-06-07 Thomas Schwinge <thomas@codesourcery.com> + Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.c-c++-common/target-present-1.c: Run code + also for non-offload_device targets; check that it runs + successfully for those and for all until a checkpoint for all + * testsuite/libgomp.c-c++-common/target-present-2.c: Likewise. + * testsuite/libgomp.c-c++-common/target-present-3.c: Likewise. + * testsuite/libgomp.fortran/target-present-1.f90: Likewise. + * testsuite/libgomp.fortran/target-present-3.f90: Likewise. + * testsuite/libgomp.fortran/target-present-2.f90: Likewise; + add missing vars to map clause. + +2023-06-06 Tobias Burnus <tobias@codesourcery.com> + + * plugin/plugin-gcn.c (GOMP_OFFLOAD_get_num_devices): Regard + unified_address requirement as supported. + * libgomp.texi (OpenMP 5.0, AMD Radeon, nvptx): Remove + 'unified_address' from the not-supported requirements. + +2023-06-06 Kwok Cheung Yeung <kcy@codesourcery.com> + Tobias Burnus <tobias@codesourcery.com> + + * libgomp.texi (OpenMP 5.1 Impl. status): Set 'present' support for + defaultmap to 'Y', add 'Y' entry for 'present' on to/from/map clauses. + * target.c (gomp_to_device_kind_p): Add map kinds with 'present' + modifier. + (gomp_map_vars_existing): Use new GOMP_MAP_FORCE_P macro. + (gomp_map_vars_internal, gomp_update, gomp_target_rev): + Emit runtime error if memory region not present. + * testsuite/libgomp.c-c++-common/target-present-1.c: New test. + * testsuite/libgomp.c-c++-common/target-present-2.c: New test. + * testsuite/libgomp.c-c++-common/target-present-3.c: New test. + * testsuite/libgomp.fortran/target-present-1.f90: New test. + * testsuite/libgomp.fortran/target-present-2.f90: New test. + * testsuite/libgomp.fortran/target-present-3.f90: New test. + +2023-06-02 Thomas Schwinge <thomas@codesourcery.com> + + PR testsuite/66005 + * testsuite/lib/libgomp.exp: 'flock' through stdout. + * testsuite/flock: New. + * configure.ac (FLOCK): Point to that if no 'flock' available, but + 'perl' is. + * configure: Regenerate. + +2023-06-02 Thomas Schwinge <thomas@codesourcery.com> + + * configure.ac (PERL): Remove. + * configure: Regenerate. + * Makefile.in: Likewise. + * testsuite/Makefile.in: Likewise. + +2023-06-01 Tobias Burnus <tobias@codesourcery.com> + + * libgomp.texi (OpenMP 5.2): Mark pure-directive handling as 'Y'. + +2023-05-26 Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.fortran/allocate-4.f90: Update dg-error. + +2023-05-21 Tobias Burnus <tobias@codesourcery.com> + + PR libgomp/109875 + * config/gcn/target.c (GOMP_teams4): Honor nteams-var ICV. + * config/nvptx/target.c (GOMP_teams4): Likewise. + * testsuite/libgomp.c-c++-common/teams-nteams-icv-1.c: New test. + * testsuite/libgomp.c-c++-common/teams-nteams-icv-2.c: New test. + * testsuite/libgomp.c-c++-common/teams-nteams-icv-3.c: New test. + * testsuite/libgomp.c-c++-common/teams-nteams-icv-4.c: New test. + +2023-05-19 Jakub Jelinek <jakub@redhat.com> + + PR libgomp/109904 + * configure.ac (link_gomp): Include also $DL_LIBS. + * configure: Regenerated. + +2023-05-17 Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.fortran/target-enter-data-3.f90: Uncomment + 'target exit data'. + * testsuite/libgomp.fortran/target-enter-data-4.f90: New test. + * testsuite/libgomp.fortran/target-enter-data-5.f90: New test. + * testsuite/libgomp.fortran/target-enter-data-6.f90: New test. + * testsuite/libgomp.fortran/target-enter-data-7.f90: New test. + +2023-05-15 Thomas Schwinge <thomas@codesourcery.com> + + PR testsuite/66005 + * configure.ac: Look for 'flock'. + * testsuite/Makefile.am (gcc_test_parallel_slots): Enable parallel testing. + * testsuite/config/default.exp: Don't 'load_lib "standard.exp"' here... + * testsuite/lib/libgomp.exp: ... but here, instead. + (libgomp_load): Override for parallel testing. + * testsuite/libgomp-site-extra.exp.in (FLOCK): Set. + * configure: Regenerate. + * Makefile.in: Regenerate. + * testsuite/Makefile.in: Regenerate. + +2023-05-15 Rainer Orth <ro@CeBiTec.Uni-Bielefeld.DE> + Thomas Schwinge <thomas@codesourcery.com> + + PR testsuite/66005 + * testsuite/Makefile.am (PWD_COMMAND): New variable. + (%/site.exp): New target. + (check_p_numbers0, check_p_numbers1, check_p_numbers2) + (check_p_numbers3, check_p_numbers4, check_p_numbers5) + (check_p_numbers6, check_p_numbers, gcc_test_parallel_slots) + (check_p_subdirs) + (check_DEJAGNU_libgomp_targets): New variables. + ($(check_DEJAGNU_libgomp_targets)): New target. + ($(check_DEJAGNU_libgomp_targets)): New dependency. + (check-DEJAGNU $(check_DEJAGNU_libgomp_targets)): New targets. + * testsuite/Makefile.in: Regenerate. + * testsuite/lib/libgomp.exp: For parallel testing, + 'load_file ../libgomp-test-support.exp'. + +2023-05-15 Thomas Schwinge <thomas@codesourcery.com> + + PR testsuite/91884 + * configure.ac: 'AC_SUBST(CXX)'. + * configure: Regenerate. + * Makefile.in: Likewise. + * testsuite/Makefile.in: Likewise. + * testsuite/libgomp-site-extra.exp.in (GXX_UNDER_TEST) + (GFORTRAN_UNDER_TEST): Set. + * testsuite/lib/libgomp.exp (libgomp_init): Adjust. + * testsuite/libgomp.c++/c++.exp: Use 'GXX_UNDER_TEST'. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + * testsuite/libgomp.fortran/fortran.exp: Use + 'GFORTRAN_UNDER_TEST'. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + +2023-05-15 Thomas Schwinge <thomas@codesourcery.com> + + PR testsuite/91884 + * testsuite/lib/libgomp.exp (libgomp_target_compile): Don't + specify compiler. + * testsuite/libgomp.c++/c++.exp (ALWAYS_CFLAGS): Specify compiler. + * testsuite/libgomp.c/c.exp (ALWAYS_CFLAGS): Likewise. + * testsuite/libgomp.fortran/fortran.exp (ALWAYS_CFLAGS): Likewise. + * testsuite/libgomp.graphite/graphite.exp (ALWAYS_CFLAGS): + Likewise. + * testsuite/libgomp.oacc-c++/c++.exp (ALWAYS_CFLAGS): Likewise. + * testsuite/libgomp.oacc-c/c.exp (ALWAYS_CFLAGS): Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp (ALWAYS_CFLAGS): + Likewise. + +2023-05-12 Tobias Burnus <tobias@codesourcery.com> + + PR libstdc++/109816 + * testsuite/libgomp.c++/target-map-class-1.C: New test. + * testsuite/libgomp.c++/target-map-class-2.C: New test. + +2023-05-12 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/lib/libgomp.exp (libgomp_target_compile): Generalize + 'lang_library_path' into a list of 'lang_library_paths'. + * testsuite/libgomp.c++/c++.exp: Adjust. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + * testsuite/libgomp.fortran/fortran.exp: Adjust. Use that for + libquadmath, too. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + +2023-05-12 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/lib/libgomp.exp (libgomp_target_compile): Don't look + at 'lang_test_file_found'. + * testsuite/libgomp.c++/c++.exp: Don't set and use it, and instead + 'return' early if not able to test. Simplify 'ld_library_path' setup. + * testsuite/libgomp.fortran/fortran.exp: Likewise. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + +2023-05-12 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c++/c++.exp: Resolve 'lang_test_file_found' + first. + * testsuite/libgomp.fortran/fortran.exp: Likewise. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + +2023-05-12 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c++/c++.exp: Localize 'lang_[...]' etc. + * testsuite/libgomp.c/c.exp: Likewise. + * testsuite/libgomp.fortran/fortran.exp: Likewise. + * testsuite/libgomp.graphite/graphite.exp: Likewise. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + * testsuite/libgomp.oacc-c/c.exp: Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + +2023-05-09 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c++/c++.exp: Don't set 'lang_test_file'. + * testsuite/libgomp.fortran/fortran.exp: Likewise. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + * testsuite/libgomp.oacc-fortran/fortran.exp: Likewise. + * testsuite/libgomp.c/c.exp: Unset 'lang_test_file_found' instead of + 'lang_test_file'. + * testsuite/libgomp.oacc-c/c.exp: Likewise. + * testsuite/libgomp.graphite/graphite.exp: Likewise. + * testsuite/lib/libgomp.exp (libgomp_target_compile): Look for + 'lang_test_file_found' instead of 'lang_test_file'. + +2023-05-09 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/lib/libgomp.exp (libgomp_init): Only use 'blddir' if + set. + * testsuite/libgomp.c++/c++.exp: Likewise. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + +2023-05-09 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c++/c++.exp (blddir): Don't set. + * testsuite/libgomp.oacc-c++/c++.exp (blddir): Likewise. + +2023-05-08 Thomas Schwinge <thomas@codesourcery.com> + + * testsuite/libgomp.c++/c++.exp: Use 'lang_include_flags' instead + of 'libstdcxx_includes'. + * testsuite/libgomp.oacc-c++/c++.exp: Likewise. + +2023-05-08 Thomas Schwinge <thomas@codesourcery.com> + + * target.c (gomp_target_rev): Instead of 'dev_to_host_cpy', + 'host_to_dev_cpy', 'token', take a single 'goacc_asyncqueue'. + * libgomp.h (gomp_target_rev): Adjust. + * libgomp-plugin.c (GOMP_PLUGIN_target_rev): Adjust. + * libgomp-plugin.h (GOMP_PLUGIN_target_rev): Adjust. + * plugin/plugin-gcn.c (process_reverse_offload): Adjust. + * plugin/plugin-nvptx.c (rev_off_dev_to_host_cpy) + (rev_off_host_to_dev_cpy): Remove. + (GOMP_OFFLOAD_run): Adjust. + +2023-05-04 Julian Brown <julian@codesourcery.com> + + PR fortran/109622 + * testsuite/libgomp.fortran/pr109622.f90: Move test... + * testsuite/libgomp.oacc-fortran/pr109622.f90: ...to here. Ignore + vector length warning. + * testsuite/libgomp.fortran/pr109622-2.f90: Move test... + * testsuite/libgomp.oacc-fortran/pr109622-2.f90: ...to here. Add + missing copyin/copyout variable. Ignore vector length warnings. + * testsuite/libgomp.fortran/pr109622-3.f90: Move test... + * testsuite/libgomp.oacc-fortran/pr109622-3.f90: ...to here. Ignore + vector length warnings. + * testsuite/libgomp.oacc-fortran/pr109622-4.f90: New test. + +2023-04-28 Julian Brown <julian@codesourcery.com> + + PR fortran/109622 + * testsuite/libgomp.fortran/pr109622.f90: New test. + * testsuite/libgomp.fortran/pr109622-2.f90: New test. + * testsuite/libgomp.fortran/pr109622-3.f90: New test. + +2023-04-25 Tobias Burnus <tobias@codesourcery.com> + + * testsuite/libgomp.c-c++-common/scan-1.c: New test. + * testsuite/libgomp.c/scan-23.c: New test. + * testsuite/libgomp.fortran/scan-2.f90: New test. + 2023-03-28 Rainer Orth <ro@CeBiTec.Uni-Bielefeld.DE> * testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c: Add diff --git a/libgomp/Makefile.in b/libgomp/Makefile.in index 2c81cca..3ef05e6 100644 --- a/libgomp/Makefile.in +++ b/libgomp/Makefile.in @@ -368,6 +368,7 @@ CFLAGS = @CFLAGS@ CPP = @CPP@ CPPFLAGS = @CPPFLAGS@ CPU_COUNT = @CPU_COUNT@ +CXX = @CXX@ CYGPATH_W = @CYGPATH_W@ DEFS = @DEFS@ DEPDIR = @DEPDIR@ @@ -382,6 +383,7 @@ EXEEXT = @EXEEXT@ FC = @FC@ FCFLAGS = @FCFLAGS@ FGREP = @FGREP@ +FLOCK = @FLOCK@ GREP = @GREP@ INSTALL = @INSTALL@ INSTALL_DATA = @INSTALL_DATA@ @@ -428,7 +430,6 @@ PACKAGE_TARNAME = @PACKAGE_TARNAME@ PACKAGE_URL = @PACKAGE_URL@ PACKAGE_VERSION = @PACKAGE_VERSION@ PATH_SEPARATOR = @PATH_SEPARATOR@ -PERL = @PERL@ RANLIB = @RANLIB@ SECTION_LDFLAGS = @SECTION_LDFLAGS@ SED = @SED@ diff --git a/libgomp/config/gcn/target.c b/libgomp/config/gcn/target.c index c6691fd..ea5eb1f 100644 --- a/libgomp/config/gcn/target.c +++ b/libgomp/config/gcn/target.c @@ -48,7 +48,9 @@ GOMP_teams4 (unsigned int num_teams_lower, unsigned int num_teams_upper, multiple times at least for some workgroups. */ (void) num_teams_lower; if (!num_teams_upper || num_teams_upper >= num_workgroups) - num_teams_upper = num_workgroups; + num_teams_upper = ((GOMP_ADDITIONAL_ICVS.nteams > 0 + && num_workgroups > GOMP_ADDITIONAL_ICVS.nteams) + ? GOMP_ADDITIONAL_ICVS.nteams : num_workgroups); else if (workgroup_id >= num_teams_upper) return false; gomp_num_teams_var = num_teams_upper - 1; diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c index f102d7d..125d92a 100644 --- a/libgomp/config/nvptx/target.c +++ b/libgomp/config/nvptx/target.c @@ -55,7 +55,9 @@ GOMP_teams4 (unsigned int num_teams_lower, unsigned int num_teams_upper, = thread_limit > INT_MAX ? UINT_MAX : thread_limit; } if (!num_teams_upper) - num_teams_upper = num_blocks; + num_teams_upper = ((GOMP_ADDITIONAL_ICVS.nteams > 0 + && num_blocks > GOMP_ADDITIONAL_ICVS.nteams) + ? GOMP_ADDITIONAL_ICVS.nteams : num_blocks); else if (num_blocks < num_teams_lower) num_teams_upper = num_teams_lower; else if (num_blocks < num_teams_upper) diff --git a/libgomp/configure b/libgomp/configure index fd0e337..e4e79c5 100755 --- a/libgomp/configure +++ b/libgomp/configure @@ -656,6 +656,7 @@ tmake_file XLDFLAGS XCFLAGS config_path +FLOCK CPU_COUNT LIBGOMP_BUILD_VERSIONED_SHLIB_SUN_FALSE LIBGOMP_BUILD_VERSIONED_SHLIB_SUN_TRUE @@ -678,6 +679,7 @@ libtool_VERSION ac_ct_FC FCFLAGS FC +CXX MAINT MAINTAINER_MODE_FALSE MAINTAINER_MODE_TRUE @@ -703,7 +705,6 @@ SED LIBTOOL BUILD_INFO_FALSE BUILD_INFO_TRUE -PERL RANLIB AR am__fastdepCC_FALSE @@ -4774,47 +4775,6 @@ else RANLIB="$ac_cv_prog_RANLIB" fi -# Extract the first word of "perl", so it can be a program name with args. -set dummy perl; ac_word=$2 -{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5 -$as_echo_n "checking for $ac_word... " >&6; } -if ${ac_cv_path_PERL+:} false; then : - $as_echo_n "(cached) " >&6 -else - case $PERL in - [\\/]* | ?:[\\/]*) - ac_cv_path_PERL="$PERL" # Let the user override the test with a path. - ;; - *) - as_save_IFS=$IFS; IFS=$PATH_SEPARATOR -for as_dir in $PATH -do - IFS=$as_save_IFS - test -z "$as_dir" && as_dir=. - for ac_exec_ext in '' $ac_executable_extensions; do - if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then - ac_cv_path_PERL="$as_dir/$ac_word$ac_exec_ext" - $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5 - break 2 - fi -done - done -IFS=$as_save_IFS - - test -z "$ac_cv_path_PERL" && ac_cv_path_PERL="perl-not-found-in-path-error" - ;; -esac -fi -PERL=$ac_cv_path_PERL -if test -n "$PERL"; then - { $as_echo "$as_me:${as_lineno-$LINENO}: result: $PERL" >&5 -$as_echo "$PERL" >&6; } -else - { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5 -$as_echo "no" >&6; } -fi - - { $as_echo "$as_me:${as_lineno-$LINENO}: checking whether ${MAKE-make} sets \$(MAKE)" >&5 $as_echo_n "checking whether ${MAKE-make} sets \$(MAKE)... " >&6; } set x ${MAKE-make} @@ -11418,7 +11378,7 @@ else lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2 lt_status=$lt_dlunknown cat > conftest.$ac_ext <<_LT_EOF -#line 11421 "configure" +#line 11381 "configure" #include "confdefs.h" #if HAVE_DLFCN_H @@ -11524,7 +11484,7 @@ else lt_dlunknown=0; lt_dlno_uscore=1; lt_dlneed_uscore=2 lt_status=$lt_dlunknown cat > conftest.$ac_ext <<_LT_EOF -#line 11527 "configure" +#line 11487 "configure" #include "confdefs.h" #if HAVE_DLFCN_H @@ -11810,11 +11770,22 @@ fi +# We optionally test libgomp C++ support, and for that want to use the proper +# C++ driver, 'g++' (or 'xg++' for build-tree testing). Given that build of +# target libstdc++-v3 depends on target libgomp (see '../Makefile.def'), we +# cannot make build of target libgomp depend on target libstdc++-v3: circular +# dependency. We thus cannot instantiate 'AC_PROG_CXX' here: we'd get +# '-funconfigured-libstdc++-v3' (see '../configure.ac'). Therefore, just +# capture 'CXX', and we'll fix this up at 'make check' time (see +# 'testsuite/lib/libgomp.exp:libgomp_init'). + + # Create a spec file, so that compile/link tests don't fail test -f libgfortran.spec || touch libgfortran.spec FCFLAGS="$FCFLAGS -L." -# We need gfortran to compile parts of the library +# We need 'gfortran' to compile parts of the library, and test libgomp Fortran +# support. # We can't use AC_PROG_FC because it expects a fully working gfortran. #AC_PROG_FC(gfortran) case `echo $GFORTRAN` in @@ -16486,6 +16457,91 @@ $as_echo "unable to detect (assuming 1)" >&6; } fi +{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for flock implementation" >&5 +$as_echo "$as_me: checking for flock implementation" >&6;} +for ac_prog in flock +do + # Extract the first word of "$ac_prog", so it can be a program name with args. +set dummy $ac_prog; ac_word=$2 +{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5 +$as_echo_n "checking for $ac_word... " >&6; } +if ${ac_cv_prog_FLOCK+:} false; then : + $as_echo_n "(cached) " >&6 +else + if test -n "$FLOCK"; then + ac_cv_prog_FLOCK="$FLOCK" # Let the user override the test. +else +as_save_IFS=$IFS; IFS=$PATH_SEPARATOR +for as_dir in $PATH +do + IFS=$as_save_IFS + test -z "$as_dir" && as_dir=. + for ac_exec_ext in '' $ac_executable_extensions; do + if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then + ac_cv_prog_FLOCK="$ac_prog" + $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5 + break 2 + fi +done + done +IFS=$as_save_IFS + +fi +fi +FLOCK=$ac_cv_prog_FLOCK +if test -n "$FLOCK"; then + { $as_echo "$as_me:${as_lineno-$LINENO}: result: $FLOCK" >&5 +$as_echo "$FLOCK" >&6; } +else + { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5 +$as_echo "no" >&6; } +fi + + + test -n "$FLOCK" && break +done + +# Fallback if 'perl' is available. +if test -z "$FLOCK"; then + # Extract the first word of "perl", so it can be a program name with args. +set dummy perl; ac_word=$2 +{ $as_echo "$as_me:${as_lineno-$LINENO}: checking for $ac_word" >&5 +$as_echo_n "checking for $ac_word... " >&6; } +if ${ac_cv_prog_FLOCK+:} false; then : + $as_echo_n "(cached) " >&6 +else + if test -n "$FLOCK"; then + ac_cv_prog_FLOCK="$FLOCK" # Let the user override the test. +else +as_save_IFS=$IFS; IFS=$PATH_SEPARATOR +for as_dir in $PATH +do + IFS=$as_save_IFS + test -z "$as_dir" && as_dir=. + for ac_exec_ext in '' $ac_executable_extensions; do + if as_fn_executable_p "$as_dir/$ac_word$ac_exec_ext"; then + ac_cv_prog_FLOCK="$srcdir/testsuite/flock" + $as_echo "$as_me:${as_lineno-$LINENO}: found $as_dir/$ac_word$ac_exec_ext" >&5 + break 2 + fi +done + done +IFS=$as_save_IFS + +fi +fi +FLOCK=$ac_cv_prog_FLOCK +if test -n "$FLOCK"; then + { $as_echo "$as_me:${as_lineno-$LINENO}: result: $FLOCK" >&5 +$as_echo "$FLOCK" >&6; } +else + { $as_echo "$as_me:${as_lineno-$LINENO}: result: no" >&5 +$as_echo "no" >&6; } +fi + + +fi + # Get target configury. . ${srcdir}/configure.tgt CFLAGS="$save_CFLAGS $XCFLAGS" @@ -16732,9 +16788,9 @@ fi # which will force linkage against -lpthread (or equivalent for the system). # That's not 100% ideal, but about the best we can do easily. if test $enable_shared = yes; then - link_gomp="-lgomp %{static: $LIBS}" + link_gomp="-lgomp %{static: $LIBS${DL_LIBS:+ $DL_LIBS}}" else - link_gomp="-lgomp $LIBS" + link_gomp="-lgomp $LIBS${DL_LIBS:+ $DL_LIBS}" fi diff --git a/libgomp/configure.ac b/libgomp/configure.ac index a9b1f39..1aad83a 100644 --- a/libgomp/configure.ac +++ b/libgomp/configure.ac @@ -132,7 +132,6 @@ AS_IF([test "x$enable_werror" != "xno" && test "x$GCC" = "xyes"], # Find other programs we need. AC_CHECK_TOOL(AR, ar) AC_CHECK_TOOL(RANLIB, ranlib, ranlib-not-found-in-path-error) -AC_PATH_PROG(PERL, perl, perl-not-found-in-path-error) AC_PROG_MAKE_SET AC_PROG_INSTALL @@ -152,11 +151,22 @@ AC_SUBST(enable_static) AM_MAINTAINER_MODE +# We optionally test libgomp C++ support, and for that want to use the proper +# C++ driver, 'g++' (or 'xg++' for build-tree testing). Given that build of +# target libstdc++-v3 depends on target libgomp (see '../Makefile.def'), we +# cannot make build of target libgomp depend on target libstdc++-v3: circular +# dependency. We thus cannot instantiate 'AC_PROG_CXX' here: we'd get +# '-funconfigured-libstdc++-v3' (see '../configure.ac'). Therefore, just +# capture 'CXX', and we'll fix this up at 'make check' time (see +# 'testsuite/lib/libgomp.exp:libgomp_init'). +AC_SUBST(CXX) + # Create a spec file, so that compile/link tests don't fail test -f libgfortran.spec || touch libgfortran.spec FCFLAGS="$FCFLAGS -L." -# We need gfortran to compile parts of the library +# We need 'gfortran' to compile parts of the library, and test libgomp Fortran +# support. # We can't use AC_PROG_FC because it expects a fully working gfortran. #AC_PROG_FC(gfortran) case `echo $GFORTRAN` in @@ -339,6 +349,13 @@ fi AX_COUNT_CPUS AC_SUBST(CPU_COUNT) +AC_MSG_NOTICE([checking for flock implementation]) +AC_CHECK_PROGS(FLOCK, flock) +# Fallback if 'perl' is available. +if test -z "$FLOCK"; then + AC_CHECK_PROG(FLOCK, perl, $srcdir/testsuite/flock) +fi + # Get target configury. . ${srcdir}/configure.tgt CFLAGS="$save_CFLAGS $XCFLAGS" @@ -385,9 +402,9 @@ fi # which will force linkage against -lpthread (or equivalent for the system). # That's not 100% ideal, but about the best we can do easily. if test $enable_shared = yes; then - link_gomp="-lgomp %{static: $LIBS}" + link_gomp="-lgomp %{static: $LIBS${DL_LIBS:+ $DL_LIBS}}" else - link_gomp="-lgomp $LIBS" + link_gomp="-lgomp $LIBS${DL_LIBS:+ $DL_LIBS}" fi AC_SUBST(link_gomp) diff --git a/libgomp/env.c b/libgomp/env.c index e7a035b..f24484d 100644 --- a/libgomp/env.c +++ b/libgomp/env.c @@ -62,13 +62,14 @@ #include "secure_getenv.h" #include "environ.h" -/* Default values of ICVs according to the OpenMP standard. */ +/* Default values of ICVs according to the OpenMP standard, + except for default-device-var. */ const struct gomp_default_icv gomp_default_icv_values = { .nthreads_var = 1, .thread_limit_var = UINT_MAX, .run_sched_var = GFS_DYNAMIC, .run_sched_chunk_size = 1, - .default_device_var = 0, + .default_device_var = INT_MIN, .max_active_levels_var = 1, .bind_var = omp_proc_bind_false, .nteams_var = 0, @@ -111,6 +112,7 @@ unsigned long gomp_bind_var_list_len; void **gomp_places_list; unsigned long gomp_places_list_len; uintptr_t gomp_def_allocator = omp_default_mem_alloc; +char *gomp_def_allocator_envvar = NULL; int gomp_debug_var; unsigned int gomp_num_teams_var; int gomp_nteams_var; @@ -1232,8 +1234,12 @@ parse_affinity (bool ignore) static bool parse_allocator (const char *env, const char *val, void *const params[]) { + const char *orig_val = val; uintptr_t *ret = (uintptr_t *) params[0]; *ret = omp_default_mem_alloc; + bool memspace = false; + size_t ntraits = 0; + omp_alloctrait_t *traits; if (val == NULL) return false; @@ -1242,28 +1248,169 @@ parse_allocator (const char *env, const char *val, void *const params[]) ++val; if (0) ; -#define C(v) \ +#define C(v, m) \ else if (strncasecmp (val, #v, sizeof (#v) - 1) == 0) \ { \ *ret = v; \ val += sizeof (#v) - 1; \ - } - C (omp_default_mem_alloc) - C (omp_large_cap_mem_alloc) - C (omp_const_mem_alloc) - C (omp_high_bw_mem_alloc) - C (omp_low_lat_mem_alloc) - C (omp_cgroup_mem_alloc) - C (omp_pteam_mem_alloc) - C (omp_thread_mem_alloc) + memspace = m; \ + } + C (omp_default_mem_alloc, false) + C (omp_large_cap_mem_alloc, false) + C (omp_const_mem_alloc, false) + C (omp_high_bw_mem_alloc, false) + C (omp_low_lat_mem_alloc, false) + C (omp_cgroup_mem_alloc, false) + C (omp_pteam_mem_alloc, false) + C (omp_thread_mem_alloc, false) + C (omp_default_mem_space, true) + C (omp_large_cap_mem_space, true) + C (omp_const_mem_space, true) + C (omp_high_bw_mem_space, true) + C (omp_low_lat_mem_space, true) #undef C else - val = "X"; + goto invalid; + if (memspace && *val == ':') + { + ++val; + const char *cp = val; + while (*cp != '\0') + { + if (*cp == '=') + ++ntraits; + ++cp; + } + traits = gomp_alloca (ntraits * sizeof (omp_alloctrait_t)); + size_t n = 0; + while (*val != '\0') + { +#define C(v) \ + else if (strncasecmp (val, #v "=", sizeof (#v)) == 0) \ + { \ + val += sizeof (#v); \ + traits[n].key = omp_atk_ ## v; +#define V(v) \ + else if (strncasecmp (val, #v, sizeof (#v) - 1) == 0) \ + { \ + val += sizeof (#v) - 1; \ + traits[n].value = omp_atv_ ## v; \ + } + if (0) + ; + C (sync_hint) + if (0) + ; + V (contended) + V (uncontended) + V (serialized) + V (private) + else + goto invalid; + } + C (alignment) + char *end; + errno = 0; + traits[n].value = strtol (val, &end, 10); + if (errno || end == val || traits[n].value <= 0) + goto invalid; + val = end; + } + C (access) + if (0) + ; + V (all) + V (cgroup) + V (pteam) + V (thread) + else + goto invalid; + } + C (pool_size) + char *end; + errno = 0; + traits[n].value = strtol (val, &end, 10); + if (errno || end == val || traits[n].value <= 0) + goto invalid; + val = end; + } + C (fallback) + if (0) + ; + V (default_mem_fb) + V (null_fb) + V (abort_fb) + V (allocator_fb) + else + goto invalid; + } + /* Ignore fb_data, which expects an allocator handle. */ + C (pinned) + if (0) + ; + V (true) + V (false) + else + goto invalid; + } + C (partition) + if (0) + ; + V (environment) + V (nearest) + V (blocked) + V (interleaved) + else + goto invalid; + } + else + goto invalid; + if (*val != ',') + break; + ++val; + ++n; + if (*val == '\0') + goto invalid; + } +#undef C +#undef V + } + else if (memspace) + switch (*ret) + { + case omp_default_mem_space: *ret = omp_default_mem_alloc; break; + case omp_large_cap_mem_space: *ret = omp_large_cap_mem_alloc; break; + case omp_const_mem_space: *ret = omp_const_mem_alloc; break; + case omp_high_bw_mem_space: *ret = omp_high_bw_mem_alloc; break; + case omp_low_lat_mem_space: *ret = omp_low_lat_mem_alloc; break; + default: __builtin_unreachable (); + } while (isspace ((unsigned char) *val)) ++val; if (*val == '\0') - return true; - print_env_var_error (env, val); + { + if (ntraits) + { + *ret = omp_init_allocator (*ret, ntraits, traits); + if (*ret == omp_null_allocator) + { + gomp_error ("Allocator of environment variable %.*s cannot be " + "created, using omp_default_mem_alloc instead", + (int) (orig_val - env - 1), env); + *ret = omp_default_mem_alloc; + } + else + gomp_def_allocator_envvar = strdup (orig_val); + } + return true; + } +invalid: + int len = (orig_val - env - 1); + if (*val == '\0') + gomp_error ("Missing value at the end of environment variable %s", env); + else + gomp_error ("Invalid value for environment variable %.*s when parsing: %s", + len, env, val); *ret = omp_default_mem_alloc; return false; } @@ -1614,6 +1761,10 @@ omp_display_env (int verbose) struct gomp_icv_list *none = gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_NO_SUFFIX); + if (none->icvs.default_device_var == INT_MIN) + /* This implies OMP_TARGET_OFFLOAD=mandatory. */ + gomp_init_targets_once (); + fputs ("\nOPENMP DISPLAY ENVIRONMENT BEGIN\n", stderr); fputs (" _OPENMP = '201511'\n", stderr); @@ -1779,7 +1930,11 @@ omp_display_env (int verbose) C (omp_pteam_mem_alloc) C (omp_thread_mem_alloc) #undef C - default: break; + /* For an OMP_ALLOCATOR with traits, '' will be output. */ + default: + if (gomp_def_allocator_envvar) + fputs (gomp_def_allocator_envvar, stderr); + break; } fputs ("'\n", stderr); @@ -2031,6 +2186,16 @@ startswith (const char *str, const char *prefix) return strncmp (str, prefix, strlen (prefix)) == 0; } +static void __attribute__((destructor)) +cleanup_env (void) +{ + if (gomp_def_allocator_envvar != NULL) + { + free (gomp_def_allocator_envvar); + omp_destroy_allocator (gomp_def_allocator); + } +} + static void __attribute__((constructor)) initialize_env (void) { @@ -2213,6 +2378,10 @@ initialize_env (void) gomp_global_icv.max_active_levels_var = gomp_supported_active_levels; } + if (gomp_global_icv.default_device_var == INT_MIN + && gomp_target_offload_var != GOMP_TARGET_OFFLOAD_MANDATORY) + none->icvs.default_device_var = gomp_global_icv.default_device_var = 0; + /* Process GOMP_* variables and dependencies between parsed ICVs. */ parse_int_secure ("GOMP_DEBUG", &gomp_debug_var, true); diff --git a/libgomp/icv-device.c b/libgomp/icv-device.c index a2bbedc..b48ea3b 100644 --- a/libgomp/icv-device.c +++ b/libgomp/icv-device.c @@ -27,6 +27,7 @@ expected to replace. */ #include "libgomp.h" +#include <limits.h> void omp_set_default_device (int device_num) @@ -41,6 +42,9 @@ int omp_get_default_device (void) { struct gomp_task_icv *icv = gomp_icv (false); + if (icv->default_device_var == INT_MIN) + /* This implies OMP_TARGET_OFFLOAD=mandatory. */ + gomp_init_targets_once (); return icv->default_device_var; } diff --git a/libgomp/libgomp-plugin.c b/libgomp/libgomp-plugin.c index 27e7c94..d696515 100644 --- a/libgomp/libgomp-plugin.c +++ b/libgomp/libgomp-plugin.c @@ -82,11 +82,8 @@ GOMP_PLUGIN_fatal (const char *msg, ...) void GOMP_PLUGIN_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, uint64_t sizes_ptr, uint64_t kinds_ptr, int dev_num, - void (*dev_to_host_cpy) (void *, const void *, size_t, - void *), - void (*host_to_dev_cpy) (void *, const void *, size_t, - void *), void *token) + struct goacc_asyncqueue *aq) { gomp_target_rev (fn_ptr, mapnum, devaddrs_ptr, sizes_ptr, kinds_ptr, dev_num, - dev_to_host_cpy, host_to_dev_cpy, token); + aq); } diff --git a/libgomp/libgomp-plugin.h b/libgomp/libgomp-plugin.h index 28267f7..42ee3d6 100644 --- a/libgomp/libgomp-plugin.h +++ b/libgomp/libgomp-plugin.h @@ -121,11 +121,7 @@ extern void GOMP_PLUGIN_fatal (const char *, ...) __attribute__ ((noreturn, format (printf, 1, 2))); extern void GOMP_PLUGIN_target_rev (uint64_t, uint64_t, uint64_t, uint64_t, - uint64_t, int, - void (*) (void *, const void *, size_t, - void *), - void (*) (void *, const void *, size_t, - void *), void *); + uint64_t, int, struct goacc_asyncqueue *); /* Prototypes for functions implemented by libgomp plugins. */ extern const char *GOMP_OFFLOAD_get_name (void); diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index ba8fe34..4d2bfab 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -1130,10 +1130,7 @@ extern void gomp_init_targets_once (void); extern int gomp_get_num_devices (void); extern bool gomp_target_task_fn (void *); extern void gomp_target_rev (uint64_t, uint64_t, uint64_t, uint64_t, uint64_t, - int, - void (*) (void *, const void *, size_t, void *), - void (*) (void *, const void *, size_t, void *), - void *); + int, struct goacc_asyncqueue *); /* Splay tree definitions. */ typedef struct splay_tree_node_s *splay_tree_node; diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi index dc6b4ac..db8b1f1 100644 --- a/libgomp/libgomp.texi +++ b/libgomp/libgomp.texi @@ -192,8 +192,7 @@ The OpenMP 4.5 specification is fully supported. env variable @tab Y @tab @item Nested-parallel changes to @emph{max-active-levels-var} ICV @tab Y @tab @item @code{requires} directive @tab P - @tab complete but no non-host devices provides @code{unified_address} or - @code{unified_shared_memory} + @tab complete but no non-host devices provides @code{unified_shared_memory} @item @code{teams} construct outside an enclosing target region @tab Y @tab @item Non-rectangular loop nests @tab P @tab Full support for C/C++, partial for Fortran @item @code{!=} as relational-op in canonical loop form for C/C++ @tab Y @tab @@ -311,7 +310,7 @@ The OpenMP 4.5 specification is fully supported. @item @code{inoutset} argument to the @code{depend} clause @tab Y @tab @item @code{private} and @code{firstprivate} argument to @code{default} clause in C and C++ @tab Y @tab -@item @code{present} argument to @code{defaultmap} clause @tab N @tab +@item @code{present} argument to @code{defaultmap} clause @tab Y @tab @item @code{omp_set_num_teams}, @code{omp_set_teams_thread_limit}, @code{omp_get_max_teams}, @code{omp_get_teams_thread_limit} runtime routines @tab Y @tab @@ -353,6 +352,8 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab @item Optional comma between directive and clause in the @code{#pragma} form @tab Y @tab @item @code{indirect} clause in @code{declare target} @tab N @tab @item @code{device_type(nohost)}/@code{device_type(host)} for variables @tab N @tab +@item @code{present} modifier to the @code{map}, @code{to} and @code{from} + clauses @tab Y @tab @end multitable @@ -383,12 +384,12 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab @item @code{declare mapper} with iterator and @code{present} modifiers @tab N @tab @item If a matching mapped list item is not found in the data environment, the - pointer retains its original value @tab N @tab + pointer retains its original value @tab Y @tab @item New @code{enter} clause as alias for @code{to} on declare target directive @tab Y @tab @item Deprecation of @code{to} clause on declare target directive @tab N @tab @item Extended list of directives permitted in Fortran pure procedures - @tab N @tab + @tab Y @tab @item New @code{allocators} directive for Fortran @tab N @tab @item Deprecation of @code{allocate} directive for Fortran allocatables/pointers @tab N @tab @@ -422,7 +423,7 @@ to address of matching mapped list item per 5.1, Sect. 2.21.7.2 @tab N @tab @item Conforming device numbers and @code{omp_initial_device} and @code{omp_invalid_device} enum/PARAMETER @tab Y @tab @item Initial value of @emph{default-device-var} ICV with - @code{OMP_TARGET_OFFLOAD=mandatory} @tab N @tab + @code{OMP_TARGET_OFFLOAD=mandatory} @tab Y @tab @item @emph{interop_types} in any position of the modifier list for the @code{init} clause of the @code{interop} construct @tab N @tab @end multitable @@ -1936,7 +1937,10 @@ section 4 of the OpenMP specification in version 4.5, while those beginning with @env{GOMP_} are GNU extensions. @menu +* OMP_ALLOCATOR:: Set the default allocator +* OMP_AFFINITY_FORMAT:: Set the format string used for affinity display * OMP_CANCELLATION:: Set whether cancellation is activated +* OMP_DISPLAY_AFFINITY:: Display thread affinity information * OMP_DISPLAY_ENV:: Show OpenMP version and environment variables * OMP_DEFAULT_DEVICE:: Set the device used in target regions * OMP_DYNAMIC:: Dynamic adjustment of threads @@ -1961,6 +1965,125 @@ beginning with @env{GOMP_} are GNU extensions. @end menu +@node OMP_ALLOCATOR +@section @env{OMP_ALLOCATOR} -- Set the default allocator +@cindex Environment Variable +@table @asis +@item @emph{Description}: +Sets the default allocator that is used when no allocator has been specified +in the @code{allocate} or @code{allocator} clause or if an OpenMP memory +routine is invoked with the @code{omp_null_allocator} allocator. +If unset, @code{omp_default_mem_alloc} is used. + +The value can either be a predefined allocator or a predefined memory space +or a predefined memory space followed by a colon and a comma-separated list +of memory trait and value pairs, separated by @code{=}. + +@multitable @columnfractions .45 .45 +@headitem Predefined allocators @tab Predefined memory spaces +@item omp_default_mem_alloc @tab omp_default_mem_space +@item omp_large_cap_mem_alloc @tab omp_large_cap_mem_space +@item omp_const_mem_alloc @tab omp_const_mem_space +@item omp_high_bw_mem_alloc @tab omp_high_bw_mem_space +@item omp_low_lat_mem_alloc @tab omp_low_lat_mem_space +@item omp_cgroup_mem_alloc @tab -- +@item omp_pteam_mem_alloc @tab -- +@item omp_thread_mem_alloc @tab -- +@end multitable + +@multitable @columnfractions .30 .60 +@headitem Trait @tab Allowed values +@item @code{sync_hint} @tab @code{contended}, @code{uncontended}, + @code{serialized}, @code{private} +@item @code{alignment} @tab Positive integer being a power of two +@item @code{access} @tab @code{all}, @code{cgroup}, + @code{pteam}, @code{thread} +@item @code{pool_size} @tab Positive integer +@item @code{fallback} @tab @code{default_mem_fb}, @code{null_fb}, + @code{abort_fb}, @code{allocator_fb} +@item @code{fb_data} @tab @emph{unsupported as it needs an allocator handle} +@item @code{pinned} @tab @code{true}, @code{false} +@item @code{partition} @tab @code{environment}, @code{nearest}, + @code{blocked}, @code{interleaved} +@end multitable + +Examples: +@smallexample +OMP_ALLOCATOR=omp_high_bw_mem_alloc +OMP_ALLOCATOR=omp_large_cap_mem_space +OMP_ALLOCATR=omp_low_lat_mem_space:pinned=true,partition=nearest +@end smallexample + +@c @item @emph{See also}: + +@item @emph{Reference}: +@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 6.21 +@end table + + + +@node OMP_AFFINITY_FORMAT +@section @env{OMP_AFFINITY_FORMAT} -- Set the format string used for affinity display +@cindex Environment Variable +@table @asis +@item @emph{Description}: +Sets the format string used when displaying OpenMP thread affinity information. +Special values are output using @code{%} followed by an optional size +specification and then either the single-character field type or its long +name enclosed in curly braces; using @code{%%} will display a literal percent. +The size specification consists of an optional @code{0.} or @code{.} followed +by a positive integer, specifing the minimal width of the output. With +@code{0.} and numerical values, the output is padded with zeros on the left; +with @code{.}, the output is padded by spaces on the left; otherwise, the +output is padded by spaces on the right. If unset, the value is +``@code{level %L thread %i affinity %A}''. + +Supported field types are: + +@multitable @columnfractions .10 .25 .60 +@item t @tab team_num @tab value returned by @code{omp_get_team_num} +@item T @tab num_teams @tab value returned by @code{omp_get_num_teams} +@item L @tab nesting_level @tab value returned by @code{omp_get_level} +@item n @tab thread_num @tab value returned by @code{omp_get_thread_num} +@item N @tab num_threads @tab value returned by @code{omp_get_num_threads} +@item a @tab ancestor_tnum + @tab value returned by + @code{omp_get_ancestor_thread_num(omp_get_level()-1)} +@item H @tab host @tab name of the host that executes the thread +@item P @tab process_id @tab process identifier +@item i @tab native_thread_id @tab native thread identifier +@item A @tab thread_affinity + @tab comma separated list of integer values or ranges, representing the + processors on which a process might execute, subject to affinity + mechanisms +@end multitable + +For instance, after setting + +@smallexample +OMP_AFFINITY_FORMAT="%0.2a!%n!%.4L!%N;%.2t;%0.2T;%@{team_num@};%@{num_teams@};%A" +@end smallexample + +with either @code{OMP_DISPLAY_AFFINITY} being set or when calling +@code{omp_display_affinity} with @code{NULL} or an empty string, the program +might display the following: + +@smallexample +00!0! 1!4; 0;01;0;1;0-11 +00!3! 1!4; 0;01;0;1;0-11 +00!2! 1!4; 0;01;0;1;0-11 +00!1! 1!4; 0;01;0;1;0-11 +@end smallexample + +@item @emph{See also}: +@ref{OMP_DISPLAY_AFFINITY} + +@item @emph{Reference}: +@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 6.14 +@end table + + + @node OMP_CANCELLATION @section @env{OMP_CANCELLATION} -- Set whether cancellation is activated @cindex Environment Variable @@ -1978,6 +2101,26 @@ if unset, cancellation is disabled and the @code{cancel} construct is ignored. +@node OMP_DISPLAY_AFFINITY +@section @env{OMP_DISPLAY_AFFINITY} -- Display thread affinity information +@cindex Environment Variable +@table @asis +@item @emph{Description}: +If set to @code{FALSE} or if unset, affinity displaying is disabled. +If set to @code{TRUE}, the runtime will display affinity information about +OpenMP threads in a parallel region upon entering the region and every time +any change occurs. + +@item @emph{See also}: +@ref{OMP_AFFINITY_FORMAT} + +@item @emph{Reference}: +@uref{https://www.openmp.org, OpenMP specification v5.0}, Section 6.13 +@end table + + + + @node OMP_DISPLAY_ENV @section @env{OMP_DISPLAY_ENV} -- Show OpenMP version and environment variables @cindex Environment Variable @@ -2005,6 +2148,8 @@ Set to choose the device which is used in a @code{target} region, unless the value is overridden by @code{omp_set_default_device} or by a @code{device} clause. The value shall be the nonnegative device number. If no device with the given device number exists, the code is executed on the host. If unset, +@env{OMP_TARGET_OFFLOAD} is @code{mandatory} and no non-host devices are +available, it is set to @code{omp_invalid_device}. Otherwise, if unset, device number 0 will be used. @@ -4458,7 +4603,7 @@ The implementation remark: @code{device(ancestor:1)}) are processed serially per @code{target} region such that the next reverse offload region is only executed after the previous one returned. -@item OpenMP code that has a requires directive with @code{unified_address} or +@item OpenMP code that has a @code{requires} directive with @code{unified_shared_memory} will remove any GCN device from the list of available devices (``host fallback''). @item The available stack size can be changed using the @code{GCN_STACK_SIZE} @@ -4520,8 +4665,8 @@ The implementation remark: Per device, reverse offload regions are processed serially such that the next reverse offload region is only executed after the previous one returned. -@item OpenMP code that has a requires directive with @code{unified_address} - or @code{unified_shared_memory} will remove any nvptx device from the +@item OpenMP code that has a @code{requires} directive with + @code{unified_shared_memory} will remove any nvptx device from the list of available devices (``host fallback''). @end itemize diff --git a/libgomp/plugin/plugin-gcn.c b/libgomp/plugin/plugin-gcn.c index 3478037..ef22d48 100644 --- a/libgomp/plugin/plugin-gcn.c +++ b/libgomp/plugin/plugin-gcn.c @@ -1949,7 +1949,7 @@ process_reverse_offload (uint64_t fn, uint64_t mapnum, uint64_t hostaddrs, { int dev_num = dev_num64; GOMP_PLUGIN_target_rev (fn, mapnum, hostaddrs, sizes, kinds, dev_num, - NULL, NULL, NULL); + NULL); } /* Output any data written to console output from the kernel. It is expected @@ -3231,7 +3231,9 @@ GOMP_OFFLOAD_get_num_devices (unsigned int omp_requires_mask) /* Return -1 if no omp_requires_mask cannot be fulfilled but devices were present. */ if (hsa_context.agent_count > 0 - && (omp_requires_mask & ~GOMP_REQUIRES_REVERSE_OFFLOAD) != 0) + && ((omp_requires_mask + & ~(GOMP_REQUIRES_UNIFIED_ADDRESS + | GOMP_REQUIRES_REVERSE_OFFLOAD)) != 0)) return -1; return hsa_context.agent_count; } diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c index b3481c4..ffc8e2d 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -56,6 +56,7 @@ #include <unistd.h> #include <assert.h> #include <errno.h> +#include <stdlib.h> /* An arbitrary fixed limit (128MB) for the size of the OpenMP soft stacks block to cache between kernel invocations. For soft-stacks blocks bigger @@ -1625,11 +1626,11 @@ GOMP_OFFLOAD_openacc_cuda_set_stream (struct goacc_asyncqueue *aq, void *stream) return 1; } -struct goacc_asyncqueue * -GOMP_OFFLOAD_openacc_async_construct (int device __attribute__((unused))) +static struct goacc_asyncqueue * +nvptx_goacc_asyncqueue_construct (unsigned int flags) { CUstream stream = NULL; - CUDA_CALL_ERET (NULL, cuStreamCreate, &stream, CU_STREAM_DEFAULT); + CUDA_CALL_ERET (NULL, cuStreamCreate, &stream, flags); struct goacc_asyncqueue *aq = GOMP_PLUGIN_malloc (sizeof (struct goacc_asyncqueue)); @@ -1637,14 +1638,26 @@ GOMP_OFFLOAD_openacc_async_construct (int device __attribute__((unused))) return aq; } -bool -GOMP_OFFLOAD_openacc_async_destruct (struct goacc_asyncqueue *aq) +struct goacc_asyncqueue * +GOMP_OFFLOAD_openacc_async_construct (int device __attribute__((unused))) +{ + return nvptx_goacc_asyncqueue_construct (CU_STREAM_DEFAULT); +} + +static bool +nvptx_goacc_asyncqueue_destruct (struct goacc_asyncqueue *aq) { CUDA_CALL_ERET (false, cuStreamDestroy, aq->cuda_stream); free (aq); return true; } +bool +GOMP_OFFLOAD_openacc_async_destruct (struct goacc_asyncqueue *aq) +{ + return nvptx_goacc_asyncqueue_destruct (aq); +} + int GOMP_OFFLOAD_openacc_async_test (struct goacc_asyncqueue *aq) { @@ -1658,14 +1671,20 @@ GOMP_OFFLOAD_openacc_async_test (struct goacc_asyncqueue *aq) return -1; } -bool -GOMP_OFFLOAD_openacc_async_synchronize (struct goacc_asyncqueue *aq) +static bool +nvptx_goacc_asyncqueue_synchronize (struct goacc_asyncqueue *aq) { CUDA_CALL_ERET (false, cuStreamSynchronize, aq->cuda_stream); return true; } bool +GOMP_OFFLOAD_openacc_async_synchronize (struct goacc_asyncqueue *aq) +{ + return nvptx_goacc_asyncqueue_synchronize (aq); +} + +bool GOMP_OFFLOAD_openacc_async_serialize (struct goacc_asyncqueue *aq1, struct goacc_asyncqueue *aq2) { @@ -1925,22 +1944,6 @@ nvptx_stacks_acquire (struct ptx_device *ptx_dev, size_t size, int num) void -rev_off_dev_to_host_cpy (void *dest, const void *src, size_t size, - CUstream stream) -{ - CUDA_CALL_ASSERT (cuMemcpyDtoHAsync, dest, (CUdeviceptr) src, size, stream); - CUDA_CALL_ASSERT (cuStreamSynchronize, stream); -} - -void -rev_off_host_to_dev_cpy (void *dest, const void *src, size_t size, - CUstream stream) -{ - CUDA_CALL_ASSERT (cuMemcpyHtoDAsync, (CUdeviceptr) dest, src, size, stream); - CUDA_CALL_ASSERT (cuStreamSynchronize, stream); -} - -void GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) { struct targ_fn_descriptor *tgt_fn_desc @@ -1973,9 +1976,17 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) } nvptx_adjust_launch_bounds (tgt_fn, ptx_dev, &teams, &threads); - size_t stack_size = nvptx_stacks_size (); bool reverse_offload = ptx_dev->rev_data != NULL; - CUstream copy_stream = NULL; + struct goacc_asyncqueue *reverse_offload_aq = NULL; + if (reverse_offload) + { + reverse_offload_aq + = nvptx_goacc_asyncqueue_construct (CU_STREAM_NON_BLOCKING); + if (!reverse_offload_aq) + exit (EXIT_FAILURE); + } + + size_t stack_size = nvptx_stacks_size (); pthread_mutex_lock (&ptx_dev->omp_stacks.lock); void *stacks = nvptx_stacks_acquire (ptx_dev, stack_size, teams * threads); @@ -1989,8 +2000,6 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) GOMP_PLUGIN_debug (0, " %s: kernel %s: launch" " [(teams: %u), 1, 1] [(lanes: 32), (threads: %u), 1]\n", __FUNCTION__, fn_name, teams, threads); - if (reverse_offload) - CUDA_CALL_ASSERT (cuStreamCreate, ©_stream, CU_STREAM_NON_BLOCKING); r = CUDA_CALL_NOCHECK (cuLaunchKernel, function, teams, 1, 1, 32, threads, 1, 0, NULL, NULL, config); if (r != CUDA_SUCCESS) @@ -2013,17 +2022,15 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) GOMP_PLUGIN_target_rev (rev_data->fn, rev_data->mapnum, rev_data->addrs, rev_data->sizes, rev_data->kinds, rev_data->dev_num, - rev_off_dev_to_host_cpy, - rev_off_host_to_dev_cpy, copy_stream); - CUDA_CALL_ASSERT (cuStreamSynchronize, copy_stream); + reverse_offload_aq); + if (!nvptx_goacc_asyncqueue_synchronize (reverse_offload_aq)) + exit (EXIT_FAILURE); __atomic_store_n (&rev_data->fn, 0, __ATOMIC_RELEASE); } usleep (1); } else r = CUDA_CALL_NOCHECK (cuCtxSynchronize, ); - if (reverse_offload) - CUDA_CALL_ASSERT (cuStreamDestroy, copy_stream); if (r == CUDA_ERROR_LAUNCH_FAILED) GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s %s\n", cuda_error (r), maybe_abort_msg); @@ -2031,6 +2038,12 @@ GOMP_OFFLOAD_run (int ord, void *tgt_fn, void *tgt_vars, void **args) GOMP_PLUGIN_fatal ("cuCtxSynchronize error: %s", cuda_error (r)); pthread_mutex_unlock (&ptx_dev->omp_stacks.lock); + + if (reverse_offload) + { + if (!nvptx_goacc_asyncqueue_destruct (reverse_offload_aq)) + exit (EXIT_FAILURE); + } } /* TODO: Implement GOMP_OFFLOAD_async_run. */ diff --git a/libgomp/target.c b/libgomp/target.c index b30c6a50..80c25a1 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -138,6 +138,10 @@ gomp_get_num_devices (void) static struct gomp_device_descr * resolve_device (int device_id, bool remapped) { + /* Get number of devices and thus ensure that 'gomp_init_targets_once' was + called, which must be done before using default_device_var. */ + int num_devices = gomp_get_num_devices (); + if (remapped && device_id == GOMP_DEVICE_ICV) { struct gomp_task_icv *icv = gomp_icv (false); @@ -150,7 +154,11 @@ resolve_device (int device_id, bool remapped) if (device_id == (remapped ? GOMP_DEVICE_HOST_FALLBACK : omp_initial_device)) return NULL; - if (device_id == omp_invalid_device) + if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY + && num_devices == 0) + gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, " + "but only the host device is available"); + else if (device_id == omp_invalid_device) gomp_fatal ("omp_invalid_device encountered"); else if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY) gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, " @@ -158,10 +166,10 @@ resolve_device (int device_id, bool remapped) return NULL; } - else if (device_id >= gomp_get_num_devices ()) + else if (device_id >= num_devices) { if (gomp_target_offload_var == GOMP_TARGET_OFFLOAD_MANDATORY - && device_id != num_devices_openmp) + && device_id != num_devices) gomp_fatal ("OMP_TARGET_OFFLOAD is set to MANDATORY, " "but device not found"); @@ -358,6 +366,8 @@ gomp_to_device_kind_p (int kind) case GOMP_MAP_FORCE_ALLOC: case GOMP_MAP_FORCE_FROM: case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_FORCE_PRESENT: return false; default: return true; @@ -593,7 +603,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep, else tgt_var->length = newn->host_end - newn->host_start; - if ((kind & GOMP_MAP_FLAG_FORCE) + if (GOMP_MAP_FORCE_P (kind) /* For implicit maps, old contained in new is valid. */ || !(implicit_subset /* Otherwise, new contained inside old is considered valid. */ @@ -1143,7 +1153,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, if (!n) { tgt->list[i].key = NULL; - tgt->list[i].offset = OFFSET_POINTER; + tgt->list[i].offset = OFFSET_INLINED; continue; } } @@ -1697,20 +1707,26 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep, i = j - 1; break; case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: { /* We already looked up the memory region above and it was missing. */ size_t size = k->host_end - k->host_start; gomp_mutex_unlock (&devicep->lock); #ifdef HAVE_INTTYPES_H - gomp_fatal ("present clause: !acc_is_present (%p, " - "%"PRIu64" (0x%"PRIx64"))", - (void *) k->host_start, - (uint64_t) size, (uint64_t) size); + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), " + "dev: %d)", (void *) k->host_start, + (uint64_t) size, (uint64_t) size, + devicep->target_id); #else - gomp_fatal ("present clause: !acc_is_present (%p, " - "%lu (0x%lx))", (void *) k->host_start, - (unsigned long) size, (unsigned long) size); + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %lu (0x%lx), dev: %d)", + (void *) k->host_start, + (unsigned long) size, (unsigned long) size, + devicep->target_id); #endif } break; @@ -2124,6 +2140,29 @@ gomp_update (struct gomp_device_descr *devicep, size_t mapnum, void **hostaddrs, gomp_copy_dev2host (devicep, NULL, hostaddr, devaddr, size); } } + else + { + int kind = get_kind (short_mapkind, kinds, i); + + if (GOMP_MAP_PRESENT_P (kind)) + { + /* We already looked up the memory region above and it + was missing. */ + gomp_mutex_unlock (&devicep->lock); +#ifdef HAVE_INTTYPES_H + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %"PRIu64" (0x%"PRIx64"), " + "dev: %d)", (void *) hostaddrs[i], + (uint64_t) sizes[i], (uint64_t) sizes[i], + devicep->target_id); +#else + gomp_fatal ("present clause: not present on the device " + "(addr: %p, size: %lu (0x%lx), dev: %d)", + (void *) hostaddrs[i], (unsigned long) sizes[i], + (unsigned long) sizes[i], devicep->target_id); +#endif + } + } } gomp_mutex_unlock (&devicep->lock); } @@ -3299,9 +3338,7 @@ gomp_map_cdata_lookup (struct cpy_data *d, uint64_t *devaddrs, void gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, uint64_t sizes_ptr, uint64_t kinds_ptr, int dev_num, - void (*dev_to_host_cpy) (void *, const void *, size_t, void*), - void (*host_to_dev_cpy) (void *, const void *, size_t, void*), - void *token) + struct goacc_asyncqueue *aq) { /* Return early if there is no offload code. */ if (sizeof (OFFLOAD_PLUGINS) == sizeof ("")) @@ -3343,26 +3380,17 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, devaddrs = (uint64_t *) gomp_malloc (mapnum * sizeof (uint64_t)); sizes = (uint64_t *) gomp_malloc (mapnum * sizeof (uint64_t)); kinds = (unsigned short *) gomp_malloc (mapnum * sizeof (unsigned short)); - if (dev_to_host_cpy) - { - dev_to_host_cpy (devaddrs, (const void *) (uintptr_t) devaddrs_ptr, - mapnum * sizeof (uint64_t), token); - dev_to_host_cpy (sizes, (const void *) (uintptr_t) sizes_ptr, - mapnum * sizeof (uint64_t), token); - dev_to_host_cpy (kinds, (const void *) (uintptr_t) kinds_ptr, - mapnum * sizeof (unsigned short), token); - } - else - { - gomp_copy_dev2host (devicep, NULL, devaddrs, - (const void *) (uintptr_t) devaddrs_ptr, - mapnum * sizeof (uint64_t)); - gomp_copy_dev2host (devicep, NULL, sizes, - (const void *) (uintptr_t) sizes_ptr, - mapnum * sizeof (uint64_t)); - gomp_copy_dev2host (devicep, NULL, kinds, (const void *) (uintptr_t) kinds_ptr, - mapnum * sizeof (unsigned short)); - } + gomp_copy_dev2host (devicep, aq, devaddrs, + (const void *) (uintptr_t) devaddrs_ptr, + mapnum * sizeof (uint64_t)); + gomp_copy_dev2host (devicep, aq, sizes, + (const void *) (uintptr_t) sizes_ptr, + mapnum * sizeof (uint64_t)); + gomp_copy_dev2host (devicep, aq, kinds, + (const void *) (uintptr_t) kinds_ptr, + mapnum * sizeof (unsigned short)); + if (aq && !devicep->openacc.async.synchronize_func (aq)) + exit (EXIT_FAILURE); } size_t tgt_align = 0, tgt_size = 0; @@ -3389,13 +3417,14 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, if (devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) memcpy (tgt + tgt_size, (void *) (uintptr_t) devaddrs[i], (size_t) sizes[i]); - else if (dev_to_host_cpy) - dev_to_host_cpy (tgt + tgt_size, (void *) (uintptr_t) devaddrs[i], - (size_t) sizes[i], token); else - gomp_copy_dev2host (devicep, NULL, tgt + tgt_size, - (void *) (uintptr_t) devaddrs[i], - (size_t) sizes[i]); + { + gomp_copy_dev2host (devicep, aq, tgt + tgt_size, + (void *) (uintptr_t) devaddrs[i], + (size_t) sizes[i]); + if (aq && !devicep->openacc.async.synchronize_func (aq)) + exit (EXIT_FAILURE); + } devaddrs[i] = (uint64_t) (uintptr_t) tgt + tgt_size; tgt_size = tgt_size + sizes[i]; if ((devicep->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM) @@ -3432,7 +3461,8 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, case GOMP_MAP_DELETE: case GOMP_MAP_RELEASE: case GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION: - /* Assume it is present; look it up - but ignore otherwise. */ + /* Assume it is present; look it up - but ignore unless the + present clause is there. */ case GOMP_MAP_ALLOC: case GOMP_MAP_FROM: case GOMP_MAP_FORCE_ALLOC: @@ -3444,6 +3474,10 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_ALWAYS_TO: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: case GOMP_MAP_ZERO_LEN_ARRAY_SECTION: cdata[i].devaddr = devaddrs[i]; bool zero_len = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION @@ -3464,7 +3498,23 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, devaddrs[i] + sizes[i], zero_len); cdata[i].present = n2 != NULL; } - if (!cdata[i].present + if (!cdata[i].present && GOMP_MAP_PRESENT_P (kind)) + { + gomp_mutex_unlock (&devicep->lock); +#ifdef HAVE_INTTYPES_H + gomp_fatal ("present clause: no corresponding data on " + "parent device at %p with size %"PRIu64, + (void *) (uintptr_t) devaddrs[i], + (uint64_t) sizes[i]); +#else + gomp_fatal ("present clause: no corresponding data on " + "parent device at %p with size %lu", + (void *) (uintptr_t) devaddrs[i], + (unsigned long) sizes[i]); +#endif + break; + } + else if (!cdata[i].present && kind != GOMP_MAP_DELETE && kind != GOMP_MAP_RELEASE && kind != GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION) @@ -3482,18 +3532,17 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, && (kind == GOMP_MAP_TO || kind == GOMP_MAP_TOFROM)) || kind == GOMP_MAP_FORCE_TO || kind == GOMP_MAP_FORCE_TOFROM - || kind == GOMP_MAP_ALWAYS_TO - || kind == GOMP_MAP_ALWAYS_TOFROM) + || GOMP_MAP_ALWAYS_TO_P (kind)) { - if (dev_to_host_cpy) - dev_to_host_cpy ((void *) (uintptr_t) devaddrs[i], - (void *) (uintptr_t) cdata[i].devaddr, - sizes[i], token); - else - gomp_copy_dev2host (devicep, NULL, - (void *) (uintptr_t) devaddrs[i], - (void *) (uintptr_t) cdata[i].devaddr, - sizes[i]); + gomp_copy_dev2host (devicep, aq, + (void *) (uintptr_t) devaddrs[i], + (void *) (uintptr_t) cdata[i].devaddr, + sizes[i]); + if (aq && !devicep->openacc.async.synchronize_func (aq)) + { + gomp_mutex_unlock (&devicep->lock); + exit (EXIT_FAILURE); + } } if (struct_cpy) struct_cpy--; @@ -3560,15 +3609,15 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, devaddrs[i] = (uint64_t) (uintptr_t) gomp_aligned_alloc (align, sizes[i]); - if (dev_to_host_cpy) - dev_to_host_cpy ((void *) (uintptr_t) devaddrs[i], - (void *) (uintptr_t) cdata[i].devaddr, - sizes[i], token); - else - gomp_copy_dev2host (devicep, NULL, - (void *) (uintptr_t) devaddrs[i], - (void *) (uintptr_t) cdata[i].devaddr, - sizes[i]); + gomp_copy_dev2host (devicep, aq, + (void *) (uintptr_t) devaddrs[i], + (void *) (uintptr_t) cdata[i].devaddr, + sizes[i]); + if (aq && !devicep->openacc.async.synchronize_func (aq)) + { + gomp_mutex_unlock (&devicep->lock); + exit (EXIT_FAILURE); + } } for (j = i + 1; j < mapnum; j++) { @@ -3668,19 +3717,21 @@ gomp_target_rev (uint64_t fn_ptr, uint64_t mapnum, uint64_t devaddrs_ptr, case GOMP_MAP_FORCE_TOFROM: case GOMP_MAP_ALWAYS_FROM: case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: copy = true; /* FALLTHRU */ case GOMP_MAP_FROM: case GOMP_MAP_TOFROM: - if (copy && host_to_dev_cpy) - host_to_dev_cpy ((void *) (uintptr_t) cdata[i].devaddr, - (void *) (uintptr_t) devaddrs[i], - sizes[i], token); - else if (copy) - gomp_copy_host2dev (devicep, NULL, - (void *) (uintptr_t) cdata[i].devaddr, - (void *) (uintptr_t) devaddrs[i], - sizes[i], false, NULL); + if (copy) + { + gomp_copy_host2dev (devicep, aq, + (void *) (uintptr_t) cdata[i].devaddr, + (void *) (uintptr_t) devaddrs[i], + sizes[i], false, NULL); + if (aq && !devicep->openacc.async.synchronize_func (aq)) + exit (EXIT_FAILURE); + } default: break; } @@ -5141,6 +5192,15 @@ gomp_target_init (void) if (devs[i].capabilities & GOMP_OFFLOAD_CAP_OPENACC_200) goacc_register (&devs[i]); } + if (gomp_global_icv.default_device_var == INT_MIN) + { + /* This implies OMP_TARGET_OFFLOAD=mandatory. */ + struct gomp_icv_list *none; + none = gomp_get_initial_icv_item (GOMP_DEVICE_NUM_FOR_NO_SUFFIX); + gomp_global_icv.default_device_var = (num_devs_openmp + ? 0 : omp_invalid_device); + none->icvs.default_device_var = gomp_global_icv.default_device_var; + } num_devices = num_devs; num_devices_openmp = num_devs_openmp; diff --git a/libgomp/testsuite/Makefile.am b/libgomp/testsuite/Makefile.am index 655a413..0cc91cc 100644 --- a/libgomp/testsuite/Makefile.am +++ b/libgomp/testsuite/Makefile.am @@ -12,6 +12,8 @@ _RUNTEST = $(shell if test -f $(top_srcdir)/../dejagnu/runtest; then \ echo $(top_srcdir)/../dejagnu/runtest; else echo runtest; fi) RUNTESTDEFAULTFLAGS = --tool $$tool --srcdir $$srcdir +PWD_COMMAND = $${PWDCMD-pwd} + EXTRA_DEJAGNU_SITE_CONFIG = libgomp-site-extra.exp # Instead of directly in ../testsuite/libgomp-test-support.exp.in, the @@ -25,17 +27,6 @@ libgomp-test-support.exp: libgomp-test-support.pt.exp Makefile 'set offload_additional_lib_paths "$(offload_additional_lib_paths)"' mv $@.tmp $@ -check-DEJAGNU: site.exp - srcdir='$(srcdir)'; export srcdir; \ - EXPECT=$(EXPECT); export EXPECT; \ - if $(SHELL) -c "$(_RUNTEST) --version" > /dev/null 2>&1; then \ - exit_status=0; l='$(PACKAGE)'; for tool in $$l; do \ - if $(_RUNTEST) $(AM_RUNTESTFLAGS) $(RUNTESTDEFAULTFLAGS) $(RUNTESTFLAGS); \ - then :; else exit_status=1; fi; \ - done; \ - else echo "WARNING: could not find '$(_RUNTEST)'" 1>&2; :;\ - fi; \ - exit $$exit_status site.exp: Makefile $(EXTRA_DEJAGNU_SITE_CONFIG) @echo 'Making a new site.exp file ...' @echo '## these variables are automatically generated by make ##' >site.tmp @@ -63,6 +54,72 @@ site.exp: Makefile $(EXTRA_DEJAGNU_SITE_CONFIG) @test ! -f site.exp || mv site.exp site.bak @mv site.tmp site.exp +%/site.exp: site.exp + -@test -d $* || mkdir $* + @srcdir=`cd $(srcdir); ${PWD_COMMAND}`; + @objdir=`${PWD_COMMAND}`/$*; \ + sed -e "s|^set srcdir .*$$|set srcdir $$srcdir|" \ + -e "s|^set objdir .*$$|set objdir $$objdir|" \ + site.exp > $*/site.exp.tmp + @-rm -f $*/site.bak + @test ! -f $*/site.exp || mv $*/site.exp $*/site.bak + @mv $*/site.exp.tmp $*/site.exp + +check_p_numbers0:=1 2 3 4 5 6 7 8 9 +check_p_numbers1:=0 $(check_p_numbers0) +check_p_numbers2:=$(foreach i,$(check_p_numbers0),$(addprefix $(i),$(check_p_numbers1))) +check_p_numbers3:=$(addprefix 0,$(check_p_numbers1)) $(check_p_numbers2) +check_p_numbers4:=$(foreach i,$(check_p_numbers0),$(addprefix $(i),$(check_p_numbers3))) +check_p_numbers5:=$(addprefix 0,$(check_p_numbers3)) $(check_p_numbers4) +check_p_numbers6:=$(foreach i,$(check_p_numbers0),$(addprefix $(i),$(check_p_numbers5))) +check_p_numbers:=$(check_p_numbers0) $(check_p_numbers2) $(check_p_numbers4) $(check_p_numbers6) +# If unable to serialize execution testing, use just one parallel slot. +gcc_test_parallel_slots:=$(if $(FLOCK),$(if $(GCC_TEST_PARALLEL_SLOTS),$(GCC_TEST_PARALLEL_SLOTS),19),1) +check_p_subdirs=$(wordlist 1,$(gcc_test_parallel_slots),$(check_p_numbers)) +check_DEJAGNU_libgomp_targets = $(addprefix check-DEJAGNUlibgomp,$(check_p_subdirs)) +$(check_DEJAGNU_libgomp_targets): check-DEJAGNUlibgomp%: libgomp%/site.exp + +check-DEJAGNU $(check_DEJAGNU_libgomp_targets): check-DEJAGNU%: site.exp + $(if $*,@)AR="$(AR)"; export AR; \ + RANLIB="$(RANLIB)"; export RANLIB; \ + if [ -z "$*" ] && [ -n "$(filter -j%, $(MFLAGS))" ]; then \ + rm -rf libgomp-parallel || true; \ + mkdir libgomp-parallel; \ + $(MAKE) $(AM_MAKEFLAGS) $(check_DEJAGNU_libgomp_targets); \ + rm -rf libgomp-parallel || true; \ + for idx in $(check_p_subdirs); do \ + if [ -d libgomp$$idx ]; then \ + mv -f libgomp$$idx/libgomp.sum libgomp$$idx/libgomp.sum.sep; \ + mv -f libgomp$$idx/libgomp.log libgomp$$idx/libgomp.log.sep; \ + fi; \ + done; \ + $(SHELL) $(srcdir)/../../contrib/dg-extract-results.sh \ + libgomp[0-9]*/libgomp.sum.sep > libgomp.sum; \ + $(SHELL) $(srcdir)/../../contrib/dg-extract-results.sh -L \ + libgomp[0-9]*/libgomp.log.sep > libgomp.log; \ + exit 0; \ + fi; \ + srcdir=`$(am__cd) $(srcdir) && pwd`; export srcdir; \ + EXPECT=$(EXPECT); export EXPECT; \ + runtest=$(_RUNTEST); \ + if [ -z "$$runtest" ]; then runtest=runtest; fi; \ + tool=libgomp; \ + if [ -n "$*" ]; then \ + if [ -f libgomp-parallel/finished ]; then rm -rf "$*"; exit 0; fi; \ + GCC_RUNTEST_PARALLELIZE_DIR=`${PWD_COMMAND}`/libgomp-parallel; \ + export GCC_RUNTEST_PARALLELIZE_DIR; \ + cd "$*"; \ + fi; \ + if $(SHELL) -c "$$runtest --version" > /dev/null 2>&1; then \ + $$runtest $(AM_RUNTESTFLAGS) $(RUNTESTDEFAULTFLAGS) \ + $(RUNTESTFLAGS); \ + if [ -n "$*" ]; then \ + touch $$GCC_RUNTEST_PARALLELIZE_DIR/finished; \ + fi; \ + else \ + echo "WARNING: could not find \`runtest'" 1>&2; :;\ + fi + distclean-DEJAGNU: -rm -f site.exp site.bak -l='$(PACKAGE)'; for tool in $$l; do \ diff --git a/libgomp/testsuite/Makefile.in b/libgomp/testsuite/Makefile.in index 7a88f0f..d743e46 100644 --- a/libgomp/testsuite/Makefile.in +++ b/libgomp/testsuite/Makefile.in @@ -147,6 +147,7 @@ CFLAGS = @CFLAGS@ CPP = @CPP@ CPPFLAGS = @CPPFLAGS@ CPU_COUNT = @CPU_COUNT@ +CXX = @CXX@ CYGPATH_W = @CYGPATH_W@ DEFS = @DEFS@ DEPDIR = @DEPDIR@ @@ -161,6 +162,7 @@ EXEEXT = @EXEEXT@ FC = @FC@ FCFLAGS = @FCFLAGS@ FGREP = @FGREP@ +FLOCK = @FLOCK@ GREP = @GREP@ INSTALL = @INSTALL@ INSTALL_DATA = @INSTALL_DATA@ @@ -207,7 +209,6 @@ PACKAGE_TARNAME = @PACKAGE_TARNAME@ PACKAGE_URL = @PACKAGE_URL@ PACKAGE_VERSION = @PACKAGE_VERSION@ PATH_SEPARATOR = @PATH_SEPARATOR@ -PERL = @PERL@ RANLIB = @RANLIB@ SECTION_LDFLAGS = @SECTION_LDFLAGS@ SED = @SED@ @@ -299,7 +300,20 @@ _RUNTEST = $(shell if test -f $(top_srcdir)/../dejagnu/runtest; then \ echo $(top_srcdir)/../dejagnu/runtest; else echo runtest; fi) RUNTESTDEFAULTFLAGS = --tool $$tool --srcdir $$srcdir +PWD_COMMAND = $${PWDCMD-pwd} EXTRA_DEJAGNU_SITE_CONFIG = libgomp-site-extra.exp +check_p_numbers0 := 1 2 3 4 5 6 7 8 9 +check_p_numbers1 := 0 $(check_p_numbers0) +check_p_numbers2 := $(foreach i,$(check_p_numbers0),$(addprefix $(i),$(check_p_numbers1))) +check_p_numbers3 := $(addprefix 0,$(check_p_numbers1)) $(check_p_numbers2) +check_p_numbers4 := $(foreach i,$(check_p_numbers0),$(addprefix $(i),$(check_p_numbers3))) +check_p_numbers5 := $(addprefix 0,$(check_p_numbers3)) $(check_p_numbers4) +check_p_numbers6 := $(foreach i,$(check_p_numbers0),$(addprefix $(i),$(check_p_numbers5))) +check_p_numbers := $(check_p_numbers0) $(check_p_numbers2) $(check_p_numbers4) $(check_p_numbers6) +# If unable to serialize execution testing, use just one parallel slot. +gcc_test_parallel_slots := $(if $(FLOCK),$(if $(GCC_TEST_PARALLEL_SLOTS),$(GCC_TEST_PARALLEL_SLOTS),19),1) +check_p_subdirs = $(wordlist 1,$(gcc_test_parallel_slots),$(check_p_numbers)) +check_DEJAGNU_libgomp_targets = $(addprefix check-DEJAGNUlibgomp,$(check_p_subdirs)) all: all-am .SUFFIXES: @@ -474,17 +488,6 @@ libgomp-test-support.exp: libgomp-test-support.pt.exp Makefile 'set offload_additional_lib_paths "$(offload_additional_lib_paths)"' mv $@.tmp $@ -check-DEJAGNU: site.exp - srcdir='$(srcdir)'; export srcdir; \ - EXPECT=$(EXPECT); export EXPECT; \ - if $(SHELL) -c "$(_RUNTEST) --version" > /dev/null 2>&1; then \ - exit_status=0; l='$(PACKAGE)'; for tool in $$l; do \ - if $(_RUNTEST) $(AM_RUNTESTFLAGS) $(RUNTESTDEFAULTFLAGS) $(RUNTESTFLAGS); \ - then :; else exit_status=1; fi; \ - done; \ - else echo "WARNING: could not find '$(_RUNTEST)'" 1>&2; :;\ - fi; \ - exit $$exit_status site.exp: Makefile $(EXTRA_DEJAGNU_SITE_CONFIG) @echo 'Making a new site.exp file ...' @echo '## these variables are automatically generated by make ##' >site.tmp @@ -512,6 +515,59 @@ site.exp: Makefile $(EXTRA_DEJAGNU_SITE_CONFIG) @test ! -f site.exp || mv site.exp site.bak @mv site.tmp site.exp +%/site.exp: site.exp + -@test -d $* || mkdir $* + @srcdir=`cd $(srcdir); ${PWD_COMMAND}`; + @objdir=`${PWD_COMMAND}`/$*; \ + sed -e "s|^set srcdir .*$$|set srcdir $$srcdir|" \ + -e "s|^set objdir .*$$|set objdir $$objdir|" \ + site.exp > $*/site.exp.tmp + @-rm -f $*/site.bak + @test ! -f $*/site.exp || mv $*/site.exp $*/site.bak + @mv $*/site.exp.tmp $*/site.exp +$(check_DEJAGNU_libgomp_targets): check-DEJAGNUlibgomp%: libgomp%/site.exp + +check-DEJAGNU $(check_DEJAGNU_libgomp_targets): check-DEJAGNU%: site.exp + $(if $*,@)AR="$(AR)"; export AR; \ + RANLIB="$(RANLIB)"; export RANLIB; \ + if [ -z "$*" ] && [ -n "$(filter -j%, $(MFLAGS))" ]; then \ + rm -rf libgomp-parallel || true; \ + mkdir libgomp-parallel; \ + $(MAKE) $(AM_MAKEFLAGS) $(check_DEJAGNU_libgomp_targets); \ + rm -rf libgomp-parallel || true; \ + for idx in $(check_p_subdirs); do \ + if [ -d libgomp$$idx ]; then \ + mv -f libgomp$$idx/libgomp.sum libgomp$$idx/libgomp.sum.sep; \ + mv -f libgomp$$idx/libgomp.log libgomp$$idx/libgomp.log.sep; \ + fi; \ + done; \ + $(SHELL) $(srcdir)/../../contrib/dg-extract-results.sh \ + libgomp[0-9]*/libgomp.sum.sep > libgomp.sum; \ + $(SHELL) $(srcdir)/../../contrib/dg-extract-results.sh -L \ + libgomp[0-9]*/libgomp.log.sep > libgomp.log; \ + exit 0; \ + fi; \ + srcdir=`$(am__cd) $(srcdir) && pwd`; export srcdir; \ + EXPECT=$(EXPECT); export EXPECT; \ + runtest=$(_RUNTEST); \ + if [ -z "$$runtest" ]; then runtest=runtest; fi; \ + tool=libgomp; \ + if [ -n "$*" ]; then \ + if [ -f libgomp-parallel/finished ]; then rm -rf "$*"; exit 0; fi; \ + GCC_RUNTEST_PARALLELIZE_DIR=`${PWD_COMMAND}`/libgomp-parallel; \ + export GCC_RUNTEST_PARALLELIZE_DIR; \ + cd "$*"; \ + fi; \ + if $(SHELL) -c "$$runtest --version" > /dev/null 2>&1; then \ + $$runtest $(AM_RUNTESTFLAGS) $(RUNTESTDEFAULTFLAGS) \ + $(RUNTESTFLAGS); \ + if [ -n "$*" ]; then \ + touch $$GCC_RUNTEST_PARALLELIZE_DIR/finished; \ + fi; \ + else \ + echo "WARNING: could not find \`runtest'" 1>&2; :;\ + fi + distclean-DEJAGNU: -rm -f site.exp site.bak -l='$(PACKAGE)'; for tool in $$l; do \ diff --git a/libgomp/testsuite/config/default.exp b/libgomp/testsuite/config/default.exp index b7afc82..01569e6 100644 --- a/libgomp/testsuite/config/default.exp +++ b/libgomp/testsuite/config/default.exp @@ -13,5 +13,3 @@ # You should have received a copy of the GNU General Public License # along with this program; see the file COPYING3. If not see # <http://www.gnu.org/licenses/>. - -load_lib "standard.exp" diff --git a/libgomp/testsuite/flock b/libgomp/testsuite/flock new file mode 100755 index 0000000..71878b1 --- /dev/null +++ b/libgomp/testsuite/flock @@ -0,0 +1,17 @@ +#!/usr/bin/env perl + +use strict; +use warnings; + +# Only arguments '--exclusive 1' exactly are supported. +(@ARGV == 2) or die; +my $mode = shift; +($mode eq "--exclusive") or die; +my $fd = shift; +($fd eq "1") or die; + +use Fcntl ':flock'; + +open(my $fh, '>&=', 1) or die "open: $!"; + +flock($fh, LOCK_EX) or die "flock: $!"; diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp index e12236e..2f9e538 100644 --- a/libgomp/testsuite/lib/libgomp.exp +++ b/libgomp/testsuite/lib/libgomp.exp @@ -9,6 +9,7 @@ proc load_gcc_lib { filename } { } load_lib dg.exp +load_lib standard.exp # Required to use gcc-dg.exp - however, the latter should NOT be # loaded until ${tool}_target_compile is defined since it uses that @@ -42,14 +43,15 @@ load_gcc_lib torture-options.exp load_gcc_lib fortran-modules.exp # Try to load a test support file, built during libgomp configuration. -load_file libgomp-test-support.exp +# Search in '..' vs. '.' to support parallel vs. sequential testing. +if [info exists ::env(GCC_RUNTEST_PARALLELIZE_DIR)] { + load_file ../libgomp-test-support.exp +} else { + load_file libgomp-test-support.exp +} set dg-do-what-default run -# -# GCC_UNDER_TEST is the compiler under test. -# - set libgomp_compile_options "" # @@ -66,12 +68,11 @@ proc libgomp_init { args } { global srcdir blddir objdir tool_root_dir global libgomp_initialized global tmpdir - global blddir global gluefile wrap_flags global ALWAYS_CFLAGS global CFLAGS global TOOL_EXECUTABLE TOOL_OPTIONS - global GCC_UNDER_TEST + global GCC_UNDER_TEST GXX_UNDER_TEST GFORTRAN_UNDER_TEST global TESTING_IN_BUILD_TREE global target_triplet global always_ld_library_path @@ -90,12 +91,45 @@ proc libgomp_init { args } { setenv LANG C.ASCII } + if { $blddir != "" } { + # Fix up '-funconfigured-libstdc++-v3' in 'GXX_UNDER_TEST' (see + # '../../configure.ac'). + set flags_file "${blddir}/../libstdc++-v3/scripts/testsuite_flags" + if { [file exists $flags_file] } { + set flags [exec sh $flags_file --build-includes] + verbose -log "GXX_UNDER_TEST = $GXX_UNDER_TEST" + set GXX_UNDER_TEST [string map [list \ + " -funconfigured-libstdc++-v3 " " $flags " \ + ] $GXX_UNDER_TEST] + verbose -log "GXX_UNDER_TEST = $GXX_UNDER_TEST" + } + } if ![info exists GCC_UNDER_TEST] then { if [info exists TOOL_EXECUTABLE] { set GCC_UNDER_TEST $TOOL_EXECUTABLE } else { set GCC_UNDER_TEST "[find_gcc]" } + # Only if we're guessing 'GCC_UNDER_TEST', we're also going to guess + # 'GXX_UNDER_TEST', 'GFORTRAN_UNDER_TEST'. + if ![info exists GXX_UNDER_TEST] then { + if [info exists TOOL_EXECUTABLE] { + set GXX_UNDER_TEST $TOOL_EXECUTABLE + } else { + set GXX_UNDER_TEST "[find_g++]" + } + } else { + error "GXX_UNDER_TEST set but not GCC_UNDER_TEST" + } + if ![info exists GFORTRAN_UNDER_TEST] then { + if [info exists TOOL_EXECUTABLE] { + set GFORTRAN_UNDER_TEST $TOOL_EXECUTABLE + } else { + set GFORTRAN_UNDER_TEST "[find_gfortran]" + } + } else { + error "GFORTRAN_UNDER_TEST set but not GCC_UNDER_TEST" + } } if ![info exists tmpdir] { @@ -118,7 +152,7 @@ proc libgomp_init { args } { } # Compute what needs to be put into LD_LIBRARY_PATH - set always_ld_library_path ".:${blddir}/.libs" + set always_ld_library_path "." global offload_additional_lib_paths if { $offload_additional_lib_paths != "" } { @@ -157,6 +191,8 @@ proc libgomp_init { args } { lappend ALWAYS_CFLAGS "additional_flags=-B${blddir}/.libs" lappend ALWAYS_CFLAGS "additional_flags=-I${blddir}" lappend ALWAYS_CFLAGS "ldflags=-L${blddir}/.libs" + + append always_ld_library_path ":${blddir}/.libs" } # The top-level include directory, for gomp-constants.h. lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/../../include" @@ -233,25 +269,26 @@ proc libgomp_target_compile { source dest type options } { global gluefile wrap_flags global ALWAYS_CFLAGS global GCC_UNDER_TEST - global lang_test_file - global lang_library_path - global lang_link_flags - global lang_include_flags - global lang_source_re - - if { [info exists lang_test_file] } { - if { $blddir != "" } { - # Some targets use libgfortran.a%s in their specs, so they need - # a -B option for uninstalled testing. - lappend options "additional_flags=-B${blddir}/${lang_library_path}" - lappend options "ldflags=-L${blddir}/${lang_library_path}" - } - lappend options "ldflags=${lang_link_flags}" - if { [info exists lang_include_flags] \ - && [regexp ${lang_source_re} ${source}] } { - lappend options "additional_flags=${lang_include_flags}" + + global lang_source_re lang_include_flags + if { [info exists lang_include_flags] \ + && [regexp ${lang_source_re} ${source}] } { + lappend options "additional_flags=${lang_include_flags}" + } + + global lang_library_paths + if { [info exists lang_library_paths] } { + foreach lang_library_path $lang_library_paths { + # targets that use lib[...].a%s in their specs need a -B option + # for uninstalled testing. + lappend options "additional_flags=-B${blddir}/${lang_library_path}" + lappend options "ldflags=-L${blddir}/${lang_library_path}" } } + global lang_link_flags + if { [info exists lang_link_flags] } { + lappend options "ldflags=${lang_link_flags}" + } if { [target_info needs_status_wrapper] != "" && [info exists gluefile] } { lappend options "libs=${gluefile}" @@ -260,7 +297,6 @@ proc libgomp_target_compile { source dest type options } { lappend options "additional_flags=[libio_include_flags]" lappend options "timeout=[timeout_value]" - lappend options "compiler=$GCC_UNDER_TEST" set options [concat $libgomp_compile_options $options] @@ -292,6 +328,36 @@ proc libgomp_option_proc { option } { } } +if ![info exists ::env(GCC_RUNTEST_PARALLELIZE_DIR)] { + # No parallel testing. +} elseif { $FLOCK == "" } { + # Using just one parallel slot. +} else { + # Using several parallel slots. Override DejaGnu + # 'standard.exp:${tool}_load'... + rename libgomp_load standard_libgomp_load + proc libgomp_load { program args } { + # ... in order to serialize execution testing via an exclusive lock. + # We use stdout, as per <https://perldoc.perl.org/functions/flock> + # "[...] FILEHANDLE [...] be open with write intent to use LOCK_EX". + set lock_file ../lock + set lock_kind --exclusive + set lock_fd [open $lock_file a+] + set lock_clock_begin [clock seconds] + global FLOCK + exec $FLOCK $lock_kind 1 >@ $lock_fd + set lock_clock_end [clock seconds] + verbose -log "Got ${FLOCK}('$lock_file', '$lock_kind') at [clock format $lock_clock_end] after [expr $lock_clock_end - $lock_clock_begin] s" 2 + + set result [standard_libgomp_load $program $args] + + # Unlock (implicit with 'close'). + close $lock_fd + + return $result + } +} + # Translate offload target to OpenACC device type. Return the empty string if # not supported, and 'host' for offload target 'disable'. proc offload_target_to_openacc_device_type { offload_target } { diff --git a/libgomp/testsuite/libgomp-site-extra.exp.in b/libgomp/testsuite/libgomp-site-extra.exp.in index c0d2666..c824d85 100644 --- a/libgomp/testsuite/libgomp-site-extra.exp.in +++ b/libgomp/testsuite/libgomp-site-extra.exp.in @@ -1 +1,4 @@ +set FLOCK {@FLOCK@} set GCC_UNDER_TEST {@CC@} +set GXX_UNDER_TEST {@CXX@} +set GFORTRAN_UNDER_TEST {@FC@} diff --git a/libgomp/testsuite/libgomp.c++/c++.exp b/libgomp/testsuite/libgomp.c++/c++.exp index f4884e2..8b4563b 100644 --- a/libgomp/testsuite/libgomp.c++/c++.exp +++ b/libgomp/testsuite/libgomp.c++/c++.exp @@ -1,15 +1,21 @@ load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp -global shlib_ext - -set shlib_ext [get_shlib_extension] -set lang_link_flags "-lstdc++" -set lang_test_file_found 0 -set lang_library_path "../libstdc++-v3/src/.libs" -if [info exists lang_include_flags] then { - unset lang_include_flags +if { $blddir != "" } { + set libstdc++_library_path "../libstdc++-v3/src/.libs" + set shlib_ext [get_shlib_extension] + if { ![file exists "${blddir}/${libstdc++_library_path}/libstdc++.a"] + && ![file exists "${blddir}/${libstdc++_library_path}/libstdc++.${shlib_ext}"] } { + verbose -log "No libstdc++ library found, will not execute c++ tests" + unset libstdc++_library_path + return + } + lappend lang_library_paths ${libstdc++_library_path} +} elseif { ![info exists GXX_UNDER_TEST] } { + verbose -log "GXX_UNDER_TEST not defined, will not execute c++ tests" + return } +lappend ALWAYS_CFLAGS "compiler=$GXX_UNDER_TEST" # If a testcase doesn't have special options, use these. if ![info exists DEFAULT_CFLAGS] then { @@ -22,61 +28,25 @@ dg-init # Turn on OpenMP. lappend ALWAYS_CFLAGS "additional_flags=-fopenmp" -# Switch into C++ mode. Otherwise, the libgomp.c-c++-common/*.c -# files would be compiled as C files. -set SAVE_GCC_UNDER_TEST "$GCC_UNDER_TEST" -set GCC_UNDER_TEST "$GCC_UNDER_TEST -x c++" - -set blddir [lookfor_file [get_multilibs] libgomp] - +# Gather a list of all tests. +set tests [lsort [concat \ + [find $srcdir/$subdir *.C] \ + [find $srcdir/$subdir/../libgomp.c-c++-common *.c]]] +set ld_library_path $always_ld_library_path if { $blddir != "" } { - # Look for a static libstdc++ first. - if [file exists "${blddir}/${lang_library_path}/libstdc++.a"] { - set lang_test_file "${lang_library_path}/libstdc++.a" - set lang_test_file_found 1 - # We may have a shared only build, so look for a shared libstdc++. - } elseif [file exists "${blddir}/${lang_library_path}/libstdc++.${shlib_ext}"] { - set lang_test_file "${lang_library_path}/libstdc++.${shlib_ext}" - set lang_test_file_found 1 - } else { - puts "No libstdc++ library found, will not execute c++ tests" - } -} elseif { [info exists GXX_UNDER_TEST] } { - set lang_test_file_found 1 - # Needs to exist for libgomp.exp. - set lang_test_file "" -} else { - puts "GXX_UNDER_TEST not defined, will not execute c++ tests" + append ld_library_path ":${blddir}/${libstdc++_library_path}" } +append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] +set_ld_library_path_env_vars -if { $lang_test_file_found } { - # Gather a list of all tests. - set tests [lsort [concat \ - [find $srcdir/$subdir *.C] \ - [find $srcdir/$subdir/../libgomp.c-c++-common *.c]]] +# Main loop. +dg-runtest $tests "" $DEFAULT_CFLAGS - if { $blddir != "" } { - set ld_library_path "$always_ld_library_path:${blddir}/${lang_library_path}" - } else { - set ld_library_path "$always_ld_library_path" - } - append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] - set_ld_library_path_env_vars - - set flags_file "${blddir}/../libstdc++-v3/scripts/testsuite_flags" - if { [file exists $flags_file] } { - set libstdcxx_includes [exec sh $flags_file --build-includes] - } else { - set libstdcxx_includes "" - } - - # Main loop. - dg-runtest $tests "" "$libstdcxx_includes $DEFAULT_CFLAGS" +if { $blddir != "" } { + unset libstdc++_library_path + unset lang_library_paths } -# See above. -set GCC_UNDER_TEST "$SAVE_GCC_UNDER_TEST" - # All done. dg-finish diff --git a/libgomp/testsuite/libgomp.c++/target-18.C b/libgomp/testsuite/libgomp.c++/target-18.C index f1085b1..a21ed4e 100644 --- a/libgomp/testsuite/libgomp.c++/target-18.C +++ b/libgomp/testsuite/libgomp.c++/target-18.C @@ -20,7 +20,9 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -37,7 +39,9 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -55,7 +59,9 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -91,7 +97,8 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -110,7 +117,8 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -130,7 +138,8 @@ foo (int *&p, int *&q, int *&r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) diff --git a/libgomp/testsuite/libgomp.c++/target-19.C b/libgomp/testsuite/libgomp.c++/target-19.C index afa6e68..7bae31d 100644 --- a/libgomp/testsuite/libgomp.c++/target-19.C +++ b/libgomp/testsuite/libgomp.c++/target-19.C @@ -1,3 +1,8 @@ +/* { dg-additional-options "-O0" } */ +/* Disable optimization to ensure that the compiler does not exploit that + S::r + t will never be NULL due to int (&r) and (&t). */ + + extern "C" void abort (); struct S { char a[64]; int (&r)[2]; char b[64]; }; @@ -19,7 +24,9 @@ foo (S s, int (&t)[3], int z) #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0]) { if (sep) - err = s.r != (int *) 0 || t != (int *) 0; + /* Since OpenMP 5.2, if no matching mapped list it has been found, + pointers retain their original value. */ + err = s.r == (int *) 0 || t == (int *) 0; else err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; } @@ -28,7 +35,9 @@ foo (S s, int (&t)[3], int z) #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z]) { if (sep) - err = s.r != (int *) 0 || t != (int *) 0; + /* Since OpenMP 5.2, if no matching mapped list it has been found, + pointers retain their original value. */ + err = s.r == (int *) 0 || t == (int *) 0; else err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7; } diff --git a/libgomp/testsuite/libgomp.c++/target-map-class-1.C b/libgomp/testsuite/libgomp.c++/target-map-class-1.C new file mode 100644 index 0000000..ad4802d --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-map-class-1.C @@ -0,0 +1,98 @@ +/* PR middle-end/109816 */ + +/* This variant: without -flto, see target-map-class-2.C for -flto. */ + +/* iostream.h adds 'globl _ZSt21ios_base_library_initv' with _GLIBCXX_SYMVER_GNU, + but it shouldn't end up in the offload assembly but only in the host assembly. */ + +/* Example based on sollve_vv's test_target_data_map_classes.cpp; however, + relevant is only the 'include' and not the actual executable code. */ + +#include <iostream> +#include <omp.h> + +using namespace std; + +#define N 1000 + +struct A +{ + int *h_array; + int size, sum; + + A (int *array, const int s) : h_array(array), size(s), sum(0) { } + ~A() { h_array = NULL; } +}; + +void +test_map_tofrom_class_heap () +{ + int *array = new int[N]; + A *obj = new A (array, N); + + #pragma omp target map(from: array[:N]) map(tofrom: obj[:1]) + { + int *tmp_h_array = obj->h_array; + obj->h_array = array; + int tmp = 0; + for (int i = 0; i < N; ++i) + { + obj->h_array[i] = 4*i; + tmp += 3; + } + obj->h_array = tmp_h_array; + obj->sum = tmp; + } + + for (int i = 0; i < N; ++i) + if (obj->h_array[i] != 4*i) + __builtin_abort (); + + if (3*N != obj->sum) + { + std::cout << "sum: " << obj->sum << std::endl; + __builtin_abort (); + } + + delete obj; + delete[] array; +} + +void +test_map_tofrom_class_stack () +{ + int array[N]; + A obj(array, N); + + #pragma omp target map(from: array[:N]) map(tofrom: obj) + { + int *tmp_h_array = obj.h_array; + obj.h_array = array; + int tmp = 0; + for (int i = 0; i < N; ++i) + { + obj.h_array[i] = 7*i; + tmp += 5; + } + obj.h_array = tmp_h_array; + obj.sum = tmp; + } + + for (int i = 0; i < N; ++i) + if (obj.h_array[i] != 7*i) + __builtin_abort (); + + if (5*N != obj.sum) + { + std::cout << "sum: " << obj.sum << std::endl; + __builtin_abort (); + } +} + +int +main() +{ + test_map_tofrom_class_heap(); + test_map_tofrom_class_stack(); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c++/target-map-class-2.C b/libgomp/testsuite/libgomp.c++/target-map-class-2.C new file mode 100644 index 0000000..1ef20f7 --- /dev/null +++ b/libgomp/testsuite/libgomp.c++/target-map-class-2.C @@ -0,0 +1,6 @@ +/* { dg-additional-options "-flto" } */ +/* PR middle-end/109816 */ + +/* This variant: with -flto, see target-map-class-1.C for without -flto. */ + +#include "target-map-class-1.C" diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c new file mode 100644 index 0000000..bff0a6b --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-1.c @@ -0,0 +1,74 @@ +/* PR libgomp/109837 */ + +#include <assert.h> +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#pragma omp requires unified_address + +#define N 15 + +void +test_device (int dev) +{ + struct st { + int *ptr; + int n; + }; + struct st s; + + s.n = 10; + s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev); + int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev); + assert (s.ptr != NULL); + assert (ptr1 != NULL); + + int q[4] = {1,2,3,4}; + int *qptr; + #pragma omp target enter data map(q) device(device_num: dev) + #pragma omp target data use_device_addr(q) device(device_num: dev) + qptr = q; + + #pragma omp target map(to:s) device(device_num: dev) + for (int i = 0; i < s.n; i++) + s.ptr[i] = 23*i; + + int *ptr2 = &s.ptr[3]; + + #pragma omp target firstprivate(qptr) map(tofrom:ptr2) device(device_num: dev) + for (int i = 0; i < 4; i++) + *(qptr++) = ptr2[i]; + + #pragma omp target exit data map(q) device(device_num: dev) + for (int i = 0; i < 4; i++) + q[i] = 23 * (i+3); + + #pragma omp target map(to: ptr1) device(device_num: dev) + for (int i = 0; i < N; i++) + ptr1[i] = 11*i; + + int *ptr3 = (int *) malloc (sizeof (int)*N); + assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0, + omp_get_initial_device(), dev)); + for (int i = 0; i < N; i++) + assert (ptr3[i] == 11*i); + + free (ptr3); + omp_target_free (ptr1, dev); + omp_target_free (s.ptr, dev); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + if (ntgts) + fprintf (stderr, "Offloading devices exist\n"); /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */ + else + fprintf (stderr, "Only host fallback\n"); /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */ + + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c new file mode 100644 index 0000000..3b5dcd3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/requires-unified-addr-2.c @@ -0,0 +1,85 @@ +/* PR middle-end/110270 */ + +/* OpenMP 5.2's 'defaultmap(default : pointer) for C/C++ pointers retains the + pointer value instead of setting it to NULL if the pointer cannot be found. + Contrary to requires-unified-addr-1.c which is valid OpenMP 5.0/5.1/5.2, + this testcase is only valid since OpenMP 5.2. */ + +/* This is kind of a follow-up to the requires-unified-addr-1.c testcase + and PR libgomp/109837 */ + + +#include <assert.h> +#include <omp.h> +#include <stdio.h> +#include <stdlib.h> + +#pragma omp requires unified_address + +#define N 15 + +void +test_device (int dev) +{ + struct st { + int *ptr; + int n; + }; + struct st s; + + s.n = 10; + s.ptr = (int *) omp_target_alloc (sizeof (int)*s.n, dev); + int *ptr1 = (int *) omp_target_alloc (sizeof (int)*N, dev); + assert (s.ptr != NULL); + assert (ptr1 != NULL); + + int q[4] = {1,2,3,4}; + int *qptr; + #pragma omp target enter data map(q) device(device_num: dev) + #pragma omp target data use_device_addr(q) device(device_num: dev) + qptr = q; + + #pragma omp target map(to:s) device(device_num: dev) + for (int i = 0; i < s.n; i++) + s.ptr[i] = 23*i; + + int *ptr2 = &s.ptr[3]; + + /* s.ptr is not mapped (but omp_target_alloc'ed) thus ptr2 shall retain its value. */ + #pragma omp target device(device_num: dev) /* implied: defaultmap(default : pointer) */ + for (int i = 0; i < 4; i++) + *(qptr++) = ptr2[i]; + + #pragma omp target exit data map(q) device(device_num: dev) + for (int i = 0; i < 4; i++) + q[i] = 23 * (i+3); + + /* ptr1 retains the value as it is not mapped (but it is omp_target_alloc'ed). */ + #pragma omp target defaultmap(default : pointer) device(device_num: dev) + for (int i = 0; i < N; i++) + ptr1[i] = 11*i; + + int *ptr3 = (int *) malloc (sizeof (int)*N); + assert (0 == omp_target_memcpy(ptr3, ptr1, N * sizeof(int), 0, 0, + omp_get_initial_device(), dev)); + for (int i = 0; i < N; i++) + assert (ptr3[i] == 11*i); + + free (ptr3); + omp_target_free (ptr1, dev); + omp_target_free (s.ptr, dev); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + if (ntgts) + fprintf (stderr, "Offloading devices exist\n"); /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */ + else + fprintf (stderr, "Only host fallback\n"); /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */ + + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/scan-1.c b/libgomp/testsuite/libgomp.c-c++-common/scan-1.c new file mode 100644 index 0000000..d1951a3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/scan-1.c @@ -0,0 +1,68 @@ +#define N 30 +#define M 3 + +int a[N][M], b[N][M], c[N][M]; + +int +main() +{ + int x, y, shift; + int j = 0; + for (int i = 0; i < N; i++) + { + a[i][0] = (i+1)*32; + a[i][1] = (i+1)*17; + a[i][2] = (i+1)*11; + b[i][0] = (i+1)*7; + b[i][1] = (i+1)*5; + b[i][2] = (i+1)*3; + } + + x = 0; + #pragma omp parallel for simd collapse(2) reduction(inscan,+: x) private(shift) + for (int i = 0; i < N; i++) + for (int j = 0; j < M; j++) + { + x += a[i][j]; + x += b[i][j]; + #pragma omp scan inclusive(x) + shift = i + 29*j; + c[i][j] = x + shift; + } + + y = 0; + for (int i = 0; i < N; i++) + for (int j = 0; j < M; j++) + { + y += a[i][j] + b[i][j]; + if (c[i][j] != y + i + 29*j) + __builtin_abort (); + } + if (x != y) + __builtin_abort (); + + x = 0; + #pragma omp parallel for simd collapse(2) reduction(inscan,+: x) private(shift) + for (int i = 0; i < N; i++) + for (int j = 0; j < M; j++) + { + shift = i + 29*j; + c[i][j] = x + shift; + #pragma omp scan exclusive(x) + x += a[i][j]; + x += b[i][j]; + } + + y = 0; + for (int i = 0; i < N; i++) + for (int j = 0; j < M; j++) + { + if (c[i][j] != y + i + 29*j) + __builtin_abort (); + y += a[i][j] + b[i][j]; + } + if (x != y) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c new file mode 100644 index 0000000..863cf0e --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-3.c @@ -0,0 +1,105 @@ +/* PR middle-end/110270 */ + +/* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2 + semantic, i.e. keeping the pointer value even if not mapped; + before OpenMP 5.0/5.1 required that it is NULL, causing issues + especially with unified-shared memory but also the code below + shows why that's not a good idea. */ + +#include <stdio.h> +#include <stdint.h> +#include <omp.h> + +/* 'unified_address' is required by the OpenMP spec as only then + 'is_device_ptr' can be left out. All our devices support this + while remote offloading would not. However, in practice it is + sufficient that the host and device pointer size is the same + (or the device pointer is smaller) - and then a device pointer is + representable and omp_target_alloc can return a bare device pointer. + + We here assume that this weaker condition holds and do not + require: #pragma omp requires unified_address */ + +void +test_device (int dev) +{ + int *p1 = (int*) 0x12345; + int *p1a = (int*) 0x67890; + int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev); + int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev); + intptr_t ip = (intptr_t) p2; + intptr_t ipa = (intptr_t) p2a; + + int A[3] = {1,2,3}; + int B[5] = {4,5,6,7,8}; + int *p3 = &A[0]; + int *p3a = &B[0]; + + #pragma omp target enter data map(to:A) device(dev) + + #pragma omp target device(dev) /* defaultmap(default:pointer) */ + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1 != 0x12345) __builtin_abort (); + if ((intptr_t) p2 != ip) __builtin_abort (); + for (int i = 0; i < 5; i++) + p2[i] = 13*i; + + for (int i = 0; i < 10; i++) + ((int *)ipa)[i] = 7*i; + + /* OpenMP: Mapped => must point to the corresponding device storage of 'A' */ + if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3) + __builtin_abort (); + p3[0] = -11; p3[1] = -22; p3[2] = -33; + } + #pragma omp target exit data map(from:A) device(dev) + + if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33) + __builtin_abort (); + + // With defaultmap: + + #pragma omp target enter data map(to:B) device(dev) + + #pragma omp target device(dev) defaultmap(default:pointer) + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1a != 0x67890) __builtin_abort (); + if ((intptr_t) p2a != ipa) __builtin_abort (); + + for (int i = 0; i < 5; i++) + ((int *)ip)[i] = 13*i; + + for (int i = 0; i < 10; i++) + p2a[i] = 7*i; + + /* OpenMP: Mapped => must point to the corresponding device storage of 'B' */ + if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8) + __builtin_abort (); + p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88; + } + #pragma omp target exit data map(from:B) device(dev) + + if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88) + __builtin_abort (); + + omp_target_free (p2, dev); + omp_target_free (p2a, dev); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + if (ntgts) + fprintf (stderr, "Offloading devices exist\n"); /* { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } */ + else + fprintf (stderr, "Only host fallback\n"); /* { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } */ + + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c new file mode 100644 index 0000000..d0b0cd1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-implicit-map-4.c @@ -0,0 +1,159 @@ +/* PR middle-end/110270 */ + +/* Same as target-implicit-map-3.c but uses the following requiement + and for not mapping the stack variables 'A' and 'B' (not mapped + but accessible -> USM makes this tested feature even more important.) */ + +#pragma omp requires unified_shared_memory + +/* Ensure that defaultmap(default : pointer) uses correct OpenMP 5.2 + semantic, i.e. keeping the pointer value even if not mapped; + before OpenMP 5.0/5.1 required that it is NULL. */ + +#include <stdio.h> +#include <stdlib.h> +#include <stdint.h> +#include <omp.h> + +void +test_device (int dev) +{ + int *p1 = (int*) 0x12345; + int *p1a = (int*) 0x67890; + int *p2 = (int*) omp_target_alloc (sizeof (int) * 5, dev); + int *p2a = (int*) omp_target_alloc (sizeof (int) * 10, dev); + intptr_t ip = (intptr_t) p2; + intptr_t ipa = (intptr_t) p2a; + + int A[3] = {1,2,3}; + int B[5] = {4,5,6,7,8}; + int *p3 = &A[0]; + int *p3a = &B[0]; + + const omp_alloctrait_t traits[] + = { { omp_atk_alignment, 128 }, + { omp_atk_pool_size, 1024 }}; + omp_allocator_handle_t a = omp_init_allocator (omp_default_mem_space, 2, traits); + + int *p4 = (int*) malloc (sizeof (int) * 5); + int *p4a = (int*) omp_alloc (sizeof (int) * 10, a); + intptr_t ip4 = (intptr_t) p4; + intptr_t ip4a = (intptr_t) p4a; + + for (int i = 0; i < 5; i++) + p4[i] = -31*i; + + for (int i = 0; i < 10; i++) + p4a[i] = -43*i; + + /* Note: 'A' is not mapped but USM accessible. */ + #pragma omp target device(dev) /* defaultmap(default:pointer) */ + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1 != 0x12345) abort (); + if ((intptr_t) p2 != ip) abort (); + for (int i = 0; i < 5; i++) + p2[i] = 13*i; + + for (int i = 0; i < 10; i++) + ((int *)ipa)[i] = 7*i; + + /* OpenMP: Points to 'A'. */ + if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3) + abort (); + p3[0] = -11; p3[1] = -22; p3[2] = -33; + + /* USM accesible allocated host memory. */ + if ((intptr_t) p4 != ip4) + abort (); + for (int i = 0; i < 5; i++) + if (p4[i] != -31*i) + abort (); + for (int i = 0; i < 10; i++) + if (((int *)ip4a)[i] != -43*i) + abort (); + for (int i = 0; i < 5; i++) + p4[i] = 9*i; + for (int i = 0; i < 10; i++) + ((int *)ip4a)[i] = 18*i; + } + + if (p3[0] != -11 || p3[1] != -22 || p3[2] != -33) + abort (); + + for (int i = 0; i < 5; i++) + if (p4[i] != 9*i) + abort (); + for (int i = 0; i < 10; i++) + if (p4a[i] != 18*i) + abort (); + for (int i = 0; i < 5; i++) + p4[i] = -77*i; + for (int i = 0; i < 10; i++) + p4a[i] = -65*i; + + // With defaultmap: + + /* Note: 'B' is not mapped but USM accessible. */ + #pragma omp target device(dev) defaultmap(default:pointer) + { + /* The pointees aren't mapped. */ + /* OpenMP 5.2 -> same value as before the target region. */ + if ((intptr_t) p1a != 0x67890) abort (); + if ((intptr_t) p2a != ipa) abort (); + + for (int i = 0; i < 5; i++) + ((int *)ip)[i] = 13*i; + + for (int i = 0; i < 10; i++) + p2a[i] = 7*i; + + /* USM accesible allocated host memory. */ + if ((intptr_t) p4a != ip4a) abort (); + + /* OpenMP: Points to 'B'. */ + if (p3a[0] != 4 || p3a[1] != 5 || p3a[2] != 6 || p3a[3] != 7 || p3a[4] != 8) + abort (); + p3a[0] = -44; p3a[1] = -55; p3a[2] = -66; p3a[3] = -77; p3a[4] = -88; + + /* USM accesible allocated host memory. */ + if ((intptr_t) p4a != ip4a) + abort (); + for (int i = 0; i < 5; i++) + if (((int *)ip4)[i] != -77*i) + abort (); + for (int i = 0; i < 10; i++) + if (p4a[i] != -65*i) + abort (); + for (int i = 0; i < 5; i++) + p4[i] = 36*i; + for (int i = 0; i < 10; i++) + ((int *)ip4a)[i] = 4*i; + } + + if (p3a[0] != -44 || p3a[1] != -55 || p3a[2] != -66 || p3a[3] != -77 || p3a[4] != -88) + abort (); + + for (int i = 0; i < 5; i++) + if (p4[i] != 36*i) + abort (); + for (int i = 0; i < 10; i++) + if (p4a[i] != 4*i) + abort (); + + omp_target_free (p2, dev); + omp_target_free (p2a, dev); + free (p4); + omp_free (p4a, a); + omp_destroy_allocator (a); +} + +int +main() +{ + int ntgts = omp_get_num_devices(); + for (int i = 0; i <= ntgts; i++) + test_device (i); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c new file mode 100644 index 0000000..5eaa9cd --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-1.c @@ -0,0 +1,40 @@ +#include <stdio.h> + +#define N 100 + +int main (void) +{ + int a[N], b[N], c[N], d[N]; + + for (int i = 0; i < N; i++) { + a[i] = i * 2; + b[i] = i * 3 + 1; + d[i] = i * 5; + } + + #pragma omp target enter data map (alloc: c, d) map(to: a) + #pragma omp target map (present, always, to: d) + for (int i = 0; i < N; i++) + if (d[i] != i * 5) + __builtin_abort (); + + /* a has already been mapped and 'c' allocated so this should be okay. */ + #pragma omp target map (present, to: a) map(present, always, from: c) + for (int i = 0; i < N; i++) + c[i] = a[i]; + + for (int i = 0; i < N; i++) + if (c[i] != i * 2) + __builtin_abort (); + + fprintf (stderr, "CheCKpOInT\n"); + /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ + + /* b has not been allocated, so this should result in an error. */ + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ + /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ + #pragma omp target map (present, to: b) + for (int i = 0; i < N; i++) + c[i] += b[i]; + #pragma omp target exit data map (from: c) +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c new file mode 100644 index 0000000..07ae90b --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-2.c @@ -0,0 +1,30 @@ +#include <stdio.h> + +#define N 100 + +int main (void) +{ + int a[N], b[N], c[N]; + + for (int i = 0; i < N; i++) { + a[i] = i * 2; + b[i] = i * 3 + 1; + } + + #pragma omp target enter data map (alloc: a, c) + /* a and c have already been allocated, so this should be okay. */ + #pragma omp target defaultmap (present) + for (int i = 0; i < N; i++) + c[i] = a[i]; + + fprintf (stderr, "CheCKpOInT\n"); + /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ + + /* b has not been allocated, so this should result in an error. */ + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ + /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ + #pragma omp target defaultmap (present) + for (int i = 0; i < N; i++) + c[i] += b[i]; + #pragma omp target exit data map (from: c) +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c new file mode 100644 index 0000000..582247d --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-present-3.c @@ -0,0 +1,41 @@ +#include <stdio.h> + +#define N 100 + +int main (void) +{ + int a[N], b[N], c[N]; + + for (int i = 0; i < N; i++) { + a[i] = i * 2; + b[i] = i * 3 + 1; + } + + #pragma omp target enter data map (alloc: a, c) + + /* This should work as a has already been allocated. */ + #pragma omp target update to (present: a) + + #pragma omp target map(present,alloc: a, c) + for (int i = 0; i < N; i++) { + if (a[i] != i * 2) + __builtin_abort (); + c[i] = 23*i; + } + + #pragma omp target update from(present : c) + for (int i = 0; i < N; i++) { + if (c[i] != 23*i) + __builtin_abort (); + } + + fprintf (stderr, "CheCKpOInT\n"); + /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ + + /* This should fail as b has not been allocated. */ + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ + /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ + #pragma omp target update to (present: b) + + #pragma omp target exit data map (from: c) +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-1.c b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-1.c new file mode 100644 index 0000000..c3c2109 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-1.c @@ -0,0 +1,198 @@ +/* Check that the nteams ICV is honored. */ +/* PR libgomp/109875 */ + +/* This base version of testcases is supposed to be run with all + OMP_NUM_TEAMS* env vars being unset. + + The variants teams-nteams-icv-{2,3,4}.c test it by setting the + various OMP_NUM_TEAMS* env vars and #define MY_... for checking. + + Currently, only <num> 0,1,2 is supported for the envar via #define + and with remote execution, dg-set-target-env-var does not work with + DejaGNU, hence, gcc/testsuite/lib/gcc-dg.exp marks those tests as + UNSUPPORTED. */ + +#define MY_MAX_DEVICES 3 + +/* OpenMP currently has: + - nteams-var ICV is initialized to 0; one ICV per device + - OMP_NUM_TEAMS(_DEV(_<dev-num>)) overrides it + OMP_NUM_TEAMS_ALL overrides it + - Number of teams is: + -> the value specific by num_teams([lower:]upper) + with lower := upper if unspecified + -> Otherwise, if nteams-var ICV > 0, #teams <= nteams-var ICV + -> Otherwise, if nteams-var ICV <= 0, #teams > 1 + GCC uses 3 as default on the host and 1 for host fallback. + For offloading, it is device specific >> 1. */ + +#include <omp.h> + +int +main () +{ + int num_teams_env = -1, num_teams_env_dev = -1; + int num_teams_env_devs[MY_MAX_DEVICES]; + +#ifdef MY_OMP_NUM_TEAMS_ALL + num_teams_env = num_teams_env_dev = MY_OMP_NUM_TEAMS_ALL; +#endif + +#ifdef MY_OMP_NUM_TEAMS + num_teams_env = MY_OMP_NUM_TEAMS; +#endif + +#ifdef MY_OMP_NUM_TEAMS_DEV + num_teams_env_dev = MY_OMP_NUM_TEAMS_DEV; +#endif + +#if MY_MAX_DEVICES != 3 + #error "Currently strictly assuming MY_MAX_DEVICES = 3" +#endif + +#if defined(MY_OMP_NUM_TEAMS_DEV_4) || defined(MY_OMP_NUM_TEAMS_DEV_5) + #error "Currently strictly assuming MY_MAX_DEVICES = 3" +#endif + +#ifdef MY_OMP_NUM_TEAMS_DEV_0 + num_teams_env_devs[0] = MY_OMP_NUM_TEAMS_DEV_0; +#else + num_teams_env_devs[0] = num_teams_env_dev; +#endif + +#ifdef MY_OMP_NUM_TEAMS_DEV_1 + num_teams_env_devs[1] = MY_OMP_NUM_TEAMS_DEV_1; +#else + num_teams_env_devs[1] = num_teams_env_dev; +#endif + +#ifdef MY_OMP_NUM_TEAMS_DEV_2 + num_teams_env_devs[2] = MY_OMP_NUM_TEAMS_DEV_2; +#else + num_teams_env_devs[2] = num_teams_env_dev; +#endif + + /* Check that the number of teams (initial device and in target) is + >= 1 and, if omp_get_max_teams() > 0, it does not + exceed omp_get_max_teams (). */ + + int nteams, num_teams; + + /* Assume that omp_get_max_teams (); returns the ICV, i.e. 0 as default init + and not the number of teams that would be run; hence: '>='. */ + nteams = omp_get_max_teams (); + if (nteams < 0 || (num_teams_env >= 0 && nteams != num_teams_env)) + __builtin_abort (); + num_teams = -1; + + #pragma omp teams + if (omp_get_team_num () == 0) + num_teams = omp_get_num_teams (); + if (num_teams < 1 || (nteams > 0 && num_teams > nteams)) + __builtin_abort (); + + /* GCC hard codes 3 teams - check for it. */ + if (nteams <= 0 && num_teams != 3) + __builtin_abort (); + + /* For each device, including host fallback. */ + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + { + int num_teams_icv = num_teams_env_dev; + if (dev == omp_get_num_devices ()) + num_teams_icv = num_teams_env; + else if (dev < MY_MAX_DEVICES) + num_teams_icv = num_teams_env_devs[dev]; + + nteams = -1; + #pragma omp target device(dev) map(from: nteams) + nteams = omp_get_max_teams (); + if (nteams < 0 || (num_teams_icv >= 0 && nteams != num_teams_icv)) + __builtin_abort (); + + num_teams = -1; + #pragma omp target teams device(dev) map(from: num_teams) + if (omp_get_team_num () == 0) + num_teams = omp_get_num_teams (); + + if (num_teams < 1 || (nteams > 0 && num_teams > nteams)) + __builtin_abort (); + + /* GCC hard codes 1 team for host fallback - check for it. */ + if (dev == omp_get_num_devices () && num_teams != 1) + __builtin_abort (); + } + + /* Now set the nteams-var ICV and check that omp_get_max_teams() + returns the set value and that the following holds: + num_teams >= 1 and num_teams <= nteams-var ICV. + + Additionally, implementation defined, assume: + - num_teams == (not '<=') nteams-var ICV, except: + - num_teams == 1 for host fallback. */ + + omp_set_num_teams (5); + + nteams = omp_get_max_teams (); + if (nteams != 5) + __builtin_abort (); + num_teams = -1; + + #pragma omp teams + if (omp_get_team_num () == 0) + num_teams = omp_get_num_teams (); + if (num_teams != 5) + __builtin_abort (); + + /* For each device, including host fallback. */ + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + { + #pragma omp target device(dev) firstprivate(dev) + omp_set_num_teams (7 + dev); + + #pragma omp target device(dev) map(from: nteams) + nteams = omp_get_max_teams (); + if (nteams != 7 + dev) + __builtin_abort (); + + num_teams = -1; + #pragma omp target teams device(dev) map(from: num_teams) + if (omp_get_team_num () == 0) + num_teams = omp_get_num_teams (); + + if (dev == omp_get_num_devices ()) + { + if (num_teams != 1) + __builtin_abort (); + } + else + { + if (num_teams != 7 + dev) + __builtin_abort (); + } + } + + /* Now use the num_teams clause explicitly. */ + + num_teams = -1; + #pragma omp teams num_teams(6) + if (omp_get_team_num () == 0) + num_teams = omp_get_num_teams (); + if (num_teams != 6) + __builtin_abort (); + + /* For each device, including host fallback. */ + for (int dev = 0; dev <= omp_get_num_devices (); dev++) + { + num_teams = -1; + #pragma omp target teams device(dev) map(from: num_teams) num_teams(dev+3) + if (omp_get_team_num () == 0) + num_teams = omp_get_num_teams (); + + /* This must match the set value, also with host fallback. */ + if (num_teams != 3 + dev) + __builtin_abort (); + } + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-2.c b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-2.c new file mode 100644 index 0000000..7bd80de --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-2.c @@ -0,0 +1,8 @@ +/* PR libgomp/109875 */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL 9 } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV 7 } */ + +#define MY_OMP_NUM_TEAMS_ALL 9 +#define MY_OMP_NUM_TEAMS_DEV 7 + +#include "teams-nteams-icv-1.c" diff --git a/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-3.c b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-3.c new file mode 100644 index 0000000..10a1cdc --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-3.c @@ -0,0 +1,8 @@ +/* PR libgomp/109875 */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL 7 } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS 8 } */ + +#define MY_OMP_NUM_TEAMS_ALL 7 +#define MY_OMP_NUM_TEAMS 8 + +#include "teams-nteams-icv-1.c" diff --git a/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-4.c b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-4.c new file mode 100644 index 0000000..c5d6577 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/teams-nteams-icv-4.c @@ -0,0 +1,14 @@ +/* PR libgomp/109875 */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_ALL 7 } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS 4 } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV 8 } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_0 5 } */ +/* { dg-set-target-env-var OMP_NUM_TEAMS_DEV_1 11 } */ + +#define MY_OMP_NUM_TEAMS_ALL 7 +#define MY_OMP_NUM_TEAMS 4 +#define MY_OMP_NUM_TEAMS_DEV 8 +#define MY_OMP_NUM_TEAMS_DEV_0 5 +#define MY_OMP_NUM_TEAMS_DEV_1 11 + +#include "teams-nteams-icv-1.c" diff --git a/libgomp/testsuite/libgomp.c/allocator-1.c b/libgomp/testsuite/libgomp.c/allocator-1.c new file mode 100644 index 0000000..2757792 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocator-1.c @@ -0,0 +1,15 @@ +/* { dg-set-target-env-var OMP_ALLOCATOR "omp_large_cap_mem_alloc" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".\\\[host\\\] OMP_ALLOCATOR = 'omp_large_cap_mem_alloc'.*" } */ + +#include <omp.h> + +int +main () +{ + omp_allocator_handle_t m = omp_get_default_allocator (); + if (m != omp_large_cap_mem_alloc) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocator-2.c b/libgomp/testsuite/libgomp.c/allocator-2.c new file mode 100644 index 0000000..ac680e5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocator-2.c @@ -0,0 +1,17 @@ +/* { dg-set-target-env-var OMP_ALLOCATOR "omp_large_cap_mem_space" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* Expect omp_large_cap_mem_alloc as allocator for omp_large_cap_mem_space. */ +/* { dg-output ".\\\[host\\\] OMP_ALLOCATOR = 'omp_large_cap_mem_alloc'.*" } */ +#include <omp.h> + +int +main () +{ + omp_allocator_handle_t m = omp_get_default_allocator (); + /* Without traits, omp_large_cap_mem_space implies + omp_large_cap_mem_alloc. */ + if (m != omp_large_cap_mem_alloc) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocator-3.c b/libgomp/testsuite/libgomp.c/allocator-3.c new file mode 100644 index 0000000..a28f4e75 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocator-3.c @@ -0,0 +1,27 @@ +/* { dg-set-target-env-var OMP_ALLOCATOR " omp_default_mem_space:alignment=512,pinned=false,access=all " } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* We copied the environment string; hence, it may contain white space. */ +/* { dg-output ".\\\[host\\\] OMP_ALLOCATOR = ' omp_default_mem_space:alignment=512,pinned=false,access=all '.*" } */ + +#include <stdint.h> +#include <omp.h> + +int +main () +{ + int *a, *b; + a = omp_alloc (sizeof (int) * 1024, omp_null_allocator); + + omp_allocator_handle_t m = omp_get_default_allocator (); + b = omp_alloc (sizeof (int) * 1024, m); + + if ((uintptr_t) a % 512 != 0) + __builtin_abort (); + + if ((uintptr_t) b % 512 != 0) + __builtin_abort (); + omp_free (a, omp_null_allocator); + omp_free (b, m); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocator-4.c b/libgomp/testsuite/libgomp.c/allocator-4.c new file mode 100644 index 0000000..e5d0b09 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocator-4.c @@ -0,0 +1,15 @@ +/* { dg-set-target-env-var OMP_ALLOCATOR "omp_const_mem_space:alignment=3,pinned=" } */ + +/* { dg-output ".*libgomp: Missing value at the end of environment variable OMP_ALLOCATOR=omp_const_mem_space:alignment=3,pinned=.*" } */ +/* OMP_ALLOCATOR syntax error -> use omp_default_mem_alloc. */ + +#include <omp.h> + +int +main () +{ + omp_allocator_handle_t m = omp_get_default_allocator (); + if (m != omp_default_mem_alloc) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocator-5.c b/libgomp/testsuite/libgomp.c/allocator-5.c new file mode 100644 index 0000000..0f900cf --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocator-5.c @@ -0,0 +1,15 @@ +/* { dg-set-target-env-var OMP_ALLOCATOR "omp_const_mem_space:access=none,pinned=false" } */ + +/* { dg-output ".*libgomp: Invalid value for environment variable OMP_ALLOCATOR when parsing: none,pinned=false.*" } */ +/* OMP_ALLOCATOR syntax error -> use omp_default_mem_alloc. */ + +#include <omp.h> + +int +main () +{ + omp_allocator_handle_t m = omp_get_default_allocator (); + if (m != omp_default_mem_alloc) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/allocator-6.c b/libgomp/testsuite/libgomp.c/allocator-6.c new file mode 100644 index 0000000..7c99e0e --- /dev/null +++ b/libgomp/testsuite/libgomp.c/allocator-6.c @@ -0,0 +1,15 @@ +/* { dg-set-target-env-var OMP_ALLOCATOR "omp_default_mem_space:alignment=3" } */ + +/* { dg-output ".*libgomp: Allocator of environment variable OMP_ALLOCATOR cannot be created, using omp_default_mem_alloc instead.*" } */ +/* OMP_ALLOCATOR's alignment is not power of 2 -> use omp_default_mem_alloc. */ + +#include <omp.h> + +int +main () +{ + omp_allocator_handle_t m = omp_get_default_allocator (); + if (m != omp_default_mem_alloc) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/c.exp b/libgomp/testsuite/libgomp.c/c.exp index 31bdd57..aae2824 100644 --- a/libgomp/testsuite/libgomp.c/c.exp +++ b/libgomp/testsuite/libgomp.c/c.exp @@ -1,17 +1,8 @@ -if [info exists lang_library_path] then { - unset lang_library_path - unset lang_link_flags -} -if [info exists lang_test_file] then { - unset lang_test_file -} -if [info exists lang_include_flags] then { - unset lang_include_flags -} - load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp +lappend ALWAYS_CFLAGS "compiler=$GCC_UNDER_TEST" + # If a testcase doesn't have special options, use these. if ![info exists DEFAULT_CFLAGS] then { set DEFAULT_CFLAGS "-O2" diff --git a/libgomp/testsuite/libgomp.c/scan-23.c b/libgomp/testsuite/libgomp.c/scan-23.c new file mode 100644 index 0000000..e7681ad --- /dev/null +++ b/libgomp/testsuite/libgomp.c/scan-23.c @@ -0,0 +1,121 @@ +/* { dg-require-effective-target size32plus } */ +/* Same as scan-9.c, instead of using { ... } it simply uses multiple + executable stmt before 'omp scan'. */ + +extern void abort (void); +int r, a[1024], b[1024], x, y, z; + +__attribute__((noipa)) void +foo (int *a, int *b) +{ + #pragma omp for reduction (inscan, +:r) lastprivate (conditional: z) firstprivate (x) private (y) + for (int i = 0; i < 1024; i++) + { + y = a[i]; + r += y + x + 12; + #pragma omp scan inclusive(r) + b[i] = r; + if ((i & 1) == 0 && i < 937) + z = r; + } +} + +__attribute__((noipa)) int +bar (void) +{ + int s = 0; + #pragma omp parallel + #pragma omp for reduction (inscan, +:s) firstprivate (x) private (y) lastprivate (z) + for (int i = 0; i < 1024; i++) + { + y = 2 * a[i]; s += y; z = y; + #pragma omp scan inclusive(s) + y = s; b[i] = y + x + 12; + } + return s; +} + +__attribute__((noipa)) void +baz (int *a, int *b) +{ + #pragma omp parallel for reduction (inscan, +:r) firstprivate (x) lastprivate (x) + for (int i = 0; i < 1024; i++) + { + r += a[i] + x + 12; if (i == 1023) x = 29; + #pragma omp scan inclusive(r) + b[i] = r; + } +} + +__attribute__((noipa)) int +qux (void) +{ + int s = 0; + #pragma omp parallel for reduction (inscan, +:s) lastprivate (conditional: x, y) + for (int i = 0; i < 1024; i++) + { + s += 2 * a[i]; if ((a[i] & 1) == 1 && i < 825) x = a[i]; + #pragma omp scan inclusive(s) + b[i] = s; if ((a[i] & 1) == 0 && i < 829) y = a[i]; + } + return s; +} + +int +main () +{ + int s = 0; + x = -12; + for (int i = 0; i < 1024; ++i) + { + a[i] = i; + b[i] = -1; + asm ("" : "+g" (i)); + } + #pragma omp parallel + foo (a, b); + if (r != 1024 * 1023 / 2 || x != -12 || z != b[936]) + abort (); + for (int i = 0; i < 1024; ++i) + { + s += i; + if (b[i] != s) + abort (); + else + b[i] = 25; + } + if (bar () != 1024 * 1023 || x != -12 || z != 2 * 1023) + abort (); + s = 0; + for (int i = 0; i < 1024; ++i) + { + s += 2 * i; + if (b[i] != s) + abort (); + else + b[i] = -1; + } + r = 0; + baz (a, b); + if (r != 1024 * 1023 / 2 || x != 29) + abort (); + s = 0; + for (int i = 0; i < 1024; ++i) + { + s += i; + if (b[i] != s) + abort (); + else + b[i] = -25; + } + if (qux () != 1024 * 1023 || x != 823 || y != 828) + abort (); + s = 0; + for (int i = 0; i < 1024; ++i) + { + s += 2 * i; + if (b[i] != s) + abort (); + } + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/simd-math-1.c b/libgomp/testsuite/libgomp.c/simd-math-1.c index 947bf60..dd2077c 100644 --- a/libgomp/testsuite/libgomp.c/simd-math-1.c +++ b/libgomp/testsuite/libgomp.c/simd-math-1.c @@ -4,7 +4,6 @@ /* { dg-do run } */ /* { dg-options "-O2 -ftree-vectorize -fno-math-errno" } */ /* { dg-additional-options -foffload-options=amdgcn-amdhsa=-mstack-size=3000000 { target offload_target_amdgcn } } */ -/* { dg-additional-options -foffload-options=-lm } */ #undef PRINT_RESULT #define VERBOSE 0 diff --git a/libgomp/testsuite/libgomp.c/target-19.c b/libgomp/testsuite/libgomp.c/target-19.c index 2505caf..dac7c56 100644 --- a/libgomp/testsuite/libgomp.c/target-19.c +++ b/libgomp/testsuite/libgomp.c/target-19.c @@ -20,7 +20,9 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -37,7 +39,9 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -55,7 +59,9 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; if (sep) { - if (q != (int *) 0 || r != (int *) 0) + /* Since OpenMP 5.2, 'q'/'r' are no longer set to NULL if pointing to + unmapped storage. */ + if (q == (int *) 0 || r == (int *) 0) err = 1; } else if (p + 8 != q || r != s) @@ -91,7 +97,8 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -110,7 +117,8 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) @@ -130,7 +138,8 @@ foo (int *p, int *q, int *r, int n, int m) err = 1; else if (sep) { - if (r != (int *) 0) + /* Since OpenMP 5.2, 'r' is no longer set to NULL if *r is unmapped.*/ + if (r == (int *) 0) err = 1; } else if (r != q + 1) diff --git a/libgomp/testsuite/libgomp.c/target-48.c b/libgomp/testsuite/libgomp.c/target-48.c new file mode 100644 index 0000000..8e95c1c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-48.c @@ -0,0 +1,31 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices; + omp_invalid_device == -4 with GCC. */ + +/* { dg-do run { target { ! offload_device } } } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '-4'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +#include <omp.h> + +int +main () +{ + if (omp_get_default_device () != omp_invalid_device) + __builtin_abort (); + + omp_set_default_device (omp_initial_device); + + /* The spec is a bit unclear whether the line above sets the device number + (a) to -1 (= omp_initial_device) or + (b) to omp_get_initial_device() == omp_get_num_devices(). Therefore, + we accept either value. */ + + if (omp_get_default_device() != omp_get_initial_device() + && omp_get_default_device() != omp_initial_device) + __builtin_abort (); + + omp_display_env (0); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-49.c b/libgomp/testsuite/libgomp.c/target-49.c new file mode 100644 index 0000000..970cb91 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-49.c @@ -0,0 +1,18 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices, + which is enforced by using -foffload=disable. */ + +/* { dg-do run } */ +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* See comment in target-50.c/target-50.c for why default-device-var can be '0'. */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '-4'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target { ! offload_device } } } */ +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" { target offload_device } } */ + +int +main () +{ + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-50.c b/libgomp/testsuite/libgomp.c/target-50.c new file mode 100644 index 0000000..6f15569 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-50.c @@ -0,0 +1,27 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices; + here with using -foffload=disable. + As default-device-var is set to 0 (= host in this case), it should not fail. */ + +/* Note that -foffload=disable will still find devices on the system and only + when trying to use them, it will fail as no binary data has been produced. + The "target offload_device" case is checked for in 'target-50a.c'. */ + +/* { dg-do run { target { ! offload_device } } } */ + +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "0" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-50a.c b/libgomp/testsuite/libgomp.c/target-50a.c new file mode 100644 index 0000000..0835cb5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-50a.c @@ -0,0 +1,43 @@ +/* Check OMP_TARGET_OFFLOAD on systems with non-host devices but no executable + code due to -foffload=disable. + + Note: While one might expect that -foffload=disable implies no non-host + devices, libgomp actually detects the devices and only fails when trying to + run as no executable code is availale for that device. + (Without MANDATORY it simply uses host fallback, which should usually be fine + but might have issues in corner cases.) + + We have default-device-var = 0 (default but also explicitly set), which will + fail at runtime. For -foffload=disable without non-host devices, see + target-50.c testcase. */ + +/* { dg-do run { target offload_device } } */ + +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DEFAULT_DEVICE "0" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +#include <omp.h> + +int +main () +{ + int x; + /* We know that there are non-host devices. With GCC, we still find them as + available devices, hence, check for it. */ + if (omp_get_num_devices() <= 0) + __builtin_abort (); + + /* But due to -foffload=disable, there are no binary code for (default) device '0' */ + + /* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but device cannot be used for offloading.*" } */ + /* { dg-shouldfail "OMP_TARGET_OFFLOAD=mandatory and no binary code for a non-host device" } */ + #pragma omp target map(tofrom:x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-51.c b/libgomp/testsuite/libgomp.c/target-51.c new file mode 100644 index 0000000..7ff8122 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-51.c @@ -0,0 +1,23 @@ +/* Check OMP_TARGET_OFFLOAD on systems with no available non-host devices, + which is enforced by using -foffload=disable. */ + +/* { dg-do run } */ +/* { dg-additional-options "-foffload=disable" } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ + +/* { dg-shouldfail "OMP_TARGET_OFFLOAD=mandatory and no available device" } */ + +/* See comment in target-50.c/target-50.c for why the output differs. */ + +/* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but .*" } */ + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-52.c b/libgomp/testsuite/libgomp.c/target-52.c new file mode 100644 index 0000000..809380c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-52.c @@ -0,0 +1,25 @@ +/* Only run this with available non-host devices; in that case, GCC sets + the default-device-var to 0. */ + +/* { dg-do run { target { offload_device } } } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'MANDATORY'.*" } */ + +#include <omp.h> + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5 + omp_is_initial_device (); + + if (x != 5) + __builtin_abort (); + + if (0 != omp_get_default_device()) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-53.c b/libgomp/testsuite/libgomp.c/target-53.c new file mode 100644 index 0000000..866e896 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-53.c @@ -0,0 +1,22 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "disabled" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '\[0-9\]+'.*OMP_TARGET_OFFLOAD = 'DISABLED'.*" } */ + +#include <omp.h> + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5 + omp_is_initial_device (); + + if (x != 5+1) + __builtin_abort (); + + if (omp_get_default_device() != omp_get_initial_device()) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-54.c b/libgomp/testsuite/libgomp.c/target-54.c new file mode 100644 index 0000000..bc4e69b --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-54.c @@ -0,0 +1,20 @@ +/* { dg-do run } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "default" } */ +/* { dg-set-target-env-var OMP_DISPLAY_ENV "true" } */ + +/* { dg-output ".*OMP_DEFAULT_DEVICE = '0'.*OMP_TARGET_OFFLOAD = 'DEFAULT'.*" } */ + +#include <omp.h> + +int +main () +{ + int x; + #pragma omp target map(tofrom:x) + x = 5 + omp_is_initial_device (); + + if (x != 5 + (omp_get_default_device() == omp_get_initial_device())) + __builtin_abort (); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-55.c b/libgomp/testsuite/libgomp.c/target-55.c new file mode 100644 index 0000000..1314b3c --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-55.c @@ -0,0 +1,20 @@ +/* { dg-do run { target { offload_device } } } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ + +/* Should pass - see target-55a.c for !offload_device */ + +/* Check OMP_TARGET_OFFLOAD - it shall run on systems with offloading + devices available and fail otherwise. Note that this did always + fail - as the device handling wasn't initialized before doing the + mandatory checking. */ + +int +main () +{ + int x = 1; + #pragma omp target map(tofrom: x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.c/target-55a.c b/libgomp/testsuite/libgomp.c/target-55a.c new file mode 100644 index 0000000..53978c3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c/target-55a.c @@ -0,0 +1,23 @@ +/* { dg-do run { target { ! offload_device } } } */ +/* { dg-set-target-env-var OMP_TARGET_OFFLOAD "mandatory" } */ + +/* Should fail - see target-55a.c for offload_device */ + +/* { dg-shouldfail "omp_invalid_device" } */ +/* { dg-output ".*libgomp: OMP_TARGET_OFFLOAD is set to MANDATORY, but only the host device is available.*" } */ + +/* Check OMP_TARGET_OFFLOAD - it shall run on systems with offloading + devices available and fail otherwise. Note that this did always + fail - as the device handling wasn't initialized before doing the + mandatory checking. */ + +int +main () +{ + int x = 1; + #pragma omp target map(tofrom: x) + x = 5; + if (x != 5) + __builtin_abort (); + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/allocate-4.f90 b/libgomp/testsuite/libgomp.fortran/allocate-4.f90 index ddb507b..1f833b6 100644 --- a/libgomp/testsuite/libgomp.fortran/allocate-4.f90 +++ b/libgomp/testsuite/libgomp.fortran/allocate-4.f90 @@ -16,27 +16,27 @@ integer, parameter :: cnst(2) = [64, 101] !$omp parallel allocate( allocator (omp_high_bw_mem_alloc) : x) firstprivate(x) ! { dg-error "Expected integer expression of the 'omp_allocator_handle_kind' kind" } !$omp end parallel -!$omp parallel allocate( align (q) : x) firstprivate(x) ! { dg-error "32:ALIGN modifier requires at \\(1\\) a scalar positive constant integer alignment expression that is a power of two" } +!$omp parallel allocate( align (q) : x) firstprivate(x) ! { dg-error "32:ALIGN requires a scalar positive constant integer alignment expression at \\(1\\) that is a power of two" } !$omp end parallel !$omp parallel allocate( align (32) : x) firstprivate(x) ! OK !$omp end parallel -!$omp parallel allocate( align(q) : x) firstprivate(x) ! { dg-error "31:ALIGN modifier requires at \\(1\\) a scalar positive constant integer alignment expression that is a power of two" } +!$omp parallel allocate( align(q) : x) firstprivate(x) ! { dg-error "31:ALIGN requires a scalar positive constant integer alignment expression at \\(1\\) that is a power of two" } !$omp end parallel !$omp parallel allocate( align(cnst(1)) : x ) firstprivate(x) ! OK !$omp end parallel -!$omp parallel allocate( align(cnst(2)) : x) firstprivate(x) ! { dg-error "31:ALIGN modifier requires at \\(1\\) a scalar positive constant integer alignment expression that is a power of two" } +!$omp parallel allocate( align(cnst(2)) : x) firstprivate(x) ! { dg-error "31:ALIGN requires a scalar positive constant integer alignment expression at \\(1\\) that is a power of two" } !$omp end parallel -!$omp parallel allocate( align( 31) :x) firstprivate(x) ! { dg-error "32:ALIGN modifier requires at \\(1\\) a scalar positive constant integer alignment expression that is a power of two" } +!$omp parallel allocate( align( 31) :x) firstprivate(x) ! { dg-error "32:ALIGN requires a scalar positive constant integer alignment expression at \\(1\\) that is a power of two" } !$omp end parallel -!$omp parallel allocate( align (32.0): x) firstprivate(x) ! { dg-error "32:ALIGN modifier requires at \\(1\\) a scalar positive constant integer alignment expression that is a power of two" } +!$omp parallel allocate( align (32.0): x) firstprivate(x) ! { dg-error "32:ALIGN requires a scalar positive constant integer alignment expression at \\(1\\) that is a power of two" } !$omp end parallel -!$omp parallel allocate( align(cnst ) : x ) firstprivate(x) ! { dg-error "31:ALIGN modifier requires at \\(1\\) a scalar positive constant integer alignment expression that is a power of two" } +!$omp parallel allocate( align(cnst ) : x ) firstprivate(x) ! { dg-error "31:ALIGN requires a scalar positive constant integer alignment expression at \\(1\\) that is a power of two" } !$omp end parallel end diff --git a/libgomp/testsuite/libgomp.fortran/fortran-torture_execute_math.f90 b/libgomp/testsuite/libgomp.fortran/fortran-torture_execute_math.f90 new file mode 100644 index 0000000..2d0caa6 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/fortran-torture_execute_math.f90 @@ -0,0 +1,3 @@ +! { dg-do run } + +include '../../../gcc/testsuite/gfortran.fortran-torture/execute/math.f90' diff --git a/libgomp/testsuite/libgomp.fortran/fortran.exp b/libgomp/testsuite/libgomp.fortran/fortran.exp index eb70131..32e4bb2 100644 --- a/libgomp/testsuite/libgomp.fortran/fortran.exp +++ b/libgomp/testsuite/libgomp.fortran/fortran.exp @@ -2,18 +2,33 @@ load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp load_gcc_lib gfortran-dg.exp -global shlib_ext -global ALWAYS_CFLAGS +if { $blddir != "" } { + set libgfortran_library_path "../libgfortran/.libs" + set shlib_ext [get_shlib_extension] + if { ![file exists "${blddir}/${libgfortran_library_path}/libgfortran.a"] + && ![file exists "${blddir}/${libgfortran_library_path}/libgfortran.${shlib_ext}"] } { + verbose -log "No libgfortran library found, will not execute fortran tests" + unset libgfortran_library_path + return + } + lappend lang_library_paths $libgfortran_library_path -set shlib_ext [get_shlib_extension] -set lang_library_path "../libgfortran/.libs" -set lang_link_flags "-lgfortran -foffload=-lgfortran" -if [info exists lang_include_flags] then { - unset lang_include_flags + set libquadmath_library_path "../libquadmath/.libs" + if { [file exists "${blddir}/${libquadmath_library_path}/libquadmath.a"] + || [file exists "${blddir}/${libquadmath_library_path}/libquadmath.${shlib_ext}"] } { + lappend lang_library_paths $libquadmath_library_path + } else { + set libquadmath_library_path "" + } +} elseif { ![info exists GFORTRAN_UNDER_TEST] } { + verbose -log "GFORTRAN_UNDER_TEST not defined, will not execute fortran tests" + return } -set lang_test_file_found 0 -set quadmath_library_path "../libquadmath/.libs" - +if { $blddir != "" } { + set lang_source_re {^.*\.[fF](|90|95|03|08)$} + set lang_include_flags "-fintrinsic-modules-path=${blddir}" +} +lappend ALWAYS_CFLAGS "compiler=$GFORTRAN_UNDER_TEST" # Initialize dg. dg-init @@ -21,58 +36,31 @@ dg-init # Turn on OpenMP. lappend ALWAYS_CFLAGS "additional_flags=-fopenmp" +# Gather a list of all tests. +set tests [lsort [find $srcdir/$subdir *.\[fF\]{,90,95,03,08}]] + +set ld_library_path $always_ld_library_path if { $blddir != "" } { - set lang_source_re {^.*\.[fF](|90|95|03|08)$} - set lang_include_flags "-fintrinsic-modules-path=${blddir}" - # Look for a static libgfortran first. - if [file exists "${blddir}/${lang_library_path}/libgfortran.a"] { - set lang_test_file "${lang_library_path}/libgfortran.a" - set lang_test_file_found 1 - # We may have a shared only build, so look for a shared libgfortran. - } elseif [file exists "${blddir}/${lang_library_path}/libgfortran.${shlib_ext}"] { - set lang_test_file "${lang_library_path}/libgfortran.${shlib_ext}" - set lang_test_file_found 1 - } else { - puts "No libgfortran library found, will not execute fortran tests" + append ld_library_path ":${blddir}/${libgfortran_library_path}" + + if { $libquadmath_library_path != "" } { + append ld_library_path ":${blddir}/${libquadmath_library_path}" } -} elseif [info exists GFORTRAN_UNDER_TEST] { - set lang_test_file_found 1 - # Needs to exist for libgomp.exp. - set lang_test_file "" -} else { - puts "GFORTRAN_UNDER_TEST not defined, will not execute fortran tests" } +append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] +set_ld_library_path_env_vars -if { $lang_test_file_found } { - # Gather a list of all tests. - set tests [lsort [find $srcdir/$subdir *.\[fF\]{,90,95,03,08}]] +# For Fortran we're doing torture testing, as Fortran has far more tests +# with arrays etc. that testing just -O0 or -O2 is insufficient, that is +# typically not the case for C/C++. +gfortran-dg-runtest $tests "" "" - if { $blddir != "" } { - if { [file exists "${blddir}/${quadmath_library_path}/libquadmath.a"] - || [file exists "${blddir}/${quadmath_library_path}/libquadmath.${shlib_ext}"] } { - lappend ALWAYS_CFLAGS "ldflags=-L${blddir}/${quadmath_library_path}/" - # Allow for spec subsitution. - lappend ALWAYS_CFLAGS "additional_flags=-B${blddir}/${quadmath_library_path}/" - set ld_library_path "$always_ld_library_path:${blddir}/${lang_library_path}:${blddir}/${quadmath_library_path}" - append lang_link_flags " -lquadmath" - } else { - set ld_library_path "$always_ld_library_path:${blddir}/${lang_library_path}" - } - } else { - set ld_library_path "$always_ld_library_path" - if { [check_no_compiler_messages has_libquadmath executable { - int main() {return 0;} - } "-lgfortran -lquadmath"] } then { - append lang_link_flags " -lquadmath" - } - } - append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] - set_ld_library_path_env_vars - - # For Fortran we're doing torture testing, as Fortran has far more tests - # with arrays etc. that testing just -O0 or -O2 is insufficient, that is - # typically not the case for C/C++. - gfortran-dg-runtest $tests "" "" +if { $blddir != "" } { + unset lang_source_re + unset lang_include_flags + unset libgfortran_library_path + unset libquadmath_library_path + unset lang_library_paths } # All done. diff --git a/libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90 b/libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90 new file mode 100644 index 0000000..f5a5adf --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/requires-unified-addr-1.f90 @@ -0,0 +1,111 @@ +! PR libgomp/109837 + +program main + use iso_c_binding + use iso_fortran_env + use omp_lib + implicit none (external, type) + !$omp requires unified_address + + integer(c_intptr_t), parameter :: N = 15 + integer :: i, ntgts + + ntgts = omp_get_num_devices(); + if (ntgts > 0) then + write (ERROR_UNIT, '(a)') "Offloading devices exist" ! { dg-output "Offloading devices exist(\n|\r\n|\r)" { target offload_device } } + else + write (ERROR_UNIT, '(a)') "Only host fallback" ! { dg-output "Only host fallback(\n|\r\n|\r)" { target { ! offload_device } } } + endif + + do i = 0, ntgts + call test_device (i); + end do + +contains + + subroutine test_device (dev) + integer, value, intent(in) :: dev + + type t + integer(c_intptr_t) :: n, m + integer, pointer :: fptr(:) + type(c_ptr) :: cptr + end type t + type(t) :: s + type(c_ptr) :: cptr, qptr, cptr2, cptr2a + integer, target :: q(4) + integer, pointer :: fptr(:) + integer(c_intptr_t) :: i + + s%n = 10; + s%m = 23; + s%cptr = omp_target_alloc (s%n * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, dev); + cptr = omp_target_alloc (s%m * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, dev); + if (.not. c_associated(s%cptr)) stop 1 + if (.not. c_associated(cptr)) stop 2 + call c_f_pointer (cptr, s%fptr, [s%m]) + + cptr = omp_target_alloc (N * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, dev); + if (.not. c_associated(cptr)) stop 3 + + q = [1, 2, 3, 4] + !$omp target enter data map(q) device(device_num: dev) + !$omp target data use_device_addr(q) device(device_num: dev) + qptr = c_loc(q) + !$omp end target data + + !$omp target map(to:s) device(device_num: dev) + block + integer, pointer :: iptr(:) + call c_f_pointer(s%cptr, iptr, [s%n]) + do i = 1, s%n + iptr(i) = 23 * int(i) + end do + do i = 1, s%m + s%fptr(i) = 35 * int(i) + end do + end block + + cptr2 = c_loc(s%fptr(4)) + cptr2a = s%cptr + + !$omp target firstprivate(qptr) map(tofrom: cptr2) map(to :cptr2a) device(device_num: dev) + block + integer, pointer :: iptr(:), iptr2(:), qvar(:) + call c_f_pointer(cptr2, iptr, [4]) + call c_f_pointer(cptr2a, iptr2, [4]) + call c_f_pointer(qptr, qvar, [4]) + qvar = iptr + iptr2 + end block + + !$omp target exit data map(q) device(device_num: dev) + do i = 1, 4 + if (q(i) /= 23 * int(i) + 35 * (int(i) + 4 - 1)) stop 4 + end do + + !$omp target map(to: cptr) device(device_num: dev) + block + integer, pointer :: p(:) + call c_f_pointer(cptr, p, [N]) + do i = 1, N + p(i) = 11 * int(i) + end do + end block + + allocate(fptr(N)) + if (0 /= omp_target_memcpy (c_loc(fptr), cptr, & + N * NUMERIC_STORAGE_SIZE/CHARACTER_STORAGE_SIZE, & + 0_c_intptr_t, 0_c_intptr_t, & + omp_get_initial_device(), dev)) & + stop 5 + + do i = 1, N + if (fptr(i) /= 11 * int(i)) stop 6 + end do + + deallocate (fptr); + call omp_target_free (cptr, dev); + call omp_target_free (s%cptr, dev); + call omp_target_free (c_loc(s%fptr), dev); + end +end diff --git a/libgomp/testsuite/libgomp.fortran/scan-2.f90 b/libgomp/testsuite/libgomp.fortran/scan-2.f90 new file mode 100644 index 0000000..2815ec7 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/scan-2.f90 @@ -0,0 +1,59 @@ +implicit none +integer, parameter :: N = 30 +integer, parameter :: M = 3 + +integer :: a(M,N), b(M,N), c(M,N) +integer :: x, y, shift +integer :: i, j + +do i = 1, N + a(1,i) = i*32 + a(2,i) = i*17 + a(3,i) = i*11 + b(1,i) = i*7 + b(2,i) = i*5 + b(3,i) = i*3 +end do + +x = 0 +!$omp parallel do simd collapse(2) reduction(inscan,+: x) private(shift) +do i = 1, N + do j = 1, M + x = x + a(j,i) + x = x + b(j,i) + !$omp scan inclusive(x) + shift = i + 29*j + c(j,i) = x + shift; + end do +end do + +y = 0 +do i = 1, N + do j = 1, M + y = y + a(j,i) + b(j,i) + if (c(j,i) /= y + i + 29*j) error stop 1 + end do +end do +if (x /= y) error stop 2 + +x = 0 +!$omp parallel do simd collapse(2) reduction(inscan,+: x) private(shift) +do i = 1, N + do j = 1, M + shift = i + 29*j + c(j,i) = x + shift; + !$omp scan exclusive(x) + x = x + a(j,i) + x = x + b(j,i) + end do +end do + +y = 0 +do i = 1, N + do j = 1, M + if (c(j,i) /= y + i + 29*j) error stop 1 + y = y + a(j,i) + b(j,i) + end do +end do +if (x /= y) error stop 2 +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 index 5d97566..1b3cdf9 100644 --- a/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 @@ -17,6 +17,6 @@ var%p2 = [46,679,54] if (any (var%p1 /= [22,53,28,6,4])) stop 3 if (any (var%p2 /= [46,679,54])) stop 4 !$omp end target -!!$omp target exit data map(from:var%p1, var%p2) +!$omp target exit data map(from:var%p1, var%p2) end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 new file mode 100644 index 0000000..6192bf2 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 @@ -0,0 +1,540 @@ +! Check that 'map(alloc:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5), ic2 + character(len=11) :: ccstr(3:4), ccstr2 + character(len=11,kind=4) :: cc4str(3:7), cc4str2 + integer, pointer :: pc(:), pc2 + character(len=:), pointer :: pcstr(:), pcstr2 + character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +end type t + +type(t) :: dt + +integer :: ii(5), ii2 +character(len=11) :: clstr(-1:1), clstr2 +character(len=11,kind=4) :: cl4str(0:3), cl4str2 +integer, pointer :: ip(:), ip2 +integer, allocatable :: ia(:), ia2 +character(len=:), pointer :: pstr(:), pstr2 +character(len=:), allocatable :: astr(:), astr2 +character(len=:,kind=4), pointer :: p4str(:), p4str2 +character(len=:,kind=4), allocatable :: a4str(:), a4str2 + +allocate(dt%pc(5), dt%pc2) +allocate(character(len=2) :: dt%pcstr(2)) +allocate(character(len=4) :: dt%pcstr2) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +allocate(character(len=5,kind=4) :: dt%pc4str2) + +allocate(ip(5), ip2, ia(8), ia2) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=4) :: pstr2) +allocate(character(len=6) :: astr(3:5)) +allocate(character(len=8) :: astr2) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=5,kind=4) :: p4str2) +allocate(character(len=7,kind=4) :: a4str(-2:3)) +allocate(character(len=9,kind=4) :: a4str2) + + +! integer :: ic(2:5), ic2 + +!$omp target enter data map(alloc: dt%ic) +!$omp target map(alloc: dt%ic) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: dt%ic2) +!$omp target map(alloc: dt%ic2) + dt%ic2 = 42 +!$omp end target +!$omp target exit data map(from: dt%ic2) +if (dt%ic2 /= 42) error stop + + +! character(len=11) :: ccstr(3:4), ccstr2 + +!$omp target enter data map(alloc: dt%ccstr) +!$omp target map(alloc: dt%ccstr) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr = ["12345678901", "abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + +!$omp target enter data map(alloc: dt%ccstr2) +!$omp target map(alloc: dt%ccstr2) + if (len(dt%ccstr2) /= 11) error stop + dt%ccstr2 = "ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%ccstr2) +if (len(dt%ccstr2) /= 11) error stop +if (dt%ccstr2 /= "ABCDEFGHIJK") error stop + + +! character(len=11,kind=4) :: cc4str(3:7), cc4str2 + +! Value check fails +!$omp target enter data map(alloc: dt%cc4str) +!$omp target map(alloc: dt%cc4str) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop + +!$omp target enter data map(alloc: dt%cc4str2) +!$omp target map(alloc: dt%cc4str2) + if (len(dt%cc4str2) /= 11) error stop + dt%cc4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%cc4str2) +if (len(dt%cc4str2) /= 11) error stop +if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop + + +! integer, pointer :: pc(:), pc2 +! allocate(dt%pc(5), dt%pc2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(alloc: dt%pc) +!$omp target map(alloc: dt%pc) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: dt%pc2) +!$omp target map(alloc: dt%pc2) + if (.not. associated(dt%pc2)) error stop + dt%pc2 = 99 +!$omp end target +!$omp target exit data map(from: dt%pc2) +if (dt%pc2 /= 99) error stop +if (.not. associated(dt%pc2)) error stop + + +! character(len=:), pointer :: pcstr(:), pcstr2 +! allocate(character(len=2) :: dt%pcstr(2)) +! allocate(character(len=4) :: dt%pcstr2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(alloc: dt%pcstr) +!$omp target map(alloc: dt%pcstr) + if (.not. associated(dt%pcstr)) error stop + if (len(dt%pcstr) /= 2) error stop + if (size(dt%pcstr) /= 2) error stop + if (lbound(dt%pcstr, 1) /= 1) error stop + if (ubound(dt%pcstr, 1) /= 2) error stop + dt%pcstr = ["01", "jk"] +!$omp end target +!$omp target exit data map(from: dt%pcstr) +if (.not. associated(dt%pcstr)) error stop +if (len(dt%pcstr) /= 2) error stop +if (size(dt%pcstr) /= 2) error stop +if (lbound(dt%pcstr, 1) /= 1) error stop +if (ubound(dt%pcstr, 1) /= 2) error stop +if (any (dt%pcstr /= ["01", "jk"])) error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(alloc: dt%pcstr2) +!$omp target map(alloc: dt%pcstr2) + if (.not. associated(dt%pcstr2)) error stop + if (len(dt%pcstr2) /= 4) error stop + dt%pcstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: dt%pcstr2) +if (.not. associated(dt%pcstr2)) error stop +if (len(dt%pcstr2) /= 4) error stop +if (dt%pcstr2 /= "HIJK") error stop + + +! character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +! allocate(character(len=5,kind=4) :: dt%pc4str2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 +! structure element when other mapped elements from the same structure weren't mapped together with it +!$omp target enter data map(alloc: dt%pc4str) +!$omp target map(alloc: dt%pc4str) + if (.not. associated(dt%pc4str)) error stop + if (len(dt%pc4str) /= 3) error stop + if (size(dt%pc4str) /= 2) error stop + if (lbound(dt%pc4str, 1) /= 2) error stop + if (ubound(dt%pc4str, 1) /= 3) error stop + dt%pc4str = [4_"456", 4_"tzu"] +!$omp end target +!$omp target exit data map(from: dt%pc4str) +if (.not. associated(dt%pc4str)) error stop +if (len(dt%pc4str) /= 3) error stop +if (size(dt%pc4str) /= 2) error stop +if (lbound(dt%pc4str, 1) /= 2) error stop +if (ubound(dt%pc4str, 1) /= 3) error stop +if (dt%pc4str(2) /= 4_"456") error stop +if (dt%pc4str(3) /= 4_"tzu") error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(alloc: dt%pc4str2) +!$omp target map(alloc: dt%pc4str2) + if (.not. associated(dt%pc4str2)) error stop + if (len(dt%pc4str2) /= 5) error stop + dt%pc4str2 = 4_"98765" +!$omp end target +!$omp target exit data map(from: dt%pc4str2) +if (.not. associated(dt%pc4str2)) error stop +if (len(dt%pc4str2) /= 5) error stop +if (dt%pc4str2 /= 4_"98765") error stop + + +! integer :: ii(5), ii2 + +!$omp target enter data map(alloc: ii) +!$omp target map(alloc: ii) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii = [-1, -2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + +!$omp target enter data map(alloc: ii2) +!$omp target map(alloc: ii2) + ii2 = -410 +!$omp end target +!$omp target exit data map(from: ii2) +if (ii2 /= -410) error stop + + +! character(len=11) :: clstr(-1:1), clstr2 + +!$omp target enter data map(alloc: clstr) +!$omp target map(alloc: clstr) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +!$omp target enter data map(alloc: clstr2) +!$omp target map(alloc: clstr2) + if (len(clstr2) /= 11) error stop + clstr2 = "ABCDEFghijk" +!$omp end target +!$omp target exit data map(from: clstr2) +if (len(clstr2) /= 11) error stop +if (clstr2 /= "ABCDEFghijk") error stop + + +! character(len=11,kind=4) :: cl4str(0:3), cl4str2 + +!$omp target enter data map(alloc: cl4str) +!$omp target map(alloc: cl4str) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + +!$omp target enter data map(alloc: cl4str2) +!$omp target map(alloc: cl4str2) + if (len(cl4str2) /= 11) error stop + cl4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: cl4str2) +if (len(cl4str2) /= 11) error stop +if (cl4str2 /= 4_"ABCDEFGHIJK") error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(alloc: ip) +!$omp target map(alloc: ip) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: ip2) +!$omp target map(alloc: ip2) + if (.not. associated(ip2)) error stop + ip2 = 99 +!$omp end target +!$omp target exit data map(from: ip2) +if (ip2 /= 99) error stop +if (.not. associated(ip2)) error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(alloc: ia) +!$omp target map(alloc: ia) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia = [1,2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + +!$omp target enter data map(alloc: ia2) +!$omp target map(alloc: ia2) + if (.not. allocated(ia2)) error stop + ia2 = 102 +!$omp end target +!$omp target exit data map(from: ia2) +if (ia2 /= 102) error stop +if (.not. allocated(ia2)) error stop + + +! character(len=:), pointer :: pstr(:), pstr2 +! allocate(character(len=2) :: pstr(-2:0)) +! allocate(character(len=4) :: pstr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(alloc: pstr) +!$omp target map(alloc: pstr) + if (.not. associated(pstr)) error stop + if (len(pstr) /= 2) error stop + if (size(pstr) /= 3) error stop + if (lbound(pstr, 1) /= -2) error stop + if (ubound(pstr, 1) /= 0) error stop + pstr = ["01", "jk", "aq"] +!$omp end target +!$omp target exit data map(from: pstr) +if (.not. associated(pstr)) error stop +if (len(pstr) /= 2) error stop +if (size(pstr) /= 3) error stop +if (lbound(pstr, 1) /= -2) error stop +if (ubound(pstr, 1) /= 0) error stop +if (any (pstr /= ["01", "jk", "aq"])) error stop + +!$omp target enter data map(alloc: pstr2) +!$omp target map(alloc: pstr2) + if (.not. associated(pstr2)) error stop + if (len(pstr2) /= 4) error stop + pstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: pstr2) +if (.not. associated(pstr2)) error stop +if (len(pstr2) /= 4) error stop +if (pstr2 /= "HIJK") error stop + + +! character(len=:), allocatable :: astr(:), astr2 +! allocate(character(len=6) :: astr(3:5)) +! allocate(character(len=8) :: astr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(alloc: astr) +!$omp target map(alloc: astr) + if (.not. allocated(astr)) error stop + if (len(astr) /= 6) error stop + if (size(astr) /= 3) error stop + if (lbound(astr, 1) /= 3) error stop + if (ubound(astr, 1) /= 5) error stop + astr = ["01db45", "jk$D%S", "zutg47"] +!$omp end target +!$omp target exit data map(from: astr) +if (.not. allocated(astr)) error stop +if (len(astr) /= 6) error stop +if (size(astr) /= 3) error stop +if (lbound(astr, 1) /= 3) error stop +if (ubound(astr, 1) /= 5) error stop +if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(alloc: astr2) +!$omp target map(alloc: astr2) + if (.not. allocated(astr2)) error stop + if (len(astr2) /= 8) error stop + astr2 = "HIJKhijk" +!$omp end target +!$omp target exit data map(from: astr2) +if (.not. allocated(astr2)) error stop +if (len(astr2) /= 8) error stop +if (astr2 /= "HIJKhijk") error stop + + +! character(len=:,kind=4), pointer :: p4str(:), p4str2 +! allocate(character(len=3,kind=4) :: p4str(2:4)) +! allocate(character(len=5,kind=4) :: p4str2) + +! FAILS with value check + +!$omp target enter data map(alloc: p4str) +!$omp target map(alloc: p4str) + if (.not. associated(p4str)) error stop + if (len(p4str) /= 3) error stop + if (size(p4str) /= 3) error stop + if (lbound(p4str, 1) /= 2) error stop + if (ubound(p4str, 1) /= 4) error stop + p4str(:) = [4_"f85", 4_"8af", 4_"A%F"] +!$omp end target +!$omp target exit data map(from: p4str) +if (.not. associated(p4str)) error stop +if (len(p4str) /= 3) error stop +if (size(p4str) /= 3) error stop +if (lbound(p4str, 1) /= 2) error stop +if (ubound(p4str, 1) /= 4) error stop +if (p4str(2) /= 4_"f85") error stop +if (p4str(3) /= 4_"8af") error stop +if (p4str(4) /= 4_"A%F") error stop + +!$omp target enter data map(alloc: p4str2) +!$omp target map(alloc: p4str2) + if (.not. associated(p4str2)) error stop + if (len(p4str2) /= 5) error stop + p4str2 = 4_"9875a" +!$omp end target +!$omp target exit data map(from: p4str2) +if (.not. associated(p4str2)) error stop +if (len(p4str2) /= 5) error stop +if (p4str2 /= 4_"9875a") error stop + + +! character(len=:,kind=4), allocatable :: a4str(:), a4str2 +! allocate(character(len=7,kind=4) :: a4str(-2:3)) +! allocate(character(len=9,kind=4) :: a4str2) + +! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped + +!$omp target enter data map(alloc: a4str) +!$omp target map(alloc: a4str) + if (.not. allocated(a4str)) error stop + if (len(a4str) /= 7) error stop + if (size(a4str) /= 6) error stop + if (lbound(a4str, 1) /= -2) error stop + if (ubound(a4str, 1) /= 3) error stop + ! See PR fortran/107508 why '(:)' is required + a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!$omp end target +!$omp target exit data map(from: a4str) +if (.not. allocated(a4str)) error stop +if (len(a4str) /= 7) error stop +if (size(a4str) /= 6) error stop +if (lbound(a4str, 1) /= -2) error stop +if (ubound(a4str, 1) /= 3) error stop +if (a4str(-2) /= 4_"sf456aq") error stop +if (a4str(-1) /= 4_"3dtzu24") error stop +if (a4str(0) /= 4_"_4fh7sm") error stop +if (a4str(1) /= 4_"=ff85s7") error stop +if (a4str(2) /= 4_"j=8af4d") error stop +if (a4str(3) /= 4_".,A%Fsz") error stop + +!$omp target enter data map(alloc: a4str2) +!$omp target map(alloc: a4str2) + if (.not. allocated(a4str2)) error stop + if (len(a4str2) /= 9) error stop + a4str2 = 4_"98765a23d" +!$omp end target +!$omp target exit data map(from: a4str2) +if (.not. allocated(a4str2)) error stop +if (len(a4str2) /= 9) error stop +if (a4str2 /= 4_"98765a23d") error stop + + +deallocate(dt%pc, dt%pc2) +deallocate(dt%pcstr) +deallocate(dt%pcstr2) + +deallocate(dt%pc4str) +deallocate(dt%pc4str2) + +deallocate(ip, ip2, ia, ia2) +deallocate(pstr) +deallocate(pstr2) +deallocate(astr) +deallocate(astr2) + +deallocate(p4str) +deallocate(p4str2) +deallocate(a4str) +deallocate(a4str2) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 new file mode 100644 index 0000000..cf75934 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 @@ -0,0 +1,540 @@ +! Check that 'map((to)from:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5), ic2 + character(len=11) :: ccstr(3:4), ccstr2 + character(len=11,kind=4) :: cc4str(3:7), cc4str2 + integer, pointer :: pc(:), pc2 + character(len=:), pointer :: pcstr(:), pcstr2 + character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +end type t + +type(t) :: dt + +integer :: ii(5), ii2 +character(len=11) :: clstr(-1:1), clstr2 +character(len=11,kind=4) :: cl4str(0:3), cl4str2 +integer, pointer :: ip(:), ip2 +integer, allocatable :: ia(:), ia2 +character(len=:), pointer :: pstr(:), pstr2 +character(len=:), allocatable :: astr(:), astr2 +character(len=:,kind=4), pointer :: p4str(:), p4str2 +character(len=:,kind=4), allocatable :: a4str(:), a4str2 + +allocate(dt%pc(5), dt%pc2) +allocate(character(len=2) :: dt%pcstr(2)) +allocate(character(len=4) :: dt%pcstr2) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +allocate(character(len=5,kind=4) :: dt%pc4str2) + +allocate(ip(5), ip2, ia(8), ia2) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=4) :: pstr2) +allocate(character(len=6) :: astr(3:5)) +allocate(character(len=8) :: astr2) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=5,kind=4) :: p4str2) +allocate(character(len=7,kind=4) :: a4str(-2:3)) +allocate(character(len=9,kind=4) :: a4str2) + + +! integer :: ic(2:5), ic2 + +!$omp target enter data map(tofrom: dt%ic) +!$omp target map(from: dt%ic) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +!$omp target enter data map(tofrom: dt%ic2) +!$omp target map(from: dt%ic2) + dt%ic2 = 42 +!$omp end target +!$omp target exit data map(from: dt%ic2) +if (dt%ic2 /= 42) error stop + + +! character(len=11) :: ccstr(3:4), ccstr2 + +!$omp target enter data map(tofrom: dt%ccstr) +!$omp target map(from: dt%ccstr) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr = ["12345678901", "abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + +!$omp target enter data map(tofrom: dt%ccstr2) +!$omp target map(from: dt%ccstr2) + if (len(dt%ccstr2) /= 11) error stop + dt%ccstr2 = "ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%ccstr2) +if (len(dt%ccstr2) /= 11) error stop +if (dt%ccstr2 /= "ABCDEFGHIJK") error stop + + +! character(len=11,kind=4) :: cc4str(3:7), cc4str2 + +! Value check fails +!$omp target enter data map(tofrom: dt%cc4str) +!$omp target map(from: dt%cc4str) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop + +!$omp target enter data map(tofrom: dt%cc4str2) +!$omp target map(from: dt%cc4str2) + if (len(dt%cc4str2) /= 11) error stop + dt%cc4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%cc4str2) +if (len(dt%cc4str2) /= 11) error stop +if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop + + +! integer, pointer :: pc(:), pc2 +! allocate(dt%pc(5), dt%pc2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(tofrom: dt%pc) +!$omp target map(from: dt%pc) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(tofrom: dt%pc2) +!$omp target map(from: dt%pc2) + if (.not. associated(dt%pc2)) error stop + dt%pc2 = 99 +!$omp end target +!$omp target exit data map(from: dt%pc2) +if (dt%pc2 /= 99) error stop +if (.not. associated(dt%pc2)) error stop + + +! character(len=:), pointer :: pcstr(:), pcstr2 +! allocate(character(len=2) :: dt%pcstr(2)) +! allocate(character(len=4) :: dt%pcstr2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(tofrom: dt%pcstr) +!$omp target map(from: dt%pcstr) + if (.not. associated(dt%pcstr)) error stop + if (len(dt%pcstr) /= 2) error stop + if (size(dt%pcstr) /= 2) error stop + if (lbound(dt%pcstr, 1) /= 1) error stop + if (ubound(dt%pcstr, 1) /= 2) error stop + dt%pcstr = ["01", "jk"] +!$omp end target +!$omp target exit data map(from: dt%pcstr) +if (.not. associated(dt%pcstr)) error stop +if (len(dt%pcstr) /= 2) error stop +if (size(dt%pcstr) /= 2) error stop +if (lbound(dt%pcstr, 1) /= 1) error stop +if (ubound(dt%pcstr, 1) /= 2) error stop +if (any (dt%pcstr /= ["01", "jk"])) error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(tofrom: dt%pcstr2) +!$omp target map(from: dt%pcstr2) + if (.not. associated(dt%pcstr2)) error stop + if (len(dt%pcstr2) /= 4) error stop + dt%pcstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: dt%pcstr2) +if (.not. associated(dt%pcstr2)) error stop +if (len(dt%pcstr2) /= 4) error stop +if (dt%pcstr2 /= "HIJK") error stop + + +! character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +! allocate(character(len=5,kind=4) :: dt%pc4str2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(tofrom: dt%pc4str) +!$omp target map(from: dt%pc4str) + if (.not. associated(dt%pc4str)) error stop + if (len(dt%pc4str) /= 3) error stop + if (size(dt%pc4str) /= 2) error stop + if (lbound(dt%pc4str, 1) /= 2) error stop + if (ubound(dt%pc4str, 1) /= 3) error stop + dt%pc4str = [4_"456", 4_"tzu"] +!$omp end target +!$omp target exit data map(from: dt%pc4str) +if (.not. associated(dt%pc4str)) error stop +if (len(dt%pc4str) /= 3) error stop +if (size(dt%pc4str) /= 2) error stop +if (lbound(dt%pc4str, 1) /= 2) error stop +if (ubound(dt%pc4str, 1) /= 3) error stop +if (dt%pc4str(2) /= 4_"456") error stop +if (dt%pc4str(3) /= 4_"tzu") error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(tofrom: dt%pc4str2) +!$omp target map(from: dt%pc4str2) + if (.not. associated(dt%pc4str2)) error stop + if (len(dt%pc4str2) /= 5) error stop + dt%pc4str2 = 4_"98765" +!$omp end target +!$omp target exit data map(from: dt%pc4str2) +if (.not. associated(dt%pc4str2)) error stop +if (len(dt%pc4str2) /= 5) error stop +if (dt%pc4str2 /= 4_"98765") error stop + + +! integer :: ii(5), ii2 + +!$omp target enter data map(tofrom: ii) +!$omp target map(from: ii) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii = [-1, -2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + +!$omp target enter data map(tofrom: ii2) +!$omp target map(from: ii2) + ii2 = -410 +!$omp end target +!$omp target exit data map(from: ii2) +if (ii2 /= -410) error stop + + +! character(len=11) :: clstr(-1:1), clstr2 + +!$omp target enter data map(tofrom: clstr) +!$omp target map(from: clstr) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +!$omp target enter data map(tofrom: clstr2) +!$omp target map(from: clstr2) + if (len(clstr2) /= 11) error stop + clstr2 = "ABCDEFghijk" +!$omp end target +!$omp target exit data map(from: clstr2) +if (len(clstr2) /= 11) error stop +if (clstr2 /= "ABCDEFghijk") error stop + + +! character(len=11,kind=4) :: cl4str(0:3), cl4str2 + +!$omp target enter data map(tofrom: cl4str) +!$omp target map(from: cl4str) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + +!$omp target enter data map(tofrom: cl4str2) +!$omp target map(from: cl4str2) + if (len(cl4str2) /= 11) error stop + cl4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: cl4str2) +if (len(cl4str2) /= 11) error stop +if (cl4str2 /= 4_"ABCDEFGHIJK") error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(tofrom: ip) +!$omp target map(from: ip) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(tofrom: ip2) +!$omp target map(from: ip2) + if (.not. associated(ip2)) error stop + ip2 = 99 +!$omp end target +!$omp target exit data map(from: ip2) +if (ip2 /= 99) error stop +if (.not. associated(ip2)) error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(tofrom: ia) +!$omp target map(from: ia) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia = [1,2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + +!$omp target enter data map(tofrom: ia2) +!$omp target map(from: ia2) + if (.not. allocated(ia2)) error stop + ia2 = 102 +!$omp end target +!$omp target exit data map(from: ia2) +if (ia2 /= 102) error stop +if (.not. allocated(ia2)) error stop + + +! character(len=:), pointer :: pstr(:), pstr2 +! allocate(character(len=2) :: pstr(-2:0)) +! allocate(character(len=4) :: pstr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(tofrom: pstr) +!$omp target map(from: pstr) + if (.not. associated(pstr)) error stop + if (len(pstr) /= 2) error stop + if (size(pstr) /= 3) error stop + if (lbound(pstr, 1) /= -2) error stop + if (ubound(pstr, 1) /= 0) error stop + pstr = ["01", "jk", "aq"] +!$omp end target +!$omp target exit data map(from: pstr) +if (.not. associated(pstr)) error stop +if (len(pstr) /= 2) error stop +if (size(pstr) /= 3) error stop +if (lbound(pstr, 1) /= -2) error stop +if (ubound(pstr, 1) /= 0) error stop +if (any (pstr /= ["01", "jk", "aq"])) error stop + +!$omp target enter data map(tofrom: pstr2) +!$omp target map(from: pstr2) + if (.not. associated(pstr2)) error stop + if (len(pstr2) /= 4) error stop + pstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: pstr2) +if (.not. associated(pstr2)) error stop +if (len(pstr2) /= 4) error stop +if (pstr2 /= "HIJK") error stop + + +! character(len=:), allocatable :: astr(:), astr2 +! allocate(character(len=6) :: astr(3:5)) +! allocate(character(len=8) :: astr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(tofrom: astr) +!$omp target map(from: astr) + if (.not. allocated(astr)) error stop + if (len(astr) /= 6) error stop + if (size(astr) /= 3) error stop + if (lbound(astr, 1) /= 3) error stop + if (ubound(astr, 1) /= 5) error stop + astr = ["01db45", "jk$D%S", "zutg47"] +!$omp end target +!$omp target exit data map(from: astr) +if (.not. allocated(astr)) error stop +if (len(astr) /= 6) error stop +if (size(astr) /= 3) error stop +if (lbound(astr, 1) /= 3) error stop +if (ubound(astr, 1) /= 5) error stop +if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(tofrom: astr2) +!$omp target map(from: astr2) + if (.not. allocated(astr2)) error stop + if (len(astr2) /= 8) error stop + astr2 = "HIJKhijk" +!$omp end target +!$omp target exit data map(from: astr2) +if (.not. allocated(astr2)) error stop +if (len(astr2) /= 8) error stop +if (astr2 /= "HIJKhijk") error stop + + +! character(len=:,kind=4), pointer :: p4str(:), p4str2 +! allocate(character(len=3,kind=4) :: p4str(2:4)) +! allocate(character(len=5,kind=4) :: p4str2) + +! FAILS with value check + +!$omp target enter data map(tofrom: p4str) +!$omp target map(from: p4str) + if (.not. associated(p4str)) error stop + if (len(p4str) /= 3) error stop + if (size(p4str) /= 3) error stop + if (lbound(p4str, 1) /= 2) error stop + if (ubound(p4str, 1) /= 4) error stop + p4str(:) = [4_"f85", 4_"8af", 4_"A%F"] +!$omp end target +!$omp target exit data map(from: p4str) +if (.not. associated(p4str)) error stop +if (len(p4str) /= 3) error stop +if (size(p4str) /= 3) error stop +if (lbound(p4str, 1) /= 2) error stop +if (ubound(p4str, 1) /= 4) error stop +if (p4str(2) /= 4_"f85") error stop +if (p4str(3) /= 4_"8af") error stop +if (p4str(4) /= 4_"A%F") error stop + +!$omp target enter data map(tofrom: p4str2) +!$omp target map(from: p4str2) + if (.not. associated(p4str2)) error stop + if (len(p4str2) /= 5) error stop + p4str2 = 4_"9875a" +!$omp end target +!$omp target exit data map(from: p4str2) +if (.not. associated(p4str2)) error stop +if (len(p4str2) /= 5) error stop +if (p4str2 /= 4_"9875a") error stop + + +! character(len=:,kind=4), allocatable :: a4str(:), a4str2 +! allocate(character(len=7,kind=4) :: a4str(-2:3)) +! allocate(character(len=9,kind=4) :: a4str2) + +! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped + +!$omp target enter data map(tofrom: a4str) +!$omp target map(from: a4str) + if (.not. allocated(a4str)) error stop + if (len(a4str) /= 7) error stop + if (size(a4str) /= 6) error stop + if (lbound(a4str, 1) /= -2) error stop + if (ubound(a4str, 1) /= 3) error stop + ! See PR fortran/107508 why '(:)' is required + a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!$omp end target +!$omp target exit data map(from: a4str) +if (.not. allocated(a4str)) error stop +if (len(a4str) /= 7) error stop +if (size(a4str) /= 6) error stop +if (lbound(a4str, 1) /= -2) error stop +if (ubound(a4str, 1) /= 3) error stop +if (a4str(-2) /= 4_"sf456aq") error stop +if (a4str(-1) /= 4_"3dtzu24") error stop +if (a4str(0) /= 4_"_4fh7sm") error stop +if (a4str(1) /= 4_"=ff85s7") error stop +if (a4str(2) /= 4_"j=8af4d") error stop +if (a4str(3) /= 4_".,A%Fsz") error stop + +!$omp target enter data map(tofrom: a4str2) +!$omp target map(from: a4str2) + if (.not. allocated(a4str2)) error stop + if (len(a4str2) /= 9) error stop + a4str2 = 4_"98765a23d" +!$omp end target +!$omp target exit data map(from: a4str2) +if (.not. allocated(a4str2)) error stop +if (len(a4str2) /= 9) error stop +if (a4str2 /= 4_"98765a23d") error stop + + +deallocate(dt%pc, dt%pc2) +deallocate(dt%pcstr) +deallocate(dt%pcstr2) + +deallocate(dt%pc4str) +deallocate(dt%pc4str2) + +deallocate(ip, ip2, ia, ia2) +deallocate(pstr) +deallocate(pstr2) +deallocate(astr) +deallocate(astr2) + +deallocate(p4str) +deallocate(p4str2) +deallocate(a4str) +deallocate(a4str2) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 new file mode 100644 index 0000000..80d30ed --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 @@ -0,0 +1,392 @@ +! Check that 'map(alloc:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5) + character(len=11) :: ccstr(3:4) + character(len=11,kind=4) :: cc4str(3:7) + integer, pointer :: pc(:) + character(len=:), pointer :: pcstr(:) + character(len=:,kind=4), pointer :: pc4str(:) +end type t + +type(t) :: dt + +integer :: ii(5) +character(len=11) :: clstr(-1:1) +character(len=11,kind=4) :: cl4str(0:3) +integer, pointer :: ip(:) +integer, allocatable :: ia(:) +character(len=:), pointer :: pstr(:) +character(len=:), allocatable :: astr(:) +character(len=:,kind=4), pointer :: p4str(:) +character(len=:,kind=4), allocatable :: a4str(:) + +allocate(dt%pc(5)) +allocate(character(len=2) :: dt%pcstr(2)) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) + +allocate(ip(5), ia(8)) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=6) :: astr(3:5)) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=7,kind=4) :: a4str(-2:3)) + + +! integer :: ic(2:5) + +!$omp target enter data map(alloc: dt%ic(3:5)) +dt%ic(2) = 22 +!$omp target map(alloc: dt%ic(3:5)) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic(3:5) = [33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic(3:5)) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +! character(len=11) :: ccstr(3:4) + +!$omp target enter data map(alloc: dt%ccstr(4:4)) +dt%ccstr(3) = "12345678901" +!$omp target map(alloc: dt%ccstr(4:4)) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr(4:4) = ["abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr(4:4)) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + + +! character(len=11,kind=4) :: cc4str(3:7) + +! Value check fails +!$omp target enter data map(alloc: dt%cc4str(4:7)) +dt%cc4str(3) = 4_"12345678901" +!$omp target map(alloc: dt%cc4str(4:7)) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str(4:7) = [4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str(4:7)) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop + +! integer, pointer :: pc(:) +! allocate(dt%pc(5)) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(alloc: dt%pc(2:5)) +dt%pc(1) = 11 +!$omp target map(alloc: dt%pc(2:5)) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc(2:5) = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc(2:5)) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +! character(len=:), pointer :: pcstr(:) +! allocate(character(len=2) :: dt%pcstr(2)) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +! FIXME: Disabled befause of PR108837 +! +!!$omp target enter data map(alloc: dt%pcstr(2:2)) +!dt%pcstr(1) = "01" +!!$omp target map(alloc: dt%pcstr(2:2)) +! if (.not. associated(dt%pcstr)) error stop +! if (len(dt%pcstr) /= 2) error stop +! if (size(dt%pcstr) /= 2) error stop +! if (lbound(dt%pcstr, 1) /= 1) error stop +! if (ubound(dt%pcstr, 1) /= 2) error stop +! dt%pcstr(2:2) = ["jk"] +!!$omp end target +!!$omp target exit data map(from: dt%pcstr(2:2)) +!if (.not. associated(dt%pcstr)) error stop +!if (len(dt%pcstr) /= 2) error stop +!if (size(dt%pcstr) /= 2) error stop +!if (lbound(dt%pcstr, 1) /= 1) error stop +!if (ubound(dt%pcstr, 1) /= 2) error stop +!if (any (dt%pcstr /= ["01", "jk"])) error stop + + +! character(len=:,kind=4), pointer :: pc4str(:) +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 +! structure element when other mapped elements from the same structure weren't mapped together with it + +! FIXME: Disabled befause of PR108837 +! +!!$omp target enter data map(alloc: dt%pc4str(3:3)) +!dt%pc4str(2) = 4_"456" +!!$omp target map(alloc: dt%pc4str(3:3)) +! if (.not. associated(dt%pc4str)) error stop +! if (len(dt%pc4str) /= 3) error stop +! if (size(dt%pc4str) /= 2) error stop +! if (lbound(dt%pc4str, 1) /= 2) error stop +! if (ubound(dt%pc4str, 1) /= 3) error stop +! dt%pc4str(3:3) = [4_"tzu"] +!!$omp end target +!!$omp target exit data map(from: dt%pc4str(3:3)) +!if (.not. associated(dt%pc4str)) error stop +!if (len(dt%pc4str) /= 3) error stop +!if (size(dt%pc4str) /= 2) error stop +!if (lbound(dt%pc4str, 1) /= 2) error stop +!if (ubound(dt%pc4str, 1) /= 3) error stop +!if (dt%pc4str(2) /= 4_"456") error stop +!if (dt%pc4str(3) /= 4_"tzu") error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +! integer :: ii(5) + +!$omp target enter data map(alloc: ii(2:5)) +ii(1) = -1 +!$omp target map(alloc: ii(2:5)) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii(2:5) = [-2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii(2:5)) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + + +! character(len=11) :: clstr(-1:1) + +!$omp target enter data map(alloc: clstr(0:1)) +clstr(-1) = "12345678901" +!$omp target map(alloc: clstr(0:1)) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr(0:1) = ["abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr(0:1)) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +! character(len=11,kind=4) :: cl4str(0:3) + +!$omp target enter data map(alloc: cl4str(1:3)) +cl4str(0) = 4_"12345678901" +!$omp target map(alloc: cl4str(1:3)) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str(1:3) = [4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str(1:3)) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + + +! allocate(ip(5), ia(8)) + +!$omp target enter data map(alloc: ip(2:5)) +ip(1) = 11 +!$omp target map(alloc: ip(2:5)) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip(2:5) = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip(2:5)) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +! allocate(ip(5), ia(8)) + +!$omp target enter data map(alloc: ia(2:8)) +ia(1) = 1 +!$omp target map(alloc: ia(2:8)) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia(2:8) = [2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia(2:8)) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + + +! character(len=:), pointer :: pstr(:) +! allocate(character(len=2) :: pstr(-2:0)) + +! libgomp: nvptx_alloc error: out of memory + +! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR +! +!!$omp target enter data map(alloc: pstr(-1:0)) +!pstr(-2) = "01" +!!$omp target map(alloc: pstr(-1:0)) +! if (.not. associated(pstr)) error stop +! if (len(pstr) /= 2) error stop +! if (size(pstr) /= 3) error stop +! if (lbound(pstr, 1) /= -2) error stop +! if (ubound(pstr, 1) /= 0) error stop +! pstr(-1:0) = ["jk", "aq"] +!!$omp end target +!!$omp target exit data map(from: pstr(-1:0)) +!if (.not. associated(pstr)) error stop +!if (len(pstr) /= 2) error stop +!if (size(pstr) /= 3) error stop +!if (lbound(pstr, 1) /= -2) error stop +!if (ubound(pstr, 1) /= 0) error stop +!if (any (pstr /= ["01", "jk", "aq"])) error stop + + +! character(len=:), allocatable :: astr(:) +! allocate(character(len=6) :: astr(3:5)) + +! libgomp: nvptx_alloc error: out of memory + +! FIXME +!!$omp target enter data map(alloc: astr(4:5)) +!astr(3) = "01db45" +!!$omp target map(alloc: astr(4:5)) +! if (.not. allocated(astr)) error stop +! if (len(astr) /= 6) error stop +! if (size(astr) /= 3) error stop +! if (lbound(astr, 1) /= 3) error stop +! if (ubound(astr, 1) /= 5) error stop +!!! astr(4:5) = ["jk$D%S", "zutg47"] +!!$omp end target +!!!$omp target exit data map(from: astr(4:5)) +!!if (.not. allocated(astr)) error stop +!!!if (len(astr) /= 6) error stop +!if (size(astr) /= 3) error stop +!if (lbound(astr, 1) /= 3) error stop +!if (ubound(astr, 1) /= 5) error stop +!if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop +! + +! character(len=:,kind=4), pointer :: p4str(:) +! allocate(character(len=3,kind=4) :: p4str(2:4)) + +! FAILS with value check + +! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR +! +!!$omp target enter data map(alloc: p4str(3:4)) +!p4str(2) = 4_"f85" +!!$omp target map(alloc: p4str(3:4)) +! if (.not. associated(p4str)) error stop +! if (len(p4str) /= 3) error stop +! if (size(p4str) /= 3) error stop +! if (lbound(p4str, 1) /= 2) error stop +! if (ubound(p4str, 1) /= 4) error stop +! p4str(3:4) = [4_"8af", 4_"A%F"] +!!$omp end target +!!$omp target exit data map(from: p4str(3:4)) +!if (.not. associated(p4str)) error stop +!if (len(p4str) /= 3) error stop +!if (size(p4str) /= 3) error stop +!if (lbound(p4str, 1) /= 2) error stop +!if (ubound(p4str, 1) /= 4) error stop +!if (p4str(2) /= 4_"f85") error stop +!if (p4str(3) /= 4_"8af") error stop +!if (p4str(4) /= 4_"A%F") error stop + +! character(len=:,kind=4), allocatable :: a4str(:) +! allocate(character(len=7,kind=4) :: a4str(-2:3)) + +! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped + +! FIXME: Disabled befause of PR108838 +!!$omp target enter data map(alloc: a4str(-1:3)) +!!a4str(-2) = 4_"sf456aq" +!!$omp target map(alloc: a4str(-1:3)) +! if (.not. allocated(a4str)) error stop +! if (len(a4str) /= 7) error stop +! if (size(a4str) /= 6) error stop +! if (lbound(a4str, 1) /= -2) error stop +! if (ubound(a4str, 1) /= 3) error stop +! a4str(-1:3) = [4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!!$omp end target +!!$omp target exit data map(from: a4str(-1:3)) +!if (.not. allocated(a4str)) error stop +!if (len(a4str) /= 7) error stop +!if (size(a4str) /= 6) error stop +!if (lbound(a4str, 1) /= -2) error stop +!if (ubound(a4str, 1) /= 3) error stop +!if (a4str(-2) /= 4_"sf456aq") error stop +!if (a4str(-1) /= 4_"3dtzu24") error stop +!if (a4str(0) /= 4_"_4fh7sm") error stop +!if (a4str(1) /= 4_"=ff85s7") error stop +!if (a4str(2) /= 4_"j=8af4d") error stop +!if (a4str(3) /= 4_".,A%Fsz") error stop + +deallocate(dt%pc) +deallocate(dt%pcstr) + +deallocate(dt%pc4str) + +deallocate(ip, ia) +deallocate(pstr) +deallocate(astr) + +deallocate(p4str) +deallocate(a4str) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 new file mode 100644 index 0000000..f129f55 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 @@ -0,0 +1,78 @@ +module m + implicit none + character(len=:), allocatable :: strA(:), strA2 + character(len=:), pointer :: strP(:), strP2 + !$omp declare target enter(strA,strA2,strP,strP2) +contains + subroutine opt_map(str1, str2, str3) + character(len=:), allocatable :: str1, str2, str3, str4 + optional :: str2, str3 + + if (.not.present(str2)) error stop + if (present(str3)) error stop + + !$omp target map(str1,str2,str3,str4) + if (allocated(str1)) error stop + if (allocated(str2)) error stop + if (present(str3)) error stop + if (allocated(str4)) error stop + !$omp end target + end + subroutine call_opt() + character(len=:), allocatable :: str1, str2 + call opt_map(str1, str2) + end + subroutine test + !$omp declare target + if (.not. allocated(strA)) error stop + !if (.not. allocated(strA2)) error stop + if (.not. associated(strP)) error stop + !if (.not. associated(strP2)) error stop + + ! ensure length was updated as well + if (len(strA) /= 3) error stop + if (len(strA2) /= 5) error stop + if (len(strP) /= 4) error stop + if (len(strP2) /= 8) error stop +! if (any (strA /= ['Hav', 'e f', 'un!'])) error stop +! if (strA2 /= 'Hello') error stop +! if (any (strP /= ['abcd', 'efgh', 'ijkl'])) error stop +! if (strP2 /= 'TestCase') error stop +! +! strA = ['123', '456', '789'] +! strA2 = 'World' +! strP = ['ABCD', 'EFGH', 'IJKL'] +! strP2 = 'Passed!!' + end +end + +program main + use m + implicit none + call call_opt + + strA = ['Hav', 'e f', 'un!'] + strA2 = 'Hello' + allocate(character(len=4) :: strP(3)) + strP = ['abcd', 'efgh', 'ijkl'] + allocate(character(len=8) :: strP2) + strP2 = 'TestCase' + + !$omp target enter data map(always, to: strA, strA2) + !$omp target enter data map(to: strP, strP2) + !$omp target + call test() + !$omp end target + !$omp target exit data map(always, from: strA, strA2, strP, strP2) + + if (len(strA) /= 3) error stop + if (len(strA2) /= 5) error stop + if (len(strP) /= 4) error stop + if (len(strP2) /= 8) error stop +! if (any (strA /= ['123', '456', '789'])) error stop +! if (strA2 /= 'World') error stop +! if (any(strP /= ['ABCD', 'EFGH', 'IJKL'])) error stop +! if (strP2 /= 'Passed!!') error stop + +! deallocate(strP, strP2, strA, strA2) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-present-1.f90 b/libgomp/testsuite/libgomp.fortran/target-present-1.f90 new file mode 100644 index 0000000..fc13609 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-present-1.f90 @@ -0,0 +1,31 @@ +program main + implicit none + integer, parameter :: N = 100 + integer :: a(N), b(N), c(N), i + + do i = 1, N + a(i) = i * 2 + b(i) = i * 3 + 1 + end do + + !$omp target enter data map (alloc: a) + ! a has already been allocated, so this should be okay. + !$omp target map (present, to: a) + do i = 1, N + c(i) = a(i) + end do + !$omp end target + + print *, "CheCKpOInT" + ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + + ! b has not been allocated, so this should result in an error. + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } + ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } + !$omp target map (present, to: b) + do i = 1, N + c(i) = c(i) + b(i) + end do + !$omp end target + !$omp target exit data map (from: c) +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-present-2.f90 b/libgomp/testsuite/libgomp.fortran/target-present-2.f90 new file mode 100644 index 0000000..524d01d --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-present-2.f90 @@ -0,0 +1,31 @@ +program main + implicit none + integer, parameter :: N = 100 + integer :: a(N), b(N), c(N), i + + do i = 1, N + a(i) = i * 2 + b(i) = i * 3 + 1 + end do + + !$omp target enter data map (alloc: a, c, i) + ! a, c, and i have already been allocated, so this should be okay. + !$omp target defaultmap (present) + do i = 1, N + c(i) = a(i) + end do + !$omp end target + + print *, "CheCKpOInT" + ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + + ! b has not been allocated, so this should result in an error. + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } + ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } + !$omp target defaultmap (present) + do i = 1, N + c(i) = c(i) + b(i) + end do + !$omp end target +!$omp target exit data map (from: c) +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-present-3.f90 b/libgomp/testsuite/libgomp.fortran/target-present-3.f90 new file mode 100644 index 0000000..dd4af4c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-present-3.f90 @@ -0,0 +1,36 @@ +program main + implicit none + integer, parameter :: N = 100 + integer :: a(N), b(N), c(N), i + + do i = 1, N + a(i) = i * 2 + b(i) = i * 3 + 1 + end do + + !$omp target enter data map (alloc: a, c) + + ! This should work as a has already been allocated. + !$omp target update to (present: a) + + !$omp target map(present, alloc: a, c) + do i = 1, N + if (a(i) /= i * 2) stop 1 + c(i) = 23 * i + end do + !$omp end target + + !$omp target update from (present: c) + do i = 1, N + if (c(i) /= 23 * i) stop 1 + end do + + print *, "CheCKpOInT" + ! { dg-output "CheCKpOInT(\n|\r\n|\r).*" } + + ! This should fail as b has not been allocated. + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } + ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } + !$omp target update to (present: b) + !$omp target exit data map (from: c) +end program diff --git a/libgomp/testsuite/libgomp.graphite/graphite.exp b/libgomp/testsuite/libgomp.graphite/graphite.exp index 1260dc9..bc3a82d 100644 --- a/libgomp/testsuite/libgomp.graphite/graphite.exp +++ b/libgomp/testsuite/libgomp.graphite/graphite.exp @@ -14,17 +14,6 @@ # along with GCC; see the file COPYING3. If not see # <http://www.gnu.org/licenses/>. -if [info exists lang_library_path] then { - unset lang_library_path - unset lang_link_flags -} -if [info exists lang_test_file] then { - unset lang_test_file -} -if [info exists lang_include_flags] then { - unset lang_include_flags -} - load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp @@ -36,6 +25,7 @@ if ![check_effective_target_fgraphite] { return } +lappend ALWAYS_CFLAGS "compiler=$GCC_UNDER_TEST" # Flags for force-parallel-*.c testcases. set PARALLEL_CFLAGS "-ansi -pedantic-errors -O2 \ -ftree-parallelize-loops=4 -floop-parallelize-all \ diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp index 42e0395..79df401 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp +++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp @@ -11,15 +11,21 @@ proc check_effective_target_c++ { } { return 1 } -global shlib_ext - -set shlib_ext [get_shlib_extension] -set lang_link_flags "-lstdc++" -set lang_test_file_found 0 -set lang_library_path "../libstdc++-v3/src/.libs" -if [info exists lang_include_flags] then { - unset lang_include_flags +if { $blddir != "" } { + set libstdc++_library_path "../libstdc++-v3/src/.libs" + set shlib_ext [get_shlib_extension] + if { ![file exists "${blddir}/${libstdc++_library_path}/libstdc++.a"] + && ![file exists "${blddir}/${libstdc++_library_path}/libstdc++.${shlib_ext}"] } { + verbose -log "No libstdc++ library found, will not execute c++ tests" + unset libstdc++_library_path + return + } + lappend lang_library_paths ${libstdc++_library_path} +} elseif { ![info exists GXX_UNDER_TEST] } { + verbose -log "GXX_UNDER_TEST not defined, will not execute c++ tests" + return } +lappend ALWAYS_CFLAGS "compiler=$GXX_UNDER_TEST" # Initialize dg. dg-init @@ -28,136 +34,96 @@ torture-init # Turn on OpenACC. lappend ALWAYS_CFLAGS "additional_flags=-fopenacc" -# Switch into C++ mode. Otherwise, the libgomp.oacc-c-c++-common/*.c -# files would be compiled as C files. -set SAVE_GCC_UNDER_TEST "$GCC_UNDER_TEST" -set GCC_UNDER_TEST "$GCC_UNDER_TEST -x c++" - -set blddir [lookfor_file [get_multilibs] libgomp] - +# Gather a list of all tests. +set tests [lsort [concat \ + [find $srcdir/$subdir *.C] \ + [find $srcdir/$subdir/../libgomp.oacc-c-c++-common *.c]]] +set ld_library_path $always_ld_library_path if { $blddir != "" } { - # Look for a static libstdc++ first. - if [file exists "${blddir}/${lang_library_path}/libstdc++.a"] { - set lang_test_file "${lang_library_path}/libstdc++.a" - set lang_test_file_found 1 - # We may have a shared only build, so look for a shared libstdc++. - } elseif [file exists "${blddir}/${lang_library_path}/libstdc++.${shlib_ext}"] { - set lang_test_file "${lang_library_path}/libstdc++.${shlib_ext}" - set lang_test_file_found 1 - } else { - puts "No libstdc++ library found, will not execute c++ tests" - } -} elseif { [info exists GXX_UNDER_TEST] } { - set lang_test_file_found 1 - # Needs to exist for libgomp.exp. - set lang_test_file "" -} else { - puts "GXX_UNDER_TEST not defined, will not execute c++ tests" + append ld_library_path ":${blddir}/${libstdc++_library_path}" } - -if { $lang_test_file_found } { - # Gather a list of all tests. - set tests [lsort [concat \ - [find $srcdir/$subdir *.C] \ - [find $srcdir/$subdir/../libgomp.oacc-c-c++-common *.c]]] - - if { $blddir != "" } { - set ld_library_path "$always_ld_library_path:${blddir}/${lang_library_path}" - } else { - set ld_library_path "$always_ld_library_path" - } - append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] - set_ld_library_path_env_vars - - set flags_file "${blddir}/../libstdc++-v3/scripts/testsuite_flags" - if { [file exists $flags_file] } { - set libstdcxx_includes [exec sh $flags_file --build-includes] - } else { - set libstdcxx_includes "" - } - - # Test with all available offload targets, and with offloading disabled. - foreach offload_target [concat [split $offload_targets ","] "disable"] { - global openacc_device_type - set openacc_device_type [offload_target_to_openacc_device_type $offload_target] - set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" - - switch $openacc_device_type { - "" { - unsupported "$subdir $offload_target offloading" +append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] +set_ld_library_path_env_vars + +# Test with all available offload targets, and with offloading disabled. +foreach offload_target [concat [split $offload_targets ","] "disable"] { + global openacc_device_type + set openacc_device_type [offload_target_to_openacc_device_type $offload_target] + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" + + switch $openacc_device_type { + "" { + unsupported "$subdir $offload_target offloading" + continue + } + host { + set acc_mem_shared 1 + } + nvidia { + if { ![check_effective_target_openacc_nvidia_accel_present] } { + # Don't bother; execution testing is going to FAIL. + untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue } - host { - set acc_mem_shared 1 - } - nvidia { - if { ![check_effective_target_openacc_nvidia_accel_present] } { - # Don't bother; execution testing is going to FAIL. - untested "$subdir $offload_target offloading: supported, but hardware not accessible" - continue - } - # Copy ptx file (TEMPORARY) - remote_download host $srcdir/libgomp.oacc-c-c++-common/subr.ptx + # Copy ptx file (TEMPORARY) + remote_download host $srcdir/libgomp.oacc-c-c++-common/subr.ptx - # Where timer.h lives - lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common" + # Where timer.h lives + lappend ALWAYS_CFLAGS "additional_flags=-I${srcdir}/libgomp.oacc-c-c++-common" - set acc_mem_shared 0 - } - radeon { - if { ![check_effective_target_openacc_radeon_accel_present] } { - # Don't bother; execution testing is going to FAIL. - untested "$subdir $offload_target offloading: supported, but hardware not accessible" - continue - } - - set acc_mem_shared 0 - } - default { - error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)" - } + set acc_mem_shared 0 } - set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared" - - # To avoid compilation overhead, and to keep simple '-foffload=[...]' - # handling in test cases, by default only build for the offload target - # that we're actually going to test. - set tagopt "$tagopt -foffload=$offload_target" - # Force usage of the corresponding OpenACC device type. - setenv ACC_DEVICE_TYPE $openacc_device_type - - # To get better test coverage for device-specific code that is only - # ever used in offloading configurations, we'd like more thorough - # testing for test cases that deal with offloading, which most of all - # OpenACC test cases are. We enable torture testing, but limit it to - # -O0 and -O2 only, to avoid testing times exploding too much, under - # the assumption that between -O0 and -O[something] there is the - # biggest difference in the overall structure of the generated code. - switch -glob $offload_target { - disable { - set-torture-options [list \ - { -O2 } ] - } - default { - set-torture-options [list \ - { -O0 } \ - { -O2 } ] + radeon { + if { ![check_effective_target_openacc_radeon_accel_present] } { + # Don't bother; execution testing is going to FAIL. + untested "$subdir $offload_target offloading: supported, but hardware not accessible" + continue } - } - gcc-dg-runtest $tests "$tagopt" "$libstdcxx_includes" + set acc_mem_shared 0 + } + default { + error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)" + } + } + set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared" + + # To avoid compilation overhead, and to keep simple '-foffload=[...]' + # handling in test cases, by default only build for the offload target + # that we're actually going to test. + set tagopt "$tagopt -foffload=$offload_target" + # Force usage of the corresponding OpenACC device type. + setenv ACC_DEVICE_TYPE $openacc_device_type + + # To get better test coverage for device-specific code that is only + # ever used in offloading configurations, we'd like more thorough + # testing for test cases that deal with offloading, which most of all + # OpenACC test cases are. We enable torture testing, but limit it to + # -O0 and -O2 only, to avoid testing times exploding too much, under + # the assumption that between -O0 and -O[something] there is the + # biggest difference in the overall structure of the generated code. + switch -glob $offload_target { + disable { + set-torture-options [list \ + { -O2 } ] + } + default { + set-torture-options [list \ + { -O0 } \ + { -O2 } ] + } } - unset offload_target -} else { - # Call this once, which placates the subsequent torture-finish. - set-torture-options [list \ - { INVALID } ] + + gcc-dg-runtest $tests "$tagopt" "" } +unset offload_target -# See above. -set GCC_UNDER_TEST "$SAVE_GCC_UNDER_TEST" +if { $blddir != "" } { + unset libstdc++_library_path + unset lang_library_paths +} # All done. torture-finish diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c index 61c8109..02fbfda 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/present-1.c @@ -48,5 +48,5 @@ main (int argc, char **argv) } /* { dg-output "CheCKpOInT(\n|\r\n|\r).*" } */ -/* { dg-output "present clause: !acc_is_present" } */ +/* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" } */ /* { dg-shouldfail "" } */ diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp index 4bb2b2a..8ca0c81 100644 --- a/libgomp/testsuite/libgomp.oacc-c/c.exp +++ b/libgomp/testsuite/libgomp.oacc-c/c.exp @@ -1,16 +1,5 @@ # This whole file adapted from libgomp.c/c.exp. -if [info exists lang_library_path] then { - unset lang_library_path - unset lang_link_flags -} -if [info exists lang_test_file] then { - unset lang_test_file -} -if [info exists lang_include_flags] then { - unset lang_include_flags -} - load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp load_gcc_lib torture-options.exp @@ -26,6 +15,8 @@ proc check_effective_target_c++ { } { dg-init torture-init +lappend ALWAYS_CFLAGS "compiler=$GCC_UNDER_TEST" + # Turn on OpenACC. lappend ALWAYS_CFLAGS "additional_flags=-fopenacc" diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran-torture_execute_math.f90 b/libgomp/testsuite/libgomp.oacc-fortran/fortran-torture_execute_math.f90 new file mode 100644 index 0000000..edc940c --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran-torture_execute_math.f90 @@ -0,0 +1,4 @@ +! { dg-do run } +!TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} } + +include '../../../gcc/testsuite/gfortran.fortran-torture/execute/math.f90' diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp index 7365b32..e5844ad 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp +++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp @@ -4,18 +4,33 @@ load_lib libgomp-dg.exp load_gcc_lib gcc-dg.exp load_gcc_lib gfortran-dg.exp -global shlib_ext -global ALWAYS_CFLAGS +if { $blddir != "" } { + set libgfortran_library_path "../libgfortran/.libs" + set shlib_ext [get_shlib_extension] + if { ![file exists "${blddir}/${libgfortran_library_path}/libgfortran.a"] + && ![file exists "${blddir}/${libgfortran_library_path}/libgfortran.${shlib_ext}"] } { + verbose -log "No libgfortran library found, will not execute fortran tests" + unset libgfortran_library_path + return + } + lappend lang_library_paths $libgfortran_library_path -set shlib_ext [get_shlib_extension] -set lang_library_path "../libgfortran/.libs" -set lang_link_flags "-lgfortran -foffload=-lgfortran" -if [info exists lang_include_flags] then { - unset lang_include_flags + set libquadmath_library_path "../libquadmath/.libs" + if { [file exists "${blddir}/${libquadmath_library_path}/libquadmath.a"] + || [file exists "${blddir}/${libquadmath_library_path}/libquadmath.${shlib_ext}"] } { + lappend lang_library_paths $libquadmath_library_path + } else { + set libquadmath_library_path "" + } +} elseif { ![info exists GFORTRAN_UNDER_TEST] } { + verbose -log "GFORTRAN_UNDER_TEST not defined, will not execute fortran tests" + return } -set lang_test_file_found 0 -set quadmath_library_path "../libquadmath/.libs" - +if { $blddir != "" } { + set lang_source_re {^.*\.[fF](|90|95|03|08)$} + set lang_include_flags "-fintrinsic-modules-path=${blddir}" +} +lappend ALWAYS_CFLAGS "compiler=$GFORTRAN_UNDER_TEST" # Initialize dg. dg-init @@ -23,105 +38,79 @@ dg-init # Turn on OpenACC. lappend ALWAYS_CFLAGS "additional_flags=-fopenacc" + +# Gather a list of all tests. +set tests [lsort [find $srcdir/$subdir *.\[fF\]{,90,95,03,08}]] + +set ld_library_path $always_ld_library_path if { $blddir != "" } { - set lang_source_re {^.*\.[fF](|90|95|03|08)$} - set lang_include_flags "-fintrinsic-modules-path=${blddir}" - # Look for a static libgfortran first. - if [file exists "${blddir}/${lang_library_path}/libgfortran.a"] { - set lang_test_file "${lang_library_path}/libgfortran.a" - set lang_test_file_found 1 - # We may have a shared only build, so look for a shared libgfortran. - } elseif [file exists "${blddir}/${lang_library_path}/libgfortran.${shlib_ext}"] { - set lang_test_file "${lang_library_path}/libgfortran.${shlib_ext}" - set lang_test_file_found 1 - } else { - puts "No libgfortran library found, will not execute fortran tests" + append ld_library_path ":${blddir}/${libgfortran_library_path}" + + if { $libquadmath_library_path != "" } { + append ld_library_path ":${blddir}/${libquadmath_library_path}" } -} elseif [info exists GFORTRAN_UNDER_TEST] { - set lang_test_file_found 1 - # Needs to exist for libgomp.exp. - set lang_test_file "" -} else { - puts "GFORTRAN_UNDER_TEST not defined, will not execute fortran tests" } - -if { $lang_test_file_found } { - # Gather a list of all tests. - set tests [lsort [find $srcdir/$subdir *.\[fF\]{,90,95,03,08}]] - - if { $blddir != "" } { - if { [file exists "${blddir}/${quadmath_library_path}/libquadmath.a"] - || [file exists "${blddir}/${quadmath_library_path}/libquadmath.${shlib_ext}"] } { - lappend ALWAYS_CFLAGS "ldflags=-L${blddir}/${quadmath_library_path}/" - # Allow for spec subsitution. - lappend ALWAYS_CFLAGS "additional_flags=-B${blddir}/${quadmath_library_path}/" - set ld_library_path "$always_ld_library_path:${blddir}/${lang_library_path}:${blddir}/${quadmath_library_path}" - append lang_link_flags " -lquadmath" - } else { - set ld_library_path "$always_ld_library_path:${blddir}/${lang_library_path}" +append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] +set_ld_library_path_env_vars + +# Test with all available offload targets, and with offloading disabled. +foreach offload_target [concat [split $offload_targets ","] "disable"] { + global openacc_device_type + set openacc_device_type [offload_target_to_openacc_device_type $offload_target] + set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" + + switch $openacc_device_type { + "" { + unsupported "$subdir $offload_target offloading" + continue } - } else { - set ld_library_path "$always_ld_library_path" - if { [check_no_compiler_messages has_libquadmath executable { - int main() {return 0;} - } "-lgfortran -lquadmath"] } then { - append lang_link_flags " -lquadmath" - } - } - append ld_library_path [gcc-set-multilib-library-path $GCC_UNDER_TEST] - set_ld_library_path_env_vars - - # Test with all available offload targets, and with offloading disabled. - foreach offload_target [concat [split $offload_targets ","] "disable"] { - global openacc_device_type - set openacc_device_type [offload_target_to_openacc_device_type $offload_target] - set tagopt "-DACC_DEVICE_TYPE_$openacc_device_type=1" - - switch $openacc_device_type { - "" { - unsupported "$subdir $offload_target offloading" + host { + set acc_mem_shared 1 + } + nvidia { + if { ![check_effective_target_openacc_nvidia_accel_present] } { + # Don't bother; execution testing is going to FAIL. + untested "$subdir $offload_target offloading: supported, but hardware not accessible" continue } - host { - set acc_mem_shared 1 - } - nvidia { - if { ![check_effective_target_openacc_nvidia_accel_present] } { - # Don't bother; execution testing is going to FAIL. - untested "$subdir $offload_target offloading: supported, but hardware not accessible" - continue - } - - set acc_mem_shared 0 - } - radeon { - if { ![check_effective_target_openacc_radeon_accel_present] } { - # Don't bother; execution testing is going to FAIL. - untested "$subdir $offload_target offloading: supported, but hardware not accessible" - continue - } - - set acc_mem_shared 0 - } - default { - error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)" + + set acc_mem_shared 0 + } + radeon { + if { ![check_effective_target_openacc_radeon_accel_present] } { + # Don't bother; execution testing is going to FAIL. + untested "$subdir $offload_target offloading: supported, but hardware not accessible" + continue } + + set acc_mem_shared 0 + } + default { + error "Unknown OpenACC device type: $openacc_device_type (offload target: $offload_target)" } - set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared" - - # To avoid compilation overhead, and to keep simple '-foffload=[...]' - # handling in test cases, by default only build for the offload target - # that we're actually going to test. - set tagopt "$tagopt -foffload=$offload_target" - # Force usage of the corresponding OpenACC device type. - setenv ACC_DEVICE_TYPE $openacc_device_type - - # For Fortran we're doing torture testing, as Fortran has far more tests - # with arrays etc. that testing just -O0 or -O2 is insufficient, that is - # typically not the case for C/C++. - gfortran-dg-runtest $tests "$tagopt" "" } - unset offload_target + set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared" + + # To avoid compilation overhead, and to keep simple '-foffload=[...]' + # handling in test cases, by default only build for the offload target + # that we're actually going to test. + set tagopt "$tagopt -foffload=$offload_target" + # Force usage of the corresponding OpenACC device type. + setenv ACC_DEVICE_TYPE $openacc_device_type + + # For Fortran we're doing torture testing, as Fortran has far more tests + # with arrays etc. that testing just -O0 or -O2 is insufficient, that is + # typically not the case for C/C++. + gfortran-dg-runtest $tests "$tagopt" "" +} +unset offload_target + +if { $blddir != "" } { + unset lang_source_re + unset lang_include_flags + unset libgfortran_library_path + unset libquadmath_library_path + unset lang_library_paths } # All done. diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90 new file mode 100644 index 0000000..d3cbebe --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-2.f90 @@ -0,0 +1,35 @@ +! { dg-do run } + +implicit none + +type t +integer :: foo +integer, pointer :: bar +end type t + +type(t) :: var +integer, target :: tgt + +var%bar => tgt + +var%foo = 99 +tgt = 199 + +!$acc enter data copyin(var, tgt) + +!$acc enter data attach(var%bar) + +!$acc serial +! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 } +var%foo = 5 +var%bar = 7 +!$acc end serial + +!$acc exit data detach(var%bar) + +!$acc exit data copyout(var, tgt) + +if (var%foo.ne.5) stop 1 +if (tgt.ne.7) stop 2 + +end diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90 new file mode 100644 index 0000000..a25b1a8 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-3.f90 @@ -0,0 +1,35 @@ +! { dg-do run } + +implicit none + +type t +integer :: foo +integer, pointer :: bar(:) +end type t + +type(t) :: var +integer, target :: tgt(20) + +var%bar => tgt + +var%foo = 99 +tgt = 199 + +!$acc enter data copyin(var, tgt) + +!$acc enter data attach(var%bar) + +!$acc serial +! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 } +var%foo = 5 +var%bar = 7 +!$acc end serial + +!$acc exit data detach(var%bar) + +!$acc exit data copyout(var, tgt) + +if (var%foo.ne.5) stop 1 +if (any(tgt.ne.7)) stop 2 + +end diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90 new file mode 100644 index 0000000..3198a0b --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622-4.f90 @@ -0,0 +1,47 @@ +! { dg-do run } + +use openacc +implicit none + +type t +integer :: foo +character(len=8), pointer :: bar +character(len=4), allocatable :: qux +end type t + +type(t) :: var +character(len=8), target :: tgt + +allocate(var%qux) + +var%bar => tgt + +var%foo = 99 +tgt = "Octopus!" +var%qux = "Fish" + +!$acc enter data copyin(var, tgt) + +! Avoid automatic attach (i.e. with "enter data") +call acc_copyin (var%qux) + +!$acc enter data attach(var%bar, var%qux) + +!$acc serial +! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 } +var%foo = 5 +var%bar = "Plankton" +var%qux = "Pond" +!$acc end serial + +!$acc exit data detach(var%bar, var%qux) + +call acc_copyout (var%qux) + +!$acc exit data copyout(var, tgt) + +if (var%foo.ne.5) stop 1 +if (tgt.ne."Plankton") stop 2 +if (var%qux.ne."Pond") stop 3 + +end diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90 new file mode 100644 index 0000000..a17c4f6 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr109622.f90 @@ -0,0 +1,35 @@ +! { dg-do run } + +implicit none + +type t +integer :: value +type(t), pointer :: chain +end type t + +type(t), target :: var, var2 + +var%value = 99 +var2%value = 199 + +var%chain => var2 +nullify(var2%chain) + +!$acc enter data copyin(var, var2) + +!$acc enter data attach(var%chain) + +!$acc serial +! { dg-warning "using .vector_length \\(32\\)., ignoring 1" "" { target openacc_nvidia_accel_selected } .-1 } +var%value = 5 +var%chain%value = 7 +!$acc end serial + +!$acc exit data detach(var%chain) + +!$acc exit data copyout(var, var2) + +if (var%value.ne.5) stop 1 +if (var2%value.ne.7) stop 2 + +end |