[gcc/devel/omp/gcc-16] OpenACC: array reductions bug fixes
Sandra Loosemore via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:c710bedf1c8618f762d9eedb50bc725899f413b9 commit c710bedf1c8618f762d9eedb50bc725899f413b9 Author: Chung-Lin Tang <[email protected]> Date: Fri May 2 00:33:07 2025 +0000 OpenACC: array reductions bug fixes This is a merge of the v4 to v5 diff patch from: https://gcc.gnu.org/pipermail/gcc-patches/2025-March/679682.html This patch fixes issues found for NVPTX sm_70 testing, and another issue related to copying to reduction buffer for worker/vector mode. gcc/ChangeLog: * config/gcn/gcn-tree.cc (gcn_goacc_reduction_setup): Fix array case copy source into reduction buffer. * config/nvptx/nvptx.cc (nvptx_expand_shared_addr): Move default size init setting place. (enum nvptx_builtins): Add NVPTX_BUILTIN_BAR_WARPSYNC. (nvptx_init_builtins): Add DEF() of nvptx_builtin_bar_warpsync. (nvptx_expand_builtin): Expand NVPTX_BUILTIN_BAR_WARPSYNC. (nvptx_goacc_reduction_setup): Fix array case copy source into reduction buffer. (nvptx_goacc_reduction_fini): Add bar.warpsync for at end of vector-mode reductions for sm_70 and above. libgomp/ChangeLog: * testsuite/libgomp.oacc-c-c++-common/reduction-arrays-2.c: Adjust test. * testsuite/libgomp.oacc-c-c++-common/reduction-arrays-3.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-arrays-4.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/reduction-arrays-5.c: Likewise. Diff: --- gcc/config/gcn/gcn-tree.cc | 3 +- gcc/config/nvptx/nvptx.cc | 32 ++++++++++++++++------ .../libgomp.oacc-c-c++-common/reduction-arrays-2.c | 24 ++++++++++++++++ .../libgomp.oacc-c-c++-common/reduction-arrays-3.c | 24 ++++++++++++++++ .../libgomp.oacc-c-c++-common/reduction-arrays-4.c | 24 ++++++++++++++++ .../libgomp.oacc-c-c++-common/reduction-arrays-5.c | 24 ++++++++++++++++ 6 files changed, 122 insertions(+), 9 deletions(-) diff --git a/gcc/config/gcn/gcn-tree.cc b/gcc/config/gcn/gcn-tree.cc index 9e9ba9251232..d31b0230b3b3 100644 --- a/gcc/config/gcn/gcn-tree.cc +++ b/gcc/config/gcn/gcn-tree.cc @@ -750,13 +750,14 @@ gcn_goacc_reduction_setup (gcall *call) tree offset = gimple_call_arg (call, 5); if (array_p) { + tree copy_src = !integer_zerop (ref_to_res) ? ref_to_res : array_addr; tree decl = gcn_goacc_get_worker_array_reduction_buffer (array_type, array_max_idx, &seq); tree ptr = make_ssa_name (TREE_TYPE (array_addr)); gimplify_assign (ptr, build_fold_addr_expr (decl), &seq); /* Store incoming value to worker reduction buffer. */ - oacc_build_array_copy (ptr, array_addr, array_max_idx, &seq); + oacc_build_array_copy (ptr, copy_src, array_max_idx, &seq); } else { diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc index 86389bbb69d7..c4b9650ecb62 100644 --- a/gcc/config/nvptx/nvptx.cc +++ b/gcc/config/nvptx/nvptx.cc @@ -6515,16 +6515,16 @@ nvptx_expand_shared_addr (tree exp, rtx target, if (TREE_CONSTANT (size_expr)) size = TREE_INT_CST_LOW (size_expr); + /* Default size for unknown size expression. */ + if (size == 0) + size = 256; + if (vector) { offload_attrs oa; populate_offload_attrs (&oa); - /* Default size for unknown size expression. */ - if (size == 0) - size = 256; - unsigned int psize = ROUND_UP (size + offset, align); unsigned int pnum = nvptx_mach_max_workers (); vector_red_partition = MAX (vector_red_partition, psize); @@ -6620,6 +6620,7 @@ enum nvptx_builtins NVPTX_BUILTIN_BAR_RED_AND, NVPTX_BUILTIN_BAR_RED_OR, NVPTX_BUILTIN_BAR_RED_POPC, + NVPTX_BUILTIN_BAR_WARPSYNC, NVPTX_BUILTIN_BREV, NVPTX_BUILTIN_BREVLL, NVPTX_BUILTIN_COND_UNI, @@ -6752,6 +6753,8 @@ nvptx_init_builtins (void) DEF (BAR_RED_POPC, "bar_red_popc", (UINT, UINT, UINT, UINT, UINT, NULL_TREE)); + DEF (BAR_WARPSYNC, "bar_warpsync", (VOID, VOID, NULL_TREE)); + DEF (BREV, "brev", (UINT, UINT, NULL_TREE)); DEF (BREVLL, "brevll", (LLUINT, LLUINT, NULL_TREE)); @@ -6802,6 +6805,10 @@ nvptx_expand_builtin (tree exp, rtx target, rtx ARG_UNUSED (subtarget), case NVPTX_BUILTIN_BAR_RED_POPC: return nvptx_expand_bar_red (exp, target, mode, ignore); + case NVPTX_BUILTIN_BAR_WARPSYNC: + emit_insn (gen_nvptx_warpsync ()); + return NULL_RTX; + case NVPTX_BUILTIN_BREV: case NVPTX_BUILTIN_BREVLL: return nvptx_expand_brev (exp, target, mode, ignore); @@ -7773,11 +7780,11 @@ nvptx_goacc_reduction_setup (gcall *call, offload_attrs *oa) push_gimplify_context (true); + /* Copy the receiver object. */ + tree ref_to_res = gimple_call_arg (call, 1); + if (level != GOMP_DIM_GANG) { - /* Copy the receiver object. */ - tree ref_to_res = gimple_call_arg (call, 1); - if (!integer_zerop (ref_to_res) && !array_p) { ref_to_res = nvptx_adjust_reduction_type (ref_to_res, @@ -7797,13 +7804,14 @@ nvptx_goacc_reduction_setup (gcall *call, offload_attrs *oa) tree call, ptr; if (array_p) { + tree copy_src = !integer_zerop (ref_to_res) ? ref_to_res : array_addr; tree array_elem_type = TREE_TYPE (array_type); call = nvptx_get_shared_red_addr (array_elem_type, array_max_idx, offset, level == GOMP_DIM_VECTOR); ptr = make_ssa_name (TREE_TYPE (call)); gimplify_assign (ptr, call, &seq); oacc_build_array_copy (fold_convert (TREE_TYPE (array_addr), ptr), - array_addr, array_max_idx, &seq); + copy_src, array_max_idx, &seq); } else { @@ -8037,6 +8045,14 @@ nvptx_goacc_reduction_fini (gcall *call, offload_attrs *oa) else r = nvptx_reduction_update (gimple_location (call), &gsi, accum, var, op, level); + + if (TARGET_SM70 && level == GOMP_DIM_VECTOR) + { + /* After SM70, with Independent Thread Scheduling introduced, + place a warpsync after vector-mode update of accum buffer. */ + tree fn = nvptx_builtin_decl (NVPTX_BUILTIN_BAR_WARPSYNC, true); + gimple_seq_add_stmt (&seq, gimple_build_call (fn, 0)); + } } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-2.c index 43e139f92314..db8b3749959f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-2.c @@ -24,6 +24,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel + #pragma acc loop gang reduction(+:a[1:2]) + ARRAY_BODY (a, 1, 2) + ARRAY_BODY (o, 1, 2) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a[3:2]) #pragma acc loop reduction(+:a[3:2]) ARRAY_BODY (a, 3, 2) @@ -32,6 +40,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a[3:2]) + #pragma acc loop worker reduction(+:a[3:2]) + ARRAY_BODY (a, 3, 2) + ARRAY_BODY (o, 3, 2) + for (int i = 0; i < 6; i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a) #pragma acc loop reduction(+:a[0:5]) ARRAY_BODY (a, 0, 5) @@ -40,6 +56,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a) + #pragma acc loop vector reduction(+:a[0:5]) + ARRAY_BODY (a, 0, 5) + ARRAY_BODY (o, 0, 5) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel #pragma acc loop reduction(+:a) ARRAY_BODY (a, 4, 1) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-3.c index aeae2e06252c..0f023b7c4607 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-3.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-3.c @@ -30,6 +30,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel + #pragma acc loop gang reduction(+:a[one:2]) + ARRAY_BODY (a, one, 2) + ARRAY_BODY (o, one, 2) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a[three:2]) #pragma acc loop reduction(+:a[three:2]) ARRAY_BODY (a, three, 2) @@ -38,6 +46,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a[three:2]) + #pragma acc loop worker reduction(+:a[three:2]) + ARRAY_BODY (a, three, 2) + ARRAY_BODY (o, three, 2) + for (int i = 0; i < 6; i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a) #pragma acc loop reduction(+:a[zero:5]) ARRAY_BODY (a, zero, 5) @@ -46,6 +62,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a) + #pragma acc loop vector reduction(+:a[zero:5]) + ARRAY_BODY (a, zero, 5) + ARRAY_BODY (o, zero, 5) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel #pragma acc loop reduction(+:a) ARRAY_BODY (a, four, 1) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-4.c index c095284e5660..94dd4c473d57 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-4.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-4.c @@ -31,6 +31,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel + #pragma acc loop gang reduction(+:a[one:two]) + ARRAY_BODY (a, one, two) + ARRAY_BODY (o, one, two) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a[three:two]) #pragma acc loop reduction(+:a[three:two]) ARRAY_BODY (a, three, two) @@ -39,6 +47,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a[three:two]) + #pragma acc loop worker reduction(+:a[three:two]) + ARRAY_BODY (a, three, two) + ARRAY_BODY (o, three, two) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a) #pragma acc loop reduction(+:a[zero:five]) ARRAY_BODY (a, zero, five) @@ -47,6 +63,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a) + #pragma acc loop vector reduction(+:a[zero:five]) + ARRAY_BODY (a, zero, five) + ARRAY_BODY (o, zero, five) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel #pragma acc loop reduction(+:a) ARRAY_BODY (a, four, one) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-5.c index 4794350e8621..56ae020d0fc1 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-5.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/reduction-arrays-5.c @@ -32,6 +32,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel + #pragma acc loop gang reduction(+:a[one:two]) + ARRAY_BODY (a, one, two) + ARRAY_BODY (o, one, two) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a[three:two]) #pragma acc loop reduction(+:a[three:two]) ARRAY_BODY (a, three, two) @@ -40,6 +48,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a[three:two]) + #pragma acc loop worker reduction(+:a[three:two]) + ARRAY_BODY (a, three, two) + ARRAY_BODY (o, three, two) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel copy(a) #pragma acc loop reduction(+:a[zero:five]) ARRAY_BODY (a, zero, five) @@ -48,6 +64,14 @@ int main (void) if (a[i] != o[i]) __builtin_abort (); + #pragma acc parallel copy(a) + #pragma acc loop vector reduction(+:a[zero:five]) + ARRAY_BODY (a, zero, five) + ARRAY_BODY (o, zero, five) + for (int i = 0; i < sizeof (a) / sizeof (int); i++) + if (a[i] != o[i]) + __builtin_abort (); + #pragma acc parallel #pragma acc loop reduction(+:a) ARRAY_BODY (a, four, one)