summaryrefslogtreecommitdiff
path: root/libgomp
AgeCommit message (Collapse)Author
2020-04-29[OpenACC] Set 'acc_device_current = -1'Thomas Schwinge
There's no point in using value '-3', and even though not directly related, value '-1' does match 'GOMP_DEVICE_ICV'. libgomp/ * config/accel/openacc.f90 (acc_device_current): Set to '-1'. * openacc.f90 (acc_device_current): Likewise. * openacc.h (acc_device_current): Likewise. * openacc_lib.h (acc_device_current): Likewise.
2020-04-29[gcn] Set 'UI_NONE' for 'TARGET_EXCEPT_UNWIND_INFO' [PR94282]Thomas Schwinge
In libgomp offloading testing, this resolves all the 'ld: error: undefined symbol: __gxx_personality_v0' FAILs. gcc/ PR target/94282 * common/config/gcn/gcn-common.c (gcn_except_unwind_info): New function. (TARGET_EXCEPT_UNWIND_INFO): Define. libgomp/ PR target/94282 * testsuite/libgomp.c-c++-common/function-not-offloaded.c: Remove 'dg-allow-blank-lines-in-output'.
2020-04-29[gcn] Use 'radeon' for the environment variable 'ACC_DEVICE_TYPE'Thomas Schwinge
..., per OpenACC 3.0, A.1.2. "AMD GPU Targets". This complements commit 6687d13a87c42dddc7d1c7adade38d31ba0d1401 "Rename acc_device_gcn to acc_device_radeon". libgomp/ * oacc-init.c (get_openacc_name): Handle 'gcn'. * testsuite/lib/libgomp.exp (offload_target_to_openacc_device_type) [amdgcn*]: Return 'radeon'. Adjust all users. (check_effective_target_openacc_amdgcn_accel_present): Rename to... (check_effective_target_openacc_radeon_accel_present): ... this. Adjust all users. (check_effective_target_openacc_amdgcn_accel_selected): Rename to... (check_effective_target_openacc_radeon_accel_selected): ... this. Adjust all users.
2020-04-29Torture testing: 'libgomp.fortran/use_device_ptr-optional-2.f90'Thomas Schwinge
Fix-up for commit a2c26c50310a336361d8129ecdd43d3001d6cb3a (r278046) "Fortran] Support absent optional args with use_device_{ptr,addr}". libgomp/ * testsuite/libgomp.fortran/use_device_ptr-optional-2.f90: Add 'dg-do run'.
2020-04-23amdgcn: Check HSA return codes [PR94629]Andrew Stubbs
Ensure that the returned status values are not ignored. The old code was not broken, but this is both safer and satisfies static analysis. 2020-04-23 Andrew Stubbs <ams@codesourcery.com> PR other/94629 libgomp/ * plugin/plugin-gcn.c (init_hsa_context): Check return value from hsa_iterate_agents. (GOMP_OFFLOAD_init_device): Check return values from both calls to hsa_agent_iterate_regions.
2020-04-20Add 'dg-do run' to 'libgomp.fortran/target-enter-data-2.F90'Thomas Schwinge
Fix-up for commit af557050fd011a03d21dc26b31959033061a0443 "[OpenMP] Fix 'omp exit data' for Fortran arrays (PR 94635)". libgomp/ PR middle-end/94635 * testsuite/libgomp.fortran/target-enter-data-2.F90: Add 'dg-do run'.
2020-04-20Fix declare copyout in libgomp.oacc-c++/declare-pr94120.CTobias Burnus
Testing on the host does not make sense for 'declare copyout' for a same-scope stack-allocated variable. Once the copyout is done, the variable is gone. Hence, test the variable on the device. This can be revisit after the OpenACC semantic has been fixed; but with that fix, the test PASSes again with devices. PR middle-end/94120 * testsuite/libgomp.oacc-c++/declare-pr94120.C: Fix 'declare copy(out)' test case.
2020-04-17[OpenMP] Fix 'omp exit data' for Fortran arrays (PR 94635)Tobias Burnus
PR middle-end/94635 * gimplify.c (gimplify_scan_omp_clauses): Turn MAP_TO_PSET to MAP_DELETE. PR middle-end/94635 * testsuite/libgomp.fortran/target-enter-data-2.F90: New.
2020-04-13Rename 'libgomp.oacc-c-c++-common/static-dynamic-lifetimes-*' to ↵Thomas Schwinge
'libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-*' [PR92843] Fix-up for commit be9862dd96945772ae0692bc95b37ec6dbcabda0 "Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]": it's "structured", not "static" data lifetimes/reference counters. libgomp/ PR libgomp/92843 * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-1.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-2.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-3.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-4.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-5.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-6.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-7.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8-lib.c: ... this. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c:: Rename to... * testsuite/libgomp.oacc-c-c++-common/structured-dynamic-lifetimes-8.c: ... this.
2020-04-10Test cases for mixed structured/dynamic data lifetimes with OpenACC [PR92843]Julian Brown
libgomp/ PR libgomp/92843 * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1-lib.c: New file. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-dynamic-lifetimes-8.c: Likewise.
2020-04-10Add 'dg-do run' to 'libgomp.fortran/target-enter-data-1.f90'Thomas Schwinge
Fix-up for commit 689418b97e5eb6a221871a2439bca3e6283ac579 "libgomp – fix handling of 'target enter data'". libgomp/ * testsuite/libgomp.fortran/target-enter-data-1.f90: Add 'dg-do run'.
2020-04-08[C/C++, OpenACC] Reject vars of different scope in acc declare (PR94120)Tobias Burnus
gcc/c/ PR middle-end/94120 * c-decl.c (c_check_in_current_scope): New function. * c-tree.h (c_check_in_current_scope): Declare it. * c-parser.c (c_parser_oacc_declare): Add check that variables are declared in the same scope as the directive. Fix handling of namespace vars. gcc/cp/ PR middle-end/94120 * paser.c (cp_parser_oacc_declare): Add check that variables are declared in the same scope as the directive. gcc/testsuite/ PR middle-end/94120 * c-c++-common/goacc/declare-pr94120.c: New. * g++.dg/declare-pr94120.C: New. libgomp/testsuite/ PR middle-end/94120 * libgomp.oacc-c++/declare-pr94120.C: New.
2020-04-06libgomp/test: Remove a build sysroot fix regressionMaciej W. Rozycki
Fix a problem with commit c8e759b4215b ("libgomp/test: Fix compilation for build sysroot") that caused a regression in some standalone test environments where testsuite/libgomp-test-support.exp is used, but the compiler is expected to be determined by `[find_gcc]', and set the GCC_UNDER_TEST TCL variable in testsuite/libgomp-site-extra.exp instead. libgomp/ * configure.ac: Add testsuite/libgomp-site-extra.exp to output files. * configure: Regenerate. * testsuite/libgomp-site-extra.exp.in: New file. * testsuite/libgomp-test-support.exp.in (GCC_UNDER_TEST): Remove variable. * testsuite/Makefile.am (EXTRA_DEJAGNU_SITE_CONFIG): New variable. * testsuite/Makefile.in: Regenerate.
2020-04-03Revert "[nvptx, libgomp] Update pr85381-{2,4}.c test-cases" [PR89713, PR94392]Thomas Schwinge
In response to PR94392 commit 75efe9cb1f8938a713ce540dc3b27bc2afcd3fae "c/94392 - only enable -ffinite-loops for C++", this reverts PR89713 commit 00908992f2a78f213d227aea8dbab014a1361df0, as apparently now again "empty oacc loops are" no longer "removed before expand". libgomp/ PR tree-optimization/89713 PR c/94392 * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Again expect 'bar.sync'. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
2020-03-31libgomp – fix handling of 'target enter data'Tobias Burnus
* target.c (GOMP_target_enter_exit_data): Handle PSET/MAP_POINTER. * testsuite/libgomp.fortran/target-enter-data-1.f90: New.
2020-03-24Fix OpenMP offload handling for target-link variables for nvptx (PR81689)Tobias Burnus
PR libgomp/81689 * lto.c (offload_handle_link_vars): Propagate TREE_PUBLIC state. PR libgomp/81689 * omp-offload.c (omp_finish_file): Fix target-link handling if targetm_common.have_named_sections is false. PR libgomp/81689 * testsuite/libgomp.c/target-link-1.c: Remove xfail.
2020-03-23libgomp – fix declare target link handling (PR94251)Tobias Burnus
PR libgomp/94251 * target.c (gomp_load_image_to_device): Fix link variable handling.
2020-03-19c++: Fix up handling of captured vars in lambdas in OpenMP clauses [PR93931]Jakub Jelinek
Without the parser.c change we were ICEing on the testcase, because while the uses of the captured vars inside of the constructs were replaced with capture proxy decls, we didn't do that for decls in OpenMP clauses. With that fixed, we don't ICE anymore, but the testcase is miscompiled and FAILs at runtime. This is because the capture proxy decls have DECL_VALUE_EXPR and during gimplification we were gimplifying those to their DECL_VALUE_EXPRs. That is fine for shared vars, but for privatized ones we must not do that. So that is what the cp-gimplify.c changes do. Had to add a DECL_CONTEXT check before calling is_capture_proxy because some VAR_DECLs don't have DECL_CONTEXT set (yet) and is_capture_proxy relies on that being non-NULL always. 2020-03-19 Jakub Jelinek <jakub@redhat.com> PR c++/93931 * parser.c (cp_parser_omp_var_list_no_open): Call process_outer_var_ref on outer_automatic_var_p decls. * cp-gimplify.c (cxx_omp_disregard_value_expr): Return true also for capture proxy decls. * testsuite/libgomp.c++/pr93931.C: New test.
2020-03-19libgomp/testsuite: ignore blank-line output for function-not-offloaded.cTobias Burnus
* testsuite/libgomp.c-c++-common/function-not-offloaded.c: Add dg-allow-blank-lines-in-output.
2020-03-18Fix libgomp.oacc-fortran/atomic_capture-1.f90Tobias Burnus
2020-03-18 Julian Brown <julian@codesourcery.com> Tobias Burnus <tobias@codesourcery.com> * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90: Really make it work concurrently.
2020-03-18libgomp testsuite - disable long double for AMDGCNTobias Burnus
* testsuite/libgomp.oacc-c++/firstprivate-mappings-1.C: Add #define DO_LONG_DOUBLE; set to 1, except for nvidia + gcn. * libgomp.oacc-c-c++-common/firstprivate-mappings-1.c: Likewise. * g++.dg/goacc/firstprivate-mappings-1.C: Only set DO_LONG_DOUBLE if not defined; update comments. * c-c++-common/goacc/firstprivate-mappings-1.c: Likewise.
2020-03-15tree-nested: Fix handling of *reduction clauses with C array sections [PR93566]Jakub Jelinek
tree-nested.c didn't handle C array sections in {,task_,in_}reduction clauses. 2020-03-14 Jakub Jelinek <jakub@redhat.com> PR middle-end/93566 * tree-nested.c (convert_nonlocal_omp_clauses, convert_local_omp_clauses): Handle {,in_,task_}reduction clauses with C/C++ array sections. * testsuite/libgomp.c/pr93566.c: New test.
2020-02-21Adapt libgomp acc_get_property.f90 testFrederik Harwath
The commit r10-6721-g8d1a1cb1b816381bf60cb1211c93b8eba1fe1472 has changed the name of the type that is used for the return value of the Fortran acc_get_property function without adapting the test acc_get_property.f90. 2020-02-21 Frederik Harwath <frederik@codesourcery.com> * testsuite/libgomp.oacc-fortran/acc_get_property.f90: Adapt to changes from 2020-02-19, i.e. use integer(c_size_t) instead of integer(acc_device_property) for the type of the return value of acc_get_property.
2020-02-19libgomp: Fixes + cleanup for OpenACC's Fortran module + openacc_lib.hTobias Burnus
2020-02-19 Tobias Burnus <tobias@codesourcery.com> * .gitattributes: New; whitespace handling for Fortran's openacc_lib.h. * config/accel/openacc.f90 (openacc_kinds): Add acc_device_current. (openacc_internal, acc_on_device_h): Fix argument name; minor cleanup. * libgomp.texi (Enabling OpenACC): No longer mark as experimental. (acc_set_device_num): Fix Fortran argument name, use same name for C. (acc_get_property): Update Fortran interface to post-OpenACC 3.0 corrections; add note about the previous interface and named constant. (OpenACC library and environment variables): Fix two typos. * openacc.f90: Use for all procedures the argument names from the spec as for …_h they are user visible. (openacc_kinds): Rename acc_device_property to acc_device_property_kinds and change value to int32 ; and update users. Re-add acc_device_property for for backward compatibility. (acc_get_property_string_h): Clean up as acc_device_property_kind changed. (acc_get_property_h): Likewise and return c_size_t instead of acc_device_property. (openacc): Also export acc_device_property_kinds. (acc_async_test_h, acc_async_test_all_h, acc_on_device_h, acc_is_present_32_h, acc_is_present_64_h): Simplify logical-return-value handling; check against /= 0 instead of == 1 to match C. * openacc_lib.h: Use for all procedures the argument names from the spec as for …_h they are user visible. Place !GCC$ into the first column to be active also for fixed-form souce form. (acc_device_current, acc_device_property_kind, acc_device_property, acc_property_memory, acc_property_free_memory, acc_property_name, acc_property_vendor, acc_property_driver): New named constants. (acc_get_property, acc_get_property_string): New generic interface.
2020-02-13Add ChangeLog entry for my last commitFrederik Harwath
2020-02-13openmp: ignore nowait if async execution is unsupported [PR93481]Frederik Harwath
An OpenMP "nowait" clause on a target construct currently leads to a call to GOMP_OFFLOAD_async_run in the plugin that is used for offloading at execution time. The nvptx plugin contains only a stub of this function that always produces a fatal error if called. This commit changes the "nowait" implementation to ignore the clause if the executing device's plugin does not implement GOMP_OFFLOAD_async_run. The stub in the nvptx plugin is removed which effectively means that programs containing "nowait" can now be executed with nvptx offloading as if the clause had not been used. This behavior is consistent with the OpenMP specification which says that "[...] execution of the target task *may* be deferred" (emphasis added), cf. OpenMP 5.0, page 172. libgomp/ * plugin/plugin-nvptx.c: Remove GOMP_OFFLOAD_async_run stub. * target.c (gomp_load_plugin_for_device): Make "async_run" loading optional. (gomp_target_task_fn): Assert "devicep->async_run_func". (clear_unsupported_flags): New function to remove unsupported flags (right now only GOMP_TARGET_FLAG_NOWAIT) that can be be ignored. (GOMP_target_ext): Apply clear_unsupported_flags to flags. * testsuite/libgomp.c/target-33.c: Remove xfail for offload_target_nvptx. * testsuite/libgomp.c/target-34.c: Likewise.
2020-02-10Add xfails to libgomp tests target-{33,34}.c, target-link-1.cFrederik Harwath
Add xfails for nvptx offloading because "no GOMP_OFFLOAD_async_run implemented in plugin-nvptx.c" (https://gcc.gnu.org/PR81688) and because "omp target link not implemented for nvptx" (https://gcc.gnu.org/PR81689). libgomp/ * testsuite/libgomp.c/target-33.c: Add xfail for execution on offload_target_nvptx, cf. https://gcc.gnu.org/PR81688. * testsuite/libgomp.c/target-34.c: Likewise. * testsuite/libgomp.c/target-link-1.c: Add xfail for offload_target_nvptx, cf. https://gcc.gnu.org/PR81689.
2020-02-09openmp: Optimize DECL_IN_CONSTANT_POOL vars in target regionsJakub Jelinek
DECL_IN_CONSTANT_POOL are shared and thus don't really get emitted in the BLOCK where they are used, so for OpenMP target regions that have initializers gimplified into copying from them we actually map them at runtime from host to offload devices. This patch instead marks them as "omp declare target", so that they are on the target device from the beginning and don't need to be copied there. 2020-02-09 Jakub Jelinek <jakub@redhat.com> * gimplify.c (gimplify_adjust_omp_clauses_1): Promote DECL_IN_CONSTANT_POOL variables into "omp declare target" to avoid copying them around between host and target. * testsuite/libgomp.c/target-38.c: New test.
2020-02-06openmp: Fix handling of non-addressable shared scalars in parallel nested ↵Jakub Jelinek
inside of target [PR93515] As the following testcase shows, we need to consider even target to be a construct that forces not to use copy in/out for shared on parallel inside of the target. E.g. for parallel nested inside another parallel or host teams, we already avoid copy in/out and we need to treat target the same. 2020-02-06 Jakub Jelinek <jakub@redhat.com> PR libgomp/93515 * omp-low.c (use_pointer_for_field): For nested constructs, also look for map clauses on target construct. (scan_omp_1_stmt) <case GIMPLE_OMP_TARGET>: Bump temporarily taskreg_nesting_level. * testsuite/libgomp.c-c++-common/pr93515.c: New test.
2020-02-05[libgomp] – Fix check_effective_target_offload_target_nvptx for remote ↵Tobias Burnus
execution * testsuite/lib/libgomp.exp (check_effective_target_offload_target_nvptx): Pass flags as 'options' and not as 'source' argument to libgomp_target_compile.
2020-02-03Remove gfx801 "carrizo" supportAndrew Stubbs
2020-02-03 Andrew Stubbs <ams@codesourcery.com> gcc/ * config.gcc: Remove "carrizo" support. * config/gcn/gcn-opts.h (processor_type): Likewise. * config/gcn/gcn.c (gcn_omp_device_kind_arch_isa): Likewise. * config/gcn/gcn.opt (gpu_type): Likewise. * config/gcn/t-omp-device: Likewise. libgomp/ * plugin/plugin-gcn.c (EF_AMDGPU_MACH_AMDGCN_GFX801): Remove. (gcn_gfx801_s): Remove. (isa_hsa_name): Remove gfx801. (isa_gcc_name): Remove gfx801/carizzo. (isa_code): Remove gfx801.
2020-02-03[OpenACC] bump version for 2.6 plus libgomp.texi updateTobias Burnus
2020-02-03 Julian Brown <julian@codesourcery.com> Tobias Burnus <tobias@codesourcery.com> gcc/c-family/ * c-cppbuiltin.c (c_cpp_builtins): Update _OPENACC define to 201711. gcc/ * doc/invoke.texi: Update mention of OpenACC version to 2.6. gcc/fortran/ * cpp.c (cpp_define_builtins): Update _OPENACC define to 201711. * intrinsic.texi: Update mentions of OpenACC version to 2.6. * gfortran.texi: Likewise. Remove experimental disclamer for OpenACC. * invoke.texi: Remove experimental disclamer for OpenACC. gcc/testsuite/ * c-c++-common/cpp/openacc-define-3.c: Update expected value for _OPENACC define. * gfortran.dg/openacc-define-3.f90: Likewise. libgomp/ * libgomp.texi (OpenACC Runtime Library Routines): Document *_async and *_finalize variants; document acc_attach and acc_detach; update references from OpenACC 2.0 to 2.6. * openacc.f90 (openacc_version): Update to 201711. * openacc_lib.h (openacc_version): Update to 201711. * testsuite/libgomp.oacc-fortran/openacc_version-1.f: Update expected openacc_version to 201711. * testsuite/libgomp.oacc-fortran/openacc_version-2.f90: Likewise.
2020-01-31[amdgcn] Scale number of threads/workers with VGPR usageKwok Cheung Yeung
2020-01-31 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/ * config/gcn/mkoffload.c (process_asm): Add sgpr_count and vgpr_count to definition of hsa_kernel_description. Parse assembly to find SGPR and VGPR count of kernel and store in hsa_kernel_description. libgomp/ * plugin/plugin-gcn.c (struct hsa_kernel_description): Add sgpr_count and vgpr_count fields. (struct kernel_info): Add a field for a hsa_kernel_description. (run_kernel): Reduce the number of threads/workers if the requested number would require too many VGPRs. (init_basic_kernel_info): Initialize description field with the hsa_kernel_description entry for the kernel.
2020-01-30Skip plugin-{gcn,hsa} for (-m)x32 (PR bootstrap/93409)Tobias Burnus
PR bootstrap/93409 * plugin/configfrag.ac (enable_offload_targets): Skip HSA and GCN plugin besides -m32 also for -mx32. * configure: Regenerate.
2020-01-29Add acc_device_radeon to name_of_acc_device_t functionFrederik Harwath
libgomp/ * oacc-init.c (name_of_acc_device_t): Handle acc_device_radeon. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
2020-01-29Adjust formatting of acc_get_property testsFrederik Harwath
libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_get_property.c: Adjust to GNU coding style. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-nvptx.c: Likewise.
2020-01-29Add OpenACC acc_get_property support for AMD GCNFrederik Harwath
Add full support for the OpenACC 2.6 acc_get_property and acc_get_property_string functions to the libgomp GCN plugin. libgomp/ * plugin-gcn.c (struct agent_info): Add fields "name" and "vendor_name" ... (GOMP_OFFLOAD_init_device): ... and init from here. (struct hsa_context_info): Add field "driver_version_s" ... (init_hsa_contest): ... and init from here. (GOMP_OFFLOAD_openacc_get_property): Replace stub with a proper implementation. * testsuite/libgomp.oacc-c-c++-common/acc_get_property.c: Enable test execution for amdgcn and host offloading targets. * testsuite/libgomp.oacc-fortran/acc_get_property.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c (expect_device_properties): Split function into ... (expect_device_string_properties): ... this new function ... (expect_device_memory): ... and this new function. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-gcn.c: Add test.
2020-01-28Don't allow mixed component and non-component accesses for OpenACC/FortranJulian Brown
gcc/fortran/ * gfortran.h (gfc_symbol): Add comp_mark bitfield. * openmp.c (resolve_omp_clauses): Disallow mixed component and full-derived-type accesses to the same variable within a single directive. libgomp/ * testsuite/libgomp.oacc-fortran/deep-copy-2.f90: Remove test from here. * testsuite/libgomp.oacc-fortran/deep-copy-3.f90: Don't use mixed component/non-component variable refs in a single directive. * testsuite/libgomp.oacc-fortran/classtypes-1.f95: Likewise. gcc/testsuite/ * gfortran.dg/goacc/deep-copy-2.f90: Move test here (from libgomp testsuite). Make a compilation test, and expect rejection of mixed component/non-component accesses. * gfortran.dg/goacc/mapping-tests-1.f90: New test.
2020-01-24Add `--with-toolexeclibdir=' configuration optionMaciej W. Rozycki
Provide means, in the form of a `--with-toolexeclibdir=' configuration option, to override the default installation directory for target libraries, otherwise known as $toolexeclibdir. This is so that it is possible to get newly-built libraries, particularly the shared ones, installed in a common place, so that they can be readily used by the target system as their host libraries, possibly over NFS, without a need to manually copy them over from the currently hardcoded location they would otherwise be installed in. In the presence of the `--enable-version-specific-runtime-libs' option and for configurations building native GCC the option is ignored. config/ * toolexeclibdir.m4: New file. gcc/ * doc/install.texi (Cross-Compiler-Specific Options): Document `--with-toolexeclibdir' option. libada/ * Makefile.in (configure_deps): Add `toolexeclibdir.m4'. * configure.ac: Handle `--with-toolexeclibdir='. * configure: Regenerate. libatomic/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * testsuite/Makefile.in: Regenerate. libffi/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * include/Makefile.in: Regenerate. * man/Makefile.in: Regenerate. * testsuite/Makefile.in: Regenerate. libgcc/ * Makefile.in (configure_deps): Add `toolexeclibdir.m4'. * configure.ac: Handle `--with-toolexeclibdir='. * configure: Regenerate. libgfortran/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. libgomp/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * testsuite/Makefile.in: Regenerate. libhsail-rt/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. libitm/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * testsuite/Makefile.in: Regenerate. libobjc/ * Makefile.in (aclocal_deps): Add `toolexeclibdir.m4'. * aclocal.m4: Include `toolexeclibdir.m4'. * configure.ac: Handle `--with-toolexeclibdir='. * configure: Regenerate. liboffloadmic/ * plugin/configure.ac: Handle `--with-toolexeclibdir='. * plugin/Makefile.in: Regenerate. * plugin/aclocal.m4: Regenerate. * plugin/configure: Regenerate. * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. libphobos/ * m4/druntime.m4: Handle `--with-toolexeclibdir='. * m4/Makefile.in: Regenerate. * libdruntime/Makefile.in: Regenerate. * src/Makefile.in: Regenerate. * testsuite/Makefile.in: Regenerate. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. libquadmath/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. libsanitizer/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * asan/Makefile.in: Regenerate. * interception/Makefile.in: Regenerate. * libbacktrace/Makefile.in: Regenerate. * lsan/Makefile.in: Regenerate. * sanitizer_common/Makefile.in: Regenerate. * tsan/Makefile.in: Regenerate. * ubsan/Makefile.in: Regenerate. libssp/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. libstdc++-v3/ * acinclude.m4: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * doc/Makefile.in: Regenerate. * include/Makefile.in: Regenerate. * libsupc++/Makefile.in: Regenerate. * po/Makefile.in: Regenerate. * python/Makefile.in: Regenerate. * src/Makefile.in: Regenerate. * src/c++11/Makefile.in: Regenerate. * src/c++17/Makefile.in: Regenerate. * src/c++98/Makefile.in: Regenerate. * src/filesystem/Makefile.in: Regenerate. * testsuite/Makefile.in: Regenerate. libvtv/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate. * testsuite/Makefile.in: Regenerate. zlib/ * configure.ac: Handle `--with-toolexeclibdir='. * Makefile.in: Regenerate. * aclocal.m4: Regenerate. * configure: Regenerate.
2020-01-24Add missing ChangeLog entry for my last commitFrederik Harwath
libgomp/ChangeLog: Add entry for commit 4bd03ed69bd.
2020-01-24Fix expectation and types in acc_get_property testsFrederik Harwath
* Weaken expectation concerning acc_property_free_memory. Do not expect the value returned by CUDA since that value might have changed in the meantime. * Use correct type for the results of calls to acc_get_property in tests. libgomp/ * testsuite/libgomp.oacc-c-c++-common/acc_get_property-aux.c (expect_device_properties): Remove "expected_free_mem" argument, change "expected_total_mem" argument type to size_t; change types of acc_get_property results to size_t, adapt format strings. * testsuite/libgomp.oacc-c-c++-common/acc_get_property.c: Use %zu instead of %zd to print size_t values. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-2.c: Adapt and rename to ... * testsuite/libgomp.oacc-c-c++-common/acc_get_property-nvptx.c: ... this. * testsuite/libgomp.oacc-c-c++-common/acc_get_property-3.c: Adapt and rename to ... * testsuite/libgomp.oacc-c-c++-common/acc_get_property-host.c: ... this. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com>
2020-01-23Fix libgomp plugin-gcn bugAndrew Stubbs
2020-01-23 Andrew Stubbs <ams@codesourcery.com> libgomp/ * plugin/plugin-gcn.c (parse_target_attributes): Use correct mask for the device id.
2020-01-21Add runtime ISA check for amdgcn offloadingFrederik Harwath
The HSA/ROCm runtime rejects binaries not built for the exact GPU device present. So far, the libgomp amdgcn plugin does not verify that the GPU ISA and the ISA specified at compile time match before handing over the binary to the runtime. In case of a mismatch, the user is confronted with an unhelpful runtime error. This commit implements a runtime ISA check. In case of an ISA mismatch, the execution is aborted with a clear error message and a hint at the correct compilation parameters for the GPU on which the execution has been attempted. libgomp/ * plugin/plugin-gcn.c (EF_AMDGPU_MACH): New enum. * (EF_AMDGPU_MACH_MASK): New constant. * (gcn_isa): New typedef. * (gcn_gfx801_s): New constant. * (gcn_gfx803_s): New constant. * (gcn_gfx900_s): New constant. * (gcn_gfx906_s): New constant. * (gcn_isa_name_len): New constant. * (elf_gcn_isa_field): New function. * (isa_hsa_name): New function. * (isa_gcc_name): New function. * (isa_code): New function. * (struct agent_info): Add field "device_isa" and remove field "gfx900_p". * (GOMP_OFFLOAD_init_device): Adapt agent init to "agent_info" field changes, fail if device has unknown ISA. * (parse_target_attributes): Replace "gfx900_p" by "device_isa". * (isa_matches_agent): New function ... * (create_and_finalize_hsa_program): ... used from here to check that the GPU ISA and the code-object ISA match.
2020-01-20Update OpenACC tests for amdgcnAndrew Stubbs
2020-01-20 Andrew Stubbs <ams@codesourcery.com> libgomp/ * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Skip test on gcn. * testsuite/libgomp.oacc-c-c++-common/loop-dim-default.c (main): Adjust test dimensions for amdgcn. * testsuite/libgomp.oacc-c-c++-common/loop-gwv-1.c (main): Adjust gang/worker/vector expectations dynamically. * testsuite/libgomp.oacc-c-c++-common/loop-red-gwv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-v-2.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-wv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-v-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-wv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c (acc_gang): Recognise acc_device_radeon. (acc_worker): Likewise. (acc_vector): Likewise. (main): Set expectations for amdgcn. * testsuite/libgomp.oacc-c-c++-common/routine-gwv-1.c (main): Adjust gang/worker/vector expectations dynamically. * testsuite/libgomp.oacc-c-c++-common/routine-v-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-1.c (main): Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Set expectations for amdgcn.
2020-01-17Rename acc_device_gcn to acc_device_radeonAndrew Stubbs
2020-01-17 Andrew Stubbs <ams@codesourcery.com> libgomp/ * config/accel/openacc.f90 (openacc_kinds): Rename acc_device_gcn to acc_device_radeon. (openacc): Likewise. * openacc.f90 (openacc_kinds): Likewise. (openacc): Likewise. * openacc.h (acc_device_t): Likewise. * openacc_lib.h: Likewise. * testsuite/lib/libgomp.exp (check_effective_target_openacc_amdgcn_accel_present): Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-init-1.c (cb_compute_construct_end): Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c (cb_enqueue_launch_start): Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-parallel-1.c (cb_enter_data_end): Likewise. (cb_exit_data_start): Likewise. (cb_exit_data_end): Likewise. (cb_compute_construct_end): Likewise. (cb_enqueue_launch_start): Likewise. (cb_enqueue_launch_end): Likewise. * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c (main): Likewise.
2020-01-10OpenACC 'acc_get_property' cleanupThomas Schwinge
include/ * gomp-constants.h (enum gomp_device_property): Remove. libgomp/ * libgomp-plugin.h (enum goacc_property): New. Adjust all users to use this instead of 'enum gomp_device_property'. (GOMP_OFFLOAD_get_property): Rename to... (GOMP_OFFLOAD_openacc_get_property): ... this. Adjust all users. * libgomp.h (struct gomp_device_descr): Move 'GOMP_OFFLOAD_openacc_get_property'... (struct acc_dispatch_t): ... here. Adjust all users. * plugin/plugin-hsa.c (GOMP_OFFLOAD_get_property): Remove. liboffloadmic/ * plugin/libgomp-plugin-intelmic.cpp (GOMP_OFFLOAD_get_property): Remove. From-SVN: r280150
2020-01-10Further changes for the OpenACC 'if_present' clause on the 'host_data' constructThomas Schwinge
gcc/ * tree.h (OMP_CLAUSE_USE_DEVICE_PTR_IF_PRESENT): New definition. * tree-core.h: Document it. * gimplify.c (gimplify_omp_workshare): Set it. * omp-low.c (lower_omp_target): Use it. * tree-pretty-print.c (dump_omp_clause): Print it. gcc/testsuite/ * c-c++-common/goacc/host_data-1.c: Extend. * gfortran.dg/goacc/host_data-tree.f95: Likewise. gcc/ * omp-low.c (lower_omp_target) <OMP_CLAUSE_USE_DEVICE_PTR etc.>: Assert that for OpenACC we always have 'GOMP_MAP_USE_DEVICE_PTR'. libgomp/ * target.c (gomp_map_vars_internal) <GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT>: Clean up/elaborate code paths. From-SVN: r280149
2020-01-10re PR libgomp/93219 (unused return value in affinity-fmt.c)Jakub Jelinek
PR libgomp/93219 * libgomp.h (gomp_print_string): Change return type from void to int. * affinity-fmt.c (gomp_print_string): Likewise. Return true if not all characters have been written. From-SVN: r280137
2020-01-10OpenACC – support "if" + "if_present" clauses with "host_data"Tobias Burnus
2020-01-10 Gergö Barany <gergo@codesourcery.com> Thomas Schwinge <thomas@codesourcery.com> Julian Brown <julian@codesourcery.com> Tobias Burnus <tobias@codesourcery.com> gcc/c/ * c-parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF and PRAGMA_OACC_CLAUSE_IF_PRESENT. gcc/cp/ * parser.c (OACC_HOST_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_IF and PRAGMA_OACC_CLAUSE_IF_PRESENT. gcc/fortran/ * openmp.c (OACC_HOST_DATA_CLAUSES): Add PRAGMA_OACC_CLAUSE_IF and PRAGMA_OACC_CLAUSE_IF_PRESENT. gcc/ * omp-low.c (lower_omp_target): Use GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT if PRAGMA_OACC_CLAUSE_IF_PRESENT exist. gcc/testsuite/ * c-c++-common/goacc/host_data-1.c: Added tests of if and if_present clauses on host_data. * gfortran.dg/goacc/host_data-tree.f95: Likewise. include/ * gomp-constants.h (enum gomp_map_kind): New enumeration constant GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT. libgomp/ * oacc-parallel.c (GOACC_data_start): Handle GOMP_MAP_USE_DEVICE_PTR_IF_PRESENT. * target.c (gomp_map_vars_async): Likewise. * testsuite/libgomp.oacc-c-c++-common/host_data-7.c: New. * testsuite/libgomp.oacc-fortran/host_data-5.F90: New. From-SVN: r280115
2020-01-08libgomp.texi: Fix typos, use https (actual change)Tobias Burnus
From-SVN: r280008
generated by cgit v1.2.3 (git 2.25.1) at 2024-05-02 02:04:31 +0000