aboutsummaryrefslogtreecommitdiff
path: root/libgomp/testsuite
AgeCommit message (Collapse)AuthorFilesLines
2021-05-04OpenMP: Support complex/float in && and || reductionTobias Burnus4-0/+770
C/C++ permit logical AND and logical OR also with floating-point or complex arguments by doing an unequal zero comparison; the result is an 'int' with value one or zero. Hence, those are also permitted as reduction variable, even though it is not the most sensible thing to do. gcc/c/ChangeLog: * c-typeck.c (c_finish_omp_clauses): Accept float + complex for || and && reductions. gcc/cp/ChangeLog: * semantics.c (finish_omp_reduction_clause): Accept float + complex for || and && reductions. gcc/ChangeLog: * omp-low.c (lower_rec_input_clauses, lower_reduction_clauses): Handle && and || with floating-point and complex arguments. gcc/testsuite/ChangeLog: * gcc.dg/gomp/clause-1.c: Use 'reduction(&:..)' instead of '...(&&:..)'. libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/reduction-1.c: New test. * testsuite/libgomp.c-c++-common/reduction-2.c: New test. * testsuite/libgomp.c-c++-common/reduction-3.c: New test.
2021-05-04OpenMP/Fortran - fix pasto + testcase in depobj [PR100397]Tobias Burnus1-6/+9
gcc/fortran/ChangeLog: PR testsuite/100397 * trans-openmp.c (gfc_trans_omp_depobj): Fix pasto in enum values. libgomp/ChangeLog: PR testsuite/100397 * testsuite/libgomp.fortran/depobj-1.f90 (dep2, dep3): Move var declaration to scope of non-'depend'-guarded assignment to avoid races.
2021-05-03[openmp, simt] Disable SIMT for user-defined reductionTom de Vries1-0/+27
The test-case included in this patch contains this target region: ... for (int i0 = 0 ; i0 < N0 ; i0++ ) counter_N0.i += 1; ... When running with nvptx accelerator, the counter variable is expected to be N0 after the region, but instead is N0 / 32. The problem is that rather than getting the result for all warp lanes, we get it for just one lane. This is caused by the implementation of SIMT being incomplete. It handles regular reductions, but appearantly not user-defined reductions. For now, handle this by disabling SIMT in this case, specifically by setting sctx->max_vf to 1. Tested libgomp on x86_64-linux with nvptx accelerator. gcc/ChangeLog: 2021-05-03 Tom de Vries <tdevries@suse.de> PR target/100321 * omp-low.c (lower_rec_input_clauses): Disable SIMT for user-defined reduction. libgomp/ChangeLog: 2021-05-03 Tom de Vries <tdevries@suse.de> PR target/100321 * testsuite/libgomp.c/target-44.c: New test.
2021-04-30modulo-sched: skip loops with strange register defs [PR100225]Roman Zhuykov1-0/+1627
PR84878 fix adds an assertion which can fail, e.g. when stack pointer is adjusted inside the loop. We have to prevent it and search earlier for any 'strange' instruction. The solution is to skip the whole loop if using 'note_stores' we found that one of hard registers is in 'df->regular_block_artificial_uses' set. Also patch properly prohibit not single-set instruction in loop body. gcc/ChangeLog: PR rtl-optimization/100225 PR rtl-optimization/84878 * modulo-sched.c (sms_schedule): Use note_stores to skip loops where we have an instruction which touches (writes) any hard register from df->regular_block_artificial_uses set. Allow not-single-set instruction only right before basic block tail. gcc/testsuite/ChangeLog: PR rtl-optimization/100225 PR rtl-optimization/84878 * gcc.dg/pr100225.c: New test. libgomp/ChangeLog: * testsuite/libgomp.oacc-c-c++-common/atomic_capture-3.c: New test.
2021-04-29[omp, simt] Handle alternative IVTom de Vries1-0/+48
Consider the test-case libgomp.c/pr81778.c added in this commit, with this core loop (note: CANARY_SIZE set to 0 for simplicity): ... int s = 1; #pragma omp target simd for (int i = N - 1; i > -1; i -= s) a[i] = 1; ... which, given that N is 32, sets a[0..31] to 1. After omp-expand, this looks like: ... <bb 5> : simduid.7 = .GOMP_SIMT_ENTER (simduid.7); .omp_simt.8 = .GOMP_SIMT_ENTER_ALLOC (simduid.7); D.3193 = -s; s.9 = s; D.3204 = .GOMP_SIMT_LANE (); D.3205 = -s.9; D.3206 = (int) D.3204; D.3207 = D.3205 * D.3206; i = D.3207 + 31; D.3209 = 0; D.3210 = -s.9; D.3211 = D.3210 - i; D.3210 = -s.9; D.3212 = D.3211 / D.3210; D.3213 = (unsigned int) D.3212; D.3213 = i >= 0 ? D.3213 : 0; <bb 19> : if (D.3209 < D.3213) goto <bb 6>; [87.50%] else goto <bb 7>; [12.50%] <bb 6> : a[i] = 1; D.3215 = -s.9; D.3219 = .GOMP_SIMT_VF (); D.3216 = (int) D.3219; D.3220 = D.3215 * D.3216; i = D.3220 + i; D.3209 = D.3209 + 1; goto <bb 19>; [100.00%] ... On nvptx, the first time bb6 is executed, i is in the 0..31 range (depending on the lane that is executing) at bb entry. So we have the following sequence: - a[0..31] is set to 1 - i is updated to -32..-1 - D.3209 is updated to 1 (being 0 initially) - bb19 is executed, and if condition (D.3209 < D.3213) == (1 < 32) evaluates to true - bb6 is once more executed, which should not happen because all the elements that needed to be handled were already handled. - consequently, elements that should not be written are written - with CANARY_SIZE == 0, we may run into a libgomp error: ... libgomp: cuCtxSynchronize error: an illegal memory access was encountered ... and with CANARY_SIZE unmodified, we run into: ... Expected 0, got 1 at base[-961] Aborted (core dumped) ... The cause of this is as follows: - because the step s is a variable rather than a constant, an alternative IV (D.3209 in our example) is generated in expand_omp_simd, and the loop condition is tested in terms of the alternative IV rather than the original IV (i in our example). - the SIMT code in expand_omp_simd works by modifying step and initial value. - The initial value fd->loop.n1 is loaded into a variable n1, which is modified by the SIMT code and then used there-after. - The step fd->loop.step is loaded into a variable step, which is modified by the SIMT code, but afterwards there are uses of both step and fd->loop.step. - There are uses of fd->loop.step in the alternative IV handling code, which should use step instead. Fix this by introducing an additional variable orig_step, which is not modified by the SIMT code and replacing all remaining uses of fd->loop.step by either step or orig_step. Build on x86_64-linux with nvptx accelerator, tested libgomp. This fixes for-5.c and for-6.c FAILs I'm currently seeing on a quadro m1200 with driver 450.66. gcc/ChangeLog: 2020-10-02 Tom de Vries <tdevries@suse.de> * omp-expand.c (expand_omp_simd): Add step_orig, and replace uses of fd->loop.step by either step or orig_step. libgomp/ChangeLog: 2020-10-02 Tom de Vries <tdevries@suse.de> * testsuite/libgomp.c/pr81778.c: New test.
2021-04-29[omp, simt] Fix expand_GOMP_SIMT_*Tom de Vries1-0/+24
When running the test-case included in this patch using an nvptx accelerator, it fails in execution. The problem is that the expansion of GOMP_SIMT_XCHG_BFLY is optimized away during pass_jump as "trivially dead insns". This is caused by this code in expand_GOMP_SIMT_XCHG_BFLY: ... class expand_operand ops[3]; create_output_operand (&ops[0], target, mode); ... expand_insn (targetm.code_for_omp_simt_xchg_bfly, 3, ops); ... which doesn't guarantee that target is assigned to by the expanded insn. F.i., if target is: ... (gdb) call debug_rtx ( target ) (subreg/s/u:QI (reg:SI 40 [ _61 ]) 0) ... then after expand_insn, we have: ... (gdb) call debug_rtx ( ops[0].value ) (reg:QI 57) ... See commit 3af3bec2e4d "internal-fn: Avoid dropping the lhs of some calls [PR94941]" for a similar problem. Fix this in the same way, by adding: ... if (!rtx_equal_p (target, ops[0].value)) emit_move_insn (target, ops[0].value); ... where applicable in the expand_GOMP_SIMT_* functions. Tested libgomp on x86_64 with nvptx accelerator. gcc/ChangeLog: 2021-04-28 Tom de Vries <tdevries@suse.de> PR target/100232 * internal-fn.c (expand_GOMP_SIMT_ENTER_ALLOC) (expand_GOMP_SIMT_LAST_LANE, expand_GOMP_SIMT_ORDERED_PRED) (expand_GOMP_SIMT_VOTE_ANY, expand_GOMP_SIMT_XCHG_BFLY) (expand_GOMP_SIMT_XCHG_IDX): Ensure target is assigned to.
2021-04-26OpenACC: Fix pattern in dg-bogus in Fortran testcases againTobias Burnus4-9/+9
It turned out that a compiler built without offloading support and one with can produce slightly different diagnostic. Offloading support implies ENABLE_OFFLOAD which implies that g->have_offload is set when offloading is actually needed. In cgraphunit.c, the latter causes flag_generate_offload = 1, which in turn affects tree.c's free_lang_data. The result is that the front-end specific diagnostic gets reset ('tree_diagnostics_defaults (global_dc)'), which affects in this case 'Warning' vs. 'warning' via the Fortran frontend. Result: 'Warning:' vs. 'warning:'. Side note: Other FE also override the diagnostic, leading to similar differences, e.g. the C++ FE outputs mangled function names differently, cf. patch thread. libgomp/ChangeLog: * testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Use [Ww]arning in dg-bogus as FE diagnostic and default diagnostic differ and the result depends on ENABLE_OFFLOAD. * testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise. * testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise. * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise. gcc/testsuite/ChangeLog: * gfortran.dg/goacc/classify-serial.f95: Use [Ww]arning in dg-bogus as FE diagnostic and default diagnostic differ and the result depends on ENABLE_OFFLOAD. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. * gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
2021-04-26OpenACC: Fix pattern in dg-bogus in Fortran testcasesTobias Burnus4-9/+9
libgomp/ChangeLog: * testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Correct spelling in dg-bogus to match -Wopenacc-parallelism. * testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise. * testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise. * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise. gcc/testsuite/ChangeLog: * gfortran.dg/goacc/classify-serial.f95: Correct spelling in dg-bogus to match -Wopenacc-parallelism. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. * gfortran.dg/goacc/routine-module-mod-1.f90: Likewise.
2021-04-26Add '-Wopenacc-parallelism'Thomas Schwinge29-2/+209
... to diagnose potentially suboptimal choices regarding OpenACC parallelism. Not enabled by default: too noisy ("*potentially* suboptimal choices"); see XFAILed 'dg-bogus'es. gcc/c-family/ * c.opt (Wopenacc-parallelism): New. gcc/fortran/ * lang.opt (Wopenacc-parallelism): New. gcc/ * omp-offload.c (oacc_validate_dims): Implement '-Wopenacc-parallelism'. * doc/invoke.texi (-Wopenacc-parallelism): Document. gcc/testsuite/ * c-c++-common/goacc/diag-parallelism-1.c: New. * c-c++-common/goacc/acc-icf.c: Specify '-Wopenacc-parallelism', and match diagnostics, as appropriate. * c-c++-common/goacc/classify-kernels-unparallelized.c: Likewise. * c-c++-common/goacc/classify-kernels.c: Likewise. * c-c++-common/goacc/classify-parallel.c: Likewise. * c-c++-common/goacc/classify-routine.c: Likewise. * c-c++-common/goacc/classify-serial.c: Likewise. * c-c++-common/goacc/kernels-decompose-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * c-c++-common/goacc/parallel-dims-1.c: Likewise. * c-c++-common/goacc/parallel-reduction.c: Likewise. * c-c++-common/goacc/pr70688.c: Likewise. * c-c++-common/goacc/routine-1.c: Likewise. * c-c++-common/goacc/routine-level-of-parallelism-2.c: Likewise. * c-c++-common/goacc/uninit-dim-clause.c: Likewise. * gfortran.dg/goacc/classify-kernels-unparallelized.f95: Likewise. * gfortran.dg/goacc/classify-kernels.f95: Likewise. * gfortran.dg/goacc/classify-parallel.f95: Likewise. * gfortran.dg/goacc/classify-routine.f95: Likewise. * gfortran.dg/goacc/classify-serial.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/routine-4.f90: Likewise. * gfortran.dg/goacc/routine-level-of-parallelism-1.f90: Likewise. * gfortran.dg/goacc/routine-module-mod-1.f90: Likewise. * gfortran.dg/goacc/routine-multiple-directives-1.f90: Likewise. * gfortran.dg/goacc/uninit-dim-clause.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/firstprivate-1.c: Specify '-Wopenacc-parallelism', and match diagnostics, as appropriate. * testsuite/libgomp.oacc-c-c++-common/loop-auto-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-red-w-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/loop-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/mode-transitions.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/private-variables.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-7.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-g-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-w-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-wv-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/static-variable-1.c: Likewise. * testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise. * testsuite/libgomp.oacc-fortran/par-reduction-2-1.f: Likewise. * testsuite/libgomp.oacc-fortran/par-reduction-2-2.f: Likewise. * testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise. * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise. * testsuite/libgomp.oacc-fortran/pr84028.f90: Likewise. * testsuite/libgomp.oacc-fortran/private-variables.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/routine-7.f90: Likewise. Co-Authored-By: Nathan Sidwell <nathan@codesourcery.com> Co-Authored-By: Tom de Vries <vries@codesourcery.com> Co-Authored-By: Julian Brown <julian@codesourcery.com> Co-Authored-By: Kwok Cheung Yeung <kcy@codesourcery.com>
2021-04-26[OpenACC] Don't compile libgomp testcases with '-w'Thomas Schwinge10-12/+1
We'd like to actually catch compiler diagnostics (and currently there aren't any). libgomp/ * testsuite/libgomp.oacc-c-c++-common/par-reduction-1.c: Don't compile with '-w'. * testsuite/libgomp.oacc-c-c++-common/par-reduction-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-reduction.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-5.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-6.c: Likewise. * testsuite/libgomp.oacc-fortran/parallel-reduction.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-5.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-6.f90: Likewise. * testsuite/libgomp.oacc-fortran/reduction-7.f90: Likewise.
2021-04-22Avoid -latomic for amdgcn offloadingRichard Biener1-1/+1
libatomic isn't built for amdgcn but reduction-16.c adds it via -foffload=-latomic when offloading for nvptx is enabled. The following avoids linker errors when offloading to amdgcn is enabled as well. 2021-04-21 Richard Biener <rguenther@suse.de> libgomp/ * testsuite/libgomp.c-c++-common/reduction-16.c: Use -latomic only on nvptx-none.
2021-04-21libgomp.fortran/depobj-1.f90: Fix omp_depend_kindTobias Burnus1-2/+1
libgomp/ * testsuite/libgomp.fortran/depobj-1.f90: Use omp_lib's omp_depend_kind instead of defining it as 16.
2021-04-21libgomp/testsuite: Fix checks for dg-excess-errorsTobias Burnus5-11/+23
For the tests modified below, the effective target line has to be effective when compiling for an offload target, except that variable-not-offloaded.c would compile with unified-share memory and pr86416-*.c if long double/float128 is supported. The previous check used a run-time device ability check. This new variant now enables those dg- lines when _compiling_ for nvptx or gcn. libgomp/ChangeLog: * testsuite/lib/libgomp.exp (offload_target_to_openacc_device_type): New, based on check_effective_target_offload_target_nvptx. (check_effective_target_offload_target_nvptx): Call it. (check_effective_target_offload_target_amdgcn): New. * testsuite/libgomp.c-c++-common/function-not-offloaded.c: Require target offload_target_nvptx || offload_target_amdgcn. * testsuite/libgomp.c-c++-common/variable-not-offloaded.c: Likewise. * testsuite/libgomp.c/pr86416-1.c: Likewise. * testsuite/libgomp.c/pr86416-2.c: Likewise.
2021-04-21Fortran/OpenMP: Add 'omp depobj' and 'depend(mutexinoutset:'Tobias Burnus1-0/+113
gcc/fortran/ChangeLog: * dump-parse-tree.c (show_omp_namelist): Handle depobj + mutexinoutset in the depend clause. (show_omp_clauses, show_omp_node, show_code_node): Handle depobj. * gfortran.h (enum gfc_statement): Add ST_OMP_DEPOBJ. (enum gfc_omp_depend_op): Add OMP_DEPEND_UNSET, OMP_DEPEND_MUTEXINOUTSET and OMP_DEPEND_DEPOBJ. (gfc_omp_clauses): Add destroy, depobj_update and depobj. (enum gfc_exec_op): Add EXEC_OMP_DEPOBJ * match.h (gfc_match_omp_depobj): Match 'omp depobj'. * openmp.c (gfc_match_omp_clauses): Add depobj + mutexinoutset to depend clause. (gfc_match_omp_depobj, resolve_omp_clauses, gfc_resolve_omp_directive): Handle 'omp depobj'. * parse.c (decode_omp_directive, next_statement, gfc_ascii_statement): Likewise. * resolve.c (gfc_resolve_code): Likewise. * st.c (gfc_free_statement): Likewise. * trans-openmp.c (gfc_trans_omp_clauses): Handle depobj + mutexinoutset in the depend clause. (gfc_trans_omp_depobj, gfc_trans_omp_directive): Handle EXEC_OMP_DEPOBJ. * trans.c (trans_code): Likewise. libgomp/ChangeLog: * testsuite/libgomp.fortran/depobj-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/depobj-1.f90: New test. * gfortran.dg/gomp/depobj-2.f90: New test.
2021-04-19[OpenACC 'kernels'] '-fopenacc-kernels=[...]' -> '--param=openacc-kernels=[...]'Thomas Schwinge4-4/+4
This configuration knob is temporary, and isn't really meant to be exposed to users. gcc/ * params.opt (-param=openacc-kernels=): Add. * omp-oacc-kernels-decompose.cc (pass_omp_oacc_kernels_decompose::gate): Use it. * doc/invoke.texi (-fopenacc-kernels=@var{mode}): Move... (--param): ... here, 'openacc-kernels'. gcc/c-family/ * c.opt (fopenacc-kernels=): Remove. gcc/fortran/ * lang.opt (fopenacc-kernels=): Remove. gcc/testsuite/ * c-c++-common/goacc/if-clause-2.c: '-fopenacc-kernels=[...]' -> '--param=openacc-kernels=[...]'. * c-c++-common/goacc/kernels-decompose-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * c-c++-common/goacc/kernels-decompose-ice-1.c: Likewise. * c-c++-common/goacc/kernels-decompose-ice-2.c: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c: '-fopenacc-kernels=[...]' -> '--param=openacc-kernels=[...]'. * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise.
2021-04-15XFAIL OpenMP/nvptx execution-time hangs for simple nested OpenMP ↵Thomas Schwinge4-3/+22
'target'/'parallel'/'task' constructs [PR99555] ... still awaiting proper resolution, of course. libgomp/ PR target/99555 * testsuite/lib/libgomp.exp (check_effective_target_offload_device_nvptx): New. * testsuite/libgomp.c/pr99555-1.c <nvptx offload device>: Until resolved, make sure that we exit quickly, with error status, XFAILed. * testsuite/libgomp.c-c++-common/task-detach-6.c: Likewise. * testsuite/libgomp.fortran/task-detach-6.f90: Likewise.
2021-04-14testsuite: Fix up libgomp.fortran/alloc-1.F90 testcase [PR100071]Jakub Jelinek1-1/+2
As can be seen under valgrind, the testcase didn't bind in the last part the fortran pointers properly to the c pointers. 2021-04-14 Jakub Jelinek <jakub@redhat.com> PR testsuite/100071 * testsuite/libgomp.fortran/alloc-1.F90: Call c_f_pointer after last cp = omp_alloc with cp, p arguments instead of cq, q and call c_f_pointer after last cq = omp_alloc with cq, q.
2021-04-11[OpenACC] Fix an ICE where a loop with GT condition is collapsed.Hafiz Abid Qadeer2-4/+28
We have seen an ICE both on trunk and devel/omp/gcc-10 branches which can be reprodued with this simple testcase. It occurs if an OpenACC loop has a collapse clause and any of the loop being collapsed uses GT or GE condition. This issue is specific to OpenACC. int main (void) { int ix, iy; int dim_x = 16, dim_y = 16; { for (iy = dim_y - 1; iy > 0; --iy) for (ix = dim_x - 1; ix > 0; --ix) ; } } The problem is caused by a failing assertion in expand_oacc_collapse_init. It checks that cond_code for fd->loop should be same as cond_code for all the loops that are being collapsed. As the cond_code for fd->loop is LT_EXPR with collapse clause (set at the end of omp_extract_for_data), this assertion forces that all the loop in collapse clause should use < operator. There does not seem to be anything in the code which demands this condition as loop with > condition works ok otherwise. I digged old mailing list a bit but could not find any discussion on this change. Looking at the code, expand_oacc_for checks that fd->loop->cond_code is either LT_EXPR or GT_EXPR. I guess the original intention was to have similar checks on the loop which are being collapsed. But the way check was written does not acheive that. I have fixed it by modifying the check in the assertion to be same as check on fd->loop->cond_code. I tested goacc and libgomp (with nvptx offloading) and did not see any regression. I have added new tests to check collapse with GT/GE condition. PR middle-end/98088 gcc/ * omp-expand.c (expand_oacc_collapse_init): Update condition in a gcc_assert. gcc/testsuite/ * c-c++-common/goacc/collapse-2.c: New. libgomp/ * testsuite/libgomp.oacc-c-c++-common/collapse-2.c: Add check for loop with GT/GE condition. * testsuite/libgomp.oacc-c-c++-common/collapse-3.c: Likewise.
2021-04-09Add 'libgomp.oacc-c-c++-common/static-variable-1.c' [PR84991, PR84992, PR90779]Thomas Schwinge1-0/+460
libgomp/ PR middle-end/84991 PR middle-end/84992 PR middle-end/90779 * testsuite/libgomp.oacc-c-c++-common/static-variable-1.c: New.
2021-03-29libgomp: Fix on_device_arch.c aux-file handling [PR99555]Tobias Burnus5-6/+6
libgomp/ChangeLog: PR target/99555 * testsuite/lib/on_device_arch.c: Move to ... * testsuite/libgomp.c-c++-common/on_device_arch.h: ... here. * testsuite/libgomp.fortran/on_device_arch.c: New file; #include on_device_arch.h. * testsuite/libgomp.c-c++-common/task-detach-6.c: #include on_device_arch.h instead of using dg-additional-source. * testsuite/libgomp.c/pr99555-1.c: Likewise. * testsuite/libgomp.fortran/task-detach-6.f90: Update to use on_device_arch.c without relative paths.
2021-03-25Avoid OpenMP/nvptx execution-time hangs for simple nested OpenMP ↵Thomas Schwinge4-0/+69
'target'/'parallel'/'task' constructs [PR99555] ... awaiting proper resolution, of course. libgomp/ PR target/99555 * testsuite/lib/on_device_arch.c: New file. * testsuite/libgomp.c/pr99555-1.c: Likewise. * testsuite/libgomp.c-c++-common/task-detach-6.c: Until resolved, skip for nvptx offloading, with error status. * testsuite/libgomp.fortran/task-detach-6.f90: Likewise.
2021-03-25'libgomp.oacc-fortran/derivedtypes-arrays-1.f90' OpenACC 'serial' construct ↵Thomas Schwinge1-0/+1
diagnostic for nvptx offloading Fixup for recent commit d28f3da11d8c0aed9b746689d723022a9b5ec04c "openacc: Fix lowering for derived-type mappings through array elements". With nvptx offloading we see the usual: [...]/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: In function 'MAIN__._omp_fn.0': [...]/libgomp.oacc-fortran/derivedtypes-arrays-1.f90:90:40: warning: using vector_length (32), ignoring 1 libgomp/ * testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: OpenACC 'serial' construct diagnostic for nvptx offloading.
2021-03-15OpenMP: Fix 'omp declare target' handling for vars [PR99509]Tobias Burnus1-0/+22
For variables with 'declare target' attribute, varpool_node::get_create marks variables as offload; however, if the node already exists, it is not updated. C/C++ may tag decl with 'declare target implicit', which may only be after varpool creation turned into 'declare target' or 'declare target link'; in this case, the tagging has to happen in the FE. gcc/c/ChangeLog: PR c++/99509 * c-decl.c (finish_decl): For 'omp declare target implicit' vars, ensure that the varpool node is marked as offloadable. gcc/cp/ChangeLog: PR c++/99509 * decl.c (cp_finish_decl): For 'omp declare target implicit' vars, ensure that the varpool node is marked as offloadable. libgomp/ChangeLog: PR c++/99509 * testsuite/libgomp.c-c++-common/declare_target-1.c: New test.
2021-03-12Fortran/OpenMP: Fix use_device_{ptr,addr} with assumed-size array [PR98858]Tobias Burnus1-0/+91
gcc/ChangeLog: PR fortran/98858 * gimplify.c (omp_add_variable): Handle NULL_TREE as size occuring for assumed-size arrays in use_device_{ptr,addr}. libgomp/ChangeLog: PR fortran/98858 * testsuite/libgomp.fortran/use_device_ptr-3.f90: New test.
2021-02-27libgomp, testsuite : Require alias support for PR96390 testcase.Iain Sandoe1-0/+1
This fails everywhere on Darwin, which does not have support for symbol aliases. Add a dg-require-alias to UNSUPPORT it. libgomp/ChangeLog: * testsuite/libgomp.c-c++-common/pr96390.c: Require alias support from the target.
2021-02-25openmp: Fix intermittent hanging of task-detach-6 libgomp tests [PR98738]Kwok Cheung Yeung22-42/+416
This adds support for the task detach clause to taskwait and taskgroup, and simplifies the handling of the detach clause by moving most of the extra handling required for detach tasks to omp_fulfill_event. 2021-02-25 Kwok Cheung Yeung <kcy@codesourcery.com> Jakub Jelinek <jakub@redhat.com> libgomp/ PR libgomp/98738 * libgomp.h (enum gomp_task_kind): Add GOMP_TASK_DETACHED. (struct gomp_task): Replace detach and completion_sem fields with union containing completion_sem and detach_team. Add deferred_p field. (struct gomp_team): Remove task_detach_queue. * task.c: Include assert.h. (gomp_init_task): Initialize deferred_p and completion_sem fields. Rearrange initialization order of fields. (task_fulfilled_p): Delete. (GOMP_task): Use address of task as the event handle. Remove initialization of detach field. Initialize deferred_p field. Use automatic local for completion_sem. Initialize detach_team field for deferred tasks. (gomp_barrier_handle_tasks): Remove handling of task_detach_queue. Set kind of suspended detach task to GOMP_TASK_DETACHED and decrement task_running_count. Move finish_cancelled block out of else branch. Relocate call to gomp_team_barrier_done. (GOMP_taskwait): Handle tasks with completion events that have not been fulfilled. (GOMP_taskgroup_end): Likewise. (omp_fulfill_event): Use address of task as event handle. Post to completion_sem for undeferred tasks. Clear detach_team if task has not finished. For finished tasks, handle post-execution tasks, call gomp_team_barrier_wake if necessary, and free task. * team.c (gomp_new_team): Remove initialization of task_detach_queue. (free_team): Remove free of task_detach_queue. * testsuite/libgomp.c-c++-common/task-detach-1.c: Fix formatting. * testsuite/libgomp.c-c++-common/task-detach-2.c: Fix formatting. * testsuite/libgomp.c-c++-common/task-detach-3.c: Fix formatting. * testsuite/libgomp.c-c++-common/task-detach-4.c: Fix formatting. * testsuite/libgomp.c-c++-common/task-detach-5.c: Fix formatting. Change data-sharing of detach events on enclosing parallel to private. * testsuite/libgomp.c-c++-common/task-detach-6.c: Likewise. Remove taskwait directive. * testsuite/libgomp.c-c++-common/task-detach-7.c: New. * testsuite/libgomp.c-c++-common/task-detach-8.c: New. * testsuite/libgomp.c-c++-common/task-detach-9.c: New. * testsuite/libgomp.c-c++-common/task-detach-10.c: New. * testsuite/libgomp.c-c++-common/task-detach-11.c: New. * testsuite/libgomp.fortran/task-detach-1.f90: Fix formatting. * testsuite/libgomp.fortran/task-detach-2.f90: Fix formatting. * testsuite/libgomp.fortran/task-detach-3.f90: Fix formatting. * testsuite/libgomp.fortran/task-detach-4.f90: Fix formatting. * testsuite/libgomp.fortran/task-detach-5.f90: Fix formatting. Change data-sharing of detach events on enclosing parallel to private. * testsuite/libgomp.fortran/task-detach-6.f90: Likewise. Remove taskwait directive. * testsuite/libgomp.fortran/task-detach-7.f90: New. * testsuite/libgomp.fortran/task-detach-8.f90: New. * testsuite/libgomp.fortran/task-detach-9.f90: New. * testsuite/libgomp.fortran/task-detach-10.f90: New. * testsuite/libgomp.fortran/task-detach-11.f90: New.
2021-02-22Fortran/OpenMP: Fix optional dummy procedures [PR99171]Tobias Burnus1-0/+393
gcc/fortran/ChangeLog: PR fortran/99171 * trans-openmp.c (gfc_omp_is_optional_argument): Regard optional dummy procs as nonoptional as no special treatment is needed. libgomp/ChangeLog: PR fortran/99171 * testsuite/libgomp.fortran/dummy-procs-1.f90: New test.
2021-02-17openacc: Strided array sections and components of derived-type arraysJulian Brown1-3/+2
This patch disallows selecting components of array sections in update directives for OpenACC, as specified in OpenACC 3.0, "2.14.4. Update Directive": In Fortran, members of variables of derived type may appear, including a subarray of a member. Members of subarrays of derived type may not appear. The diagnostic for attempting to use the same construct on other directives has also been improved. gcc/fortran/ * openmp.c (resolve_omp_clauses): Disallow selecting components of arrays of derived type. gcc/testsuite/ * gfortran.dg/goacc/array-with-dt-2.f90: Remove expected errors. * gfortran.dg/goacc/array-with-dt-6.f90: New test. * gfortran.dg/goacc/mapping-tests-2.f90: Update expected error. * gfortran.dg/goacc/ref_inquiry.f90: Update expected errors. * gfortran.dg/gomp/ref_inquiry.f90: Likewise. libgomp/ * testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: Remove expected errors.
2021-02-17openacc: Fix lowering for derived-type mappings through array elementsJulian Brown2-0/+162
This patch fixes lowering of derived-type mappings which select elements of arrays of derived types, and similar. These would previously lead to ICEs. With this change, OpenACC directives can pass through constructs that are no longer recognized by the gimplifier, hence alterations are needed there also. gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Handle element selection for arrays of derived types. gcc/ * gimplify.c (gimplify_scan_omp_clauses): Handle ATTACH_DETACH for non-decls. gcc/testsuite/ * gfortran.dg/goacc/array-with-dt-1.f90: New test. * gfortran.dg/goacc/array-with-dt-3.f90: Likewise. * gfortran.dg/goacc/array-with-dt-4.f90: Likewise. * gfortran.dg/goacc/array-with-dt-5.f90: Likewise. * gfortran.dg/goacc/derived-chartypes-1.f90: Re-enable test. * gfortran.dg/goacc/derived-chartypes-2.f90: Likewise. * gfortran.dg/goacc/derived-classtypes-1.f95: Uncomment previously-broken directives. libgomp/ * testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: New test. * testsuite/libgomp.oacc-fortran/update-dt-array.f90: Likewise.
2021-02-09openacc: Add XFAILs [PR98979]Julian Brown1-2/+3
This patch adds some XFAILs for PR98979 until the patch to fix them has been approved. See: https://gcc.gnu.org/pipermail/gcc-patches/2021-February/564711.html gcc/testsuite/ PR fortran/98979 * gfortran.dg/goacc/array-with-dt-2.f90: Add expected errors. * gfortran.dg/goacc/derived-chartypes-1.f90: Skip ICEing test. * gfortran.dg/goacc/derived-chartypes-2.f90: Likewise. libgomp/ PR fortran/98979 * testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: Add expected errors.
2021-02-04openacc: Allow strided arrays in update directivesJulian Brown1-0/+44
OpenACC 3.0 ("2.14.4. Update Directive") states: Noncontiguous subarrays may appear. It is implementation-specific whether noncontiguous regions are updated by using one transfer for each contiguous subregion, or whether the non-contiguous data is packed, transferred once, and unpacked, or whether one or more larger subarrays (no larger than the smallest contiguous region that contains the specified subarray) are updated. This patch relaxes some conditions in the Fortran front-end so that strided accesses are permitted for update directives. gcc/fortran/ * openmp.c (resolve_omp_clauses): Omit OpenACC update in contiguity check and stride-specified error. gcc/testsuite/ * gfortran.dg/goacc/array-with-dt-2.f90: New test. libgomp/ * testsuite/libgomp.oacc-fortran/array-stride-dt-1.f90: New test.
2021-01-19OpenMP/Fortran: Fixes for {use,is}_device_ptrTobias Burnus1-0/+54
gcc/fortran/ChangeLog: PR fortran/98476 * openmp.c (resolve_omp_clauses): Change use_device_ptr to use_device_addr for unless type(c_ptr); check all list item for is_device_ptr. gcc/ChangeLog: PR fortran/98476 * omp-low.c (lower_omp_target): Handle nonpointer is_device_ptr. libgomp/ChangeLog: PR fortran/98476 * testsuite/libgomp.fortran/is_device_ptr-1.f90: New test. gcc/testsuite/ChangeLog: PR fortran/98476 * gfortran.dg/gomp/map-3.f90: Update expected scan-dump-tree. * gfortran.dg/gomp/is_device_ptr-2.f90: New test. * gfortran.dg/gomp/use_device_ptr-1.f90: New test.
2021-01-16openmp: Add support for the OpenMP 5.0 task detach clauseKwok Cheung Yeung12-0/+423
2021-01-16 Kwok Cheung Yeung <kcy@codesourcery.com> gcc/ * builtin-types.def (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename to... (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR): ...this. Add extra argument. * gimplify.c (omp_default_clause): Ensure that event handle is firstprivate in a task region. (gimplify_scan_omp_clauses): Handle OMP_CLAUSE_DETACH. (gimplify_adjust_omp_clauses): Likewise. * omp-builtins.def (BUILT_IN_GOMP_TASK): Change function type to BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR. * omp-expand.c (expand_task_call): Add GOMP_TASK_FLAG_DETACH to flags if detach clause specified. Add detach argument when generating call to GOMP_task. * omp-low.c (scan_sharing_clauses): Setup data environment for detach clause. (finish_taskreg_scan): Move field for variable containing the event handle to the front of the struct. * tree-core.h (enum omp_clause_code): Add OMP_CLAUSE_DETACH. Fix ordering. * tree-nested.c (convert_nonlocal_omp_clauses): Handle OMP_CLAUSE_DETACH clause. (convert_local_omp_clauses): Handle OMP_CLAUSE_DETACH clause. * tree-pretty-print.c (dump_omp_clause): Handle OMP_CLAUSE_DETACH. * tree.c (omp_clause_num_ops): Add entry for OMP_CLAUSE_DETACH. Fix ordering. (omp_clause_code_name): Add entry for OMP_CLAUSE_DETACH. Fix ordering. (walk_tree_1): Handle OMP_CLAUSE_DETACH. gcc/c-family/ * c-pragma.h (pragma_omp_clause): Add PRAGMA_OMP_CLAUSE_DETACH. Redefine PRAGMA_OACC_CLAUSE_DETACH. gcc/c/ * c-parser.c (c_parser_omp_clause_detach): New. (c_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause. (OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH. * c-typeck.c (c_finish_omp_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH clause. Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. gcc/cp/ * parser.c (cp_parser_omp_clause_detach): New. (cp_parser_omp_all_clauses): Handle PRAGMA_OMP_CLAUSE_DETACH. (OMP_TASK_CLAUSE_MASK): Add mask for PRAGMA_OMP_CLAUSE_DETACH. * pt.c (tsubst_omp_clauses): Handle OMP_CLAUSE_DETACH clause. * semantics.c (finish_omp_clauses): Handle OMP_CLAUSE_DETACH clause. Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. gcc/fortran/ * dump-parse-tree.c (show_omp_clauses): Handle detach clause. * frontend-passes.c (gfc_code_walker): Walk detach expression. * gfortran.h (struct gfc_omp_clauses): Add detach field. (gfc_c_intptr_kind): New. * openmp.c (gfc_free_omp_clauses): Free detach clause. (gfc_match_omp_detach): New. (enum omp_mask1): Add OMP_CLAUSE_DETACH. (enum omp_mask2): Remove OMP_CLAUSE_DETACH. (gfc_match_omp_clauses): Handle OMP_CLAUSE_DETACH for OpenMP. (OMP_TASK_CLAUSES): Add OMP_CLAUSE_DETACH. (resolve_omp_clauses): Prevent use of detach with mergeable and overriding the data sharing mode of the event handle. * trans-openmp.c (gfc_trans_omp_clauses): Handle detach clause. * trans-types.c (gfc_c_intptr_kind): New. (gfc_init_kinds): Initialize gfc_c_intptr_kind. * types.def (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT): Rename to... (BT_FN_VOID_OMPFN_PTR_OMPCPYFN_LONG_LONG_BOOL_UINT_PTR_INT_PTR): ...this. Add extra argument. gcc/testsuite/ * c-c++-common/gomp/task-detach-1.c: New. * g++.dg/gomp/task-detach-1.C: New. * gcc.dg/gomp/task-detach-1.c: New. * gfortran.dg/gomp/task-detach-1.f90: New. include/ * gomp-constants.h (GOMP_TASK_FLAG_DETACH): New. libgomp/ * fortran.c (omp_fulfill_event_): New. * libgomp.h (struct gomp_task): Add detach and completion_sem fields. (struct gomp_team): Add task_detach_queue and task_detach_count fields. * libgomp.map (OMP_5.0.1): Add omp_fulfill_event and omp_fulfill_event_. * libgomp_g.h (GOMP_task): Add extra argument. * omp.h.in (enum omp_event_handle_t): New. (omp_fulfill_event): New. * omp_lib.f90.in (omp_event_handle_kind): New. (omp_fulfill_event): New. * omp_lib.h.in (omp_event_handle_kind): New. (omp_fulfill_event): Declare. * priority_queue.c (priority_tree_find): New. (priority_list_find): New. (priority_queue_find): New. * priority_queue.h (priority_queue_predicate): New. (priority_queue_find): New. * task.c (gomp_init_task): Initialize detach field. (task_fulfilled_p): New. (GOMP_task): Add detach argument. Ignore detach argument if GOMP_TASK_FLAG_DETACH not set in flags. Initialize completion_sem field. Copy address of completion_sem into detach argument and into the start of the data record. Wait for detach event if task not deferred. (gomp_barrier_handle_tasks): Queue tasks with unfulfilled events. Remove completed tasks and requeue dependent tasks. (omp_fulfill_event): New. * team.c (gomp_new_team): Initialize task_detach_queue and task_detach_count fields. (free_team): Free task_detach_queue field. * testsuite/libgomp.c-c++-common/task-detach-1.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-2.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-3.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-4.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-5.c: New testcase. * testsuite/libgomp.c-c++-common/task-detach-6.c: New testcase. * testsuite/libgomp.fortran/task-detach-1.f90: New testcase. * testsuite/libgomp.fortran/task-detach-2.f90: New testcase. * testsuite/libgomp.fortran/task-detach-3.f90: New testcase. * testsuite/libgomp.fortran/task-detach-4.f90: New testcase. * testsuite/libgomp.fortran/task-detach-5.f90: New testcase. * testsuite/libgomp.fortran/task-detach-6.f90: New testcase.
2021-01-04Update copyright years.Jakub Jelinek4-4/+4
2020-12-18openmp: Don't optimize shared to firstprivate on task with depend clauseJakub Jelinek1-0/+47
The attached testcase is miscompiled, because we optimize shared clauses to firstprivate when task body can't modify the variable even when the task has depend clause. That is wrong, because firstprivate means the variable will be copied immediately when the task is created, while with depend clause some other task might change it later before the dependencies are satisfied and the task should observe the value only after the change. 2020-12-18 Jakub Jelinek <jakub@redhat.com> * gimplify.c (struct gimplify_omp_ctx): Add has_depend member. (gimplify_scan_omp_clauses): Set it to true if OMP_CLAUSE_DEPEND appears on OMP_TASK. (gimplify_adjust_omp_clauses_1, gimplify_adjust_omp_clauses): Force GOVD_WRITTEN on shared variables if task construct has depend clause. * testsuite/libgomp.c/task-6.c: New test.
2020-12-08Fortran: Add 'omp scan' support of OpenMP 5.0Tobias Burnus1-0/+115
gcc/fortran/ChangeLog: * dump-parse-tree.c (show_omp_clauses, show_omp_node, show_code_node): Handle OMP SCAN. * gfortran.h (enum gfc_statement): Add ST_OMP_SCAN. (enum): Add OMP_LIST_SCAN_IN and OMP_LIST_SCAN_EX. (enum gfc_exec_op): Add EXEC_OMP_SCAN. * match.h (gfc_match_omp_scan): New prototype. * openmp.c (gfc_match_omp_scan): New. (gfc_match_omp_taskgroup): Cleanup. (resolve_omp_clauses, gfc_resolve_omp_do_blocks, omp_code_to_statement, gfc_resolve_omp_directive): Handle 'omp scan'. * parse.c (decode_omp_directive, next_statement, gfc_ascii_statement): Likewise. * resolve.c (gfc_resolve_code): Handle EXEC_OMP_SCAN. * st.c (gfc_free_statement): Likewise. * trans-openmp.c (gfc_trans_omp_clauses, gfc_trans_omp_do, gfc_split_omp_clauses): Handle 'omp scan'. libgomp/ChangeLog: * testsuite/libgomp.fortran/scan-1.f90: New test. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/reduction4.f90: Update; move FE some tests to ... * gfortran.dg/gomp/reduction6.f90: ... this new test and ... * gfortran.dg/gomp/reduction7.f90: ... this new test. * gfortran.dg/gomp/reduction5.f90: Add dg-error. * gfortran.dg/gomp/scan-1.f90: New test. * gfortran.dg/gomp/scan-2.f90: New test. * gfortran.dg/gomp/scan-3.f90: New test. * gfortran.dg/gomp/scan-4.f90: New test. * gfortran.dg/gomp/scan-5.f90: New test. * gfortran.dg/gomp/scan-6.f90: New test. * gfortran.dg/gomp/scan-7.f90: New test.
2020-11-25Fix templatized C++ OpenACC 'cache' directive ICEsThomas Schwinge2-1/+24
This has been broken forever, whoops... gcc/cp/ * pt.c (tsubst_omp_clauses): Handle 'OMP_CLAUSE__CACHE_'. (tsubst_expr): Handle 'OACC_CACHE'. gcc/testsuite/ * c-c++-common/goacc/cache-1.c: Update. * c-c++-common/goacc/cache-2.c: Likewise. * g++.dg/goacc/cache-1.C: New. * g++.dg/goacc/cache-2.C: Likewise. libgomp/ * testsuite/libgomp.oacc-c++/cache-1.C: New. * testsuite/libgomp.oacc-c-c++-common/cache-1.c: Update.
2020-11-25Fix atomic_capture-1.f90 testcaseAndrew Stubbs1-19/+19
The testcase had invalid assumptions about which loop iterations would run first and last. libgomp/ChangeLog * testsuite/libgomp.oacc-fortran/atomic_capture-1.f90 (main): Adjust expected results.
2020-11-24[testsuite] Avoid Tcl 8.5-specific behaviorThomas Schwinge2-0/+16
gcc/ * doc/install.texi (Prerequisites) <Tcl>: Add comment. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-1.c: Avoid Tcl 8.5-specific behavior. * c-c++-common/goacc/kernels-decompose-2.c: Likewise. * gfortran.dg/goacc/kernels-decompose-1.f95: Likewise. * gfortran.dg/goacc/kernels-decompose-2.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Avoid Tcl 8.5-specific behavior. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise. Reported-by: David Edelsohn <dje.gcc@gmail.com>
2020-11-18testsuite/libgomp.c/usleep.h: Use sleep-loop also for GCNTobias Burnus1-3/+4
As typically configured, newlib's libc.a does not build 'posix' and, hence, usleep is not available. Thus, use the same fallback as for nvptx. libgomp/ * testsuite/libgomp.c/usleep.h (fallback_usleep): Renamed from nvptx_usleep; use also for device={arch(gcn)}.
2020-11-14openmp: Add support for non-VLA {,first}private allocate on omp taskJakub Jelinek2-86/+195
This patch adds support for custom allocators on private/firstprivate clauses for task (and taskloop) constructs. Private didn't need anything special, but firstprivate if it is passed by reference needs the GOMP_alloc calls in the copyfn and GOMP_free in the task body. 2020-11-14 Jakub Jelinek <jakub@redhat.com> * gimplify.c (gimplify_omp_for): Add OMP_CLAUSE_ALLOCATE_ALLOCATOR decls as firstprivate on task clauses even when allocate clause decl is not lastprivate. * omp-low.c (install_var_field): Don't dereference omp_is_reference types if mask is 33 rather than 1. (scan_sharing_clauses): Populate allocate_map even for task constructs. For now remove it back for variables mentioned in reduction and in_reduction clauses on task/taskloop constructs or on VLA task firstprivates. For firstprivate on task construct, install the var field into field_map with by_ref and 33 instead of false and 1 if mentioned in allocate clause. (lower_private_allocate): Set TREE_THIS_NOTRAP on the created MEM_REF. (lower_rec_input_clauses): Handle allocate for task firstprivatized non-VLA variables. (create_task_copyfn): Likewise. * testsuite/libgomp.c-c++-common/allocate-1.c (struct S): New type. (foo): Add tests for non-VLA private and firstprivate clauses on omp task. (bar): Likewise. Remove taking of address from private/firstprivate variables. * testsuite/libgomp.c++/allocate-1.C (struct S): New type. (foo): Add p, q, px and s arguments. Add tests for array reductions and for non-VLA private and firstprivate clauses on omp task. (bar): Removed. (main): Adjust foo caller. Don't call bar.
2020-11-13Decompose OpenACC 'kernels' constructs into parts, a sequence of compute ↵Gergö Barany5-3/+66
constructs Not yet enabled by default: for now, the current mode of OpenACC 'kernels' constructs handling still remains '-fopenacc-kernels=parloops', but that is to change later. gcc/ * omp-oacc-kernels-decompose.cc: New. * Makefile.in (OBJS): Add it. * passes.def: Instantiate it. * tree-pass.h (make_pass_omp_oacc_kernels_decompose): Declare. * flag-types.h (enum openacc_kernels): Add. * doc/invoke.texi (-fopenacc-kernels): Document. * gimple.h (enum gf_mask): Add 'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_PARALLELIZED', 'GF_OMP_TARGET_KIND_OACC_PARALLEL_KERNELS_GANG_SINGLE', 'GF_OMP_TARGET_KIND_OACC_DATA_KERNELS'. (is_gimple_omp_oacc, is_gimple_omp_offloaded): Handle these. * gimple-pretty-print.c (dump_gimple_omp_target): Likewise. * omp-expand.c (expand_omp_target, build_omp_regions_1) (omp_make_gimple_edges): Likewise. * omp-low.c (scan_sharing_clauses, scan_omp_for) (check_omp_nesting_restrictions, lower_oacc_reductions) (lower_oacc_head_mark, lower_omp_target): Likewise. * omp-offload.c (execute_oacc_device_lower): Likewise. gcc/c-family/ * c.opt (fopenacc-kernels): Add. gcc/fortran/ * lang.opt (fopenacc-kernels): Add. gcc/testsuite/ * c-c++-common/goacc/kernels-decompose-1.c: New. * c-c++-common/goacc/kernels-decompose-2.c: New. * c-c++-common/goacc/kernels-decompose-ice-1.c: New. * c-c++-common/goacc/kernels-decompose-ice-2.c: New. * gfortran.dg/goacc/kernels-decompose-1.f95: New. * gfortran.dg/goacc/kernels-decompose-2.f95: New. * c-c++-common/goacc/if-clause-2.c: Adjust. * gfortran.dg/goacc/kernels-tree.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose-ice-1.c: New. * testsuite/libgomp.oacc-c-c++-common/declare-vla-kernels-decompose.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/declare-vla.c: Adjust. * testsuite/libgomp.oacc-fortran/pr94358-1.f90: Likewise. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2020-11-13Add 'libgomp.oacc-fortran/pr94358-1.f90' [PR94358]Gergö Barany1-0/+34
Document status quo re PR94358 "[OMP] Privatize internal array variables introduced by the Fortran FE". libgomp/ PR fortran/94358 * testsuite/libgomp.oacc-fortran/pr94358-1.f90: New. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
2020-11-13openmp: Support allocate for C/C++ array section reductionsJakub Jelinek1-6/+31
This adds allocate clause support for array section reductions. Furthermore, it fixes one bug that would cause inscan reductions with allocate to be rejected by C, and for now just ignores allocate for inscan/task reductions, that will need slightly more work. 2020-11-13 Jakub Jelinek <jakub@redhat.com> gcc/ * omp-low.c (scan_sharing_clauses): For now remove for reduction clauses with inscan or task modifiers decl from allocate_map. (lower_private_allocate): Handle TYPE_P (new_var). (lower_rec_input_clauses): Handle allocate clause for C/C++ array reductions. gcc/c/ * c-typeck.c (c_finish_omp_clauses): Don't clear OMP_CLAUSE_REDUCTION_INSCAN unless reduction_seen == -2. libgomp/ * testsuite/libgomp.c-c++-common/allocate-1.c (foo): Add tests for array reductions. (main): Adjust foo callers.
2020-11-12openmp: Implement allocate clause in omp lowering.Jakub Jelinek2-0/+448
For now, task/taskloop constructs aren't handled and C/C++ array reductions and reductions with task or inscan modifiers need further work. Instead of calling omp_alloc/omp_free (where the former doesn't have alignment argument and omp_aligned_alloc is 5.1 only feature), this calls GOMP_alloc/GOMP_free, so that the library can fail if it would fall back into NULL (exception is zero length allocations). 2020-11-12 Jakub Jelinek <jakub@redhat.com> gcc/ * builtin-types.def (BT_FN_PTR_SIZE_SIZE_PTRMODE): New function type. * omp-builtins.def (BUILT_IN_GOACC_DECLARE): Move earlier. (BUILT_IN_GOMP_ALLOC, BUILT_IN_GOMP_FREE): New builtins. * gimplify.c (gimplify_scan_omp_clauses): Force allocator into a decl if it is not NULL, INTEGER_CST or decl. (gimplify_adjust_omp_clauses): Clear GOVD_EXPLICIT on explicit clauses which are being removed. Remove allocate clauses for variables not seen if they are private, firstprivate or linear too. Call omp_notice_variable on the allocator otherwise. (gimplify_omp_for): Handle iterator vars mentioned in allocate clauses similarly to non-is_gimple_reg iterators. * omp-low.c (struct omp_context): Add allocate_map field. (delete_omp_context): Delete it. (scan_sharing_clauses): Fill it from allocate clauses. Remove it if mentioned also in shared clause. (lower_private_allocate): New function. (lower_rec_input_clauses): Handle allocate clause for privatized variables, except for task/taskloop, C/C++ array reductions for now and task/inscan variables. (lower_send_shared_vars): Don't consider variables in allocate_map as shared. * omp-expand.c (expand_omp_for_generic, expand_omp_for_static_nochunk, expand_omp_for_static_chunk): Use expand_omp_build_assign instead of gimple_build_assign + gsi_insert_after. * builtins.c (builtin_fnspec): Handle BUILTIN_GOMP_ALLOC and BUILTIN_GOMP_FREE. * tree-ssa-ccp.c (evaluate_stmt): Handle BUILTIN_GOMP_ALLOC. * tree-ssa-dce.c (mark_stmt_if_obviously_necessary): Handle BUILTIN_GOMP_ALLOC. (mark_all_reaching_defs_necessary_1): Handle BUILTIN_GOMP_ALLOC and BUILTIN_GOMP_FREE. (propagate_necessity): Likewise. gcc/fortran/ * f95-lang.c (ATTR_ALLOC_WARN_UNUSED_RESULT_SIZE_2_NOTHROW_LIST): Define. (gfc_init_builtin_functions): Add alloc_size and warn_unused_result attributes to __builtin_GOMP_alloc. * types.def (BT_PTRMODE): New primitive type. (BT_FN_VOID_PTR_PTRMODE, BT_FN_PTR_SIZE_SIZE_PTRMODE): New function types. libgomp/ * libgomp.map (GOMP_alloc, GOMP_free): Export at GOMP_5.0.1. * omp.h.in (omp_alloc): Add malloc and alloc_size attributes. * libgomp_g.h (GOMP_alloc, GOMP_free): Declare. * allocator.c (omp_aligned_alloc): New for now static function, add alignment argument and handle it. (omp_alloc): Reimplement using omp_aligned_alloc. (GOMP_alloc, GOMP_free): New functions. (omp_free): Add ialias. * testsuite/libgomp.c-c++-common/allocate-1.c: New test. * testsuite/libgomp.c++/allocate-1.C: New test.
2020-11-12Adjust 'libgomp.oacc-fortran/attach-descriptor-1.f90' for improved location ↵Thomas Schwinge1-2/+1
information Fix-up for commit b71ff8c15f5a7d6b1cc1524b4d27843f0d88dbda "Fortran: improve location data for OpenACC/OpenMP directives [PR97782]". libgomp/ PR fortran/97782 * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Adjust.
2020-11-10openmp: Implement OpenMP 5.0 base-pointer attachement and clause orderingChung-Lin Tang1-0/+82
This patch implements some parts of the target variable mapping changes specified in OpenMP 5.0, including base-pointer attachment/detachment behavior for array section list-items in map clauses, and ordering of map clauses according to map kind. 2020-11-10 Chung-Lin Tang <cltang@codesourcery.com> gcc/c-family/ChangeLog: * c-common.h (c_omp_adjust_map_clauses): New declaration. * c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses. (c_omp_adjust_map_clauses): New function. gcc/c/ChangeLog: * c-parser.c (c_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (c_parser_omp_target_enter_data): Likewise. (c_parser_omp_target_exit_data): Likewise. (c_parser_omp_target): Likewise. * c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. (c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/cp/ChangeLog: * parser.c (cp_parser_omp_target_data): Add use of new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as handled map clause kind. (cp_parser_omp_target_enter_data): Likewise. (cp_parser_omp_target_exit_data): Likewise. (cp_parser_omp_target): Likewise. * semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix interaction between reference case and attach/detach. (finish_omp_clauses): Adjust bitmap checks to allow struct decl and same struct field access to co-exist on OpenMP construct. gcc/ChangeLog: * gimplify.c (is_or_contains_p): New static helper function. (omp_target_reorder_clauses): New function. (gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to reorder clause list according to OpenMP 5.0 rules. Add handling of GOMP_MAP_ATTACH_DETACH for OpenMP cases. * omp-low.c (is_omp_target): New static helper function. (scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. (lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for OpenMP cases. gcc/testsuite/ChangeLog: * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid. * gfortran.dg/gomp/map-2.f90: Likewise. * c-c++-common/gomp/map-5.c: New testcase. libgomp/ChangeLog: * libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag usable. * oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to 'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'. (goacc_enter_datum): Likewise for call to gomp_map_vars_async. (goacc_enter_data_internal): Likewise. * target.c (gomp_map_vars_internal): Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use of gomp_attach_pointer for OpenMP cases. (gomp_exit_data): Add handling of GOMP_MAP_DETACH. (GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH. * testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.
2020-11-05openmp: Mark deprecated symbols in OpenMP 5.0Kwok Cheung Yeung26-1/+37
2020-11-05 Ulrich Drepper <drepper@redhat.com> Kwok Cheung Yeung <kcy@codesourcery.com> libgomp/ * Makefile.am (%.mod): Add -cpp and -fopenmp to compile flags. * Makefile.in: Regenerate. * fortran.c: Wrap uses of omp_set_nested and omp_get_nested with pragmas to ignore -Wdeprecated-declarations warnings. * icv.c: Likewise. * omp.h.in (__GOMP_DEPRECATED_5_0): Define. Mark omp_lock_hint_* enum values, omp_lock_hint_t, omp_set_nested, and omp_get_nested with __GOMP_DEPRECATED_5_0. * omp_lib.f90.in: Mark omp_get_nested and omp_set_nested as deprecated. * testsuite/libgomp.c++/affinity-1.C: Add -Wno-deprecated-declarations to test options. * testsuite/libgomp.c/affinity-1.c: Likewise. * testsuite/libgomp.c/affinity-2.c: Likewise. * testsuite/libgomp.c/appendix-a/a.15.1.c: Likewise. * testsuite/libgomp.c/lib-1.c: Likewise. * testsuite/libgomp.c/nested-1.c: Likewise. * testsuite/libgomp.c/nested-2.c: Likewise. * testsuite/libgomp.c/nested-3.c: Likewise. * testsuite/libgomp.c/pr32362-1.c: Likewise. * testsuite/libgomp.c/pr32362-2.c: Likewise. * testsuite/libgomp.c/pr32362-3.c: Likewise. * testsuite/libgomp.c/pr35549.c: Likewise. * testsuite/libgomp.c/pr42942.c: Likewise. * testsuite/libgomp.c/pr61200.c: Likewise. * testsuite/libgomp.c/sort-1.c: Likewise. * testsuite/libgomp.c/target-5.c: Likewise. * testsuite/libgomp.c/target-6.c: Likewise. * testsuite/libgomp.c/teams-1.c: Likewise. * testsuite/libgomp.c/thread-limit-1.c: Likewise. * testsuite/libgomp.c/thread-limit-2.c: Likewise. * testsuite/libgomp.c/thread-limit-4.c: Likewise. * testsuite/libgomp.fortran/affinity1.f90: Likewise. * testsuite/libgomp.fortran/lib1.f90: Likewise. * testsuite/libgomp.fortran/lib2.f: Likewise. * testsuite/libgomp.fortran/nested1.f90: Likewise. * testsuite/libgomp.fortran/teams1.f90: Likewise.
2020-11-02Simplify and enhance 'libgomp.oacc-c-c++-common/pr85486*.c' [PR85486]Thomas Schwinge3-97/+20
Avoid code duplication, and better test what we expect to happen. libgomp/ PR target/85486 * testsuite/libgomp.oacc-c-c++-common/pr85486-2.c: Simplify and enhance. * testsuite/libgomp.oacc-c-c++-common/pr85486-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise.
2020-11-02libgomp testsuite: tell warning from error diagnostics, etc. [PR80219, PR85303]Thomas Schwinge1-0/+6
This changes makes 'dg-warning', 'dg-error', 'dg-bogus', 'dg-message' behave as expected, and also enables use of relative line numbers as well as 'dg-line'. libgomp/ PR testsuite/80219 PR testsuite/85303 * testsuite/lib/libgomp.exp (libgomp_init): Set 'gcc_warning_prefix', 'gcc_error_prefix'.
generated by cgit v1.1 at 2025-09-25 07:24:20 +0000