[gomp4] signed nums are better for dimensions

[gomp4] signed nums are better for dimensions

[Nathan Sidwell]

[-- Attachment #1: Type: text/plain, Size: 1111 bytes --]

This patch does a number of things.
1) signed results from the DIM_SIZE and DIM_POS internal fns are better.  The 
rest of the loop machinery uses signed ints -- to show overflow is undefined. 
Returning unsigned simply caused an unsigned->signed conversion and subsequent 
confusion.

2) We really should not be getting to the expanders if there's nothing to expand 
to.  That's simply covering up lack of earlier handling.  That earlier removal 
gives optimizers more leeway.

3) Added a hook so the backend can say what the absolute dimension limits are. 
This is used by the DIM_SIZE/DIM_POS transform as a fall back.  It's main raison 
d'etre is for a min/max elision patch I've been  working on, but turned out 
useful here.

4) I'm forcing the oacc function attribute to always have dimensions.  This 
makes users of that data simpler.  However, Cesar discovered that the 
overwriting we were already doing was being broken by LTO's (reasonable) habit 
of commonizing identical tree lists.  I'll be working on correcting the 
behaviour once I have these patches flushed out.

committed to gomp4.

nathan

[-- Attachment #2: gomp4-signed-dim.patch --]
[-- Type: text/x-patch, Size: 16225 bytes --]

2015-08-11  Nathan Sidwell  <nathan@codesourcery.com>

	* internal-fn.c (expand_GOACC_FORK): Use ARG_UNUSED.  Rename
	varible sensibly. Assert we're never reached if there is no
	expansion.
	(expand_GOACC_JOIN, expand_GOACC_DIM_SIZE,
	expand_GOACC_DIM_POS): likewise.
	* targhooks.h (default_goacc_dim_limit): Declare.
	* target.def (validate_dims): Expand documentation.
	(dim_limit): New hook.
	* config/nvptx/nvptx.c (nvptx_reorg): Expect attribute to always
	have dimensions.
	(nvptx_record_offload_symbol): Likewise.
	(nvptx_validate_syms): Force dimension creation and adjust.
	(nvtpx_dim_limit): Define.
	(TARGET_GOACC_DIM_LIMIT): Override.
	* omp-low.c (expand_oacc_get_num_threads): Builtins return int not
	unsigned.
	(expand_oacc_get_thread_num, oacc_init_count_vars): Likewise.
	(oacc_xform_dim): Result is signed.
	(execute_oacc_transform): Expect dims to always be present.
	(default_goacc_validate_dims): Force dimension creation.
	(default_goacc_dim_limit): New.
	* doc/tm.texi.in (TARGET_GOACC_DIM_LIMIT): Add.
	* doc/tm.texi: Rebuilt.

Index: internal-fn.c
===================================================================
--- internal-fn.c	(revision 226770)
+++ internal-fn.c	(working copy)
@@ -1965,58 +1965,62 @@ expand_GOACC_DATA_END_WITH_ARG (gcall *s
 }
 
 static void
-expand_GOACC_FORK (gcall *stmt ATTRIBUTE_UNUSED)
+expand_GOACC_FORK (gcall *ARG_UNUSED (stmt))
 {
 #ifdef HAVE_oacc_fork
-  rtx mode = expand_normal (gimple_call_arg (stmt, 0));
+  rtx dim = expand_normal (gimple_call_arg (stmt, 0));
   
-  emit_insn (gen_oacc_fork (mode));
+  emit_insn (gen_oacc_fork (dim));
+#else
+  gcc_unreachable ();
 #endif
 }
 
 static void
-expand_GOACC_JOIN (gcall *stmt ATTRIBUTE_UNUSED)
+expand_GOACC_JOIN (gcall *ARG_UNUSED (stmt))
 {
 #ifdef HAVE_oacc_join
-  rtx mode = expand_normal (gimple_call_arg (stmt, 0));
+  rtx dim = expand_normal (gimple_call_arg (stmt, 0));
   
-  emit_insn (gen_oacc_join (mode));
+  emit_insn (gen_oacc_join (dim));
+#else
+  gcc_unreachable ();
 #endif
 }
 
 static void
-expand_GOACC_DIM_SIZE (gcall *stmt)
+expand_GOACC_DIM_SIZE (gcall *ARG_UNUSED (stmt))
 {
+#ifdef HAVE_oacc_dim_size
   tree lhs = gimple_call_lhs (stmt);
 
   if (!lhs)
     return;
   
   rtx target = expand_expr (lhs, NULL_RTX, VOIDmode, EXPAND_WRITE);
-#ifdef HAVE_oacc_dim_size
-  rtx val = expand_expr (gimple_call_arg (stmt, 0), NULL_RTX,
+  rtx dim = expand_expr (gimple_call_arg (stmt, 0), NULL_RTX,
 			 VOIDmode, EXPAND_NORMAL);
-  emit_insn (gen_oacc_dim_size (target, val));
+  emit_insn (gen_oacc_dim_size (target, dim));
 #else
-  emit_move_insn (target, const1_rtx);
+  gcc_unreachable ();
 #endif
 }
 
 static void
-expand_GOACC_DIM_POS (gcall *stmt)
+expand_GOACC_DIM_POS (gcall *ARG_UNUSED (stmt))
 {
+#ifdef HAVE_oacc_dim_pos
   tree lhs = gimple_call_lhs (stmt);
 
   if (!lhs)
     return;
   
   rtx target = expand_expr (lhs, NULL_RTX, VOIDmode, EXPAND_WRITE);
-#ifdef HAVE_oacc_dim_pos
-  rtx val = expand_expr (gimple_call_arg (stmt, 0), NULL_RTX,
+  rtx dim = expand_expr (gimple_call_arg (stmt, 0), NULL_RTX,
 			 VOIDmode, EXPAND_NORMAL);
-  emit_insn (gen_oacc_dim_pos (target, val));
+  emit_insn (gen_oacc_dim_pos (target, dim));
 #else
-  emit_move_insn (target, const0_rtx);
+  gcc_unreachable ();
 #endif
 }
 
Index: targhooks.h
===================================================================
--- targhooks.h	(revision 226770)
+++ targhooks.h	(working copy)
@@ -108,6 +108,7 @@ extern void default_finish_cost (void *,
 extern void default_destroy_cost_data (void *);
 
 extern tree default_goacc_validate_dims (tree, tree);
+extern unsigned default_goacc_dim_limit (unsigned);
 extern bool default_goacc_fork_join (bool, gimple_stmt_iterator *, gimple);
 
 /* These are here, and not in hooks.[ch], because not all users of
Index: target.def
===================================================================
--- target.def	(revision 226770)
+++ target.def	(working copy)
@@ -1648,11 +1648,19 @@ DEFHOOK
 (validate_dims,
 "This hook should check the launch dimensions provided.  It should fill\n\
 in default values and verify non-defaults.  The TREE_LIST is unshared\n\
-and may be overwritten.  Diagnostics should be issued as appropriate.",
+and may be overwritten.  If the dimension list is NULL, one should be\n\
+created.  Diagnostics should be issued as appropriate.",
 tree, (tree, tree),
 default_goacc_validate_dims)
 
 DEFHOOK
+(dim_limit,
+"This hook should return the maximum size of a particular dimension,\n\
+or zero if unbounded.",
+unsigned, (unsigned),
+default_goacc_dim_limit)
+
+DEFHOOK
 (fork_join,
 "This hook should convert IFN_GOACC_FORK and IFN_GOACC_JOIN function\n\
 calls to target-specific gimple.  It is executed during the oacc_xform\n\
Index: config/nvptx/nvptx.c
===================================================================
--- config/nvptx/nvptx.c	(revision 226770)
+++ config/nvptx/nvptx.c	(working copy)
@@ -3051,30 +3051,23 @@ nvptx_reorg (void)
   tree attr = get_oacc_fn_attrib (current_function_decl);
   if (attr)
     {
+      /* If we determined this mask before RTL expansion, we could
+	 elide emission of some levels of forks and joins.  */
       unsigned mask = 0;
       tree dims = TREE_VALUE (attr);
       unsigned ix;
 
-      for (ix = 0; ix != GOMP_DIM_MAX; ix++)
+      for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims))
 	{
-	  unsigned HOST_WIDE_INT dim = 0;
+	  HOST_WIDE_INT size = TREE_INT_CST_LOW (TREE_VALUE (dims));
 
-	  if (dims)
-	    {
-	      tree cst = TREE_VALUE (dims);
-
-	      dim = TREE_INT_CST_LOW (cst);
-	      dims = TREE_CHAIN (dims);
-	    }
-	  if (dim != 1)
+	  if (size > 1 || (!size && !TREE_PURPOSE (dims)))
 	    mask |= GOMP_DIM_MASK (ix);
 	}
       /* If there is worker neutering, there must be vector
-	 neutering.  Otherwise the hardware will fail.  This really
-	 should be dealt with earlier because it indicates faulty
-	 logic in determining launch dimensions.  */
-      if (mask & GOMP_DIM_MASK (GOMP_DIM_WORKER))
-	mask |= GOMP_DIM_MASK (GOMP_DIM_VECTOR);
+	 neutering.  Otherwise the hardware will fail.  */
+      gcc_assert (!(mask & GOMP_DIM_MASK (GOMP_DIM_WORKER))
+		  || (mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR)));
 
       parallel *pars = nvptx_discover_pars (&bb_insn_map);
       nvptx_process_pars (pars);
@@ -3175,15 +3168,10 @@ nvptx_record_offload_symbol (tree decl)
 
 	for (ix = 0; ix != GOMP_DIM_MAX; ix++, dims = TREE_CHAIN (dims))
 	  {
-	    tree cst = TREE_VALUE (dims);
+	    HOST_WIDE_INT size = TREE_INT_CST_LOW (TREE_VALUE (dims));
 
-	    /* When device_type support is added an earlier pass
-	       should have massaged the attribute to be
-	       ptx-specific.  */
-	    gcc_assert (TREE_CODE (cst) == INTEGER_CST);
-
-	    unsigned HOST_WIDE_INT dim = TREE_INT_CST_LOW (cst);
-	    fprintf (asm_out_file, ", " HOST_WIDE_INT_PRINT_HEX, dim);
+	    gcc_assert (!TREE_PURPOSE (dims));
+	    fprintf (asm_out_file, ", " HOST_WIDE_INT_PRINT_HEX, size);
 	  }
 
 	fprintf (asm_out_file, "\n");
@@ -3537,18 +3525,23 @@ nvptx_validate_dims (tree decl, tree dim
 {
   tree adims[GOMP_DIM_MAX];
   unsigned ix;
-  tree pos = dims;
+  tree *pos_ptr;
 
-  for (ix = 0; ix != GOMP_DIM_MAX; ix++)
+  for (ix = 0, pos_ptr = &dims; ix != GOMP_DIM_MAX;
+       ix++, pos_ptr = &TREE_CHAIN (*pos_ptr))
     {
-      adims[ix] = TREE_VALUE (pos);
-      pos = TREE_CHAIN (pos);
+      if (!*pos_ptr)
+	*pos_ptr = tree_cons (NULL_TREE, NULL_TREE, NULL_TREE);
+      
+      adims[ix] = TREE_VALUE (*pos_ptr);
     }
+
   /* Define vector size for known hardware.  */
 #define PTX_VECTOR_LENGTH 32
 #define PTX_WORKER_LENGTH 32
+
   /* If the worker size is not 1, the vector size must be 32.  If
-     the vector size is not 1, it must be 32.  */
+     the vector  size is not 1, it must be 32.  */
   if ((adims[GOMP_DIM_WORKER]
        && TREE_INT_CST_LOW (adims[GOMP_DIM_WORKER]) != 1)
       || (adims[GOMP_DIM_VECTOR]
@@ -3585,12 +3578,31 @@ nvptx_validate_dims (tree decl, tree dim
       adims[ix] = integer_one_node;
 
   /* Write results.  */
-  pos = dims;
-  for (ix = 0; ix != GOMP_DIM_MAX; ix++, pos = TREE_CHAIN (pos))
+  tree pos;
+  for (ix = 0, pos = dims; ix != GOMP_DIM_MAX; ix++, pos = TREE_CHAIN (pos))
     TREE_VALUE (pos) = adims[ix];
-  
+
   return dims;
 }
+
+/* Return maximum dimension size, or zero for unbounded.  */
+
+static unsigned
+nvptx_dim_limit (unsigned axis)
+{
+  switch (axis)
+    {
+    case GOMP_DIM_WORKER:
+      return PTX_WORKER_LENGTH;
+      break;
+    case GOMP_DIM_VECTOR:
+      return  PTX_VECTOR_LENGTH;
+      break;
+    default: break;
+    }
+  return 0;
+}
+
 \f
 #undef TARGET_OPTION_OVERRIDE
 #define TARGET_OPTION_OVERRIDE nvptx_option_override
@@ -3689,6 +3701,9 @@ nvptx_validate_dims (tree decl, tree dim
 #undef TARGET_GOACC_VALIDATE_DIMS
 #define TARGET_GOACC_VALIDATE_DIMS nvptx_validate_dims
 
+#undef TARGET_GOACC_DIM_LIMIT
+#define TARGET_GOACC_DIM_LIMIT nvptx_dim_limit
+
 struct gcc_target targetm = TARGET_INITIALIZER;
 
 #include "gt-nvptx.h"
Index: omp-low.c
===================================================================
--- omp-low.c	(revision 226770)
+++ omp-low.c	(working copy)
@@ -4681,13 +4681,13 @@ expand_oacc_get_num_threads (gimple_seq
   for (ix = GOMP_DIM_GANG; ix != GOMP_DIM_MAX; ix++)
     if (GOMP_DIM_MASK(ix) & gwv_bits)
       {
-	tree arg = build_int_cst (unsigned_type_node, ix);
-	tree count = create_tmp_var (unsigned_type_node);
+	tree arg = build_int_cst (integer_type_node, ix);
+	tree count = create_tmp_var (integer_type_node);
 	gimple call = gimple_build_call_internal (IFN_GOACC_DIM_SIZE, 1, arg);
 	
 	gimple_call_set_lhs (call, count);
 	gimple_seq_add_stmt (seq, call);
-	res = fold_build2 (MULT_EXPR, unsigned_type_node, res, count);
+	res = fold_build2 (MULT_EXPR, integer_type_node, res, count);
       }
   
   return res;
@@ -4713,30 +4713,30 @@ expand_oacc_get_thread_num (gimple_seq *
 	  {
 	    /* We had an outer index, so scale that by the size of
 	       this dimension.  */
-	    tree n = create_tmp_var (unsigned_type_node);
+	    tree n = create_tmp_var (integer_type_node);
 	    gimple call
 	      = gimple_build_call_internal (IFN_GOACC_DIM_SIZE, 1, arg);
 	    
 	    gimple_call_set_lhs (call, n);
 	    gimple_seq_add_stmt (seq, call);
-	    res = fold_build2 (MULT_EXPR, unsigned_type_node, res, n);
+	    res = fold_build2 (MULT_EXPR, integer_type_node, res, n);
 	  }
 
 	/* Determine index in this dimension.  */
-	tree id = create_tmp_var (unsigned_type_node);
+	tree id = create_tmp_var (integer_type_node);
 	gimple call = gimple_build_call_internal (IFN_GOACC_DIM_POS, 1, arg);
 	
 	gimple_call_set_lhs (call, id);
 	gimple_seq_add_stmt (seq, call);
 	if (res)
-	  res = fold_build2 (PLUS_EXPR, unsigned_type_node, res, id);
+	  res = fold_build2 (PLUS_EXPR, integer_type_node, res, id);
 	else
 	  res = id;
       }
 
   if (res == NULL_TREE)
-    res = build_int_cst (unsigned_type_node, 0);
-			 
+    res = build_int_cst (integer_type_node, 0);
+
   return res;
 }
 
@@ -11662,8 +11662,8 @@ oacc_init_count_vars (omp_context *ctx,
     {
       tree arg = build_int_cst (unsigned_type_node, GOMP_DIM_WORKER);
       
-      worker_var = create_tmp_var (unsigned_type_node, ".worker");
-      worker_count = create_tmp_var (unsigned_type_node, ".workercount");
+      worker_var = create_tmp_var (integer_type_node, ".worker");
+      worker_count = create_tmp_var (integer_type_node, ".workercount");
       
       gimple call1 = gimple_build_call_internal (IFN_GOACC_DIM_POS, 1, arg);
       gimple_call_set_lhs (call1, worker_var);
@@ -11675,8 +11675,8 @@ oacc_init_count_vars (omp_context *ctx,
     }
   else
     {
-      worker_var = build_int_cst (unsigned_type_node, 0);
-      worker_count = build_int_cst (unsigned_type_node, 1);
+      worker_var = build_int_cst (integer_type_node, 0);
+      worker_count = build_int_cst (integer_type_node, 1);
     }
   
   ctx->worker_var = worker_var;
@@ -14582,17 +14582,19 @@ oacc_xform_dim (gimple_stmt_iterator *gs
 {
   tree arg = gimple_call_arg (stmt, 0);
   unsigned axis = (unsigned)TREE_INT_CST_LOW (arg);
+
   while (axis--)
     dims = TREE_CHAIN (dims);
-  unsigned size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+  int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
 
-  if (size == 0)
+  if (!size)
     /* Dimension size is dynamic.  */
     return;
+  
   if (is_pos)
     {
       if (size != 1)
-	/* Size is more than 1.  */
+	/* Size is more than 1, so POS might be non-zero.  */
 	return;
       size = 0;
     }
@@ -14600,7 +14602,7 @@ oacc_xform_dim (gimple_stmt_iterator *gs
   /* Replace the internal call with a constant.  */
   tree lhs = gimple_call_lhs (stmt);
   gimple g = gimple_build_assign
-    (lhs, build_int_cst (unsigned_type_node, size));
+    (lhs, build_int_cst (integer_type_node, size));
   gsi_replace (gsi, g, false);
 }
 
@@ -14619,8 +14621,7 @@ execute_oacc_transform ()
     return 0;
 
   tree dims = TREE_VALUE (attrs);
-  if (dims)
-    dims = targetm.goacc.validate_dims (current_function_decl, dims);
+  dims = targetm.goacc.validate_dims (current_function_decl, dims);
   /* Safe to overwrite, this attribute chain is unshared.  */
   TREE_VALUE (attrs) = dims;
   
@@ -14645,13 +14646,11 @@ execute_oacc_transform ()
 		    default: break;
 
 		    case IFN_GOACC_DIM_SIZE:
-		      if (dims)
-			oacc_xform_dim (&gsi, stmt, dims, false);
+		      oacc_xform_dim (&gsi, stmt, dims, false);
 		      break;
 
 		    case IFN_GOACC_DIM_POS:
-		      if (dims)
-			oacc_xform_dim (&gsi, stmt, dims, true);
+		      oacc_xform_dim (&gsi, stmt, dims, true);
 		      break;
 
 		    case IFN_GOACC_FORK:
@@ -14681,24 +14680,42 @@ execute_oacc_transform ()
 tree
 default_goacc_validate_dims (tree ARG_UNUSED (decl), tree dims)
 {
-  tree pos = dims;
-  for (unsigned ix = GOMP_DIM_MAX; ix--; pos = TREE_CHAIN (pos))
+  tree *pos_ptr;
+  unsigned ix;
+  
+  for (ix = 0, pos_ptr = &dims;
+       ix != GOMP_DIM_MAX; ix++, pos_ptr = &TREE_CHAIN (*pos_ptr))
     {
-      tree val = TREE_VALUE (pos);
+      /* Cons up a default, if the attribue list is NULL.  This
+	 happens on 'declare' routines, as theyy do not currently set
+	 the dimensions over which the routine may be active.  */
+      if (!*pos_ptr)
+	*pos_ptr = tree_cons (NULL_TREE, NULL_TREE, NULL_TREE);
+      
+      tree val = TREE_VALUE (*pos_ptr);
       
 #ifdef ACCEL_COMPILER
       if (!val)
-	val = integer_one_node;
-#else
-      /* Set to 1 on the host. */
-      val = integer_one_node;
 #endif
-      TREE_VALUE (pos) =  val;
+	val = integer_one_node;
+      TREE_VALUE (*pos_ptr) = val;
     }
 
   return dims;
 }
 
+/* Default dimension bound is unknown on accelerator and 1 on host. */
+
+unsigned
+default_goacc_dim_limit (unsigned ARG_UNUSED (axis))
+{
+#ifdef ACCEL_COMPILER
+  return 0;
+#else
+  return 1;
+#endif
+}
+
 /* Default fork/join early expander.  Delete the function calls if
    there is no RTL expander.  */
 
Index: doc/tm.texi.in
===================================================================
--- doc/tm.texi.in	(revision 226770)
+++ doc/tm.texi.in	(working copy)
@@ -4247,6 +4247,8 @@ address;  but often a machine-dependent
 
 @hook TARGET_GOACC_VALIDATE_DIMS
 
+@hook TARGET_GOACC_DIM_LIMIT
+
 @hook TARGET_GOACC_FORK_JOIN
 
 @node Anchored Addresses
Index: doc/tm.texi
===================================================================
--- doc/tm.texi	(revision 226770)
+++ doc/tm.texi	(working copy)
@@ -5743,7 +5743,13 @@ to use it.
 @deftypefn {Target Hook} tree TARGET_GOACC_VALIDATE_DIMS (tree, @var{tree})
 This hook should check the launch dimensions provided.  It should fill
 in default values and verify non-defaults.  The TREE_LIST is unshared
-and may be overwritten.  Diagnostics should be issued as appropriate.
+and may be overwritten.  If the dimension list is NULL, one should be
+created.  Diagnostics should be issued as appropriate.
+@end deftypefn
+
+@deftypefn {Target Hook} unsigned TARGET_GOACC_DIM_LIMIT (unsigned)
+This hook should return the maximum size of a particular dimension,
+or zero if unbounded.
 @end deftypefn
 
 @deftypefn {Target Hook} bool TARGET_GOACC_FORK_JOIN (bool, gimple_stmt_iterator *@var{}, @var{gimple})

[gomp4] Work around expand_GOACC_DIM_SIZE/expand_GOACC_DIM_POS ICEs (was: signed nums are better for dimensions)

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 5723 bytes --]

Hi!

On Tue, 11 Aug 2015 13:38:34 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> 2) We really should not be getting to the expanders if there's nothing to expand 
> to.  That's simply covering up lack of earlier handling.  That earlier removal 
> gives optimizers more leeway.

> --- internal-fn.c	(revision 226770)
> +++ internal-fn.c	(working copy)

>  static void
> -expand_GOACC_FORK (gcall *stmt ATTRIBUTE_UNUSED)
> +expand_GOACC_FORK (gcall *ARG_UNUSED (stmt))
>  {
>  #ifdef HAVE_oacc_fork
> [...]
> +#else
> +  gcc_unreachable ();
>  #endif
>  }
>  
>  static void
> -expand_GOACC_JOIN (gcall *stmt ATTRIBUTE_UNUSED)
> +expand_GOACC_JOIN (gcall *ARG_UNUSED (stmt))
>  {
>  #ifdef HAVE_oacc_join
> [...]
> +#else
> +  gcc_unreachable ();
>  #endif
>  }

>  static void
> -expand_GOACC_DIM_SIZE (gcall *stmt)
> +expand_GOACC_DIM_SIZE (gcall *ARG_UNUSED (stmt))
>  {
> +#ifdef HAVE_oacc_dim_size
> [...]
>  #else
> -  emit_move_insn (target, const1_rtx);
> +  gcc_unreachable ();
>  #endif
>  }
>  
>  static void
> -expand_GOACC_DIM_POS (gcall *stmt)
> +expand_GOACC_DIM_POS (gcall *ARG_UNUSED (stmt))
>  {
> +#ifdef HAVE_oacc_dim_pos
> [...]
>  #else
> -  emit_move_insn (target, const0_rtx);
> +  gcc_unreachable ();
>  #endif
>  }

It's not an issue for expand_GOACC_FORK and expand_GOACC_JOIN, but
expand_GOACC_DIM_SIZE and expand_GOACC_DIM_POS then blow up when the
intelmic offloading compiler works through the OpenACC offloading code,
for example:

    [...]/source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/firstprivate-1.c: In function 'main._omp_fn.0':
    [...]/source-gcc/libgomp/testsuite/libgomp.oacc-c/../libgomp.oacc-c-c++-common/firstprivate-1.c:18:11: internal compiler error: in expand_GOACC_DIM_POS, at internal-fn.c:2023
       #pragma acc parallel num_gangs (n) firstprivate (a)
               ^
    0x90bce7 expand_GOACC_DIM_POS
            [...]/source-gcc/gcc/internal-fn.c:2023
    0x62945a expand_call_stmt
            [...]/source-gcc/gcc/cfgexpand.c:2279
    0x62945a expand_gimple_stmt_1
            [...]/source-gcc/gcc/cfgexpand.c:3238
    0x62945a expand_gimple_stmt
            [...]/source-gcc/gcc/cfgexpand.c:3399
    0x62a82d expand_gimple_basic_block
            [...]/source-gcc/gcc/cfgexpand.c:5411
    0x62fc86 execute
            [...]/source-gcc/gcc/cfgexpand.c:6023
    Please submit a full bug report,
    with preprocessed source if appropriate.
    Please include the complete backtrace with any bug report.
    See <http://gcc.gnu.org/bugs.html> for instructions.
    mkoffload-intelmic: fatal error: [...]/install/offload-x86_64-intelmicemul-linux-gnu/bin//x86_64-unknown-linux-gnu-accel-x86_64-intelmicemul-linux-gnu-gcc returned 1 exit status

Admittedly, compiling OpenACC offloading code for Intel MIC doesn't make
a lot of sense (currently), but it is what's being done, and has caused a
lot of regressions in my testing, so I committed the following workaround
to gomp-4_0-branch in r226804.  A different approach would have been to
invent some machinery to not compile OpenACC offloading code on the Intel
MIC offloading path, but as the latter eventually is to support OpenACC
offloading, too (see
<http://news.gmane.org/find-root.php?message_id=%3C20141112111207.GW5026%40tucnak.redhat.com%3E>,
for example), this seemed like a step into the wrong direction.
Eventually, this will need to be fixed differently/properly -- have to
implement oacc_dim_size and oacc_dim_pos for Intel MIC offloading, or can
we have a generic fallback solution not specific to the offloading
target?

commit 8a3187f17e13bd45e630ff8a587c1fc7086abece
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed Aug 12 09:56:59 2015 +0000

    Work around expand_GOACC_DIM_SIZE/expand_GOACC_DIM_POS ICEs
    
    Followup to r226783.
    
    	gcc/
    	* internal-fn.c: Include "builtins.h".
    	(expand_GOACC_DIM_SIZE, expand_GOACC_DIM_POS): Instead of
    	gcc_unreachable, expand_builtin_trap.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@226804 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp |    6 ++++++
 gcc/internal-fn.c  |    5 +++--
 2 files changed, 9 insertions(+), 2 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index ea254c5..044fdf7 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2015-08-12  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* internal-fn.c: Include "builtins.h".
+	(expand_GOACC_DIM_SIZE, expand_GOACC_DIM_POS): Instead of
+	gcc_unreachable, expand_builtin_trap.
+
 2015-08-10  Nathan Sidwell  <nathan@acm.org>
 
 	* tree-ssa-phiopt.c (minmax_replacement): Create new ssa name if
diff --git gcc/internal-fn.c gcc/internal-fn.c
index 8b8c6e1..70bffd4 100644
--- gcc/internal-fn.c
+++ gcc/internal-fn.c
@@ -47,6 +47,7 @@ along with GCC; see the file COPYING3.  If not see
 #include "stringpool.h"
 #include "tree-ssanames.h"
 #include "diagnostic-core.h"
+#include "builtins.h"
 
 /* The names of each internal function, indexed by function number.  */
 const char *const internal_fn_name_array[] = {
@@ -2003,7 +2004,7 @@ expand_GOACC_DIM_SIZE (gcall *ARG_UNUSED (stmt))
 			 VOIDmode, EXPAND_NORMAL);
   emit_insn (gen_oacc_dim_size (target, dim));
 #else
-  gcc_unreachable ();
+  expand_builtin_trap ();
 #endif
 }
 
@@ -2021,7 +2022,7 @@ expand_GOACC_DIM_POS (gcall *ARG_UNUSED (stmt))
 			 VOIDmode, EXPAND_NORMAL);
   emit_insn (gen_oacc_dim_pos (target, dim));
 #else
-  gcc_unreachable ();
+  expand_builtin_trap ();
 #endif
 }
 


Grüße,
 Thomas

[-- Attachment #2: signature.asc --]
[-- Type: application/pgp-signature, Size: 472 bytes --]

Re: [gomp4] Work around expand_GOACC_DIM_SIZE/expand_GOACC_DIM_POS ICEs

[Nathan Sidwell]

On 08/12/15 06:21, Thomas Schwinge wrote:
> Hi!
  OpenACC offloading code for Intel MIC doesn't make
> a lot of sense (currently), but it is what's being done, and has caused a
> lot of regressions in my testing, so I committed the following workaround
> to gomp-4_0-branch in r226804.  A different approach would have been to
> invent some machinery to not compile OpenACC offloading code on the Intel
> MIC offloading path, but as the latter eventually is to support OpenACC
> offloading, too (see

The right solution is probably for the default validated_dims to set the size to 
unity, just like the host.  Thereby forcing any backend that wants to provide a 
larger size to override that hook.

I.e. remove the #ifndef ACCEL_COMPILER from default_validate_dims in omp-low

nathan
-- 
Nathan Sidwell

Re: [gomp4] Work around expand_GOACC_DIM_SIZE/expand_GOACC_DIM_POS ICEs

[Nathan Sidwell]

On 08/12/15 08:46, Nathan Sidwell wrote:
> On 08/12/15 06:21, Thomas Schwinge wrote:
>> Hi!
>   OpenACC offloading code for Intel MIC doesn't make
>> a lot of sense (currently), but it is what's being done, and has caused a
>> lot of regressions in my testing, so I committed the following workaround
>> to gomp-4_0-branch in r226804.  A different approach would have been to
>> invent some machinery to not compile OpenACC offloading code on the Intel
>> MIC offloading path, but as the latter eventually is to support OpenACC
>> offloading, too (see
>
> The right solution is probably for the default validated_dims to set the size to
> unity, just like the host.  Thereby forcing any backend that wants to provide a
> larger size to override that hook.
>
> I.e. remove the #ifndef ACCEL_COMPILER from default_validate_dims in omp-low

for avoidance of doubt, I'll add that to the patch I have in progress.

nathan

-- 
Nathan Sidwell - Director, Sourcery Services - Mentor Embedded

[gomp4] dimension API

[Nathan Sidwell]

[-- Attachment #1: Type: text/plain, Size: 656 bytes --]

I've committed this patch to gomp4.  It reworks the oacc functiuon attribute 
dimension handling.  Rather than pass the TREE_LIST to the various hooks, I 
convert it to a regular C array of ints.  That makes life simpler for the 
consumers.  They return a 'changed' flag to indicate whether the attrribute 
should be rewritten.  That rewriting is done in a way that doesn;t presume the 
attribute is unshared (Cesar, your workaround should no longer be necessary).

Also  put in the change I mentioned earlier this morning about the default 
validate dims hook setting the dimensions to 1 on accelerators to.  I'll  revert 
Thomas's patch shortly.

nathan

[-- Attachment #2: gomp4-dim-api.patch --]
[-- Type: text/x-patch, Size: 10422 bytes --]

2015-08-12  Nathan Sidwell  <nathan@codesourcery.com>

	* target.def (validate_dims): Adjust API.
	* targhooks.h (default_goacc_validate_dims): Adjust.
	* omp-low.c (replace_oacc_fn_attrib): New function.
	(set_oacc_fn_attrib): Use it.
	(oacc_xform_dim): Dims is array of ints.
	(execute_oacc_transform): Create int array of dims, adjust uses.
	(default_goacc_validate_dims): Adjust API.  Force to w everywhere.
	* doc/tm.texi: Rebuild.
	* config/nvptx/nvptx.c (nvptx_validate_dims): Adjust API.

Index: gcc/config/nvptx/nvptx.c
===================================================================
--- gcc/config/nvptx/nvptx.c	(revision 226808)
+++ gcc/config/nvptx/nvptx.c	(working copy)
@@ -3538,71 +3538,46 @@ nvptx_expand_builtin (tree exp, rtx targ
   return d->expander (exp, target, mode, ignore);
 }
 \f
-/* Validate compute dimensions, fill in defaults.  */
-
-static tree
-nvptx_validate_dims (tree decl, tree dims)
-{
-  tree adims[GOMP_DIM_MAX];
-  unsigned ix;
-  tree *pos_ptr;
-
-  for (ix = 0, pos_ptr = &dims; ix != GOMP_DIM_MAX;
-       ix++, pos_ptr = &TREE_CHAIN (*pos_ptr))
-    {
-      if (!*pos_ptr)
-	*pos_ptr = tree_cons (NULL_TREE, NULL_TREE, NULL_TREE);
-      
-      adims[ix] = TREE_VALUE (*pos_ptr);
-    }
-
-  /* Define vector size for known hardware.  */
+/* Define vector size for known hardware.  */
 #define PTX_VECTOR_LENGTH 32
 #define PTX_WORKER_LENGTH 32
 
+/* Validate compute dimensions, fill in non-unity defaults.  */
+
+static bool
+nvptx_validate_dims (tree decl, int dims[])
+{
+  bool changed = false;
+
   /* If the worker size is not 1, the vector size must be 32.  If
      the vector  size is not 1, it must be 32.  */
-  if ((adims[GOMP_DIM_WORKER]
-       && TREE_INT_CST_LOW (adims[GOMP_DIM_WORKER]) != 1)
-      || (adims[GOMP_DIM_VECTOR]
-	  && TREE_INT_CST_LOW (adims[GOMP_DIM_VECTOR]) != 1))
+  if ((dims[GOMP_DIM_WORKER] > 1 || dims[GOMP_DIM_WORKER] == 0)
+      || (dims[GOMP_DIM_VECTOR] > 1 || dims[GOMP_DIM_VECTOR] == 0))
     {
-      if (!adims[GOMP_DIM_VECTOR]
-	  || TREE_INT_CST_LOW (adims[GOMP_DIM_VECTOR]) != PTX_VECTOR_LENGTH)
+      if (dims[GOMP_DIM_VECTOR] != PTX_VECTOR_LENGTH)
 	{
-	  tree use = build_int_cst (integer_type_node, PTX_VECTOR_LENGTH);
-	  if (adims[GOMP_DIM_VECTOR])
+	  if (dims[GOMP_DIM_VECTOR] >= 0)
 	    warning_at (DECL_SOURCE_LOCATION (decl), 0,
-			TREE_INT_CST_LOW (adims[GOMP_DIM_VECTOR])
-			? "using vector_length (%E), ignoring %E"
-			: "using vector_length (%E), ignoring runtime setting",
-			use, adims[GOMP_DIM_VECTOR]);
-	  adims[GOMP_DIM_VECTOR] = use;
+			dims[GOMP_DIM_VECTOR]
+			? "using vector_length (%d), ignoring %d"
+			: "using vector_length (%d), ignoring runtime setting",
+			PTX_VECTOR_LENGTH, dims[GOMP_DIM_VECTOR]);
+	  dims[GOMP_DIM_VECTOR] = PTX_VECTOR_LENGTH;
+	  changed = true;
 	}
     }
 
   /* Check the num workers is not too large.  */
-  if (adims[GOMP_DIM_WORKER]
-      && TREE_INT_CST_LOW (adims[GOMP_DIM_WORKER]) > PTX_WORKER_LENGTH)
+  if (dims[GOMP_DIM_WORKER] > PTX_WORKER_LENGTH)
     {
-      tree use = build_int_cst (integer_type_node, PTX_WORKER_LENGTH);
       warning_at (DECL_SOURCE_LOCATION (decl), 0,
-		  "using num_workers (%E), ignoring %E",
-		  use, adims[GOMP_DIM_WORKER]);
-      adims[GOMP_DIM_WORKER] = use;
+		  "using num_workers (%d), ignoring %d",
+		  PTX_WORKER_LENGTH, dims[GOMP_DIM_WORKER]);
+      dims[GOMP_DIM_WORKER] = PTX_WORKER_LENGTH;
+      changed = true;
     }
 
-  /* Set defaults.  */
-  for (ix = 0; ix != GOMP_DIM_MAX; ix++)
-    if (!adims[ix])
-      adims[ix] = integer_one_node;
-
-  /* Write results.  */
-  tree pos;
-  for (ix = 0, pos = dims; ix != GOMP_DIM_MAX; ix++, pos = TREE_CHAIN (pos))
-    TREE_VALUE (pos) = adims[ix];
-
-  return dims;
+  return changed;
 }
 
 /* Return maximum dimension size, or zero for unbounded.  */
Index: gcc/doc/tm.texi
===================================================================
--- gcc/doc/tm.texi	(revision 226808)
+++ gcc/doc/tm.texi	(working copy)
@@ -5740,11 +5740,12 @@ usable.  In that case, the smaller the n
 to use it.
 @end deftypefn
 
-@deftypefn {Target Hook} tree TARGET_GOACC_VALIDATE_DIMS (tree, @var{tree})
+@deftypefn {Target Hook} bool TARGET_GOACC_VALIDATE_DIMS (tree, int @var{[]})
 This hook should check the launch dimensions provided.  It should fill
-in default values and verify non-defaults.  The TREE_LIST is unshared
-and may be overwritten.  If the dimension list is NULL, one should be
-created.  Diagnostics should be issued as appropriate.
+in anything that needs default to non-unity and verify non-defaults.
+Defaults are represented as -1.  Diagnostics should be issuedas 
+ppropriate.  Return true if changes have been made.  You must override
+this hook to provide dimensions larger than 1.
 @end deftypefn
 
 @deftypefn {Target Hook} unsigned TARGET_GOACC_DIM_LIMIT (unsigned)
Index: gcc/omp-low.c
===================================================================
--- gcc/omp-low.c	(revision 226808)
+++ gcc/omp-low.c	(working copy)
@@ -9314,6 +9314,16 @@ oacc_launch_pack (unsigned code, tree de
 
 #define OACC_FN_ATTRIB "oacc function"
 
+/* Replace any existing oacc fn attribute with updated dimensions.  */
+
+void
+replace_oacc_fn_attrib (tree fn, tree dims)
+{
+  /* Simply cons onto the beginning of the list.  */
+  DECL_ATTRIBUTES (fn) =
+    tree_cons (get_identifier (OACC_FN_ATTRIB), dims, DECL_ATTRIBUTES (fn));
+}
+
 static void
 set_oacc_fn_attrib (tree clauses, tree fn, vec<tree> *args)
 {
@@ -9341,9 +9351,7 @@ set_oacc_fn_attrib (tree clauses, tree f
       attr = tree_cons (NULL_TREE, dim, attr);
     }
 
-  /* Add the attributes.  */
-  DECL_ATTRIBUTES (fn) =
-    tree_cons (get_identifier (OACC_FN_ATTRIB), attr, DECL_ATTRIBUTES (fn));
+  replace_oacc_fn_attrib (fn, attr);
 
   if (non_const)
     {
@@ -14581,14 +14589,11 @@ oacc_xform_on_device (gimple_stmt_iterat
 
 static void
 oacc_xform_dim (gimple_stmt_iterator *gsi, gimple stmt,
-		tree dims, bool is_pos)
+		int dims[], bool is_pos)
 {
   tree arg = gimple_call_arg (stmt, 0);
   unsigned axis = (unsigned)TREE_INT_CST_LOW (arg);
-
-  while (axis--)
-    dims = TREE_CHAIN (dims);
-  int size = TREE_INT_CST_LOW (TREE_VALUE (dims));
+  int size = dims[axis];
 
   if (!size)
     /* Dimension size is dynamic.  */
@@ -14618,15 +14623,50 @@ execute_oacc_transform ()
 {
   basic_block bb;
   tree attrs = get_oacc_fn_attrib (current_function_decl);
+  int dims[GOMP_DIM_MAX];
   
   if (!attrs)
     /* Not an offloaded function.  */
     return 0;
 
-  tree dims = TREE_VALUE (attrs);
-  dims = targetm.goacc.validate_dims (current_function_decl, dims);
-  /* Safe to overwrite, this attribute chain is unshared.  */
-  TREE_VALUE (attrs) = dims;
+  {
+    unsigned ix;
+    tree pos = TREE_VALUE (attrs);
+
+    for (ix = 0; ix != GOMP_DIM_MAX; ix++)
+      {
+	if (!pos)
+	  dims[ix] = 0;
+	else
+	  {
+	    tree val = TREE_VALUE (pos);
+	    
+	    dims[ix] = val ? TREE_INT_CST_LOW (val) : -1;
+	    pos = TREE_CHAIN (pos);
+	  }
+      }
+
+    bool changed = targetm.goacc.validate_dims (current_function_decl, dims);
+
+    /* Default anything left undefaulted to 1.  */
+    for (ix = 0; ix != GOMP_DIM_MAX; ix++)
+      if (dims[ix] < 0)
+	{
+	  dims[ix] = 1;
+	  changed = true;
+	}
+  
+    if (changed)
+      {
+	/* Replace the attribute with new values.  */
+	pos = NULL_TREE;
+	for (ix = GOMP_DIM_MAX; ix--;)
+	  pos = tree_cons (NULL_TREE,
+			   build_int_cst (integer_type_node, dims[ix]),
+			   pos);
+	replace_oacc_fn_attrib (current_function_decl, pos);
+      }
+  }
   
   FOR_ALL_BB_FN (bb, cfun)
     {
@@ -14677,34 +14717,25 @@ execute_oacc_transform ()
   return 0;
 }
 
-/* Default launch dimension validator.  Force everything to 1 on the
-   host and default to 1 otherwise.  */
+/* Default launch dimension validator.  Force everything to 1.  A
+   backend that wants to provide larger dimensions must override this
+   hook.  */
 
-tree
-default_goacc_validate_dims (tree ARG_UNUSED (decl), tree dims)
+bool
+default_goacc_validate_dims (tree ARG_UNUSED (decl), int *ARG_UNUSED (dims))
 {
-  tree *pos_ptr;
-  unsigned ix;
-  
-  for (ix = 0, pos_ptr = &dims;
-       ix != GOMP_DIM_MAX; ix++, pos_ptr = &TREE_CHAIN (*pos_ptr))
+  bool changed = false;
+
+  for (unsigned ix = 0; ix != GOMP_DIM_MAX; ix++)
     {
-      /* Cons up a default, if the attribue list is NULL.  This
-	 happens on 'declare' routines, as theyy do not currently set
-	 the dimensions over which the routine may be active.  */
-      if (!*pos_ptr)
-	*pos_ptr = tree_cons (NULL_TREE, NULL_TREE, NULL_TREE);
-      
-      tree val = TREE_VALUE (*pos_ptr);
-      
-#ifdef ACCEL_COMPILER
-      if (!val)
-#endif
-	val = integer_one_node;
-      TREE_VALUE (*pos_ptr) = val;
+      if (dims[ix] != 1)
+	{
+	  dims[ix] = 1;
+	  changed = true;
+	}
     }
 
-  return dims;
+  return changed;
 }
 
 /* Default dimension bound is unknown on accelerator and 1 on host. */
Index: gcc/target.def
===================================================================
--- gcc/target.def	(revision 226808)
+++ gcc/target.def	(working copy)
@@ -1647,10 +1647,11 @@ HOOK_VECTOR (TARGET_GOACC, goacc)
 DEFHOOK
 (validate_dims,
 "This hook should check the launch dimensions provided.  It should fill\n\
-in default values and verify non-defaults.  The TREE_LIST is unshared\n\
-and may be overwritten.  If the dimension list is NULL, one should be\n\
-created.  Diagnostics should be issued as appropriate.",
-tree, (tree, tree),
+in anything that needs default to non-unity and verify non-defaults.\n\
+Defaults are represented as -1.  Diagnostics should be issuedas \n\
+ppropriate.  Return true if changes have been made.  You must override\n\
+this hook to provide dimensions larger than 1.",
+bool, (tree, int []),
 default_goacc_validate_dims)
 
 DEFHOOK
Index: gcc/targhooks.h
===================================================================
--- gcc/targhooks.h	(revision 226808)
+++ gcc/targhooks.h	(working copy)
@@ -107,7 +107,7 @@ extern unsigned default_add_stmt_cost (v
 extern void default_finish_cost (void *, unsigned *, unsigned *, unsigned *);
 extern void default_destroy_cost_data (void *);
 
-extern tree default_goacc_validate_dims (tree, tree);
+extern bool default_goacc_validate_dims (tree, int []);
 extern unsigned default_goacc_dim_limit (unsigned);
 extern bool default_goacc_fork_join (bool, gimple_stmt_iterator *, gimple);
 

Re: [gomp4] dimension API

[Nathan Sidwell]

On 08/12/15 09:30, Nathan Sidwell wrote:
> I've committed this patch to gomp4.  It reworks the oacc functiuon attribute
> dimension handling.  Rather than pass the TREE_LIST to the various hooks, I
> convert it to a regular C array of ints.  That makes life simpler for the
> consumers.  They return a 'changed' flag to indicate whether the attrribute
> should be rewritten.  That rewriting is done in a way that doesn;t presume the
> attribute is unshared (Cesar, your workaround should no longer be necessary).

I've discovered I'd committed a slightly earlier version of what I'd tested. 
Consequently there is some breakage.  I'll be fixing it later today, after a retest.

nathan

Re: [gomp4] dimension API

[Thomas Schwinge]

[-- Attachment #1: Type: text/plain, Size: 388 bytes --]

Hi Nathan!

On Wed, 12 Aug 2015 09:30:10 -0400, Nathan Sidwell <nathan@acm.org> wrote:
> I've committed this patch to gomp4.  [...]
> 
> Also  put in the change I mentioned earlier this morning about the default 
> validate dims hook setting the dimensions to 1 on accelerators to.  I'll  revert 
> Thomas's patch shortly.

Confirmed, that works, thanks!


Grüße,
 Thomas

[-- Attachment #2: Type: application/pgp-signature, Size: 472 bytes --]