[gomp4, wip] remove references to ganglocal shared memory inside gcc

[gomp4, wip] remove references to ganglocal shared memory inside gcc

[Cesar Philippidis]

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

This patch strips out all of the references to ganglocal memory in gcc.
Unfortunately, the runtime api still takes a shared memory parameter, so
I haven't made any changes there yet. Perhaps we could still keep the
shared memory argument to GOACC_parallel, but remove all of the support
for ganglocal mappings. Then again, maybe we still need support
ganglocal mappings for legacy purposes.

With the ganglocal mapping aside, I'm in favor of leaving the shared
memory argument to GOACC_parallel, just in case we find another use for
shared memory in the future.

Nathan, what do you want to do here?

Cesar

[-- Attachment #2: ganglocal-removal.diff --]
[-- Type: text/x-patch, Size: 15141 bytes --]

2015-08-26  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* builtins.c (expand_oacc_ganglocal_ptr): Delete.
	(expand_builtin): Remove stale GOACC_GET_GANGLOCAL_PTR builtin.
	* config/nvptx/nvptx.md (ganglocal_ptr): Delete.
	* gimple.h (struct gimple_statement_omp_parallel_layout): Remove
	ganglocal_size member.
	(gimple_omp_target_ganglocal_size): Delete.
	(gimple_omp_target_set_ganglocal_size): Delete.
	* omp-builtins.def (BUILT_IN_GOACC_GET_GANGLOCAL_PTR): Delete.
	* omp-low.c (struct omp_context): Remove ganglocal_init, ganglocal_ptr,
	ganglocal_size, ganglocal_size_host, worker_var, worker_count and
	worker_sync_elt.
	(alloc_var_ganglocal): Delete.
	(install_var_ganglocal): Delete.
	(new_omp_context): Don't use ganglocal memory.
	(expand_omp_target): Likewise.
	(lower_omp_taskreg): Likewise.
	(lower_omp_target): Likewise.
	* tree-parloops.c (create_parallel_loop): Likewise.
	* tree-pretty-print.c (dump_omp_clause): Remove support for
	GOMP_MAP_FORCE_TO_GANGLOCAL

diff --git a/gcc/builtins.c b/gcc/builtins.c
index 7c3ead1..f465716 100644
--- a/gcc/builtins.c
+++ b/gcc/builtins.c
@@ -5913,25 +5913,6 @@ expand_builtin_acc_on_device (tree exp, rtx target)
   return target;
 }
 
-static rtx
-expand_oacc_ganglocal_ptr (rtx target ATTRIBUTE_UNUSED)
-{
-#ifdef HAVE_ganglocal_ptr
-  enum insn_code icode;
-  icode = CODE_FOR_ganglocal_ptr;
-  rtx tmp = target;
-  if (!REG_P (tmp) || GET_MODE (tmp) != Pmode)
-    tmp = gen_reg_rtx (Pmode);
-  rtx insn = GEN_FCN (icode) (tmp);
-  if (insn != NULL_RTX)
-    {
-      emit_insn (insn);
-      return tmp;
-    }
-#endif
-  return NULL_RTX;
-}
-
 /* Expand an expression EXP that calls a built-in function,
    with result going to TARGET if that's convenient
    (and in mode MODE if that's convenient).
@@ -7074,12 +7055,6 @@ expand_builtin (tree exp, rtx target, rtx subtarget, machine_mode mode,
 	return target;
       break;
 
-    case BUILT_IN_GOACC_GET_GANGLOCAL_PTR:
-      target = expand_oacc_ganglocal_ptr (target);
-      if (target)
-	return target;
-      break;
-
     default:	/* just do library call, if unknown builtin */
       break;
     }
diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index 3d734a8..d0d6564 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -1485,23 +1485,6 @@
   ""
   "%.\\tst.shared%u1\\t%1,%0;")
 
