diff mbox series

[OpenACC,1/3] Non-contiguous array support for OpenACC data clauses (re-submission), front-end patches

Message ID f2e23787-6255-b3b4-c571-a0635c91ed71@mentor.com
State New
Headers show
Series [OpenACC,1/3] Non-contiguous array support for OpenACC data clauses (re-submission), front-end patches | expand

Commit Message

Chung-Lin Tang Aug. 20, 2019, 11:36 a.m. UTC 
Hi Jakub, Thomas,
this is a re-submission of the patch-set from [1].

The usage of the term "dynamic arrays" didn't go well with Jakub the last time,
so this time I'm referring to this functionality as "non-contiguous arrays".

int *a[100], **b;

// re-constructs array slices on GPU and copies data in
#pragma acc parallel copyin (a[0:n][0:m], b[1:x][5:y])

The overall implementation has not changed much from the last submission,
mainly the renaming changes and rebasing to current trunk.

The first patch here are the C/C++ front-end patches.

Thanks,
Chung-Lin

[1] https://gcc.gnu.org/ml/gcc-patches/2018-10/msg00937.html

	gcc/c/
	* c-typeck.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
	parameter, adjust recursive call site, add cases for allowing
	pointer based multi-dimensional arrays for OpenACC.
	(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
	handle non-contiguous case to create dynamic array map.

	gcc/cp/
	* semantics.c (handle_omp_array_sections_1): Add 'bool &non_contiguous'
	parameter, adjust recursive call site, add cases for allowing
	pointer based multi-dimensional arrays for OpenACC.
	(handle_omp_array_sections): Adjust handle_omp_array_sections_1 call,
	handle non-contiguous case to create dynamic array map.

Comments

Thomas Schwinge Oct. 7, 2019, 1:50 p.m. UTC | #1 
Hi Chung-Lin!

Thanks for your work on this.


Please reference PR76739 in your submission/ChangeLog updates.


We'll need Jakub to review the generic code changes, but let me provide
some first review remarks, too.


On 2019-08-20T19:36:24+0800, Chung-Lin Tang <chunglin_tang@mentor.com> wrote:
> The first patch here are the C/C++ front-end patches.

As far as I'm concerned, it doesn't make sense to artificially split up
patches like that, given that the individual three pieces can only be
considered all together.

And if posting split-up for other reasonse, then please make sure that
the individual patch submission emails have a common "cover letter" email
so that they show up as one email thread.


> --- gcc/c/c-typeck.c	(revision 274618)
> +++ gcc/c/c-typeck.c	(working copy)

> @@ -13099,6 +13100,21 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
>  		    }
>  		}
>  	    }
> +
> +	  /* For OpenACC, if the low_bound/length suggest this is a subarray,
> +	     and is referenced through by a pointer, then mark this as
> +	     non-contiguous.  */

I don't directly understand this logic.  I'll have to think about it
more.

> +	  if (ort == C_ORT_ACC
> +	      && types.length () > 0
> +	      && (TREE_CODE (low_bound) != INTEGER_CST
> +		  || integer_nonzerop (low_bound)
> +		  || (length && (TREE_CODE (length) != INTEGER_CST
> +				 || !tree_int_cst_equal (size, length)))))
> +	    {
> +	      tree x = types.last ();
> +	      if (TREE_CODE (x) == POINTER_TYPE)
> +		non_contiguous = true;
> +	    }
>  	}
>        else if (length == NULL_TREE)
>  	{
> @@ -13142,7 +13158,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
>        /* If there is a pointer type anywhere but in the very first
>  	 array-section-subscript, the array section can't be contiguous.  */
>        if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
> -	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
> +	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST
> +	  && ort != C_ORT_ACC)
>  	{
>  	  error_at (OMP_CLAUSE_LOCATION (c),
>  		    "array section is not contiguous in %qs clause",
> @@ -13149,6 +13166,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<t
>  		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
>  	  return error_mark_node;
>  	}
> +      else if (TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
> +	non_contiguous = true;
>      }
>    else
>      {


> @@ -13337,6 +13365,14 @@ handle_omp_array_sections (tree c, enum c_omp_regi
>  		size = size_binop (MULT_EXPR, size, l);
>  	    }
>  	}
> +      if (non_contiguous)
> +	{
> +	  int kind = OMP_CLAUSE_MAP_KIND (c);
> +	  OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
> +	  OMP_CLAUSE_DECL (c) = t;
> +	  OMP_CLAUSE_SIZE (c) = ncarray_dims;
> +	  return false;
> +	}

I'm expecting to see front end test cases (probably
'-fdump-tree-original' scanning?) for a good number of different data
clauses/array variants, whether that flag 'GOMP_MAP_NONCONTIG_ARRAY' has
been set or not.  (That would then also document the logic presented
above, and should thus help me understand that.)


