[openacc, PR78266, committed] Fix diff_type in expand_oacc_for char iter_type

[openacc, PR78266, committed] Fix diff_type in expand_oacc_for char iter_type

[Tom de Vries]

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

Hi,

this fixes PR78266, an openacc PR.

When compiling a gang loop with an iteration variable of type 'unsigned 
char' and 256 gangs:
...
#pragma acc parallel loop num_gangs (256)
for (unsigned char j = 0; j < 5; j++)
..
we run into trouble.

The 'diff_type' in expand_oacc_for is set to 'signed char', and we 
generate f.i.:
...
   _41 = GOACC_DIM_SIZE (0);
   _29 = (signed char) _41;
...
where _41 is 256, so forwprop2 folds _29 to '0'.

The patch fixes this by ensuring that diff_type is chosen big enough in 
expand_oacc_for.

Tested libgomp on x86_64 with nvptx accelerator.

Committed.

Thanks,
- Tom

[-- Attachment #2: 0001-Fix-diff_type-in-expand_oacc_for-char-iter_type.patch --]
[-- Type: text/x-patch, Size: 3405 bytes --]

Fix diff_type in expand_oacc_for char iter_type

2017-08-07  Tom de Vries  <tom@codesourcery.com>

	PR middle-end/78266
	* omp-expand.c (expand_oacc_for): Ensure diff_type is large enough.

	* testsuite/libgomp.oacc-c-c++-common/vprop-2.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/vprop.c: Remove xfail.

---
 gcc/omp-expand.c                                   |  2 +
 .../testsuite/libgomp.oacc-c-c++-common/vprop-2.c  | 45 ++++++++++++++++++++++
 .../testsuite/libgomp.oacc-c-c++-common/vprop.c    |  1 -
 3 files changed, 47 insertions(+), 1 deletion(-)

diff --git a/gcc/omp-expand.c b/gcc/omp-expand.c
index 970e04f..1eef7c0 100644
--- a/gcc/omp-expand.c
+++ b/gcc/omp-expand.c
@@ -5328,6 +5328,8 @@ expand_oacc_for (struct omp_region *region, struct omp_for_data *fd)
     }
   if (POINTER_TYPE_P (diff_type) || TYPE_UNSIGNED (diff_type))
     diff_type = signed_type_for (diff_type);
+  if (TYPE_PRECISION (diff_type) < TYPE_PRECISION (integer_type_node))
+    diff_type = integer_type_node;
 
   basic_block entry_bb = region->entry; /* BB ending in OMP_FOR */
   basic_block exit_bb = region->exit; /* BB ending in OMP_RETURN */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop-2.c
new file mode 100644
index 0000000..046ac68
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop-2.c
@@ -0,0 +1,45 @@
+/* { dg-do run } */
+
+#include <assert.h>
+
+#define DO_PRAGMA(x) _Pragma (#x)
+
+#define test(idx,type,ngangs)                           \
+  void                                                  \
+  test_##idx ()                                         \
+  {                                                     \
+    int b[100];                                         \
+                                                        \
+    for (unsigned int i = 0; i < 100; i++)              \
+      b[i] = 0;                                         \
+                                                        \
+    DO_PRAGMA(acc parallel num_gangs (ngangs) copy (b)) \
+      {                                                 \
+        _Pragma("acc loop gang")                        \
+          for (type j = 0; j < 5; j++)                  \
+            {                                           \
+              _Pragma("acc loop vector")                \
+                for (unsigned int i = 0; i < 20; i++)   \
+                  b[j * 20 + i] = -2;                   \
+            }                                           \
+      }                                                 \
+                                                        \
+    for (unsigned int i = 0; i < 100; i++)              \
+      assert (b[i] == -2);                              \
+  }
+
+test (0, signed char, 256)
+test (1, unsigned char, 256)
+test (2, signed short, 65535)
+test (3, unsigned short, 65535)
+
+int
+main ()
+{
+  test_0 ();
+  test_1 ();
+  test_2 ();
+  test_3 ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
index 0ac0cf6..e4dd682 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/vprop.c
@@ -1,5 +1,4 @@
 /* { dg-do run } */
-/* { dg-xfail-run-if "PR78266" { openacc_nvidia_accel_selected } } */
 
 #include <assert.h>