-(define_insn "ganglocal_ptr<mode>"
-  [(set (match_operand:P 0 "nvptx_register_operand" "")
-	(unspec:P [(const_int 0)] UNSPEC_SHARED_DATA))]
-  ""
-  "%.\\tcvta.shared%t0\\t%0, sdata;")
-
-(define_expand "ganglocal_ptr"
-  [(match_operand 0 "nvptx_register_operand" "")]
-  ""
-{
-  if (Pmode == DImode)
-    emit_insn (gen_ganglocal_ptrdi (operands[0]));
-  else
-    emit_insn (gen_ganglocal_ptrsi (operands[0]));
-  DONE;
-})
-
 ;; Atomic insns.
 
 (define_expand "atomic_compare_and_swap<mode>"
diff --git a/gcc/gimple.h b/gcc/gimple.h
index d8d8742..278b49f 100644
--- a/gcc/gimple.h
+++ b/gcc/gimple.h
@@ -580,10 +580,6 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT")))
   /* [ WORD 10 ]
      Shared data argument.  */
   tree data_arg;
-
-  /* [ WORD 11 ]
-     Size of the gang-local memory to allocate.  */
-  tree ganglocal_size;
 };
 
 /* GIMPLE_OMP_PARALLEL or GIMPLE_TASK */
@@ -5232,25 +5228,6 @@ gimple_omp_target_set_data_arg (gomp_target *omp_target_stmt,
 }
 
 
-/* Return the size of gang-local data associated with OMP_TARGET GS.  */
-
-static inline tree
-gimple_omp_target_ganglocal_size (const gomp_target *omp_target_stmt)
-{
-  return omp_target_stmt->ganglocal_size;
-}
-
-
-/* Set SIZE to be the size of gang-local memory associated with OMP_TARGET
-   GS.  */
-
-static inline void
-gimple_omp_target_set_ganglocal_size (gomp_target *omp_target_stmt, tree size)
-{
-  omp_target_stmt->ganglocal_size = size;
-}
-
-
 /* Return the clauses associated with OMP_TEAMS GS.  */
 
 static inline tree
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 0d9f386..615c4e0 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -58,8 +58,6 @@ DEF_GOACC_BUILTIN_FNSPEC (BUILT_IN_GOACC_UPDATE, "GOACC_update",
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_WAIT, "GOACC_wait",
 		   BT_FN_VOID_INT_INT_VAR,
 		   ATTR_NOTHROW_LIST)