> --- gcc/cp/semantics.c	(revision 274618)
> +++ gcc/cp/semantics.c	(working copy)

Likewise.


Grüße
 Thomas
diff mbox series

Patch

Index: gcc/c/c-typeck.c
===================================================================
--- gcc/c/c-typeck.c	(revision 274618)
+++ gcc/c/c-typeck.c	(working copy)
@@ -12848,7 +12848,7 @@  c_finish_omp_cancellation_point (location_t loc, t
 static tree
 handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     bool &maybe_zero_len, unsigned int &first_non_one,
-			     enum c_omp_region_type ort)
+			     bool &non_contiguous, enum c_omp_region_type ort)
 {
   tree ret, low_bound, length, type;
   if (TREE_CODE (t) != TREE_LIST)
@@ -12933,7 +12933,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
     }
 
   ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
-				     maybe_zero_len, first_non_one, ort);
+				     maybe_zero_len, first_non_one,
+				     non_contiguous, ort);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
@@ -13099,6 +13100,21 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
 		    }
 		}
 	    }
+
+	  /* For OpenACC, if the low_bound/length suggest this is a subarray,
+	     and is referenced through by a pointer, then mark this as
+	     non-contiguous.  */
+	  if (ort == C_ORT_ACC
+	      && types.length () > 0
+	      && (TREE_CODE (low_bound) != INTEGER_CST
+		  || integer_nonzerop (low_bound)
+		  || (length && (TREE_CODE (length) != INTEGER_CST
+				 || !tree_int_cst_equal (size, length)))))
+	    {
+	      tree x = types.last ();
+	      if (TREE_CODE (x) == POINTER_TYPE)
+		non_contiguous = true;
+	    }
 	}
       else if (length == NULL_TREE)
 	{
@@ -13142,7 +13158,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
       /* If there is a pointer type anywhere but in the very first
 	 array-section-subscript, the array section can't be contiguous.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
-	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
+	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST
+	  && ort != C_ORT_ACC)
 	{
 	  error_at (OMP_CLAUSE_LOCATION (c),
 		    "array section is not contiguous in %qs clause",
@@ -13149,6 +13166,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
 		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	  return error_mark_node;
 	}
+      else if (TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
+	non_contiguous = true;
     }
   else
     {
@@ -13176,6 +13195,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
+  bool non_contiguous = false;
   auto_vec<tree, 10> types;
   tree *tp = &OMP_CLAUSE_DECL (c);
   if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -13185,7 +13205,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
     tp = &TREE_VALUE (*tp);
   tree first = handle_omp_array_sections_1 (c, *tp, types,
 					    maybe_zero_len, first_non_one,
-					    ort);
+					    non_contiguous, ort);
   if (first == error_mark_node)
     return true;
   if (first == NULL_TREE)
@@ -13218,6 +13238,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
       unsigned int num = types.length (), i;
       tree t, side_effects = NULL_TREE, size = NULL_TREE;
       tree condition = NULL_TREE;
+      tree ncarray_dims = NULL_TREE;
 
       if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
 	maybe_zero_len = true;
@@ -13241,6 +13262,13 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 	    length = fold_convert (sizetype, length);
 	  if (low_bound == NULL_TREE)
 	    low_bound = integer_zero_node;
+
+	  if (non_contiguous)
+	    {
+	      ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+	      continue;
+	    }
+
 	  if (!maybe_zero_len && i > first_non_one)
 	    {
 	      if (integer_nonzerop (low_bound))
@@ -13337,6 +13365,14 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 		size = size_binop (MULT_EXPR, size, l);
 	    }
 	}
+      if (non_contiguous)
+	{
+	  int kind = OMP_CLAUSE_MAP_KIND (c);
+	  OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+	  OMP_CLAUSE_DECL (c) = t;
+	  OMP_CLAUSE_SIZE (c) = ncarray_dims;
+	  return false;
+	}
       if (side_effects)
 	size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
       if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
Index: gcc/cp/semantics.c
===================================================================
--- gcc/cp/semantics.c	(revision 274618)
+++ gcc/cp/semantics.c	(working copy)
@@ -4626,7 +4626,7 @@  omp_privatize_field (tree t, bool shared)
 static tree
 handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
 			     bool &maybe_zero_len, unsigned int &first_non_one,
-			     enum c_omp_region_type ort)
+			     bool &non_contiguous, enum c_omp_region_type ort)
 {
   tree ret, low_bound, length, type;
   if (TREE_CODE (t) != TREE_LIST)
@@ -4711,7 +4711,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
       && TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
     TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t), false);
   ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
-				     maybe_zero_len, first_non_one, ort);
+				     maybe_zero_len, first_non_one,
+				     non_contiguous, ort);
   if (ret == error_mark_node || ret == NULL_TREE)
     return ret;
 
@@ -4889,6 +4890,21 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
 		    }
 		}
 	    }
+
+	  /* For OpenACC, if the low_bound/length suggest this is a subarray,
+	     and is referenced through by a pointer, then mark this as
+	     non-contiguous.  */
+	  if (ort == C_ORT_ACC
+	      && types.length () > 0
+	      && (TREE_CODE (low_bound) != INTEGER_CST
+		  || integer_nonzerop (low_bound)
+		  || (length && (TREE_CODE (length) != INTEGER_CST
+				 || !tree_int_cst_equal (size, length)))))
+	    {
+	      tree x = types.last ();
+	      if (TREE_CODE (x) == POINTER_TYPE)
+		non_contiguous = true;
+	    }
 	}
       else if (length == NULL_TREE)
 	{
@@ -4932,7 +4948,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
       /* If there is a pointer type anywhere but in the very first
 	 array-section-subscript, the array section can't be contiguous.  */
       if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_DEPEND
-	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
+	  && TREE_CODE (TREE_CHAIN (t)) == TREE_LIST
+	  && ort != C_ORT_ACC)
 	{
 	  error_at (OMP_CLAUSE_LOCATION (c),
 		    "array section is not contiguous in %qs clause",
@@ -4939,6 +4956,8 @@  handle_omp_array_sections_1 (tree c, tree t, vec<t
 		    omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	  return error_mark_node;
 	}
+      else if (TREE_CODE (TREE_CHAIN (t)) == TREE_LIST)
+	non_contiguous = true;
     }
   else
     {
@@ -4966,6 +4985,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 {
   bool maybe_zero_len = false;
   unsigned int first_non_one = 0;
+  bool non_contiguous = false;
   auto_vec<tree, 10> types;
   tree *tp = &OMP_CLAUSE_DECL (c);
   if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_DEPEND
@@ -4975,7 +4995,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
     tp = &TREE_VALUE (*tp);
   tree first = handle_omp_array_sections_1 (c, *tp, types,
 					    maybe_zero_len, first_non_one,
-					    ort);
+					    non_contiguous, ort);
   if (first == error_mark_node)
     return true;
   if (first == NULL_TREE)
@@ -5009,6 +5029,7 @@  handle_omp_array_sections (tree c, enum c_omp_regi
       unsigned int num = types.length (), i;
       tree t, side_effects = NULL_TREE, size = NULL_TREE;
       tree condition = NULL_TREE;
+      tree ncarray_dims = NULL_TREE;
 
       if (int_size_in_bytes (TREE_TYPE (first)) <= 0)
 	maybe_zero_len = true;
@@ -5034,6 +5055,13 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 	    length = fold_convert (sizetype, length);
 	  if (low_bound == NULL_TREE)
 	    low_bound = integer_zero_node;
+
+	  if (non_contiguous)
+	    {
+	      ncarray_dims = tree_cons (low_bound, length, ncarray_dims);
+	      continue;
+	    }
+
 	  if (!maybe_zero_len && i > first_non_one)
 	    {
 	      if (integer_nonzerop (low_bound))
@@ -5125,6 +5153,14 @@  handle_omp_array_sections (tree c, enum c_omp_regi
 	}
       if (!processing_template_decl)
 	{
+	  if (non_contiguous)
+	    {
+	      int kind = OMP_CLAUSE_MAP_KIND (c);
+	      OMP_CLAUSE_SET_MAP_KIND (c, kind | GOMP_MAP_NONCONTIG_ARRAY);
+	      OMP_CLAUSE_DECL (c) = t;
+	      OMP_CLAUSE_SIZE (c) = ncarray_dims;
+	      return false;
+	    }
 	  if (side_effects)
 	    size = build2 (COMPOUND_EXPR, sizetype, side_effects, size);
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION