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 --git a/gcc/config/gcn/gcn-tree.cc b/gcc/config/gcn/gcn-tree.cc
index a71aadc395a0af2b9e58e353c56985e988aa71dc..97f35cd6f7a065025c205fb3d7be8828eb90fc33 100644 (file)
--- 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 886a39a26e95e0951f69448ad37caf57a0affc75..8e5d9ecce70ac82ef9a7e974861ea199fb79055d 100644 (file)
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -6500,16 +6500,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);
@@ -6605,6 +6605,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,
@@ -6737,6 +6738,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));
 
@@ -6787,6 +6790,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);
@@ -7758,11 +7765,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))
        {
          if (!array_p)
@@ -7785,13 +7792,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
        {
@@ -8026,6 +8034,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 43e139f923147ff1a56c6ee594feaab2283d9114..db8b3749959ffc936584eef5383659c6739810d1 100644 (file)
--- 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 aeae2e06252c326beaad44f876c74582fefcaca3..0f023b7c4607e55c12314819ab925a7a5e1f6111 100644 (file)
--- 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 c095284e56607299401f70fcf10aba3d2f0ba696..94dd4c473d57da14de5b6c9bd8ddcc7ea04e6a4d 100644 (file)
--- 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 4794350e8621706d987a70fa95222af4709e1230..56ae020d0fc137a19ca511003dc310246bbc92e7 100644 (file)
--- 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)