-DEF_GOACC_BUILTIN (BUILT_IN_GOACC_GET_GANGLOCAL_PTR, "GOACC_get_ganglocal_ptr",
-		   BT_FN_PTR, ATTR_NOTHROW_LEAF_LIST)
 DEF_GOACC_BUILTIN (BUILT_IN_GOACC_DEVICEPTR, "GOACC_deviceptr",
 		   BT_FN_PTR_PTR, ATTR_CONST_NOTHROW_LEAF_LIST)
 DEF_GOACC_BUILTIN_COMPILER (BUILT_IN_ACC_ON_DEVICE, "acc_on_device",
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ee92141..19d66fa 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -212,21 +212,6 @@ typedef struct omp_context
      this level and above.  For parallel and kernels clauses, a mask
      indicating which of num_gangs/num_workers/num_vectors was used.  */
   int gwv_this;
-
-  gimple_seq ganglocal_init;
-  tree ganglocal_ptr;
-  tree ganglocal_size;
-  tree ganglocal_size_host;
-
-  /* For OpenACC offloaded regions, variables holding the worker id and count
-     of workers.  */
-  tree worker_var;
-  tree worker_count;
-
-  /* For offloaded regions, at runtime this variable holds a pointer
-     to the location that should be used for thread
-     synchronization.  */
-  tree worker_sync_elt;
 } omp_context;
 
 /* A structure holding the elements of:
@@ -1448,96 +1433,6 @@ align_and_expand (tree *poldsz, tree size, unsigned int align)
   return fold_build2 (PLUS_EXPR, size_type_node, oldsz, size);
 }
 
-/* Generate code at the entry to CTX to allocate SIZE bytes of gang-local
-   memory, cast this to be a pointer to VARTYPE and return the result.
-   UNDERLYING_VAR is used to choose a name, it can be NULL.
-
-   FIXME: Consider pointer to arrays.  */
-
-static tree
-alloc_var_ganglocal (tree underlying_var, tree vartype, omp_context *ctx,
-		     tree size, tree host_size = NULL)
-{
-  /* If this is a pointer mapping, then we need to create too mappings, one
-     for the pointer, and another to the data.  Add the offset for the pointer
-     to size.  */
-  bool pointer = underlying_var ? is_reference (underlying_var) : false;
-
-  if (pointer)
-    size = fold_build2 (PLUS_EXPR, TREE_TYPE (size), size,
-			TYPE_SIZE_UNIT (ptr_type_node));
-
-  if (host_size == NULL)
-    host_size = size;
-  else if (pointer)
-    {
-      host_size = fold_build2 (PLUS_EXPR, TREE_TYPE (size), host_size,
-			       TYPE_SIZE_UNIT (ptr_type_node));
-    }
-
-  if (ctx->ganglocal_ptr == NULL_TREE)
-    {
-      tree ptr = create_tmp_var (ptr_type_node, "__ganglocal_ptr");
-      tree fndecl = builtin_decl_explicit (BUILT_IN_GOACC_GET_GANGLOCAL_PTR);
-      gimple g = gimple_build_call (fndecl, 0);
-      gimple_call_set_lhs (g, ptr);
-      gimple_seq_add_stmt (&ctx->ganglocal_init, g);
-
-      ctx->ganglocal_ptr = ptr;
-    }
-  const char *name = NULL;
-  if (underlying_var && DECL_NAME (underlying_var))
-    name = IDENTIFIER_POINTER (DECL_NAME (underlying_var));
-  tree oldsz = ctx->ganglocal_size;
-  tree varptrtype = build_pointer_type (vartype);
-
-  unsigned int align = TYPE_ALIGN_UNIT (vartype);
-  tree newsz = align_and_expand (&oldsz, size, align);
-  tree gl_host = ctx->ganglocal_size_host;
-  ctx->ganglocal_size_host = align_and_expand (&gl_host, host_size, align);
-  ctx->ganglocal_size = newsz;
-
-  tree newvarptr = create_tmp_var (varptrtype, name);
-#if 0
-  DECL_CHAIN (newvarptr) = ctx->block_vars;
-  ctx->block_vars = newvarptr;
-#endif
-  tree device_off = oldsz;
-  walk_tree (&device_off, copy_tree_body_r, &ctx->cb, NULL);
-
-  tree x = fold_convert (varptrtype,
-			 fold_build_pointer_plus (ctx->ganglocal_ptr,
-						  device_off));
-  gimple g = gimple_build_assign (newvarptr, x);
-  gimple_seq_add_stmt (&ctx->ganglocal_init, g);
-
-  /* Add a target for the pointer.  */
-  if (pointer)
-    {
-      tree vtype = build_pointer_type (size_type_node);
-      tree t = create_tmp_var (vtype);
-      x = fold_convert (vtype, x);
-      g = gimple_build_assign (t, x);
-      gimple_seq_add_stmt (&ctx->ganglocal_init, g);
-
-      x = fold_build_pointer_plus (x, build_int_cst (sizetype, 8));
-      g = gimple_build_assign (build_simple_mem_ref (t),
-			       fold_convert (sizetype, x));
-      gimple_seq_add_stmt (&ctx->ganglocal_init, g);
-    }
-  return newvarptr;
-}
-
-static tree
-install_var_ganglocal (tree decl, omp_context *ctx)
-{
-  tree type = TREE_TYPE (decl);
-  tree ptr = alloc_var_ganglocal (decl, type, ctx, TYPE_SIZE_UNIT (type));
-  tree new_var = build1 (INDIRECT_REF, type, ptr);
-  insert_decl_map (&ctx->cb, decl, new_var);
-  return new_var;
-}
-
 /* Debugging dumps for parallel regions.  */
 void dump_omp_region (FILE *, struct omp_region *, int);
 void debug_omp_region (struct omp_region *);
@@ -1675,8 +1570,6 @@ new_omp_context (gimple stmt, omp_context *outer_ctx)
       ctx->cb.transform_call_graph_edges = CB_CGE_MOVE;
       ctx->depth = 1;
     }
