[patch] Avoid an unwanted decl re-map in copy_gimple_seq_and_replace_locals

[Martin Jambor <mjambor at suse dot cz>]

Hi,

I ran into an ICE when compiling the following function on the HSA branch:

foo (int n, int m, int o, int (*a)[m][o])
{
  int i, j, k;
#pragma omp target teams distribute parallel for shared(a) firstprivate(n, m, o) private(i,j,k)
  for (i = 0; i < n; i++)
    for (j = 0; j < m; j++)
      for (k = 0; k < o; k++)
        a[i][j][k] = i + j + k;
}

The problem is that when I duplicate the loop with
copy_gimple_seq_and_replace_locals, I get one extra re-mapping.
Specifically, I feed the function this:

{
  int i.2;

  #pragma omp teams shared(a) firstprivate(n) firstprivate(m) firstprivate(o) shared(m.1) shared(D.3275) shared(o.0)
    {
      #pragma omp distribute private(i.2)
      for (i.2 = 0; i.2 < n; i.2 = i.2 + 1)
        {
          #pragma omp parallel shared(a) firstprivate(n) firstprivate(m) firstprivate(o) private(i) private(j) private(k) shared(m.1) shared(D.3275) shared(o.0)
            {
              sizetype D.3286;
              long unsigned int D.3287;
              sizetype D.3288;
              sizetype D.3289;
              sizetype D.3290;
              long unsigned int D.3291;
              long unsigned int D.3292;
              int[0:D.3279][0:D.3271] * D.3293;
              int D.3294;
              int D.3295;

              #pragma omp for nowait
              for (i = 0; i < n; i = i + 1)
                {
                  j = 0;
                  goto <D.3244>;
                  <D.3243>:
                  k = 0;
                  goto <D.3241>;
                  <D.3240>:
                  D.3286 = D.3275 /[ex] 4;   <--- here I get wrog decl
                  D.3287 = (long unsigned int) i;
                  D.3288 = (sizetype) o.0;
                  D.3289 = (sizetype) m.1;
                  D.3290 = D.3288 * D.3289;
                  D.3291 = D.3287 * D.3290;
                  D.3292 = D.3291 * 4;
                  D.3293 = a + D.3292;
                  D.3294 = i + j;
                  D.3295 = D.3294 + k;
                  *D.3293[j]{lb: 0 sz: D.3286 * 4}[k] = D.3295;
                  k = k + 1;
                  <D.3241>:
                  if (k < o) goto <D.3240>; else goto <D.3242>;
                  <D.3242>:
                  j = j + 1;
                  <D.3244>:
                  if (j < m) goto <D.3243>; else goto <D.3245>;
                  <D.3245>:
                }
            }
        }
    }
}

and it replaces D.3275 with its new copy with undefined value.  The
mapping is created when an array type where the size is defined in
terms of that variable declaration is copied.  The comment in
type-remapping code says that we "use the already remaped data" but
that is not true.

My solution was to prevent declaration duplication in this case with
yet another state variable in struct copy_body_data that holds a
special value when we are running copy_gimple_seq_and_replace_locals
and another when we are within type-remapping.

I'll be happy for any suggestion how to deal with this without
cluttering copy_body_date even more but so far I have not found any.
If nobody has a better idea, is the following good for trunk?  (I am
about to commit it to the hsa branch.)  It has passed bootstrap and
testing on x86_64-linux.

Thanks,

Martin


2016-01-06  Martin Jambor  <mjambor@suse.cz>

	* tree-inline.h (copy_body_data): New field
	decl_creation_prevention_level.  Moved remap_var_for_cilk to minimize
	padding.
	* tree-inline.c (remap_decl): Return original decls if
	decl_creation_prevention_level is two or bigger.
	(remap_type_1): Increment and decrement decl_creation_prevention_level
	if appropriate.
	(copy_gimple_seq_and_replace_locals): Set
	decl_creation_prevention_level to 1.

diff --git a/gcc/tree-inline.c b/gcc/tree-inline.c
index 88a6753..2df11a2 100644
--- a/gcc/tree-inline.c
+++ b/gcc/tree-inline.c
@@ -340,8 +340,20 @@ remap_decl (tree decl, copy_body_data *id)
       return decl;
     }
 
-  /* If we didn't already have an equivalent for this declaration,
-     create one now.  */
+  /* If decl copying is forbidden (which happens when copying a type with size
+     defined outside of the copied sequence) work with the original decl. */
+  if (!n
+      && id->decl_creation_prevention_level > 1
+      && (VAR_P (decl) || TREE_CODE (decl) == PARM_DECL))
+    {
+      if (id->do_not_unshare)
+	return decl;
+      else
+	return unshare_expr (decl);
+    }
+
+  /* If we didn't already have an equivalent for this declaration, create one
+     now.  */
   if (!n)
     {
       /* Make a copy of the variable or label.  */
@@ -526,7 +538,10 @@ remap_type_1 (tree type, copy_body_data *id)
       gcc_unreachable ();
     }
 
-  /* All variants of type share the same size, so use the already remaped data.  */
+  /* All variants of type share the same size, so use the already remaped
+     data.  */
+  if (id->decl_creation_prevention_level > 0)
+    id->decl_creation_prevention_level++;
   if (TYPE_MAIN_VARIANT (new_tree) != new_tree)
     {
       gcc_checking_assert (TYPE_SIZE (type) == TYPE_SIZE (TYPE_MAIN_VARIANT (type)));
@@ -540,6 +555,8 @@ remap_type_1 (tree type, copy_body_data *id)
       walk_tree (&TYPE_SIZE (new_tree), copy_tree_body_r, id, NULL);
       walk_tree (&TYPE_SIZE_UNIT (new_tree), copy_tree_body_r, id, NULL);
     }
+  if (id->decl_creation_prevention_level > 1)
+    id->decl_creation_prevention_level--;
 
   return new_tree;
 }
@@ -5276,6 +5293,7 @@ copy_gimple_seq_and_replace_locals (gimple_seq seq)
   id.transform_return_to_modify = false;
   id.transform_parameter = false;
   id.transform_lang_insert_block = NULL;
+  id.decl_creation_prevention_level = 1;
 
   /* Walk the tree once to find local labels.  */
   memset (&wi, 0, sizeof (wi));
diff --git a/gcc/tree-inline.h b/gcc/tree-inline.h
index b8fb2a2..f737daf 100644
--- a/gcc/tree-inline.h
+++ b/gcc/tree-inline.h
@@ -140,14 +140,21 @@ struct copy_body_data
      the originals have been mapped to a value rather than to a
      variable.  */
   hash_map<tree, tree> *debug_map;
- 
-  /* Cilk keywords currently need to replace some variables that
-     ordinary nested functions do not.  */ 
-  bool remap_var_for_cilk;
 
   /* A map from the inlined functions dependence info cliques to
      equivalents in the function into which it is being inlined.  */
   hash_map<dependence_hash, unsigned short> *dependence_map;
+
+  /* Cilk keywords currently need to replace some variables that
+     ordinary nested functions do not.  */
+  bool remap_var_for_cilk;
+
+  /* When zero, do nothing.  When one or higher, increment during type
+     remapping, When two or higher, do not create new variables when remapping
+     decls.  Used when remapping types with variable size, but when the size is
+     defined outside the sequence copied by
+     copy_gimple_seq_and_replace_locals.  */
+  unsigned decl_creation_prevention_level;
 };
 
 /* Weights of constructions for estimate_num_insns.  */