[gomp4, committed] Add goacc/kernels-acc-on-device.c

[gomp4, committed] Add goacc/kernels-acc-on-device.c

[Tom de Vries]

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

Hi,

this patch adds a new kernels region test-case.

Committed to gomp-4_0-branch as trivial.

Thanks,
- Tom

[-- Attachment #2: 0001-Add-goacc-kernels-acc-on-device.c.patch --]
[-- Type: text/x-patch, Size: 1203 bytes --]

Add goacc/kernels-acc-on-device.c

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

	* c-c++-common/goacc/kernels-acc-on-device.c: New test.
---
 .../c-c++-common/goacc/kernels-acc-on-device.c     | 39 ++++++++++++++++++++++
 1 file changed, 39 insertions(+)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
new file mode 100644
index 0000000..e9e93c7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
@@ -0,0 +1,39 @@
+/* { dg-additional-options "-O2" } */
+
+#include <openacc.h>
+
+#define N 32
+
+void
+foo (float *a, float *b)
+{
+  float exp;
+  int i;
+  int n;
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
+  {
+    int ii;
+
+    for (ii = 0; ii < N; ii++)
+      {
+	if (acc_on_device (acc_device_host))
+	  b[ii] = a[ii] + 1;
+	else
+	  b[ii] = a[ii];
+      }
+  }
+
+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
+  {
+    int ii;
+
+    for (ii = 0; ii < N; ii++)
+      {
+	if (acc_on_device (acc_device_host))
+	  b[ii] = a[ii] + 2;
+	else
+	  b[ii] = a[ii];
+      }
+  }
+}
-- 
1.9.1

Re: [gomp4, committed] Add goacc/kernels-acc-on-device.c

[Thomas Schwinge]

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

Hi Tom!

On Sat, 10 Oct 2015 12:49:01 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> --- /dev/null
> +++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
> @@ -0,0 +1,39 @@
> +/* { dg-additional-options "-O2" } */
> +
> +#include <openacc.h>

That doesn't work (at least in build-tree testing), as gcc/testsuite/ is
not set up to look for header files in [target]/libgomp/:

    [...]/source-gcc/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c:3:21: fatal error: openacc.h: No such file or directory
    compilation terminated.
    compiler exited with status 1

> +
> +#define N 32
> +
> +void
> +foo (float *a, float *b)
> +{
> +  float exp;
> +  int i;
> +  int n;
> +
> +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
> +  {
> +    int ii;
> +
> +    for (ii = 0; ii < N; ii++)
> +      {
> +	if (acc_on_device (acc_device_host))

Your two options are: if that's applicable/sufficient for what you intend
to test here, use __builtin_acc_on_device with a hard-coded acc_device_*,
or duplicate part of <openacc.h> as done for example in
gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c.

> +	  b[ii] = a[ii] + 1;
> +	else
> +	  b[ii] = a[ii];
> +      }
> +  }
> +
> +#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
> +  {
> +    int ii;
> +
> +    for (ii = 0; ii < N; ii++)
> +      {
> +	if (acc_on_device (acc_device_host))
> +	  b[ii] = a[ii] + 2;
> +	else
> +	  b[ii] = a[ii];
> +      }
> +  }
> +}


Grüße,
 Thomas

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

Re: [gomp4, committed] Add goacc/kernels-acc-on-device.c

[Tom de Vries]

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

On 12/10/15 12:49, Thomas Schwinge wrote:
> Hi Tom!
>
> On Sat, 10 Oct 2015 12:49:01 +0200, Tom de Vries<Tom_deVries@mentor.com>  wrote:
>> >--- /dev/null
>> >+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
>> >@@ -0,0 +1,39 @@
>> >+/* { dg-additional-options "-O2" } */
>> >+
>> >+#include <openacc.h>

Hi Thomas,

> That doesn't work (at least in build-tree testing), as gcc/testsuite/ is
> not set up to look for header files in [target]/libgomp/:
>
>      [...]/source-gcc/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c:3:21: fatal error: openacc.h: No such file or directory
>      compilation terminated.
>      compiler exited with status 1
>

Ah, I see. I was doing 'make' followed by 'make install', and then 
build-tree testing. The build-tree testing seems to pick up the header 
file from the install directory. So for me test passed.

>> >+
>> >+#define N 32
>> >+
>> >+void
>> >+foo (float *a, float *b)
>> >+{
>> >+  float exp;
>> >+  int i;
>> >+  int n;
>> >+
>> >+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
>> >+  {
>> >+    int ii;
>> >+
>> >+    for (ii = 0; ii < N; ii++)
>> >+      {
>> >+	if (acc_on_device (acc_device_host))
> Your two options are: if that's applicable/sufficient for what you intend
> to test here, use __builtin_acc_on_device with a hard-coded acc_device_*,
> or duplicate part of <openacc.h> as done for example in
> gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c.
>

Went with second option, committed as attached.

Thanks,
- Tom


[-- Attachment #2: 0001-Remove-openacc.h-include-from-goacc-kernels-acc-on-d.patch --]
[-- Type: text/x-patch, Size: 1705 bytes --]

Remove openacc.h include from goacc/kernels-acc-on-device.c

2015-10-12  Tom de Vries  <tom@codesourcery.com>

	* c-c++-common/goacc/kernels-acc-on-device.c: Remove openacc.h include.
	(enum acc_device_t, acc_on_device): Declare.
	(foo): Remove unused vars.  Use acc_device_X.
---
 .../c-c++-common/goacc/kernels-acc-on-device.c     | 27 ++++++++++++++++------
 1 file changed, 20 insertions(+), 7 deletions(-)

diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
index e9e93c7..784c66a 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
@@ -1,23 +1,36 @@
 /* { dg-additional-options "-O2" } */
 
-#include <openacc.h>
+#if __cplusplus
+extern "C" {
+#endif
+
+#if __cplusplus >= 201103
+# define __GOACC_NOTHROW noexcept
+#elif __cplusplus
+# define __GOACC_NOTHROW throw ()
+#else /* Not C++ */
+# define __GOACC_NOTHROW __attribute__ ((__nothrow__))
+#endif
+
+typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
+int acc_on_device (int) __GOACC_NOTHROW;
+
+#if __cplusplus
+}
+#endif
 
 #define N 32
 
 void
 foo (float *a, float *b)
 {
-  float exp;
-  int i;
-  int n;
-
 #pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
   {
     int ii;
 
     for (ii = 0; ii < N; ii++)
       {
-	if (acc_on_device (acc_device_host))
+	if (acc_on_device (acc_device_X))
 	  b[ii] = a[ii] + 1;
 	else
 	  b[ii] = a[ii];
@@ -30,7 +43,7 @@ foo (float *a, float *b)
 
     for (ii = 0; ii < N; ii++)
       {
-	if (acc_on_device (acc_device_host))
+	if (acc_on_device (acc_device_X))
 	  b[ii] = a[ii] + 2;
 	else
 	  b[ii] = a[ii];
-- 
1.9.1

Re: [gomp4, committed] Add goacc/kernels-acc-on-device.c

[Tom de Vries]

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

On 12/10/15 14:52, Tom de Vries wrote:
> On 12/10/15 12:49, Thomas Schwinge wrote:
>> Hi Tom!
>>
>> On Sat, 10 Oct 2015 12:49:01 +0200, Tom de
>> Vries<Tom_deVries@mentor.com>  wrote:
>>> >--- /dev/null
>>> >+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
>>> >@@ -0,0 +1,39 @@
>>> >+/* { dg-additional-options "-O2" } */
>>> >+
>>> >+#include <openacc.h>
>
> Hi Thomas,
>
>> That doesn't work (at least in build-tree testing), as gcc/testsuite/ is
>> not set up to look for header files in [target]/libgomp/:
>>
>> [...]/source-gcc/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c:3:21:
>> fatal error: openacc.h: No such file or directory
>>      compilation terminated.
>>      compiler exited with status 1
>>
>
> Ah, I see. I was doing 'make' followed by 'make install', and then
> build-tree testing. The build-tree testing seems to pick up the header
> file from the install directory. So for me test passed.
>
>>> >+
>>> >+#define N 32
>>> >+
>>> >+void
>>> >+foo (float *a, float *b)
>>> >+{
>>> >+  float exp;
>>> >+  int i;
>>> >+  int n;
>>> >+
>>> >+#pragma acc kernels copyin(a[0:N]) copyout(b[0:N])
>>> >+  {
>>> >+    int ii;
>>> >+
>>> >+    for (ii = 0; ii < N; ii++)
>>> >+      {
>>> >+    if (acc_on_device (acc_device_host))
>> Your two options are: if that's applicable/sufficient for what you intend
>> to test here, use __builtin_acc_on_device with a hard-coded acc_device_*,
>> or duplicate part of <openacc.h> as done for example in
>> gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c.
>>
>
> Went with second option, committed as attached.

As a follow-up patch, I've factored the code into a mockup openacc.h, 
now shared by several test-cases.

Committed to gomp-4_0-branch.

Thanks,
- Tom

[-- Attachment #2: 0001-Factor-out-goacc-openacc.h.patch --]
[-- Type: text/x-patch, Size: 3405 bytes --]

Factor out goacc/openacc.h

2015-10-13  Tom de Vries  <tom@codesourcery.com>

	* c-c++-common/goacc/openacc.h: New header file, factored out of ...
	* c-c++-common/goacc/kernels-acc-on-device.c: ... here.
	* c-c++-common/goacc/acc_on_device-2-off.c: Use openacc.h.
	* c-c++-common/goacc/acc_on_device-2.c: Same.
---
 .../c-c++-common/goacc/acc_on_device-2-off.c          | 11 +----------
 gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c    | 13 +------------
 .../c-c++-common/goacc/kernels-acc-on-device.c        | 19 +------------------
 gcc/testsuite/c-c++-common/goacc/openacc.h            | 18 ++++++++++++++++++
 4 files changed, 21 insertions(+), 40 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/openacc.h

diff --git a/gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c b/gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c
index 71abe11..cce58de 100644
--- a/gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c
+++ b/gcc/testsuite/c-c++-common/goacc/acc_on_device-2-off.c
@@ -3,16 +3,7 @@
 
 /* Duplicate parts of libgomp/openacc.h, because we can't include it here.  */
 
-#if __cplusplus
-extern "C" {
-#endif
-
-typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
-extern int acc_on_device (int);
-
-#if __cplusplus
-}
-#endif
+#include "openacc.h"
 
 int
 f (void)
diff --git a/gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c b/gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c
index 243e562..19a5bd3 100644
--- a/gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c
+++ b/gcc/testsuite/c-c++-common/goacc/acc_on_device-2.c
@@ -1,18 +1,7 @@
 /* Have to enable optimizations, as otherwise builtins won't be expanded.  */
 /* { dg-additional-options "-O -fdump-rtl-expand" } */
 
-/* Duplicate parts of libgomp/openacc.h, because we can't include it here.  */
-
-#if __cplusplus
-extern "C" {
-#endif
-
-typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
-extern int acc_on_device (int);
-
-#if __cplusplus
-}
-#endif
+#include "openacc.h"
 
 int
 f (void)
diff --git a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
index 784c66a..958b65b 100644
--- a/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
@@ -1,23 +1,6 @@
 /* { dg-additional-options "-O2" } */
 
-#if __cplusplus
-extern "C" {
-#endif
-
-#if __cplusplus >= 201103
-# define __GOACC_NOTHROW noexcept
-#elif __cplusplus
-# define __GOACC_NOTHROW throw ()
-#else /* Not C++ */
-# define __GOACC_NOTHROW __attribute__ ((__nothrow__))
-#endif
-
-typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
-int acc_on_device (int) __GOACC_NOTHROW;
-
-#if __cplusplus
-}
-#endif
+#include "openacc.h"
 
 #define N 32
 
diff --git a/gcc/testsuite/c-c++-common/goacc/openacc.h b/gcc/testsuite/c-c++-common/goacc/openacc.h
new file mode 100644
index 0000000..a74a482
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/openacc.h
@@ -0,0 +1,18 @@
+#if __cplusplus
+extern "C" {
+#endif
+
+#if __cplusplus >= 201103
+# define __GOACC_NOTHROW noexcept
+#elif __cplusplus
+# define __GOACC_NOTHROW throw ()
+#else /* Not C++ */
+# define __GOACC_NOTHROW __attribute__ ((__nothrow__))
+#endif
+
+typedef enum acc_device_t { acc_device_X = 123 } acc_device_t;
+int acc_on_device (int) __GOACC_NOTHROW;
+
+#if __cplusplus
+}
+#endif
-- 
1.9.1

Re: [gomp4, committed] Add goacc/kernels-acc-on-device.c

[Thomas Schwinge]

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

Hi Tom!

On Tue, 13 Oct 2015 17:49:21 +0200, Tom de Vries <Tom_deVries@mentor.com> wrote:
> On 12/10/15 14:52, Tom de Vries wrote:
> > On 12/10/15 12:49, Thomas Schwinge wrote:
> >> On Sat, 10 Oct 2015 12:49:01 +0200, Tom de
> >> Vries<Tom_deVries@mentor.com>  wrote:
> >>> >--- /dev/null
> >>> >+++ b/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c
> >>> >@@ -0,0 +1,39 @@
> >>> >+/* { dg-additional-options "-O2" } */
> >>> >+
> >>> >+#include <openacc.h>
> >
> > Hi Thomas,
> >
> >> That doesn't work (at least in build-tree testing), as gcc/testsuite/ is
> >> not set up to look for header files in [target]/libgomp/:
> >>
> >> [...]/source-gcc/gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device.c:3:21:
> >> fatal error: openacc.h: No such file or directory
> >>      compilation terminated.
> >>      compiler exited with status 1

> As a follow-up patch, I've factored the code into a mockup openacc.h, 
> now shared by several test-cases.

> Factor out goacc/openacc.h
> 
> 2015-10-13  Tom de Vries  <tom@codesourcery.com>
> 
> 	* c-c++-common/goacc/openacc.h: New header file, factored out of ...
> 	* c-c++-common/goacc/kernels-acc-on-device.c: ... here.
> 	* c-c++-common/goacc/acc_on_device-2-off.c: Use openacc.h.
> 	* c-c++-common/goacc/acc_on_device-2.c: Same.

As a clean-up, and to move acc_on_device testing to where that OpenACC
Runtime Library routine is actually defined (libgomp's openacc.h), Nathan
has just removed this stub openacc.h header file, and also the test files
listed just above, plus the
gcc/testsuite/c-c++-common/goacc/kernels-acc-on-device-2.c that you added
later, see:
<http://news.gmane.org/find-root.php?message_id=%3C56327E51.1090507%40acm.org%3E>
(gomp-4_0-branch),
<http://news.gmane.org/find-root.php?message_id=%3C5632B856.4050509%40acm.org%3E>
(trunk).  If the kernels tests are still important (they tested ICEs, as
far as I remember), you'll have to re-instantiate these in
libgomp/testsuite/, or using __builtin_acc_on_device in gcc/testsuite/.


Grüße
 Thomas

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