-  ctx->ganglocal_size = size_zero_node;
-  ctx->ganglocal_size_host = size_zero_node;
   ctx->cb.decl_map = new hash_map<tree, tree>;
 
   return ctx;
@@ -10056,8 +9949,9 @@ expand_omp_target (struct omp_region *region)
     case BUILT_IN_GOACC_KERNELS_INTERNAL:
     case BUILT_IN_GOACC_PARALLEL:
       {
-	args.quick_push (gimple_omp_target_ganglocal_size (entry_stmt));
-	set_oacc_fn_attrib (child_fn, clauses, &args);
+	tree shared_memory = build_int_cst (integer_type_node, 0);
+	args.quick_push (shared_memory);
+	set_oacc_fn_attrib (clauses, child_fn, &args);
 	tagging = true;
       }
       /* FALLTHRU */
@@ -11935,39 +11829,6 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
     }
 }
 
-/* A subroutine of lower_omp_target.  Build variables holding the
-   worker count and index for use inside in ganglocal memory allocations.  */
-
-static void
-oacc_init_count_vars (omp_context *ctx, tree clauses ATTRIBUTE_UNUSED)
-{
-  tree worker_var, worker_count;
-  
-  if (ctx->gwv_this & GOMP_DIM_MASK (GOMP_DIM_WORKER))
-    {
-      tree arg = build_int_cst (unsigned_type_node, GOMP_DIM_WORKER);
-      
-      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);
-      gimple_seq_add_stmt (&ctx->ganglocal_init, call1);
-
-      gimple call2 = gimple_build_call_internal (IFN_GOACC_DIM_SIZE, 1, arg);
-      gimple_call_set_lhs (call2, worker_count);
-      gimple_seq_add_stmt (&ctx->ganglocal_init, call2);
-    }
-  else
-    {
-      worker_var = build_int_cst (integer_type_node, 0);
-      worker_count = build_int_cst (integer_type_node, 1);
-    }
-  
-  ctx->worker_var = worker_var;
-  ctx->worker_count = worker_count;
-}
-
 /* Lower the GIMPLE_OMP_TARGET in the current statement
    in GSI_P.  CTX holds context information for the directive.  */
 
@@ -12057,7 +11918,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    break;
 	  case GOMP_MAP_FORCE_ALLOC:
 	  case GOMP_MAP_FORCE_TO:
-	  case GOMP_MAP_FORCE_TO_GANGLOCAL:
 	  case GOMP_MAP_FORCE_FROM:
 	  case GOMP_MAP_FORCE_TOFROM:
 	  case GOMP_MAP_FORCE_PRESENT:
@@ -12127,15 +11987,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		DECL_HAS_VALUE_EXPR_P (new_var) = 1;
 	      }
 	    else
-	      {
-		/* Copy from the receiver field to gang-local memory.  */
-		gcc_assert (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
-			    && (OMP_CLAUSE_MAP_KIND (c)
-				== GOMP_MAP_FORCE_TO_GANGLOCAL)
-			    && TREE_CODE (new_var) == INDIRECT_REF);
-		gimple g = gimple_build_assign (new_var, x);
-		gimple_seq_add_stmt (&ctx->ganglocal_init, g);
-	      }
+	      gcc_unreachable ();
 	  }
 	map_cnt++;
 	break;
@@ -12153,9 +12005,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   irlist = NULL;
   orlist = NULL;
 
-  if (is_gimple_omp_oacc (stmt))
-    oacc_init_count_vars (ctx, clauses);
-
   if (is_oacc_parallel (ctx) && ctx->reductions)
     {
       lower_oacc_reductions (IFN_GOACC_REDUCTION_SETUP, GOMP_DIM_GANG,
@@ -12168,8 +12017,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 			      clauses, &orlist, ctx, true);
     }
 
