aboutsummaryrefslogtreecommitdiff
path: root/libgomp
AgeCommit message (Collapse)AuthorFilesLines
2019-01-31re PR c++/88988 (ICE: Segmentation fault (in lookup_name_real_1))Jakub Jelinek2-0/+33
PR c++/88988 * lambda.c (is_capture_proxy): Don't return true for DECL_OMP_PRIVATIZED_MEMBER artificial vars. * testsuite/libgomp.c++/pr88988.C: New test. From-SVN: r268407
2019-01-28re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)Jakub Jelinek1-43/+0
PR middle-end/89002 * gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ for lastprivate/linear IV, push gimplify context around gimplify_assign and, if it needed any temporaries, pop it into a gimple bind around the sequence. * testsuite/libgomp.c/pr89002.c: New test. From-SVN: r268346
2019-01-28re PR middle-end/89002 (ICE in scan_omp_1_op, at omp-low.c:3166)Jakub Jelinek2-0/+91
PR middle-end/89002 * gimplify.c (gimplify_omp_for): When adding OMP_CLAUSE_*_GIMPLE_SEQ for lastprivate/linear IV, push gimplify context around gimplify_assign and, if it needed any temporaries, pop it into a gimple bind around the sequence. * testsuite/libgomp.c/pr89002.c: New test. From-SVN: r268345
2019-01-28re PR testsuite/89064 (libgomp.graphite/force-parallel-5.c fails starting ↵Richard Biener2-1/+7
with r268257) 2019-01-28 Richard Biener <rguenther@suse.de> PR testsuite/89064 PR tree-optimization/86865 * testsuite/libgomp.graphite/force-parallel-5.c: XFAIL. From-SVN: r268333
2019-01-24[nvptx, libgomp] Fix memleak in GOMP_OFFLOAD_fini_deviceTom de Vries2-0/+11
I wrote a test-case: ... int main (void) { for (unsigned i = 0; i < 128; ++i) { acc_init (acc_device_nvidia); acc_shutdown (acc_device_nvidia); } return 0; } ... and ran it under valgrind. The only leak location reported with a frequency of 128, was the allocation of ptx_devices in nvptx_init. Fix this by freeing ptx_devices in GOMP_OFFLOAD_fini_device, once instantiated_devices drops to 0. 2019-01-24 Tom de Vries <tdevries@suse.de> * plugin/plugin-nvptx.c (GOMP_OFFLOAD_fini_device): Free ptx_devices once instantiated_devices drops to 0. From-SVN: r268237
2019-01-23[nvptx, libgomp] Fix cuMemAlloc with size zeroTom de Vries3-20/+51
Consider test-case: ... int main (void) { #pragma acc parallel async ; #pragma acc parallel async ; #pragma acc wait return 0; } ... This fails with: ... libgomp: cuMemAlloc error: invalid argument Segmentation fault (core dumped) ... The cuMemAlloc error is due to the fact that we're try to allocate 0 bytes. Fix this by preventing calling map_push with size zero argument in nvptx_exec. This also has the consequence that for the abort-1.c test-case, we end up calling cuMemFree during map_fini for the struct cuda_map allocated in map_init, which fails because an abort happened. Fix this by calling cuMemFree with CUDA_CALL_NOCHECK in cuda_map_destroy. 2019-01-23 Tom de Vries <tdevries@suse.de> PR target/PR88946 * plugin/plugin-nvptx.c (cuda_map_destroy): Use CUDA_CALL_NOCHECK for cuMemFree. (nvptx_exec): Don't call map_push if mapnum == 0. * testsuite/libgomp.oacc-c-c++-common/pr88946.c: New test. From-SVN: r268178
2019-01-23[nvptx, libgomp] Fix assert (!s->map->active) in map_finiTom de Vries3-2/+48
There are currently two situations where this assert triggers: ... libgomp/plugin/plugin-nvptx.c: map_fini: Assertion `!s->map->active' failed. ... First, in abort-1.c, a parallel region triggering an abort: ... int main (void) { #pragma acc parallel abort (); return 0; } ... The abort is detected in nvptx_exec as the CUDA_ERROR_ILLEGAL_INSTRUCTION return status of the cuStreamSynchronize call after kernel launch, which is then handled by calling non-returning function GOMP_PLUGIN_fatal. Consequently, the map_pop in nvptx_exec that in case of cuStreamSynchronize success would remove or inactive the element added by the map_push earlier in nvptx_exec, does not trigger. With the element no longer active, but still marked active and a member of s->map, we run into the assert during GOMP_OFFLOAD_fini_device, which is triggered from atexit handler gomp_target_fini (which is triggered by the GOMP_PLUGIN_fatal mentioned above calling exit). Second, in pr88941.c, an async parallel region without wait: ... int main (void) { #pragma acc parallel async ; /* no #pragma acc wait */ return 0; } ... Because nvptx_exec is handling an async region, it does not call map_pop for the element added by map_push, but schedules an kernel execution completion event to call map_pop. Again, we run into the assert during GOMP_OFFLOAD_fini_device, which is triggered from atexit handler gomp_target_fini, but the exit in this case is triggered by returning from main. So either the kernel is still running, or the kernel has completed but the corresponding event that is supposed to call map_pop is stuck in the event queue, waiting for an event_gc. Fix this by removing the assert, and skipping the freeing of device memory if the map is still marked active (though in the async case, this is more a workaround than an fix). 2019-01-23 Tom de Vries <tdevries@suse.de> PR target/88941 PR target/88939 * plugin/plugin-nvptx.c (cuda_map_destroy): Handle map->active case. (map_fini): Remove "assert (!s->map->active)". * testsuite/libgomp.oacc-c-c++-common/pr88941.c: New test. From-SVN: r268177
2019-01-23[nvptx, libgomp] Fix map_pushTom de Vries3-18/+97
The map field of a struct ptx_stream is a FIFO. The FIFO is implemented as a single linked list, with pop-from-the-front semantics. The function map_pop pops an element, either by: - deallocating the element, if there is more than one element - or marking the element inactive, if there's only one element The responsibility of map_push is to push an element to the back, as well as selecting the element to push, by: - allocating an element, or - reusing the element at the front if inactive and big enough, or - dropping the element at the front if inactive and not big enough, and allocating one that's big enough The current implemention gets at least the first and most basic scenario wrong: > map = cuda_map_create (size); We create an element, and assign it to map. > for (t = s->map; t->next != NULL; t = t->next) > ; We determine the last element in the fifo. > t->next = map; We append the new element. > s->map = map; But here, we throw away the rest of the FIFO, and declare the FIFO to be just the new element. This problem causes the test-case asyncwait-1.c to fail intermittently on some systems. The pr87835.c test-case added here is a a minimized and modified version of asyncwait-1.c (avoiding the kernel construct) that is more likely to fail. Fix this by rewriting map_pop more robustly, by: - seperating the function in two phases: select element, push element - when reusing or dropping an element, making sure that the element is cleanly popped from the queue - rewriting the push element part in such a way that it can handle all cases without needing if statements, such that each line is exercised for each of the three cases. 2019-01-23 Tom de Vries <tdevries@suse.de> PR target/87835 * plugin/plugin-nvptx.c (map_push): Fix adding of allocated element. * testsuite/libgomp.oacc-c-c++-common/pr87835.c: New test. From-SVN: r268176
2019-01-15[nvptx] Handle assignment to gang-level reduction variableTom de Vries2-0/+22
2019-01-15 Tom de Vries <tdevries@suse.de> PR target/80547 * config/nvptx/nvptx.c (nvptx_goacc_reduction_init): Handle lhs == NULL_TREE for gang-level reduction. * testsuite/libgomp.oacc-c-c++-common/gang-reduction-var-assignment.c: New test. From-SVN: r267934
2019-01-12[nvptx] Enable setting vector length using -fopenacc-dim -- testcasesTom de Vries5-0/+219
Add some test-cases that set vector length using -fopenacc-dim. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-5.c: New test. * testsuite/libgomp.oacc-fortran/gemm-2.f90: New test. From-SVN: r267897
2019-01-12[nvptx] Enable setting vector length using -fopenacc-dimTom de Vries2-1/+5
Enable setting vector length using -fopenacc-dim, f.i. -fopenacc-dim=::128. 2019-01-12 Tom de Vries <tdevries@suse.de> * config/nvptx/nvptx.c (nvptx_goacc_validate_dims_1): Alow setting vector length using -fopenacc-dim. * plugin/plugin-nvptx.c (nvptx_exec): Update error message. From-SVN: r267896
2019-01-12[nvptx] Add vector_length 64 test-casesTom de Vries4-0/+61
Add some test-cases using vector_length 64. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-3.c: New test. From-SVN: r267895
2019-01-12[nvptx] Force vl32 if calling vector-partitionable routines -- test-casesTom de Vries3-0/+111
Add test-cases for "[nvptx] Force vl32 if calling vector-partitionable routines". 2019-01-12 Tom de Vries <tdevries@suse.de> PR target/85486 * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: New test. From-SVN: r267894
2019-01-12[nvptx] Don't emit barriers for empty loops -- test-casesTom de Vries3-0/+48
Add test-cases for PR85381. 2019-01-12 Tom de Vries <tdevries@suse.de> PR target/85381 * testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381.c: New test. From-SVN: r267893
2019-01-12[nvptx] Enable large vectors -- reduction testcasesTom de Vries4-0/+179
Add various reduction test-cases with vector length 128. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vred2d-128.c: New test. * testsuite/libgomp.oacc-fortran/gemm.f90: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-10.c: New test. From-SVN: r267892
2019-01-12[nvptx] Enable large vectors -- test-casesTom de Vries4-0/+127
Add various test-cases with vector length 128. 2019-01-12 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-7.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-4.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-6.c: New test. From-SVN: r267891
2019-01-12[nvptx] Update insufficient launch message for variable vector_lengthTom de Vries2-8/+15
Update message in nvptx libgomp plugin about insufficient resources to launch kernel, to accommodate for the fact the vector_length can now be variable. 2019-01-12 Tom de Vries <tdevries@suse.de> * plugin/plugin-nvptx.c (nvptx_exec): Update insufficient hardware resources diagnostic. From-SVN: r267890
2019-01-12[nvptx] Enable large vectorsTom de Vries3-5/+11
Allow vector_length clauses to accept values larger than warp size. Note that this does not enable setting vector_length to values larger than warp size using -fopenacc-dim. 2019-01-12 Tom de Vries <tdevries@suse.de> * config/nvptx/nvptx.c (nvptx_goacc_validate_dims): Take larger vector lengths into account. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: Expect vector length to be 128. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Expect vector length 2097152 to be reduced to 1024 instead of 32. From-SVN: r267889
2019-01-11Better distinguish OpenACC and OpenMP sections in libgomp.texiThomas Schwinge2-7/+15
2019-01-11 Thomas Schwinge <thomas@codesourcery.com> James Norris <jnorris@codesourcery.com> * libgomp.texi: Better distinguish OpenACC and OpenMP "Runtime Library Routines", and "Environment Variables". Co-Authored-By: James Norris <jnorris@codesourcery.com> From-SVN: r267841
2019-01-11[nvptx] Don't allow vector_length 64 with num_workers 16Tom de Vries2-0/+27
When using a compiler build with: ... +#define PTX_DEFAULT_VECTOR_LENGTH PTX_CTA_SIZE ... consider a test-case: ... int main (void) { #pragma acc parallel vector_length (64) #pragma acc loop worker for (unsigned int i = 0; i < 32; i++) #pragma acc loop vector for (unsigned int j = 0; j < 64; j++) ; return 0; } ... If num_workers is 16, either because: - we add a "num_workers (16)" clause on the parallel directive, or - we set "GOMP_OPENACC_DIM=:16:", or - the libgomp plugin chooses 16 num_workers we run into an illegal instruction at runtime, because a bar.sync instruction tries to use a barrier 16. The instruction is illegal, because ptx supports only 16 barriers per CTA, and the valid range is 0..15. The problem is that with a warp-multiple vector length, we use a code generation scheme with a per-worker barrier. And because barrier zero is reserved for per-cta barrier, only the remaining 15 barriers can be used as per-worker barrier, and consequently we can't use num_workers larger than 15. This problem occurs only for vector_length 64. For vector_length 32, we use a different code generation scheme, and for vector_length >= 96, the maximum num_workers is not big enough not to trigger this problem. Also, this problem only occurs for num_workers 16. As explained above, num_workers 15 is safe to use, and 16 is already the maximum num_workers for vector_length 64. This patch fixes the problem in both the compiler (handling "num_workers (16)") and in the libgomp nvptx plugin (with and without "GOMP_OPENACC_DIM=:16:"). 2019-01-11 Tom de Vries <tdevries@suse.de> * config/nvptx/nvptx.c (PTX_CTA_NUM_BARRIERS, PTX_PER_CTA_BARRIER) (PTX_NUM_PER_CTA_BARRIER, PTX_FIRST_PER_WORKER_BARRIER) (PTX_NUM_PER_WORKER_BARRIERS): Define. (nvptx_apply_dim_limits): Prevent vector_length 64 and num_workers 16. * plugin/plugin-nvptx.c (nvptx_exec): Prevent vector_length 64 and num_workers 16. From-SVN: r267838
2019-01-11[libgomp, testsuite, openacc] Remove -foffload=-w in reduction-[1-5].cTom de Vries6-15/+9
Before the commit "[libgomp, testsuite, openacc] Don't use const int for dimensions", the "const int" construct was used to set launch dimensions in reductions-[1-5].c. In the case of -xc -O0, the const int is implemented as a variable by the C front-end. Consequently, the nvptx back-end generated warnings that vector_length was overridden to be hard-coded, rather than left to be set at runtime. The test-cases silenced these warnings by switching off all warnings in the accelerator compiler using "-foffload=-w". Given that no warnings occur anymore, remove the "-foffload=-w" setting. 2019-01-11 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/reduction-1.c: Remove -foffload=-w. * testsuite/libgomp.oacc-c-c++-common/reduction-2.c: Same. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c: Same. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c: Same. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Same. From-SVN: r267836
2019-01-11[nvptx, testsuite, openacc, libgomp] Add insufficient-resources.cTom de Vries2-0/+26
Add a test-case that tests the "insufficient resources" fatal in the nvptx libgomp plugin. 2019-01-11 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/insufficient-resources.c: New test. From-SVN: r267835
2019-01-10Add testcase from PR71959Nathan Sidwell3-0/+73
libgomp/ PR lto/71959 * testsuite/libgomp.oacc-c++/pr71959-aux.cc: New. * testsuite/libgomp.oacc-c++/pr71959.C: New. Co-Authored-By: Julian Brown <julian@codesourcery.com> From-SVN: r267806
2019-01-09libgomp: Reduce copy and paste for RTEMSSebastian Huber2-181/+7
libgomp/ * config/rtems/bar.c: Include "../linux/bar.c" and delete copy and paste code. From-SVN: r267752
2019-01-09libgomp: Avoid complex dependencies for RTEMSSebastian Huber2-0/+55
libgomp/ * config/rtems/affinity-fmt.c: New file. Include affinity-fmt.c, undefining HAVE_GETPID and HAVE_GETHOSTNAME, and mapping fwrite to write. From-SVN: r267751
2019-01-09[libgomp, testsuite, openacc] Don't use const int for dimensionsTom de Vries6-15/+25
Const int is handled differently at -O0 for -xc and -xc++, which can cause noise in testsuite/libgomp.oacc-c-c++-common test-cases (which are both run for c and c++) if const int is used for launch dimensions. Fix this by using #defines instead. 2019-01-09 Tom de Vries <tdevries@suse.de> PR target/88756 * testsuite/libgomp.oacc-c-c++-common/reduction-1.c (ng, nw, vl): Use #define instead of "const int". * testsuite/libgomp.oacc-c-c++-common/reduction-2.c (ng, nw, vl): Same. * testsuite/libgomp.oacc-c-c++-common/reduction-3.c (ng, nw, vl): Same. * testsuite/libgomp.oacc-c-c++-common/reduction-4.c (ng, nw, vl): Same. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c (ng, nw, vl): Same. From-SVN: r267747
2019-01-09[nvptx, libgomp] Don't launch with num_workers == 0Tom de Vries2-0/+6
When using a compiler build with: ... +#define PTX_DEFAULT_VECTOR_LENGTH PTX_CTA_SIZE +#define PTX_MAX_VECTOR_LENGTH PTX_CTA_SIZE ... and running the libgomp testsuite, we run into an execution failure in parallel-loop-1.c, due to a cuda launch failure: ... nvptx_exec: kernel f6_none_none$_omp_fn$0: launch gangs=480, workers=0, \ vectors=1024 libgomp: cuLaunchKernel error: invalid argument ... because workers == 0. The workers variable is set to 0 here in nvptx_exec: ... workers = blocks / actual_vectors; ... because actual_vectors is 1024, and blocks is 768: ... cuOccupancyMaxPotentialBlockSize: grid = 10, block = 768 ... Fix this by ensuring that workers is at least one. 2019-01-09 Tom de Vries <tdevries@suse.de> * plugin/plugin-nvptx.c (nvptx_exec): Make sure to launch with at least one worker. From-SVN: r267746
2019-01-07[nvptx] Fix libgomp.oacc-c-c++-common/vector-length-128-3.cTom de Vries2-1/+6
The vector-length-128-3.c test-case uses GOMP_OPENACC_DIM=-:-:128, but '-' is not yet supported on trunk. Use GOMP_OPENACC_DIM=::128 instead. 2019-01-07 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: Fix GOMP_OPENACC_DIM argument. From-SVN: r267624
2019-01-03[nvptx] Add vector_length 128 testcasesTom de Vries3-0/+86
Add a couple of test-cases using vector length 128, while checking that we override to vector length 32. 2019-01-03 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/vector-length-128-1.c: New test. * testsuite/libgomp.oacc-c-c++-common/vector-length-128-3.c: New test. From-SVN: r267559
2019-01-01Update copyright years.Jakub Jelinek134-134/+138
From-SVN: r267494
2019-01-01gcc.c (process_command): Update copyright notice dates.Jakub Jelinek2-1/+5
* gcc.c (process_command): Update copyright notice dates. * gcov-dump.c (print_version): Ditto. * gcov.c (print_version): Ditto. * gcov-tool.c (print_version): Ditto. * gengtype.c (create_file): Ditto. * doc/cpp.texi: Bump @copying's copyright year. * doc/cppinternals.texi: Ditto. * doc/gcc.texi: Ditto. * doc/gccint.texi: Ditto. * doc/gcov.texi: Ditto. * doc/install.texi: Ditto. * doc/invoke.texi: Ditto. gcc/fortran/ * gfortranspec.c (lang_specific_driver): Update copyright notice dates. * gfc-internals.texi: Bump @copying's copyright year. * gfortran.texi: Ditto. * intrinsic.texi: Ditto. * invoke.texi: Ditto. gcc/go/ * gccgo.texi: Bump @copyrights-go year. gcc/ada/ * gnat_ugn.texi: Bump @copying's copyright year. * gnat_rm.texi: Likewise. gcc/d/ * gdc.texi: Bump @copyrights-d year. libitm/ * libitm.texi: Bump @copying's copyright year. libgomp/ * libgomp.texi: Bump @copying's copyright year. libquadmath/ * libquadmath.texi: Bump @copying's copyright year. From-SVN: r267492
2018-12-28For libgomp OpenACC entry points, redefine the "device" argument to "flags"Thomas Schwinge2-19/+39
... so that we're then able to use this for other flags in addition to "GOACC_FLAG_HOST_FALLBACK". gcc/ * omp-expand.c (expand_omp_target): Restructure OpenACC vs. OpenMP code paths. Update for libgomp OpenACC entry points change. include/ * gomp-constants.h (GOACC_FLAG_HOST_FALLBACK) (GOACC_FLAGS_MARSHAL_OP, GOACC_FLAGS_UNMARSHAL): Define. libgomp/ * oacc-parallel.c (GOACC_parallel_keyed, GOACC_parallel) (GOACC_data_start, GOACC_enter_exit_data, GOACC_update) (GOACC_declare): Redefine the "device" argument to "flags". From-SVN: r267448
2018-12-28Cleanup libgomp's coalesce chunk data structuresThomas Schwinge2-21/+38
libgomp/ * target.c (struct gomp_coalesce_chunk): New structure. (struct gomp_coalesce_buf): Update the chunks member to use that type. Adjust all users. Co-Authored-By: Cesar Philippidis <cesar@codesourcery.com> From-SVN: r267446
2018-12-27aligned1.f03: Fix invalid code that now causes an error after r267415.Steven G. Kargl1-2/+2
2018-12-27 Steven G. Kargl <kargl@gcc.gnu.org> * libgomp.fortran/aligned1.f03: Fix invalid code that now causes an error after r267415. From-SVN: r267436
2018-12-19[nvptx] Commit passing pr85381-*.c test-casesTom de Vries4-0/+104
Add pr85381*.c test-cases that are already passing without the fix for PR85381. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-19 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: New test. * testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: New test. From-SVN: r267268
2018-12-19[nvptx, libgomp] Move rtl-dump test-cases to libgompTom de Vries4-0/+73
The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing during linking due to missing libgomp.spec. Move them to the libgomp testsuite. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-19 Tom de Vries <tdevries@suse.de> * gcc.dg/goacc/nvptx-merged-loop.c: Move to libgomp/testsuite/libgomp.oacc-c-c++-common. * gcc.dg/goacc/nvptx-sese-1.c: Same. * testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp. * testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from gcc/testsuite/gcc.dg/goacc. * testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same. From-SVN: r267267
2018-12-14Missing changes from "Adjust copy/copyin/copyout/create for OpenACC 2.5"Thomas Schwinge19-297/+51
Most of that patch's changes were already committed as part of r261813 "Update OpenACC data clause semantics to the 2.5 behavior", but not all of them. libgomp/ * oacc-mem.c (acc_present_or_create): Remove definition and change to alias of acc_create. (acc_present_or_copyin): Remove definition and change to alias of acc_copyin. * oacc-parallel.c (GOACC_enter_exit_data): Call acc_create instead of acc_present_or_create. * testsuite/libgomp.oacc-c-c++-common/data-already-1.c: Remove. * testsuite/libgomp.oacc-c-c++-common/data-already-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-6.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-already-8.c: Likewise. * testsuite/libgomp.oacc-fortran/data-already-1.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-2.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-3.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-4.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-5.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-6.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-7.f: Likewise. * testsuite/libgomp.oacc-fortran/data-already-8.f: Likewise. Co-Authored-By: Chung-Lin Tang <cltang@codesourcery.com> From-SVN: r267153
2018-12-14[PR88495] An OpenACC async queue is always synchronized with itselfThomas Schwinge4-139/+8
An OpenACC async queue is always synchronized with itself, so invocations like "#pragma acc wait(0) async(0)", or "acc_wait_async (0, 0)" don't make a lot of sense, but are still valid. libgomp/ PR libgomp/88495 * plugin/plugin-nvptx.c (nvptx_wait_async): Don't refuse "identical parameters". * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/lib-80.c: Remove. From-SVN: r267152
2018-12-14[PR88484] OpenACC wait directive without wait argument but with async clauseThomas Schwinge3-2/+84
We don't correctly handle "#pragma acc wait async (a)" for "a >= 0", handling as a no-op whereas it should enqueue the appropriate wait operations on "async (a)". libgomp/ PR libgomp/88484 * oacc-parallel.c (GOACC_wait): Correct handling for "async >= 0". * testsuite/libgomp.oacc-c-c++-common/asyncwait-nop-1.c: New file. From-SVN: r267151
2018-12-14[PR88407] [OpenACC] Correctly handle unseen async-argumentsThomas Schwinge11-267/+93
... which turn the operation into a no-op. libgomp/ PR libgomp/88407 * plugin/plugin-nvptx.c (nvptx_async_test, nvptx_wait) (nvptx_wait_async): Unseen async-argument is a no-op. * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Update. * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-79.c: Likewise. * testsuite/libgomp.oacc-fortran/lib-12.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-71.c: Merge into... * testsuite/libgomp.oacc-c-c++-common/lib-69.c: ... this. Update. * testsuite/libgomp.oacc-c-c++-common/lib-77.c: Merge into... * testsuite/libgomp.oacc-c-c++-common/lib-74.c: ... this. Update From-SVN: r267150
2018-12-14Revise libgomp.oacc-c-c++-common/data-2-lib.c, ↵Thomas Schwinge3-157/+125
libgomp.oacc-c-c++-common/data-2.c These are meant to be functionally equivalent (but no longer are), just using different means. Also, use the OpenACC "*_async" functions recently added. libgomp/ * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Revise. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. From-SVN: r267149
2018-12-14Correctly describe OpenACC async/wait dependenciesChung-Lin Tang4-3/+9
libgomp/ * testsuite/libgomp.oacc-c-c++-common/data-2-lib.c: Adjust. * testsuite/libgomp.oacc-c-c++-common/data-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/data-3.c: Likewise. Reviewed-by: Thomas Schwinge <thomas@codesourcery.com> From-SVN: r267148
2018-12-14[PR88370] acc_get_cuda_stream/acc_set_cuda_stream: acc_async_sync, ↵Thomas Schwinge8-20/+222
acc_async_noval Per my reading of the OpenACC specification (and as supported by secondary documentation, such as code examples, or presentations), it's valid to call "acc_get_cuda_stream"/"acc_set_cuda_stream" also with "acc_async_sync", "acc_async_noval" arguments, not just with the nonnegative values as currently implemented. libgomp/ PR libgomp/88370 * libgomp.texi (acc_get_current_cuda_context, acc_get_cuda_stream) (acc_set_cuda_stream): Clarify. * oacc-cuda.c (acc_get_cuda_stream, acc_set_cuda_stream): Use "async_valid_p". * plugin/plugin-nvptx.c (nvptx_set_cuda_stream): Refuse "async == acc_async_sync". * testsuite/libgomp.oacc-c-c++-common/acc_set_cuda_stream-1.c: New file. * testsuite/libgomp.oacc-c-c++-common/async_queue-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/lib-84.c: Update. * testsuite/libgomp.oacc-c-c++-common/lib-85.c: Likewise. From-SVN: r267147
2018-12-14[offloading] Error on missing symbolsTom de Vries6-0/+90
When compiling an OpenMP or OpenACC program containing a reference in the offloaded code to a symbol that has not been included in the offloaded code, the offloading compiler may ICE in lto1. Fix this by erroring out instead, mentioning the problematic symbol: ... error: variable 'var' has been referenced in offloaded code but hasn't been marked to be included in the offloaded code lto1: fatal error: errors during merging of translation units compilation terminated. ... Build x86_64 with nvptx accelerator and reg-tested libgomp. Build x86_64 and reg-tested libgomp. 2018-12-14 Tom de Vries <tdevries@suse.de> * lto-cgraph.c (verify_node_partition): New function. (input_overwrite_node, input_varpool_node): Use verify_node_partition. * testsuite/libgomp.c-c++-common/function-not-offloaded-aux.c: New test. * testsuite/libgomp.c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: New test. * testsuite/libgomp.oacc-c-c++-common/function-not-offloaded.c: New test. * testsuite/libgomp.oacc-c-c++-common/variable-not-offloaded.c: New test. From-SVN: r267134
2018-12-13[libgomp, nvptx] Fix libgomp.c/target-5.c compilationTom de Vries5-6/+74
Libgomp test-case libgomp.c/target-5.c is failing to compile when building for x86_64 with nvptx accelerator due to missing: - getpid - gethostname - isatty (pulled in by fwrite) in the nvptx newlib. This patch fixes the build failure by: - adding a function gomp_print_string which limits the use of fwrite to a single location (in affinity-fmt.c), and - creating an nvptx version of affinity-fmt.c, which: - overrides the configure test results HAVE_GETPID and HAVE_GETHOSTNAME, and - implements fwrite using write. Build and reg-tested on x86_64 with nvptx accelerator. 2018-12-13 Tom de Vries <tdevries@suse.de> * affinity-fmt.c (gomp_print_string): New function, factored out of ... (omp_display_affinity, gomp_display_affinity_thread): ... here, and ... * fortran.c (omp_display_affinity_): ... here. * libgomp.h (gomp_print_string): Declare. * config/nvptx/affinity-fmt.c: New file. Include affinity-fmt.c, undefining HAVE_GETPID and HAVE_GETHOSTNAME, and mapping fwrite to write. From-SVN: r267100
2018-12-13re PR libgomp/88460 ([nvptx] FAIL: libgomp.c++/for-24.C (internal compiler ↵Jakub Jelinek2-20/+26
error)) PR libgomp/88460 * testsuite/libgomp.c++/for-24.C (results): Include it in omp declare target region. (main): Use map (always, tofrom: results) instead of map (tofrom: results). From-SVN: r267093
2018-12-12re PR fortran/88463 (Rejects conforming source, OpenMP Parallel region ↵Jakub Jelinek3-0/+32
Default(None) reference to module parameter array, separate source) PR fortran/88463 * trans-openmp.c (gfc_omp_predetermined_sharing): Handle TREE_READONLY VAR_DECLs with DECL_EXTERNAL like those with TREE_STATIC. * testsuite/libgomp.fortran/pr88463-1.f90: New test. * testsuite/libgomp.fortran/pr88463-2.f90: New test. From-SVN: r267069
2018-12-12omp-builtins.def (BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START, [...]): Fix ↵Jakub Jelinek2-0/+118
up function types - remove one argument. * omp-builtins.def (BUILT_IN_GOMP_LOOP_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_MAYBE_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_LOOP_ULL_MAYBE_NONMONOTONIC_RUNTIME_START, BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_RUNTIME, BUILT_IN_GOMP_PARALLEL_LOOP_MAYBE_NONMONOTONIC_RUNTIME): Fix up function types - remove one argument. * testsuite/libgomp.c-c++-common/for-16.c: New test. From-SVN: r267067
2018-12-12Don't mark stack pointer as clobbered in asmAndreas Schwab2-2/+7
* config/linux/ia64/futex.h (sys_futex0): Don't mark r12 as clobbered. From-SVN: r267052
2018-12-09re PR libfortran/88411 (Random crashes for ASYNCHRONOUS writes (bad locking?))Thomas Koenig2-0/+35
2018-12-09 Thomas Koenig <tkoenig@gcc.gnu.org> PR fortran/88411 * io/transfer.c (dta_transfer_init): Do not treat as an asynchronous statement unless the statement has ASYNCHRONOUS="YES". (st_write_done): Likewise. (st_read_done): Do not perform async_wait for synchronous I/O on an async unit. (st_read_done): Likewise. 2018-12-09 Thomas Koenig <tkoenig@gcc.gnu.org> PR fortran/88411 * testsuite/libgomp.fortran/async_io_8.f90: New test. From-SVN: r266929
generated by cgit v1.1 at 2024-07-20 01:26:46 +0000