[PATCH] [og13] OpenACC: Vector length warning fixes for implicit mapping/declare create tests

[PATCH] [og13] OpenACC: Vector length warning fixes for implicit mapping/declare create tests

[Julian Brown]

This patch adds expected "vector length" warnings to several tests
for NVPTX.

Tested with offloading to NVPTX. I will apply (to og13) shortly.

2023-07-11  Julian Brown  <julian@codesourcery.com>

libgomp/
	* testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c: Add
	expected warning.
	* testsuite/libgomp.oacc-fortran/declare-create-1.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/declare-create-2.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/declare-create-3.f90: Likewise.
	* testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90:
	Likewise.
	* testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90:
	Likewise.
---
 libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c | 1 +
 libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90      | 1 +
 libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90      | 1 +
 libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90      | 1 +
 .../testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90 | 1 +
 .../testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90 | 1 +
 6 files changed, 6 insertions(+)

diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c
index 4825e875998..ed0ab94cd8f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/implicit-mapping-1.c
@@ -12,6 +12,7 @@ int main(void)
 #pragma acc enter data copyin(arr[30:10])
 
 #pragma acc serial
+/* { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 } */
   {
     arr[33] = 66;
   }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90
index 9e7e60f1440..057b5eb958a 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-1.f90
@@ -11,6 +11,7 @@ use m
 mint = 0
 
 !$acc serial
+! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 }
 mint = 5
 !$acc end serial
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90
index 675f6902775..dd7c9798fba 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-2.f90
@@ -13,6 +13,7 @@ allocate(mint)
 mint = 0
 
 !$acc serial
+! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 }
 mint = 5
 !$acc end serial
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90
index 16651cb1f5e..7cceaa5f8a3 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/declare-create-3.f90
@@ -13,6 +13,7 @@ allocate(mint(1:20))
 mint = 0
 
 !$acc serial
+! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 }
 mint = 5
 !$acc end serial
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90
index 4b61e1cee9b..8b173c72d88 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-1.f90
@@ -19,6 +19,7 @@ integer :: arr(*)
 !$acc enter data copyin(arr(1:10))
 
 !$acc serial
+! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 }
 arr(5) = 5
 !$acc end serial
 
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90
index daf7089915f..659fe8e3c06 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/nonlexical-assumed-size-2.f90
@@ -30,6 +30,7 @@ integer :: arr(*)
 ! overwritten.
 
 !$acc serial
+! { dg-warning {using .vector_length \(32\)., ignoring 1} "" { target openacc_nvidia_accel_selected } .-1 }
 ! This access is then done via the on-target pointer.
 arr(5) = 5
 !$acc end serial
-- 
2.41.0

[PATCH] OpenMP: Strided/rectangular 'target update' out-of-bounds array lookup fix

[Julian Brown]

This patch fixes a bug with the calculation of array bounds in the
metadata for noncontiguous 'target update' directives.  We record the
array base address, a bias and the array length to pass to libgomp --
but at present, we use the 'whole array size' for the last, which means
that at runtime we might look up an array with lower bound "base+bias"
and upper bound "base+bias+length", which for non-zero bias will overflow
the actual bounds of the array on the host and will (sometimes) return
an unrelated block instead of the correct one.

The fix is to instead calculate a size for the array that encloses the
elements to be transferred, and is guaranteed to be entirely within the
array (user errors excepted).

Tested with offloading to NVPTX.  I will apply (to og13) shortly.

2023-07-11  Julian Brown  <julian@codesourcery.com>

gcc/
	* omp-low.cc (lower_omp_target): Calculate volume enclosing
	transferred elements instead of using whole array size for
	noncontiguous 'target update' operations.
---
 gcc/omp-low.cc | 23 +++++++++++++++++++----
 1 file changed, 19 insertions(+), 4 deletions(-)

diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc
index 05ac917fb27..c7706a5921f 100644
--- a/gcc/omp-low.cc
+++ b/gcc/omp-low.cc
@@ -14444,11 +14444,13 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 
 		tree bias = size_zero_node;
 		tree volume = size_one_node;
+		tree enclosure = size_one_node;
 		for (i = dims - 1; i >= 0; i--)
 		  {
 		    tree dim = (*vdim)[i].value;
 		    tree index = (*vindex)[i].value;
 		    tree stride = (*vstride)[i].value;
+		    tree len = (*vlen)[i].value;
 
 		    /* For the bias we want, e.g.:
 
@@ -14463,6 +14465,20 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 				       size_binop (MULT_EXPR, volume,
 						   index_stride));
 		    volume = size_binop (MULT_EXPR, volume, dim);
+
+		    if (i == 0)
+		      {
+			tree elems_covered = size_binop (MINUS_EXPR, len,
+							 size_one_node);
+			elems_covered = size_binop (MULT_EXPR, elems_covered,
+						    stride);
+			elems_covered = size_binop (PLUS_EXPR, elems_covered,
+						    size_one_node);
+			enclosure = size_binop (MULT_EXPR, enclosure,
+						elems_covered);
+		      }
+		    else
+		      enclosure = volume;
 		  }
 
 		/* If we don't have a separate span size, use the element size
@@ -14470,10 +14486,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		if (!span)
 		  span = fold_convert (sizetype, elsize);
 
-		/* The size of the whole array -- to make sure we find any
-		   part of the array via splay-tree lookup that might be
-		   mapped on the target at runtime.  */
-		OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, arrsize, span);
+		/* The size of a volume enclosing the elements to be
+		   transferred.  */
+		OMP_CLAUSE_SIZE (oc) = size_binop (MULT_EXPR, enclosure, span);
 		/* And the bias of the first element we will update.  */
 		OMP_CLAUSE_SIZE (dn) = size_binop (MULT_EXPR, bias, span);
 
-- 
2.41.0