-  lower_omp (&ctx->ganglocal_init, ctx);
-
   if (offloaded)
     {
       /* Declare all the variables created by mapping and the variables
@@ -12424,7 +12271,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
       gimple_seq_add_stmt (&new_body,
 	  		   gimple_build_assign (ctx->receiver_decl, t));
     }
-  gimple_seq_add_seq (&new_body, ctx->ganglocal_init);
   gimple_seq_add_seq (&new_body, fplist);
 
   if (offloaded)
@@ -12457,14 +12303,6 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
   gimple_bind_add_stmt (bind, stmt);
   gimple_bind_add_seq (bind, olist);
 
-  gimple_seq sz_ilist = NULL;
-  tree sz = create_tmp_var (size_type_node, "__ganglocal_size");
-  gimplify_and_add (build2 (MODIFY_EXPR, size_type_node, sz,
-			    ctx->ganglocal_size_host),
-		    &sz_ilist);
-  gsi_insert_seq_before (gsi_p, sz_ilist, GSI_SAME_STMT);
-
-  gimple_omp_target_set_ganglocal_size (stmt, sz);
   pop_gimplify_context (NULL);
 }
 
diff --git a/gcc/tree-parloops.c b/gcc/tree-parloops.c
index c124b87..c6942d5 100644
--- a/gcc/tree-parloops.c
+++ b/gcc/tree-parloops.c
@@ -2080,9 +2080,6 @@ create_parallel_loop (struct loop *loop, tree loop_fn, tree data,
       gimple_omp_target_set_child_fn (stmt, child_fn);
       tree data_arg = gimple_omp_target_data_arg (kernels);
       gimple_omp_target_set_data_arg (stmt, data_arg);
-      tree ganglocal_size
-	= gimple_call_arg (goacc_kernels_internal, /* TODO */ 6);
-      gimple_omp_target_set_ganglocal_size (stmt, ganglocal_size);
 
       gimple_set_location (stmt, loc);
 
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 8da3dcd..1b52aa2 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -537,9 +537,6 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
 	case GOMP_MAP_FORCE_TO:
 	  pp_string (pp, "force_to");
 	  break;
-	case GOMP_MAP_FORCE_TO_GANGLOCAL:
-	  pp_string (pp, "force_to_ganglocal");
-	  break;
 	case GOMP_MAP_FORCE_FROM:
 	  pp_string (pp, "force_from");
 	  break;
@@ -570,21 +567,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, int flags)
 	  else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
 		   && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET)
 	    pp_string (pp, " [pointer set, len: ");
-
-	  else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
-		   && (OMP_CLAUSE_MAP_KIND (clause)
-		       == GOMP_MAP_FORCE_TO_GANGLOCAL))
-	    pp_string (pp, " [gang-local copy, len: ");
 	  else
 	    pp_string (pp, " [len: ");
 	  dump_generic_node (pp, OMP_CLAUSE_SIZE (clause),
 			     spc, flags, false);
 	  pp_right_bracket (pp);
 	}
-      else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
-	       && (OMP_CLAUSE_MAP_KIND (clause)
-		   == GOMP_MAP_FORCE_TO_GANGLOCAL))
-	pp_string (pp, " [gang-local copy]");
       pp_right_paren (pp);
       break;
 

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Nathan Sidwell]

On 08/26/15 21:37, Cesar Philippidis wrote:
> This patch strips out all of the references to ganglocal memory in gcc.
> Unfortunately, the runtime api still takes a shared memory parameter, so
> I haven't made any changes there yet. Perhaps we could still keep the
> shared memory argument to GOACC_parallel, but remove all of the support
> for ganglocal mappings. Then again, maybe we still need support
> ganglocal mappings for legacy purposes.
>
> With the ganglocal mapping aside, I'm in favor of leaving the shared
> memory argument to GOACC_parallel, just in case we find another use for
> shared memory in the future.
>
> Nathan, what do you want to do here?

We should remove the parameter.

1) the patch I posted earlier this week for trunk review doesn't have it
2) if it turns out to be needed in the future, it can be done by extending the 
tagging scheme we now have in that API
3) It's a target-specific concept and if needed I strongly suspect either 
compile time known by the target compiler (and hence emittable in the offload 
data), or deducible at runtime from other data.

WRT to the patch you've posted, I think you can totally excise 
'GOMP_MAP_FORCE_TO_GANGLOCAL' and friends from gomp-constants.h and from the 
runtime too.  (that could be a separate patch).

nathan

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Cesar Philippidis]

On 08/27/2015 06:13 AM, Nathan Sidwell wrote:
> On 08/26/15 21:37, Cesar Philippidis wrote:
>> This patch strips out all of the references to ganglocal memory in gcc.
>> Unfortunately, the runtime api still takes a shared memory parameter, so
>> I haven't made any changes there yet. Perhaps we could still keep the
>> shared memory argument to GOACC_parallel, but remove all of the support
>> for ganglocal mappings. Then again, maybe we still need support
>> ganglocal mappings for legacy purposes.
>>
>> With the ganglocal mapping aside, I'm in favor of leaving the shared
>> memory argument to GOACC_parallel, just in case we find another use for
>> shared memory in the future.
>>
>> Nathan, what do you want to do here?
> 
> We should remove the parameter.
> 
> 1) the patch I posted earlier this week for trunk review doesn't have it

I've committed this patch to gomp-4_0-branch. Do you want to apply that
patch to gomp-4_0-branch since ganglocal memory is no longer used? Just
remember that you'll need to teach expand_omp_target not to pass the
shared memory argument to the runtime.

> 2) if it turns out to be needed in the future, it can be done by
> extending the tagging scheme we now have in that API
> 3) It's a target-specific concept and if needed I strongly suspect
> either compile time known by the target compiler (and hence emittable in
> the offload data), or deducible at runtime from other data.

That sounds reasonable.

> WRT to the patch you've posted, I think you can totally excise
> 'GOMP_MAP_FORCE_TO_GANGLOCAL' and friends from gomp-constants.h and from
> the runtime too.  (that could be a separate patch).

I'll create a follow up patch for that later, probably after I finish
working on the auto-independent loop patch. In the meantime, I'm found a
bug where acc routine calls aren't being checked for compatible
parallelism. E.g.

  #pragma acc routine gang
  void foo ();

  ...

  #pragma acc parallel loop worker
  for (...)
     foo ();

The call to foo isn't being reported as an error, which it should. I'm
testing a fix for this.

Cesar

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Nathan Sidwell]

On 08/27/15 13:25, Cesar Philippidis wrote:
> On 08/27/2015 06:13 AM, Nathan Sidwell wrote:

> I'll create a follow up patch for that later, probably after I finish

ok

> working on the auto-independent loop patch. In the meantime, I'm found a
> bug where acc routine calls aren't being checked for compatible
> parallelism. E.g.
>
>    #pragma acc routine gang
>    void foo ();
>
>    ...
>
>    #pragma acc parallel loop worker
>    for (...)
>       foo ();
>
> The call to foo isn't being reported as an error, which it should. I'm
> testing a fix for this.

Yeah, I discovered  this yesterday, coincidentally when the partition merging 
optimization blew up.  thanks for fixing.

nathan

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Tom de Vries]

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

On 27/08/15 03:37, Cesar Philippidis wrote:
> -  ctx->ganglocal_size_host = align_and_expand (&gl_host, host_size, align);

I suspect this caused a bootstrap failure (align_and_expand unused). 
Worked-around as attached.

Thanks,
- Tom

[-- Attachment #2: 0002-Mark-align_and_expand-as-unused.patch --]
[-- Type: text/x-patch, Size: 646 bytes --]

Mark align_and_expand as unused

2015-09-01  Tom de Vries  <tom@codesourcery.com>

	* omp-low.c (align_and_expand): Mark as unused.
---
 gcc/omp-low.c | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a62daa2..fdca880 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -1450,7 +1450,7 @@ omp_copy_decl (tree var, copy_body_data *cb)
 
 /* Modify the old size *POLDSZ to align it up to ALIGN, and then return
    a value with SIZE added to it.  */
-static tree
+static tree ATTRIBUTE_UNUSED
 align_and_expand (tree *poldsz, tree size, unsigned int align)
 {
   tree oldsz = *poldsz;
-- 
1.9.1

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Tom de Vries]

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

On 27/08/15 03:37, Cesar Philippidis wrote:
> -      tree ganglocal_size
> -	= gimple_call_arg (goacc_kernels_internal, /* TODO */ 6);
> -      gimple_omp_target_set_ganglocal_size (stmt, ganglocal_size);

This caused a bootstrap failure. Committed as attached.

Thanks,
- Tom

[-- Attachment #2: 0003-Remove-unused-var-in-create_parallel_loop.patch --]
[-- Type: text/x-patch, Size: 838 bytes --]

Remove unused var in create_parallel_loop

2015-09-01  Tom de Vries  <tom@codesourcery.com>

	* tree-parloops.c (create_parallel_loop): Remove unused variable
	goacc_kernels_internal.
---
 gcc/tree-parloops.c | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/gcc/tree-parloops.c b/gcc/tree-parloops.c
index c6942d5..02dd6d5 100644
--- a/gcc/tree-parloops.c
+++ b/gcc/tree-parloops.c
@@ -2058,9 +2058,6 @@ create_parallel_loop (struct loop *loop, tree loop_fn, tree data,
 	 GOACC_kernels_internal call.  */
       gomp_target *kernels = as_a <gomp_target *> (gsi_stmt (gsi));
 
-      gsi_prev (&gsi);
-      gcall *goacc_kernels_internal = as_a <gcall *> (gsi_stmt (gsi));
-
       tree clauses = gimple_omp_target_clauses (kernels);
       /* FIXME: We need a more intelligent mapping onto vector, gangs,
 	 workers.  */
-- 
1.9.1

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Thomas Schwinge]

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

Hi!

On Tue, 1 Sep 2015 18:29:55 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> On 27/08/15 03:37, Cesar Philippidis wrote:
> > -  ctx->ganglocal_size_host = align_and_expand (&gl_host, host_size, align);
> 
> I suspect this caused a bootstrap failure (align_and_expand unused). 
> Worked-around as attached.

> --- a/gcc/omp-low.c
> +++ b/gcc/omp-low.c
> @@ -1450,7 +1450,7 @@ omp_copy_decl (tree var, copy_body_data *cb)
>  
>  /* Modify the old size *POLDSZ to align it up to ALIGN, and then return
>     a value with SIZE added to it.  */
> -static tree
> +static tree ATTRIBUTE_UNUSED
>  align_and_expand (tree *poldsz, tree size, unsigned int align)
>  {
>    tree oldsz = *poldsz;

If I remember correctly, this has only ever been used in the "ganglocal"
implementation -- which is now gone.  So, should align_and_expand also be
elided (Cesar)?


Grüße,
 Thomas

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

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Cesar Philippidis]

On 09/18/2015 01:39 AM, Thomas Schwinge wrote:

> On Tue, 1 Sep 2015 18:29:55 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
>> On 27/08/15 03:37, Cesar Philippidis wrote:
>>> -  ctx->ganglocal_size_host = align_and_expand (&gl_host, host_size, align);
>>
>> I suspect this caused a bootstrap failure (align_and_expand unused). 
>> Worked-around as attached.
> 
>> --- a/gcc/omp-low.c
>> +++ b/gcc/omp-low.c
>> @@ -1450,7 +1450,7 @@ omp_copy_decl (tree var, copy_body_data *cb)
>>  
>>  /* Modify the old size *POLDSZ to align it up to ALIGN, and then return
>>     a value with SIZE added to it.  */
>> -static tree
>> +static tree ATTRIBUTE_UNUSED
>>  align_and_expand (tree *poldsz, tree size, unsigned int align)
>>  {
>>    tree oldsz = *poldsz;
> 
> If I remember correctly, this has only ever been used in the "ganglocal"
> implementation -- which is now gone.  So, should align_and_expand also be
> elided (Cesar)?

Most likely. I probably overlooked it when I was working on that
ganglocal removal patch. Can you remove it please? I'm already juggling
a couple of patches right now.

Thanks,
Cesar

Re: [gomp4, wip] remove references to ganglocal shared memory inside gcc

[Thomas Schwinge]

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

Hi!

On Fri, 18 Sep 2015 06:51:18 -0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> On 09/18/2015 01:39 AM, Thomas Schwinge wrote:
> 
> > On Tue, 1 Sep 2015 18:29:55 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> >> On 27/08/15 03:37, Cesar Philippidis wrote:
> >>> -  ctx->ganglocal_size_host = align_and_expand (&gl_host, host_size, align);
> >>
> >> I suspect this caused a bootstrap failure (align_and_expand unused). 
> >> Worked-around as attached.

> > If I remember correctly, this has only ever been used in the "ganglocal"
> > implementation -- which is now gone.  So, should align_and_expand also be
> > elided (Cesar)?
> 
> Most likely. I probably overlooked it when I was working on that
> ganglocal removal patch. Can you remove it please? I'm already juggling
> a couple of patches right now.

Together with removal of printing the declarator for sdata, committed to
gomp-4_0-branch in r228038:

commit f5890b47c1b6f09134c4bfadcc7ece0d5403a1d7
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date:   Wed Sep 23 10:35:31 2015 +0000

    More "ganglocal" cleanup
    
    	gcc/
    	* config/nvptx/nvptx.c (nvptx_file_start): Don't print declaration
    	of sdata.
    	* omp-low.c (align_and_expand): Remove function.
    
    git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@228038 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/ChangeLog.gomp       |  6 ++++++
 gcc/config/nvptx/nvptx.c |  1 -
 gcc/omp-low.c            | 15 ---------------
 3 files changed, 6 insertions(+), 16 deletions(-)

diff --git gcc/ChangeLog.gomp gcc/ChangeLog.gomp
index 21c6fa0..c66f80a 100644
--- gcc/ChangeLog.gomp
+++ gcc/ChangeLog.gomp
@@ -1,3 +1,9 @@
+2015-09-23  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* config/nvptx/nvptx.c (nvptx_file_start): Don't print declaration
+	of sdata.
+	* omp-low.c (align_and_expand): Remove function.
+
 2015-09-22  Cesar Philippidis  <cesar@codesourcery.com>
 
 	* gimplify.c (oacc_default_clause): Inspect pointer types when
diff --git gcc/config/nvptx/nvptx.c gcc/config/nvptx/nvptx.c
index 5640e34..37b50a3 100644
--- gcc/config/nvptx/nvptx.c
+++ gcc/config/nvptx/nvptx.c
@@ -4063,7 +4063,6 @@ nvptx_file_start (void)
   else
     fputs ("\t.target\tsm_30\n", asm_out_file);
   fprintf (asm_out_file, "\t.address_size %d\n", GET_MODE_BITSIZE (Pmode));
-  fprintf (asm_out_file, "\t.extern .shared .u8 sdata[];\n");
   fputs ("// END PREAMBLE\n", asm_out_file);
 }
 
diff --git gcc/omp-low.c gcc/omp-low.c
index ee527d0..ec76096 100644
--- gcc/omp-low.c
+++ gcc/omp-low.c
@@ -1446,21 +1446,6 @@ omp_copy_decl (tree var, copy_body_data *cb)
   return error_mark_node;
 }
 
-/* Modify the old size *POLDSZ to align it up to ALIGN, and then return
-   a value with SIZE added to it.  */
-static tree ATTRIBUTE_UNUSED
-align_and_expand (tree *poldsz, tree size, unsigned int align)
-{
-  tree oldsz = *poldsz;
-  oldsz = fold_build2 (BIT_AND_EXPR, size_type_node,
-		       fold_build2 (PLUS_EXPR, size_type_node,
-				    oldsz, size_int (align - 1)),
-		       fold_build1 (BIT_NOT_EXPR, size_type_node,
-				    size_int (align - 1)));
-  *poldsz = oldsz;
-  return fold_build2 (PLUS_EXPR, size_type_node, oldsz, size);
-}
-
 /* Debugging dumps for parallel regions.  */
 void dump_omp_region (FILE *, struct omp_region *, int);
 void debug_omp_region (struct omp_region *);


Grüße,
 Thomas

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