OpenACC-Library-Interoperability

OpenACC-Library-Interoperability

[Vikram Singh]

We have a project that is written in gfortran.

I am trying to see whether we could use OpenACC in our project. I got
gcc5.2 installed  using instructions here.

https://github.com/olcf/OLCFHack15

Now I want to do something similar to what is stated here.

https://gcc.gnu.org/onlinedocs/libgomp/OpenACC-Library-Interoperability.html

More specifically what is stated in secion 8.3. I am trying to exactly
reproduce it using gfortran. Unfortunately I don't see how I can do it
in fortran. In the example,

d_X = acc_copyin(&h_X[0], N * sizeof (float));

This allows d_X to be directly used in s = cublasSaxpy(h, N, &alpha,
d_X, 1, d_Y, 1);

But in fortran, the acc_copyin does not return anything.


So how would I replicate the case in Fortran?

Thanks and Regards,
Vikram

Re: OpenACC-Library-Interoperability

[Thomas Schwinge]

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

Hi!

On Tue, 12 Apr 2016 18:00:02 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
> We have a project that is written in gfortran.
> 
> I am trying to see whether we could use OpenACC in our project.

Thanks for your interest!

> I got
> gcc5.2 installed

I will just note that the upcoming GCC 6 releases will offer much
extended support for OpenACC with nvptx offloading.  (If you're building
GCC from sources anyway, you may want to give the current GCC trunk or
gomp-4_0-branch a try, after resolving the problem discussed below.)

> using instructions here.
> 
> https://github.com/olcf/OLCFHack15

Nice.  

> Now I want to do something similar to what is stated here.
> 
> https://gcc.gnu.org/onlinedocs/libgomp/OpenACC-Library-Interoperability.html
> 
> More specifically what is stated in secion 8.3. I am trying to exactly
> reproduce it using gfortran. Unfortunately I don't see how I can do it
> in fortran. In the example,
> 
> d_X = acc_copyin(&h_X[0], N * sizeof (float));
> 
> This allows d_X to be directly used in s = cublasSaxpy(h, N, &alpha,
> d_X, 1, d_Y, 1);
> 
> But in fortran, the acc_copyin does not return anything.

Correct.  I'm not much of a Fortran programmer, but as far as I remember,
Fortran doesn't have a notion of a pointer as C/C++ have.  I suppose
that's why the OpenACC specification defines that in Fortran, acc_copyin
et al. are subroutines (no return value) instead of functions, and why
acc_malloc, acc_deviceptr, acc_map et al. don't even exist in Fortran.

> So how would I replicate the case in Fortran?

You would use the OpenACC host_data construct.  So, instead of:

    [...]
    d_X = acc_copyin(&h_X[0], N * sizeof (float));
    [...]
    s = cublasSaxpy(h, N, &alpha, d_X, 1, d_Y, 1);
    [...]

... you would do something like (notice X instead of d_X and h_X)
(untested):

    [...]
    acc_copyin(X, N * sizeof (float));
    [...]
    #pragma acc host_data use_device(X)
    {
      s = cublasSaxpy(h, N, &alpha, X, 1, d_Y, 1);
    }
    [...]

I'm copying Jim, who is the author of this chapter in the documentation
as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
test cases, and much of the relevant libgomp code, too, and who should
please correct me if I'm wrong.  I'll make a note for later, that we
should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
Fortran, and also replicate them using the OpenACC host_data construct
(like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
the same for the documentation you referenced.  (Vikram, please tell if
you're interested to work on these items.)


..., and now that've written all this, I think we have a show-stopper
here: as far as I remember, the OpenACC host_data construct is not yet
supported in the GCC 5 releases, and it's currently broken for Fortran,
too: <https://gcc.gnu.org/PR70598> (Chung-Lin if working on that).


Grüße
 Thomas

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

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Hi Thomas.

Thanks for the reply. Yes, I came to the conclusion that host_data
would be the only way to do it in fortran.

On the other hand, I though there were no plans to implement it in
gfortran 6 either, so I pivoted back to using PGI compilers which I
had originally not wanted to do.

> I'm copying Jim, who is the author of this chapter in the documentation
> as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
> test cases, and much of the relevant libgomp code, too, and who should
> please correct me if I'm wrong.  I'll make a note for later, that we
> should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
> Fortran, and also replicate them using the OpenACC host_data construct
> (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
> the same for the documentation you referenced.  (Vikram, please tell if
> you're interested to work on these items.)

I am not sure I understand what you want me to work on exactly. I am
not really much of a C programmer, so I wouldn't be good at it. I
would still like to help.


Thanks and Regards,
Vikram



On Fri, Apr 15, 2016 at 9:16 AM, Thomas Schwinge
<thomas@codesourcery.com> wrote:
> Hi!
>
> On Tue, 12 Apr 2016 18:00:02 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
>> We have a project that is written in gfortran.
>>
>> I am trying to see whether we could use OpenACC in our project.
>
> Thanks for your interest!
>
>> I got
>> gcc5.2 installed
>
> I will just note that the upcoming GCC 6 releases will offer much
> extended support for OpenACC with nvptx offloading.  (If you're building
> GCC from sources anyway, you may want to give the current GCC trunk or
> gomp-4_0-branch a try, after resolving the problem discussed below.)
>
>> using instructions here.
>>
>> https://github.com/olcf/OLCFHack15
>
> Nice.
>
>> Now I want to do something similar to what is stated here.
>>
>> https://gcc.gnu.org/onlinedocs/libgomp/OpenACC-Library-Interoperability.html
>>
>> More specifically what is stated in secion 8.3. I am trying to exactly
>> reproduce it using gfortran. Unfortunately I don't see how I can do it
>> in fortran. In the example,
>>
>> d_X = acc_copyin(&h_X[0], N * sizeof (float));
>>
>> This allows d_X to be directly used in s = cublasSaxpy(h, N, &alpha,
>> d_X, 1, d_Y, 1);
>>
>> But in fortran, the acc_copyin does not return anything.
>
> Correct.  I'm not much of a Fortran programmer, but as far as I remember,
> Fortran doesn't have a notion of a pointer as C/C++ have.  I suppose
> that's why the OpenACC specification defines that in Fortran, acc_copyin
> et al. are subroutines (no return value) instead of functions, and why
> acc_malloc, acc_deviceptr, acc_map et al. don't even exist in Fortran.
>
>> So how would I replicate the case in Fortran?
>
> You would use the OpenACC host_data construct.  So, instead of:
>
>     [...]
>     d_X = acc_copyin(&h_X[0], N * sizeof (float));
>     [...]
>     s = cublasSaxpy(h, N, &alpha, d_X, 1, d_Y, 1);
>     [...]
>
> ... you would do something like (notice X instead of d_X and h_X)
> (untested):
>
>     [...]
>     acc_copyin(X, N * sizeof (float));
>     [...]
>     #pragma acc host_data use_device(X)
>     {
>       s = cublasSaxpy(h, N, &alpha, X, 1, d_Y, 1);
>     }
>     [...]
>
> I'm copying Jim, who is the author of this chapter in the documentation
> as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
> test cases, and much of the relevant libgomp code, too, and who should
> please correct me if I'm wrong.  I'll make a note for later, that we
> should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
> Fortran, and also replicate them using the OpenACC host_data construct
> (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
> the same for the documentation you referenced.  (Vikram, please tell if
> you're interested to work on these items.)
>
>
> ..., and now that've written all this, I think we have a show-stopper
> here: as far as I remember, the OpenACC host_data construct is not yet
> supported in the GCC 5 releases, and it's currently broken for Fortran,
> too: <https://gcc.gnu.org/PR70598> (Chung-Lin if working on that).
>
>
> Grüße
>  Thomas

Re: OpenACC-Library-Interoperability

[Thomas Schwinge]

Hi Vikram!

On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
> Yes, I came to the conclusion that host_data
> would be the only way to do it in fortran.
> 
> On the other hand, I though there were no plans to implement it in
> gfortran 6 either

I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
fixed in time for the GCC 6.1 release.  I'll keep you posted.


> > I'm copying Jim, who is the author of this chapter in the documentation
> > as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
> > test cases, and much of the relevant libgomp code, too, and who should
> > please correct me if I'm wrong.  I'll make a note for later, that we
> > should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
> > Fortran, and also replicate them using the OpenACC host_data construct
> > (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
> > the same for the documentation you referenced.  (Vikram, please tell if
> > you're interested to work on these items.)
> 
> I am not sure I understand what you want me to work on exactly. I am
> not really much of a C programmer, so I wouldn't be good at it. I
> would still like to help.

Sorry for being unclear.  My idea/question has been whether you're
interested in helping by translating the documentation as well as the
libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran (using
the OpenACC host_data construct instead of the acc_* functions).  If yes,
then that's great, if not, then one of us will do it at some point.


Grüße
 Thomas

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Hi Thomas,

Thanks for the clarification.

I checked libgomp.oacc-c-c++-common/context-*.c, and they seem to be
exactly what I needed to start with.

If I understand it correctly, I'll anyway be trying to implement
something similar with PGI compilers, so I'll have something as a
starting point.

But again, it will need the OpenACC host_data construct to be setup
for gfortran to test. So yeah, I think I can help.

Regards,
Vikram



On Fri, Apr 15, 2016 at 11:57 AM, Thomas Schwinge
<thomas@codesourcery.com> wrote:
> Hi Vikram!
>
> On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
>> Yes, I came to the conclusion that host_data
>> would be the only way to do it in fortran.
>>
>> On the other hand, I though there were no plans to implement it in
>> gfortran 6 either
>
> I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
> fixed in time for the GCC 6.1 release.  I'll keep you posted.
>
>
>> > I'm copying Jim, who is the author of this chapter in the documentation
>> > as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
>> > test cases, and much of the relevant libgomp code, too, and who should
>> > please correct me if I'm wrong.  I'll make a note for later, that we
>> > should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
>> > Fortran, and also replicate them using the OpenACC host_data construct
>> > (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
>> > the same for the documentation you referenced.  (Vikram, please tell if
>> > you're interested to work on these items.)
>>
>> I am not sure I understand what you want me to work on exactly. I am
>> not really much of a C programmer, so I wouldn't be good at it. I
>> would still like to help.
>
> Sorry for being unclear.  My idea/question has been whether you're
> interested in helping by translating the documentation as well as the
> libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran (using
> the OpenACC host_data construct instead of the acc_* functions).  If yes,
> then that's great, if not, then one of us will do it at some point.
>
>
> Grüße
>  Thomas

Re: OpenACC-Library-Interoperabilit

[Salvatore Filippone]

You can count me in as well for testing, provided you can help me in
getting the compilation process to complete; I tried the scripts on
the gfortran openACC wiki page and could not get them to complete...

Re: OpenACC-Library-Interoperability

[Vladimír Fuka]

Hi!

 > Correct. I'm not much of a Fortran programmer, but as far as I remember,
 >Fortran doesn't have a notion of a pointer as C/C++ have.  I suppose
 >that's why the OpenACC specification defines that in Fortran, acc_copyin
 >et al. are subroutines (no return value) instead of functions, and why
 >acc_malloc, acc_deviceptr, acc_map et al. don't even exist in Fortran.

 I already commented on this when Vikram asked about this issue at
stackoverflow: http://stackoverflow.com/questions/36577628/openacc-library-interoperability-how-to-get-device-pointer


 In my opinion it is a clear defficiency of OpenACC and very likely caused due
 to wrong information similar to "Fortran doesn't have a notion of a
pointer as C/C++ have.". Fortran 2003
 has intrinsic derived type `c_ptr` which is interoperable with `void
*`. I don't know if
 the C compiler does any magic besides returning or using the device
pointer as a `void *`, but if not I can't see
 a reason why it can't be implemented in Fortran. That's why I also
suggested to Vikram to try to
 create a Fortran interface to the C API functions himself. I think it
is likely to work but I don't know if he has tried it.

   Vladimir

[PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE

[Chung-Lin Tang]

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

Hi, this patch resolves an ICE for Fortran when using the OpenACC
host_data directive.  Actually, rather than say resolve, it's more like
adjusting the front-end to same middle-end restrictions as C/C++,
namely that we only support pointers or arrays for host_data right now.

This patch contains a little bit of adjustments in fortran/openmp.c:resolve_omp_clauses(),
and some testcase adjustments. This has been tested without regressions for Fortran.

Is this okay for trunk?

Thanks,
Chung-Lin

2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>

	gcc/
	* fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
	handling to only allow pointers and arrays.

	gcc/testsuite/
	* gfortran.dg/goacc/host_data-tree.f95: Adjust to use accept pointers in use_device clause.
	* gfortran.dg/goacc/uninit-use-device-clause.f95: Likewise.
	* gfortran.dg/goacc/list.f95: Adjust to catch "neither a pointer nor an array" error messages.

	libgomp/testsuite/
	* libgomp.oacc-fortran/host_data-1.f90: New testcase.

[-- Attachment #2: pr70598.patch --]
[-- Type: text/x-patch, Size: 5694 bytes --]

Index: gcc/fortran/openmp.c
===================================================================
--- gcc/fortran/openmp.c	(revision 236020)
+++ gcc/fortran/openmp.c	(working copy)
@@ -3743,11 +3743,18 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
 			      && CLASS_DATA (n->sym)->attr.allocatable))
 			gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
 				   n->sym->name, name, &n->where);
-		      if (n->sym->attr.pointer
-			  || (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
-			      && CLASS_DATA (n->sym)->attr.class_pointer))
-			gfc_error ("POINTER object %qs in %s clause at %L",
-				   n->sym->name, name, &n->where);
+		      if (n->sym->attr.flavor == FL_VARIABLE
+			  && !n->sym->as && !n->sym->attr.pointer
+			  && !n->sym->attr.cray_pointer
+			  && !n->sym->attr.cray_pointee)
+			gfc_error ("%s clause variable %qs at %L is neither "
+				   "a pointer nor an array", name,
+				   n->sym->name, &n->where);
+		      if (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
+			  && CLASS_DATA (n->sym)->attr.class_pointer)
+			gfc_error ("POINTER object %qs of polymorphic type in "
+				   "%s clause at %L", n->sym->name, name,
+				   &n->where);
 		      if (n->sym->attr.cray_pointer)
 			gfc_error ("Cray pointer object %qs in %s clause at %L",
 				   n->sym->name, name, &n->where);
Index: gcc/testsuite/gfortran.dg/goacc/uninit-use-device-clause.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/uninit-use-device-clause.f95	(revision 236020)
+++ gcc/testsuite/gfortran.dg/goacc/uninit-use-device-clause.f95	(working copy)
@@ -2,9 +2,9 @@
 ! { dg-additional-options "-Wuninitialized" }
 
 subroutine test
-  integer :: i
+  integer, pointer :: p
 
-  !$acc host_data use_device(i) ! { dg-warning "is used uninitialized in this function" }
+  !$acc host_data use_device(p) ! { dg-warning "is used uninitialized in this function" }
   !$acc end host_data
 end subroutine test
 
Index: gcc/testsuite/gfortran.dg/goacc/list.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/list.f95	(revision 236020)
+++ gcc/testsuite/gfortran.dg/goacc/list.f95	(working copy)
@@ -76,19 +76,19 @@ program test
   !$acc parallel private (i) firstprivate (i) ! { dg-error "present on multiple clauses" }
   !$acc end parallel
 
-  !$acc host_data use_device(i)
+  !$acc host_data use_device(i) ! { dg-error "neither a pointer nor an array" }
   !$acc end host_data
 
-  !$acc host_data use_device(c, d)
+  !$acc host_data use_device(c, d) ! { dg-error "neither a pointer nor an array" }
   !$acc end host_data
 
   !$acc host_data use_device(a)
   !$acc end host_data
 
-  !$acc host_data use_device(i, j, k, l, a)
+  !$acc host_data use_device(i, j, k, l, a) ! { dg-error "neither a pointer nor an array" }
   !$acc end host_data  
 
-  !$acc host_data use_device (i) use_device (j)
+  !$acc host_data use_device (i) use_device (j) ! { dg-error "neither a pointer nor an array" }
   !$acc end host_data
 
   !$acc host_data use_device ! { dg-error "Unclassifiable OpenACC directive" }
@@ -99,13 +99,17 @@ program test
 
   !$acc host_data use_device(10) ! { dg-error "Syntax error" }
 
-  !$acc host_data use_device(/b/, /b/) ! { dg-error "present on multiple clauses" }
+  !$acc host_data use_device(/b/, /b/)
   !$acc end host_data
+  ! { dg-error "neither a pointer nor an array" "" { target *-*-* } 102 }
+  ! { dg-error "present on multiple clauses" "" { target *-*-* } 102 }
 
-  !$acc host_data use_device(i, j, i) ! { dg-error "present on multiple clauses" }
+  !$acc host_data use_device(i, j, i)
   !$acc end host_data
+  ! { dg-error "neither a pointer nor an array" "" { target *-*-* } 107 }
+  ! { dg-error "present on multiple clauses" "" { target *-*-* } 107 }
 
-  !$acc host_data use_device(p1) ! { dg-error "POINTER" }
+  !$acc host_data use_device(p1)
   !$acc end host_data
 
 end program test
Index: gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95
===================================================================
--- gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95	(revision 236020)
+++ gcc/testsuite/gfortran.dg/goacc/host_data-tree.f95	(working copy)
@@ -3,9 +3,9 @@
 
 program test
   implicit none
-  integer :: i = 1
+  integer, pointer :: p
 
-  !$acc host_data use_device(i)
+  !$acc host_data use_device(p)
   !$acc end host_data
 end program test
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(i\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 1 "original" } }
Index: libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90
===================================================================
--- libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90	(revision 0)
+++ libgomp/testsuite/libgomp.oacc-fortran/host_data-1.f90	(revision 0)
@@ -0,0 +1,35 @@
+! { dg-do run }
+! { dg-additional-options "-cpp" }
+
+! { dg-xfail-if "TODO" { *-*-* } }
+! { dg-excess-errors "TODO" }
+
+program test
+  implicit none
+
+  integer, target :: i, arr(1000)
+  integer, pointer :: ip, iph
+  integer, contiguous, pointer :: parr(:), parrh(:)
+
+  ! Assign the same targets
+  ip => i
+  parr => arr
+  iph => i
+  parrh => arr
+
+  !$acc data copyin(i, arr)
+  !$acc host_data use_device(ip, parr)
+
+  ! Test how the pointers compare inside a host_data construct
+#if ACC_MEM_SHARED
+  if (.not. associated(ip, iph)) call abort
+  if (.not. associated(parr, parrh)) call abort
+#else
+  if (associated(ip, iph)) call abort
+  if (associated(parr, parrh)) call abort
+#endif
+
+  !$acc end host_data
+  !$acc end data
+
+end program test

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE

[Bernhard Reutner-Fischer]

On May 9, 2016 4:26:50 PM GMT+02:00, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>Hi, this patch resolves an ICE for Fortran when using the OpenACC
>host_data directive.  Actually, rather than say resolve, it's more like
>adjusting the front-end to same middle-end restrictions as C/C++,
>namely that we only support pointers or arrays for host_data right now.
>
>This patch contains a little bit of adjustments in
>fortran/openmp.c:resolve_omp_clauses(),
>and some testcase adjustments. This has been tested without regressions
>for Fortran.
>
>Is this okay for trunk?
>
>Thanks,
>Chung-Lin
>
>2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
>
>	gcc/
>	* fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
>	handling to only allow pointers and arrays.

Fortran has it's own ChangeLog. The patch itself looks somewhat plausible to me, fwiw, but Jakub or a FE maintainer has the say.

Re: OpenACC-Library-Interoperability

[Thomas Schwinge]

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

Hi!

On Fri, 15 Apr 2016 13:59:53 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
> I checked libgomp.oacc-c-c++-common/context-*.c, and they seem to be
> exactly what I needed to start with.
> 
> If I understand it correctly, I'll anyway be trying to implement
> something similar with PGI compilers, so I'll have something as a
> starting point.

Any progress on that already?  I had a very quick look myself, but it's
not as easy as I thought...  A "courageous" use of "use cublas" in
Fortran code compiled with gfortran results in: "Fatal Error: Can't open
module file 'cublas.mod' for reading".  The problem is: Fortran
interfacing to C libraries (which cuBLAS is).
<http://docs.nvidia.com/cuda/cublas/index.html#appendix-b-cublas-fortran-bindings>
has some instructions how to do it.  Anyone got that to work already?

> But again, it will need the OpenACC host_data construct to be setup
> for gfortran to test.

In
<http://news.gmane.org/find-root.php?message_id=%3C2b4f59d5-be38-2814-27bb-73aa7ffb4b8f%40codesourcery.com%3E>,
Chung-Lin has now posted a patch (pending review) that should make the
OpenACC host_data construct usable in GCC Fortran.  (Problem discussed in
<http://news.gmane.org/find-root.php?message_id=%3C878u0o6wwj.fsf%40kepler.schwinge.homeip.net%3E>
before.)

For reference:

> On Fri, Apr 15, 2016 at 11:57 AM, Thomas Schwinge
> <thomas@codesourcery.com> wrote:
> > On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
> >> Yes, I came to the conclusion that host_data
> >> would be the only way to do it in fortran.
> >>
> >> On the other hand, I though there were no plans to implement it in
> >> gfortran 6 either
> >
> > I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
> > fixed in time for the GCC 6.1 release.  I'll keep you posted.
> >
> >
> >> > I'm copying Jim, who is the author of this chapter in the documentation
> >> > as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
> >> > test cases, and much of the relevant libgomp code, too, and who should
> >> > please correct me if I'm wrong.  I'll make a note for later, that we
> >> > should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
> >> > Fortran, and also replicate them using the OpenACC host_data construct
> >> > (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
> >> > the same for the documentation you referenced.  (Vikram, please tell if
> >> > you're interested to work on these items.)
> >>
> >> I am not sure I understand what you want me to work on exactly. I am
> >> not really much of a C programmer, so I wouldn't be good at it. I
> >> would still like to help.
> >
> > Sorry for being unclear.  My idea/question has been whether you're
> > interested in helping by translating the documentation as well as the
> > libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran (using
> > the OpenACC host_data construct instead of the acc_* functions).  If yes,
> > then that's great, if not, then one of us will do it at some point.


Grüße
 Thomas

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

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Well, I have already written code that I had originally wanted to do,
that is, use openacc with cublas.

And it works perfectly well with PGI compiler. I guess I will have to
check on how I would go about trying to run it using gfortran now that
host_data works using gfortran.

On Thu, May 12, 2016 at 6:34 PM, Thomas Schwinge
<thomas@codesourcery.com> wrote:
> Hi!
>
> On Fri, 15 Apr 2016 13:59:53 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
>> I checked libgomp.oacc-c-c++-common/context-*.c, and they seem to be
>> exactly what I needed to start with.
>>
>> If I understand it correctly, I'll anyway be trying to implement
>> something similar with PGI compilers, so I'll have something as a
>> starting point.
>
> Any progress on that already?  I had a very quick look myself, but it's
> not as easy as I thought...  A "courageous" use of "use cublas" in
> Fortran code compiled with gfortran results in: "Fatal Error: Can't open
> module file 'cublas.mod' for reading".  The problem is: Fortran
> interfacing to C libraries (which cuBLAS is).
> <http://docs.nvidia.com/cuda/cublas/index.html#appendix-b-cublas-fortran-bindings>
> has some instructions how to do it.  Anyone got that to work already?
>
>> But again, it will need the OpenACC host_data construct to be setup
>> for gfortran to test.
>
> In
> <http://news.gmane.org/find-root.php?message_id=%3C2b4f59d5-be38-2814-27bb-73aa7ffb4b8f%40codesourcery.com%3E>,
> Chung-Lin has now posted a patch (pending review) that should make the
> OpenACC host_data construct usable in GCC Fortran.  (Problem discussed in
> <http://news.gmane.org/find-root.php?message_id=%3C878u0o6wwj.fsf%40kepler.schwinge.homeip.net%3E>
> before.)
>
> For reference:
>
>> On Fri, Apr 15, 2016 at 11:57 AM, Thomas Schwinge
>> <thomas@codesourcery.com> wrote:
>> > On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
>> >> Yes, I came to the conclusion that host_data
>> >> would be the only way to do it in fortran.
>> >>
>> >> On the other hand, I though there were no plans to implement it in
>> >> gfortran 6 either
>> >
>> > I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
>> > fixed in time for the GCC 6.1 release.  I'll keep you posted.
>> >
>> >
>> >> > I'm copying Jim, who is the author of this chapter in the documentation
>> >> > as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
>> >> > test cases, and much of the relevant libgomp code, too, and who should
>> >> > please correct me if I'm wrong.  I'll make a note for later, that we
>> >> > should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
>> >> > Fortran, and also replicate them using the OpenACC host_data construct
>> >> > (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
>> >> > the same for the documentation you referenced.  (Vikram, please tell if
>> >> > you're interested to work on these items.)
>> >>
>> >> I am not sure I understand what you want me to work on exactly. I am
>> >> not really much of a C programmer, so I wouldn't be good at it. I
>> >> would still like to help.
>> >
>> > Sorry for being unclear.  My idea/question has been whether you're
>> > interested in helping by translating the documentation as well as the
>> > libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran (using
>> > the OpenACC host_data construct instead of the acc_* functions).  If yes,
>> > then that's great, if not, then one of us will do it at some point.
>
>
> Grüße
>  Thomas

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Hi.

How would I go about getting this patched version of gfortran?



On Thu, May 12, 2016 at 7:40 PM, Vikram Singh <vikramsingh001@gmail.com> wrote:
> Well, I have already written code that I had originally wanted to do,
> that is, use openacc with cublas.
>
> And it works perfectly well with PGI compiler. I guess I will have to
> check on how I would go about trying to run it using gfortran now that
> host_data works using gfortran.
>
> On Thu, May 12, 2016 at 6:34 PM, Thomas Schwinge
> <thomas@codesourcery.com> wrote:
>> Hi!
>>
>> On Fri, 15 Apr 2016 13:59:53 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
>>> I checked libgomp.oacc-c-c++-common/context-*.c, and they seem to be
>>> exactly what I needed to start with.
>>>
>>> If I understand it correctly, I'll anyway be trying to implement
>>> something similar with PGI compilers, so I'll have something as a
>>> starting point.
>>
>> Any progress on that already?  I had a very quick look myself, but it's
>> not as easy as I thought...  A "courageous" use of "use cublas" in
>> Fortran code compiled with gfortran results in: "Fatal Error: Can't open
>> module file 'cublas.mod' for reading".  The problem is: Fortran
>> interfacing to C libraries (which cuBLAS is).
>> <http://docs.nvidia.com/cuda/cublas/index.html#appendix-b-cublas-fortran-bindings>
>> has some instructions how to do it.  Anyone got that to work already?
>>
>>> But again, it will need the OpenACC host_data construct to be setup
>>> for gfortran to test.
>>
>> In
>> <http://news.gmane.org/find-root.php?message_id=%3C2b4f59d5-be38-2814-27bb-73aa7ffb4b8f%40codesourcery.com%3E>,
>> Chung-Lin has now posted a patch (pending review) that should make the
>> OpenACC host_data construct usable in GCC Fortran.  (Problem discussed in
>> <http://news.gmane.org/find-root.php?message_id=%3C878u0o6wwj.fsf%40kepler.schwinge.homeip.net%3E>
>> before.)
>>
>> For reference:
>>
>>> On Fri, Apr 15, 2016 at 11:57 AM, Thomas Schwinge
>>> <thomas@codesourcery.com> wrote:
>>> > On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh <vikramsingh001@gmail.com> wrote:
>>> >> Yes, I came to the conclusion that host_data
>>> >> would be the only way to do it in fortran.
>>> >>
>>> >> On the other hand, I though there were no plans to implement it in
>>> >> gfortran 6 either
>>> >
>>> > I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
>>> > fixed in time for the GCC 6.1 release.  I'll keep you posted.
>>> >
>>> >
>>> >> > I'm copying Jim, who is the author of this chapter in the documentation
>>> >> > as well as the libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
>>> >> > test cases, and much of the relevant libgomp code, too, and who should
>>> >> > please correct me if I'm wrong.  I'll make a note for later, that we
>>> >> > should translate the libgomp.oacc-c-c++-common/context-*.c test cases to
>>> >> > Fortran, and also replicate them using the OpenACC host_data construct
>>> >> > (like in libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
>>> >> > the same for the documentation you referenced.  (Vikram, please tell if
>>> >> > you're interested to work on these items.)
>>> >>
>>> >> I am not sure I understand what you want me to work on exactly. I am
>>> >> not really much of a C programmer, so I wouldn't be good at it. I
>>> >> would still like to help.
>>> >
>>> > Sorry for being unclear.  My idea/question has been whether you're
>>> > interested in helping by translating the documentation as well as the
>>> > libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran (using
>>> > the OpenACC host_data construct instead of the acc_* functions).  If yes,
>>> > then that's great, if not, then one of us will do it at some point.
>>
>>
>> Grüße
>>  Thomas

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE

[Chung-Lin Tang]

Ping.

On 2016/5/11 02:57 AM, Bernhard Reutner-Fischer wrote:
> On May 9, 2016 4:26:50 PM GMT+02:00, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>> Hi, this patch resolves an ICE for Fortran when using the OpenACC
>> host_data directive.  Actually, rather than say resolve, it's more like
>> adjusting the front-end to same middle-end restrictions as C/C++,
>> namely that we only support pointers or arrays for host_data right now.
>>
>> This patch contains a little bit of adjustments in
>> fortran/openmp.c:resolve_omp_clauses(),
>> and some testcase adjustments. This has been tested without regressions
>> for Fortran.
>>
>> Is this okay for trunk?
>>
>> Thanks,
>> Chung-Lin
>>
>> 2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
>>
>> 	gcc/
>> 	* fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
>> 	handling to only allow pointers and arrays.
> 
> Fortran has it's own ChangeLog. The patch itself looks somewhat plausible to me, fwiw, but Jakub or a FE maintainer has the say.
> 

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE (ping x2)

[Chung-Lin Tang]

Ping x2

On 2016/6/7 08:03 PM, Chung-Lin Tang wrote:
> Ping.
> 
> On 2016/5/11 02:57 AM, Bernhard Reutner-Fischer wrote:
>> On May 9, 2016 4:26:50 PM GMT+02:00, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>>> Hi, this patch resolves an ICE for Fortran when using the OpenACC
>>> host_data directive.  Actually, rather than say resolve, it's more like
>>> adjusting the front-end to same middle-end restrictions as C/C++,
>>> namely that we only support pointers or arrays for host_data right now.
>>>
>>> This patch contains a little bit of adjustments in
>>> fortran/openmp.c:resolve_omp_clauses(),
>>> and some testcase adjustments. This has been tested without regressions
>>> for Fortran.
>>>
>>> Is this okay for trunk?
>>>
>>> Thanks,
>>> Chung-Lin
>>>
>>> 2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
>>>
>>> 	gcc/
>>> 	* fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
>>> 	handling to only allow pointers and arrays.
>>
>> Fortran has it's own ChangeLog. The patch itself looks somewhat plausible to me, fwiw, but Jakub or a FE maintainer has the say.
>>
> 

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE (ping x3)

[Chung-Lin Tang]

Ping x3

On 06/21/2016 02:18 PM, Chung-Lin Tang wrote:
> Ping x2
>
> On 2016/6/7 08:03 PM, Chung-Lin Tang wrote:
>> Ping.
>>
>> On 2016/5/11 02:57 AM, Bernhard Reutner-Fischer wrote:
>>> On May 9, 2016 4:26:50 PM GMT+02:00, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>>>> Hi, this patch resolves an ICE for Fortran when using the OpenACC
>>>> host_data directive.  Actually, rather than say resolve, it's more like
>>>> adjusting the front-end to same middle-end restrictions as C/C++,
>>>> namely that we only support pointers or arrays for host_data right now.
>>>>
>>>> This patch contains a little bit of adjustments in
>>>> fortran/openmp.c:resolve_omp_clauses(),
>>>> and some testcase adjustments. This has been tested without regressions
>>>> for Fortran.
>>>>
>>>> Is this okay for trunk?
>>>>
>>>> Thanks,
>>>> Chung-Lin
>>>>
>>>> 2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
>>>>
>>>> 	gcc/
>>>> 	* fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
>>>> 	handling to only allow pointers and arrays.
>>>
>>> Fortran has it's own ChangeLog. The patch itself looks somewhat plausible to me, fwiw, but Jakub or a FE maintainer has the say.
>>>
>>
>

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE (ping x4)

[Chung-Lin Tang]

Ping x4

On 2016/7/13 7:52 PM, Chung-Lin Tang wrote:
> Ping x3
> 
> On 06/21/2016 02:18 PM, Chung-Lin Tang wrote:
>> Ping x2
>>
>> On 2016/6/7 08:03 PM, Chung-Lin Tang wrote:
>>> Ping.
>>>
>>> On 2016/5/11 02:57 AM, Bernhard Reutner-Fischer wrote:
>>>> On May 9, 2016 4:26:50 PM GMT+02:00, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>>>>> Hi, this patch resolves an ICE for Fortran when using the OpenACC
>>>>> host_data directive.  Actually, rather than say resolve, it's more like
>>>>> adjusting the front-end to same middle-end restrictions as C/C++,
>>>>> namely that we only support pointers or arrays for host_data right now.
>>>>>
>>>>> This patch contains a little bit of adjustments in
>>>>> fortran/openmp.c:resolve_omp_clauses(),
>>>>> and some testcase adjustments. This has been tested without regressions
>>>>> for Fortran.
>>>>>
>>>>> Is this okay for trunk?
>>>>>
>>>>> Thanks,
>>>>> Chung-Lin
>>>>>
>>>>> 2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
>>>>>
>>>>>     gcc/
>>>>>     * fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
>>>>>     handling to only allow pointers and arrays.
>>>>
>>>> Fortran has it's own ChangeLog. The patch itself looks somewhat plausible to me, fwiw, but Jakub or a FE maintainer has the
>>>> say.
>>>>
>>>
>>

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE (ping x4)

[Paul Richard Thomas]

Hi Chung-Lin,

I was ignoring your patch on the grounds that one of the omp gurus
should deal with it.  That says, it looks OK for trunk to me.

I presume that fortran/host_data-1.f90 should have the XFAIL removed?

Cheers

Paul

On 21 July 2016 at 11:28, Chung-Lin Tang <cltang@codesourcery.com> wrote:
> Ping x4
>
> On 2016/7/13 7:52 PM, Chung-Lin Tang wrote:
>> Ping x3
>>
>> On 06/21/2016 02:18 PM, Chung-Lin Tang wrote:
>>> Ping x2
>>>
>>> On 2016/6/7 08:03 PM, Chung-Lin Tang wrote:
>>>> Ping.
>>>>
>>>> On 2016/5/11 02:57 AM, Bernhard Reutner-Fischer wrote:
>>>>> On May 9, 2016 4:26:50 PM GMT+02:00, Chung-Lin Tang <cltang@codesourcery.com> wrote:
>>>>>> Hi, this patch resolves an ICE for Fortran when using the OpenACC
>>>>>> host_data directive.  Actually, rather than say resolve, it's more like
>>>>>> adjusting the front-end to same middle-end restrictions as C/C++,
>>>>>> namely that we only support pointers or arrays for host_data right now.
>>>>>>
>>>>>> This patch contains a little bit of adjustments in
>>>>>> fortran/openmp.c:resolve_omp_clauses(),
>>>>>> and some testcase adjustments. This has been tested without regressions
>>>>>> for Fortran.
>>>>>>
>>>>>> Is this okay for trunk?
>>>>>>
>>>>>> Thanks,
>>>>>> Chung-Lin
>>>>>>
>>>>>> 2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
>>>>>>
>>>>>>     gcc/
>>>>>>     * fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
>>>>>>     handling to only allow pointers and arrays.
>>>>>
>>>>> Fortran has it's own ChangeLog. The patch itself looks somewhat plausible to me, fwiw, but Jakub or a FE maintainer has the
>>>>> say.
>>>>>
>>>>
>>>
>



-- 
The difference between genius and stupidity is; genius has its limits.

Albert Einstein

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE

[Jakub Jelinek]

On Mon, May 09, 2016 at 10:26:50PM +0800, Chung-Lin Tang wrote:
> 2015-05-09  Chung-Lin Tang  <cltang@codesourcery.com>
> 
> 	gcc/
> 	* fortran/openmp.c (resolve_omp_clauses): Adjust use_device clause
> 	handling to only allow pointers and arrays.

As has been mentioned earlier, this should go into gcc/fortran/ChangeLog
without fortran/ prefix.

> --- gcc/fortran/openmp.c	(revision 236020)
> +++ gcc/fortran/openmp.c	(working copy)
> @@ -3743,11 +3743,18 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
>  			      && CLASS_DATA (n->sym)->attr.allocatable))
>  			gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
>  				   n->sym->name, name, &n->where);
> -		      if (n->sym->attr.pointer
> -			  || (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
> -			      && CLASS_DATA (n->sym)->attr.class_pointer))
> -			gfc_error ("POINTER object %qs in %s clause at %L",
> -				   n->sym->name, name, &n->where);
> +		      if (n->sym->attr.flavor == FL_VARIABLE
> +			  && !n->sym->as && !n->sym->attr.pointer

Better put every && on a separate line if the whole if (...) doesn't fit
on a single line.

> +			  && !n->sym->attr.cray_pointer
> +			  && !n->sym->attr.cray_pointee)

This is too ugly.  I'd instead move the if after the cray pointer/pointee
tests, i.e.
if (n->sym->attr.cray_pointer)
  gfc_error (...);
else if (n->sym->attr.cray_pointee)
  gfc_error (...);
else if (n->sym->attr.flavor == FL_VARIABLE
	 && !n->sym->as
	 && !n->sym->attr.pointer)
  gfc_error (...);

> +			gfc_error ("%s clause variable %qs at %L is neither "
> +				   "a pointer nor an array", name,
> +				   n->sym->name, &n->where);

Why pointer rather than POINTER ?

> +		      if (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
> +			  && CLASS_DATA (n->sym)->attr.class_pointer)
> +			gfc_error ("POINTER object %qs of polymorphic type in "
> +				   "%s clause at %L", n->sym->name, name,
> +				   &n->where);

And, doesn't this mean you emit 2 errors, the first one because it is not
a pointer nor array, and the second one because it is a class with
attr.class_pointer?  Also put && on the next line.

I think best would be to make it else if after the cray pointer/pointee
checks.

	Jakub

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE

[Chung-Lin Tang]

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

On 2016/7/21 07:13 PM, Jakub Jelinek wrote:
> Better put every && on a separate line if the whole if (...) doesn't fit
> on a single line.
> 
>> > +			  && !n->sym->attr.cray_pointer
>> > +			  && !n->sym->attr.cray_pointee)
> This is too ugly.  I'd instead move the if after the cray pointer/pointee
> tests, i.e.
> if (n->sym->attr.cray_pointer)
>   gfc_error (...);
> else if (n->sym->attr.cray_pointee)
>   gfc_error (...);
> else if (n->sym->attr.flavor == FL_VARIABLE
> 	 && !n->sym->as
> 	 && !n->sym->attr.pointer)
>   gfc_error (...);
> 

Hi Jakub, I've adjusted the patch like you suggested.

Patch has been re-tested and applied to gomp-4_0-branch,
okay for trunk as well?


[-- Attachment #2: x.diff --]
[-- Type: text/plain, Size: 4842 bytes --]

Index: fortran/openmp.c
===================================================================
--- fortran/openmp.c	(revision 238751)
+++ fortran/openmp.c	(working copy)
@@ -3752,17 +3752,24 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_claus
 			      && CLASS_DATA (n->sym)->attr.allocatable))
 			gfc_error ("ALLOCATABLE object %qs in %s clause at %L",
 				   n->sym->name, name, &n->where);
-		      if (n->sym->attr.pointer
-			  || (n->sym->ts.type == BT_CLASS && CLASS_DATA (n->sym)
-			      && CLASS_DATA (n->sym)->attr.class_pointer))
-			gfc_error ("POINTER object %qs in %s clause at %L",
-				   n->sym->name, name, &n->where);
+		      if (n->sym->ts.type == BT_CLASS
+			  && CLASS_DATA (n->sym)
+			  && CLASS_DATA (n->sym)->attr.class_pointer)
+			gfc_error ("POINTER object %qs of polymorphic type in "
+				   "%s clause at %L", n->sym->name, name,
+				   &n->where);
 		      if (n->sym->attr.cray_pointer)
 			gfc_error ("Cray pointer object %qs in %s clause at %L",
 				   n->sym->name, name, &n->where);
-		      if (n->sym->attr.cray_pointee)
+		      else if (n->sym->attr.cray_pointee)
 			gfc_error ("Cray pointee object %qs in %s clause at %L",
 				   n->sym->name, name, &n->where);
+		      else if (n->sym->attr.flavor == FL_VARIABLE
+			       && !n->sym->as
+			       && !n->sym->attr.pointer)
+			gfc_error ("%s clause variable %qs at %L is neither "
+				   "a POINTER nor an array", name,
+				   n->sym->name, &n->where);
 		      /* FALLTHRU */
 		  case OMP_LIST_DEVICE_RESIDENT:
 		    check_symbol_not_pointer (n->sym, n->where, name);
Index: testsuite/gfortran.dg/goacc/uninit-use-device-clause.f95
===================================================================
--- testsuite/gfortran.dg/goacc/uninit-use-device-clause.f95	(revision 238751)
+++ testsuite/gfortran.dg/goacc/uninit-use-device-clause.f95	(working copy)
@@ -2,9 +2,9 @@
 ! { dg-additional-options "-Wuninitialized" }
 
 subroutine test
-  integer :: i
+  integer, pointer :: p
 
-  !$acc host_data use_device(i) ! { dg-warning "is used uninitialized in this function" }
+  !$acc host_data use_device(p) ! { dg-warning "is used uninitialized in this function" }
   !$acc end host_data
 end subroutine test
 
Index: testsuite/gfortran.dg/goacc/list.f95
===================================================================
--- testsuite/gfortran.dg/goacc/list.f95	(revision 238751)
+++ testsuite/gfortran.dg/goacc/list.f95	(working copy)
@@ -76,19 +76,19 @@ program test
   !$acc parallel private (i) firstprivate (i) ! { dg-error "present on multiple clauses" }
   !$acc end parallel
 
-  !$acc host_data use_device(i)
+  !$acc host_data use_device(i) ! { dg-error "neither a POINTER nor an array" }
   !$acc end host_data
 
-  !$acc host_data use_device(c, d)
+  !$acc host_data use_device(c, d) ! { dg-error "neither a POINTER nor an array" }
   !$acc end host_data
 
   !$acc host_data use_device(a)
   !$acc end host_data
 
-  !$acc host_data use_device(i, j, k, l, a)
+  !$acc host_data use_device(i, j, k, l, a) ! { dg-error "neither a POINTER nor an array" }
   !$acc end host_data  
 
-  !$acc host_data use_device (i) use_device (j)
+  !$acc host_data use_device (i) use_device (j) ! { dg-error "neither a POINTER nor an array" }
   !$acc end host_data
 
   !$acc host_data use_device ! { dg-error "Unclassifiable OpenACC directive" }
@@ -99,13 +99,17 @@ program test
 
   !$acc host_data use_device(10) ! { dg-error "Syntax error" }
 
-  !$acc host_data use_device(/b/, /b/) ! { dg-error "present on multiple clauses" }
+  !$acc host_data use_device(/b/, /b/)
   !$acc end host_data
+  ! { dg-error "neither a POINTER nor an array" "" { target *-*-* } 102 }
+  ! { dg-error "present on multiple clauses" "" { target *-*-* } 102 }
 
-  !$acc host_data use_device(i, j, i) ! { dg-error "present on multiple clauses" }
+  !$acc host_data use_device(i, j, i)
   !$acc end host_data
+  ! { dg-error "neither a POINTER nor an array" "" { target *-*-* } 107 }
+  ! { dg-error "present on multiple clauses" "" { target *-*-* } 107 }
 
-  !$acc host_data use_device(p1) ! { dg-error "POINTER" }
+  !$acc host_data use_device(p1)
   !$acc end host_data
 
 end program test
Index: testsuite/gfortran.dg/goacc/host_data-tree.f95
===================================================================
--- testsuite/gfortran.dg/goacc/host_data-tree.f95	(revision 238751)
+++ testsuite/gfortran.dg/goacc/host_data-tree.f95	(working copy)
@@ -3,9 +3,9 @@
 
 program test
   implicit none
-  integer :: i = 1
+  integer, pointer :: p
 
-  !$acc host_data use_device(i)
+  !$acc host_data use_device(p)
   !$acc end host_data
 end program test
-! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(i\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "pragma acc host_data use_device_ptr\\(p\\)" 1 "original" } }

[-- Attachment #3: y.diff --]
[-- Type: text/plain, Size: 970 bytes --]

Index: testsuite/libgomp.oacc-fortran/host_data-1.f90
===================================================================
--- testsuite/libgomp.oacc-fortran/host_data-1.f90	(revision 0)
+++ testsuite/libgomp.oacc-fortran/host_data-1.f90	(revision 0)
@@ -0,0 +1,32 @@
+! { dg-do run }
+! { dg-additional-options "-cpp" }
+
+program test
+  implicit none
+
+  integer, target :: i, arr(1000)
+  integer, pointer :: ip, iph
+  integer, contiguous, pointer :: parr(:), parrh(:)
+
+  ! Assign the same targets
+  ip => i
+  parr => arr
+  iph => i
+  parrh => arr
+
+  !$acc data copyin(i, arr)
+  !$acc host_data use_device(ip, parr)
+
+  ! Test how the pointers compare inside a host_data construct
+#if ACC_MEM_SHARED
+  if (.not. associated(ip, iph)) call abort
+  if (.not. associated(parr, parrh)) call abort
+#else
+  if (associated(ip, iph)) call abort
+  if (associated(parr, parrh)) call abort
+#endif
+
+  !$acc end host_data
+  !$acc end data
+
+end program test

Re: OpenACC-Library-Interoperability

[Salvatore Filippone]

Hi there,
I tried to use the same script, but then I get
    done
/bin/sh /opt/gnu/6.1.1-acc//src/gcc/libgcc/../mkinstalldirs
/opt/gnu/6.1.1-acc//install/lib/gcc/x86_64-pc-linux-gnu/6.1.1/include
/usr/bin/install -c -m 644 unwind.h
/opt/gnu/6.1.1-acc//install/lib/gcc/x86_64-pc-linux-gnu/6.1.1/include
make[2]: Leaving directory
`/opt/gnu/6.1.1-acc/build/gcc6/x86_64-pc-linux-gnu/libgcc'
/bin/sh: line 3: cd: x86_64-pc-linux-gnu/libstdc++-v3: No such file or directory
make[1]: *** [install-target-libstdc++-v3] Error 1
make[1]: Leaving directory `/opt/gnu/6.1.1-acc/build/gcc6'
make: *** [install] Error 2

Any ideas? It appears that the gcc6 build step is not doing everything
that it is supposed to do.....
Thanks

On Tue, Aug 2, 2016 at 4:54 PM, Salvatore Filippone
<filippone.salvatore@gmail.com> wrote:
> Hi there,
> I tried to use the same script, but then I get
>     done
> /bin/sh /opt/gnu/6.1.1-acc//src/gcc/libgcc/../mkinstalldirs
> /opt/gnu/6.1.1-acc//install/lib/gcc/x86_64-pc-linux-gnu/6.1.1/include
> /usr/bin/install -c -m 644 unwind.h
> /opt/gnu/6.1.1-acc//install/lib/gcc/x86_64-pc-linux-gnu/6.1.1/include
> make[2]: Leaving directory
> `/opt/gnu/6.1.1-acc/build/gcc6/x86_64-pc-linux-gnu/libgcc'
> /bin/sh: line 3: cd: x86_64-pc-linux-gnu/libstdc++-v3: No such file or
> directory
> make[1]: *** [install-target-libstdc++-v3] Error 1
> make[1]: Leaving directory `/opt/gnu/6.1.1-acc/build/gcc6'
> make: *** [install] Error 2
>
> Any ideas? It appears that the gcc6 build step is not doing all that it is
> supposed to do.....
>
> On Mon, Aug 1, 2016 at 4:58 PM, Vikram Singh <vikramsingh001@gmail.com>
> wrote:
>>
>> I finally found some time to repeat my previous attempts, but with gcc
>> 6.1. I'll try to be a little detailed so that people don't have to waste
>> time in doing stuff I did.
>>
>> I used the same script as before
>>
>>
>> https://github.com/olcf/OLCFHack15/blob/master/GCC5OffloadTest/auto-gcc5-offload-openacc-build-install.sh
>>
>> And modified it for the paths for gcc6.1. For some reason
>>
>> x86_64-pc-linux-gnu-accel-nvptx-none-gcc
>>
>> that builds in install/bin
>>
>> does not look for directories in $LD_LIBRARY_PATH. I found this using the
>> -print-search-dirs option.
>>
>> x86_64-pc-linux-gnu-accel-nvptx-none-gcc -print-search-dirs
>>
>> So, I got a few errors for libraries not found for libraries that were
>> present in install/lib64. I just copied those libraries to
>> install/nvptx-none/lib/.
>>
>> With that I can now use gcc6.
>>
>> Onto the next task. I made a simple testcase. I'll copy paste the whole
>> thing here. Lets call this test.f90
>>
>>
>> program example_dgemm
>>
>>   use iso_c_binding
>>   implicit none
>>
>>   integer                     :: N = 8
>>   real(c_double), allocatable :: A(:,:), B(:, :), C(:, :)
>>   integer                     :: size_of_real, i, j
>>   integer*8 :: devPtrA, devPtrB, devPtrC
>>
>>   size_of_real = 16
>>
>>   allocate(A(N, N))
>>   allocate(B(N, N))
>>   allocate(C(N, N))
>>
>>   !$ACC PARALLEL COPY(A)
>>   do i = 1, N
>>       do j = 1, N
>>           A(i, j) = i + j
>>       end do
>>   end do
>>   !$ACC END PARALLEL
>>   !$ACC PARALLEL COPY(B)
>>   do i = 1, N
>>       do j = 1, N
>>           B(i, j) = j
>>       end do
>>   end do
>>   !$ACC END PARALLEL
>>
>>   call cublas_Alloc(N*N, size_of_real, devPtrA)
>>   call cublas_Alloc(N*N, size_of_real, devPtrB)
>>   call cublas_Alloc(N*N, size_of_real, devPtrC)
>>
>>   ! Copy Fixed Data to the GPU
>>   call cublas_Set_Matrix(N, N, size_of_real, A, N, devPtrA, N)
>>   call cublas_Set_Matrix(N, N, size_of_real, C, N, devPtrC, N)
>>
>>   ! Copy data to the GPU
>>   call cublas_Set_Matrix(N, N, size_of_real, B, N, devPtrB, N)
>>
>>   ! Do DGEMM on the GPU
>>   call cublas_DGEMM('N', 'N', N, N, N, &
>>        1.0_c_double, devPtrA, N, devPtrB, N, 0.0_c_double, devPtrC, N)
>>
>>   ! Copy data from the GPU
>>   call cublas_Get_matrix(N, N, size_of_real, devPtrC, N, C, N)
>>
>>   call cublas_Free(devPtrA)
>>   call cublas_Free(devPtrB)
>>   call cublas_Free(devPtrC)
>>
>>   deallocate(A)
>>   deallocate(B)
>>   deallocate(C)
>>
>> end program example_dgemm
>>
>> Basically, I can test both CUBLAS and OPENACC using this simple code. But,
>> using CUBLAS needs some extra stuff. So what I need to do is copy fortran.h,
>> fortran_common.h and fortran.c from /usr/local/cuda/src.
>>
>> Now, I do
>>
>>  ./rungcc6.sh gcc -Wall -g -I/usr/local/cuda/include -I/usr/local/cuda/src
>> -DCUBLAS_GFORTRAN -c fortran.c
>>
>> Finally, I do
>>
>> ./rungcc6.sh gfortran -Wall -g test.f90 fortran.o -fopenacc
>> -foffload=nvptx-none -foffload=-O3 -O3 -o gpu.x -L/usr/local/cuda/lib64
>> -lcublas -lcudart
>>
>> And finally nvprof ./gpu.x gives output
>>
>> ==9155== Profiling application: ./gpu.x
>> ==9155== Profiling result:
>> Time(%)      Time     Calls       Avg       Min       Max  Name
>>  23.60%  13.952us         1  13.952us  13.952us  13.952us
>> MAIN__$_omp_fn$0
>>  23.55%  13.920us         1  13.920us  13.920us  13.920us
>> MAIN__$_omp_fn$1
>>  18.57%  10.976us        16     686ns     576ns  1.0880us  [CUDA memcpy
>> HtoD]
>>  13.72%  8.1080us         2  4.0540us  2.0430us  6.0650us  [CUDA memcpy
>> HtoH]
>>  12.07%  7.1360us         1  7.1360us  7.1360us  7.1360us  void
>> gemm_kernel2x2_core<double, bool=0, bool=0, bool=0, bool=0, bool=0>(double*,
>> double const *, double const *, int, int, int, int, int, int, double*,
>> double*, double, double, int)
>>   8.50%  5.0240us         3  1.6740us  1.5680us  1.8880us  [CUDA memcpy
>> DtoH]
>>
>> Clearly openacc loops run on GPU and DGEMM kernels also run on GPU.
>>
>> So the only thing left now, is to put host_data so that the whole process
>> of cublas_alloc, cublas_set_matrix, cublas_get_matrix need not be done.
>>
>> If you can tell me how the patch is applied, I can test it with this
>> simple example.
>>
>> Regards,
>> Vikram
>>
>>
>> On Thu, May 12, 2016 at 6:34 PM, Thomas Schwinge <thomas@codesourcery.com>
>> wrote:
>>>
>>> Hi!
>>>
>>> On Fri, 15 Apr 2016 13:59:53 +0300, Vikram Singh
>>> <vikramsingh001@gmail.com> wrote:
>>> > I checked libgomp.oacc-c-c++-common/context-*.c, and they seem to be
>>> > exactly what I needed to start with.
>>> >
>>> > If I understand it correctly, I'll anyway be trying to implement
>>> > something similar with PGI compilers, so I'll have something as a
>>> > starting point.
>>>
>>> Any progress on that already?  I had a very quick look myself, but it's
>>> not as easy as I thought...  A "courageous" use of "use cublas" in
>>> Fortran code compiled with gfortran results in: "Fatal Error: Can't open
>>> module file 'cublas.mod' for reading".  The problem is: Fortran
>>> interfacing to C libraries (which cuBLAS is).
>>>
>>> <http://docs.nvidia.com/cuda/cublas/index.html#appendix-b-cublas-fortran-bindings>
>>> has some instructions how to do it.  Anyone got that to work already?
>>>
>>> > But again, it will need the OpenACC host_data construct to be setup
>>> > for gfortran to test.
>>>
>>> In
>>>
>>> <http://news.gmane.org/find-root.php?message_id=%3C2b4f59d5-be38-2814-27bb-73aa7ffb4b8f%40codesourcery.com%3E>,
>>> Chung-Lin has now posted a patch (pending review) that should make the
>>> OpenACC host_data construct usable in GCC Fortran.  (Problem discussed in
>>>
>>> <http://news.gmane.org/find-root.php?message_id=%3C878u0o6wwj.fsf%40kepler.schwinge.homeip.net%3E>
>>> before.)
>>>
>>> For reference:
>>>
>>> > On Fri, Apr 15, 2016 at 11:57 AM, Thomas Schwinge
>>> > <thomas@codesourcery.com> wrote:
>>> > > On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh
>>> > > <vikramsingh001@gmail.com> wrote:
>>> > >> Yes, I came to the conclusion that host_data
>>> > >> would be the only way to do it in fortran.
>>> > >>
>>> > >> On the other hand, I though there were no plans to implement it in
>>> > >> gfortran 6 either
>>> > >
>>> > > I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
>>> > > fixed in time for the GCC 6.1 release.  I'll keep you posted.
>>> > >
>>> > >
>>> > >> > I'm copying Jim, who is the author of this chapter in the
>>> > >> > documentation
>>> > >> > as well as the
>>> > >> > libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
>>> > >> > test cases, and much of the relevant libgomp code, too, and who
>>> > >> > should
>>> > >> > please correct me if I'm wrong.  I'll make a note for later, that
>>> > >> > we
>>> > >> > should translate the libgomp.oacc-c-c++-common/context-*.c test
>>> > >> > cases to
>>> > >> > Fortran, and also replicate them using the OpenACC host_data
>>> > >> > construct
>>> > >> > (like in
>>> > >> > libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
>>> > >> > the same for the documentation you referenced.  (Vikram, please
>>> > >> > tell if
>>> > >> > you're interested to work on these items.)
>>> > >>
>>> > >> I am not sure I understand what you want me to work on exactly. I am
>>> > >> not really much of a C programmer, so I wouldn't be good at it. I
>>> > >> would still like to help.
>>> > >
>>> > > Sorry for being unclear.  My idea/question has been whether you're
>>> > > interested in helping by translating the documentation as well as the
>>> > > libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran
>>> > > (using
>>> > > the OpenACC host_data construct instead of the acc_* functions).  If
>>> > > yes,
>>> > > then that's great, if not, then one of us will do it at some point.
>>>
>>>
>>> Grüße
>>>  Thomas
>>
>>
>

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Sorry, I realized I have to send it as plain text.

I finally found some time to repeat my previous attempts, but with gcc
6.1. I'll try to be a little detailed so that people don't have to
waste time in doing stuff I did.

I used the same script as before

https://github.com/olcf/OLCFHack15/blob/master/GCC5OffloadTest/auto-gcc5-offload-openacc-build-install.sh

And modified it for the paths for gcc6.1. For some reason

x86_64-pc-linux-gnu-accel-nvptx-none-gcc

that builds in install/bin

does not look for directories in $LD_LIBRARY_PATH. I found this using
the -print-search-dirs option.

x86_64-pc-linux-gnu-accel-nvptx-none-gcc -print-search-dirs

So, I got a few errors for libraries not found for libraries that were
present in install/lib64. I just copied those libraries to
install/nvptx-none/lib/.

With that I can now use gcc6.

Onto the next task. I made a simple testcase. I'll copy paste the
whole thing here. Lets call this test.f90


program example_dgemm

  use iso_c_binding
  implicit none

  integer                     :: N = 8
  real(c_double), allocatable :: A(:,:), B(:, :), C(:, :)
  integer                     :: size_of_real, i, j
  integer*8 :: devPtrA, devPtrB, devPtrC

  size_of_real = 16

  allocate(A(N, N))
  allocate(B(N, N))
  allocate(C(N, N))

  !$ACC PARALLEL COPY(A)
  do i = 1, N
      do j = 1, N
          A(i, j) = i + j
      end do
  end do
  !$ACC END PARALLEL
  !$ACC PARALLEL COPY(B)
  do i = 1, N
      do j = 1, N
          B(i, j) = j
      end do
  end do
  !$ACC END PARALLEL

  call cublas_Alloc(N*N, size_of_real, devPtrA)
  call cublas_Alloc(N*N, size_of_real, devPtrB)
  call cublas_Alloc(N*N, size_of_real, devPtrC)

  ! Copy Fixed Data to the GPU
  call cublas_Set_Matrix(N, N, size_of_real, A, N, devPtrA, N)
  call cublas_Set_Matrix(N, N, size_of_real, C, N, devPtrC, N)

  ! Copy data to the GPU
  call cublas_Set_Matrix(N, N, size_of_real, B, N, devPtrB, N)

  ! Do DGEMM on the GPU
  call cublas_DGEMM('N', 'N', N, N, N, &
       1.0_c_double, devPtrA, N, devPtrB, N, 0.0_c_double, devPtrC, N)

  ! Copy data from the GPU
  call cublas_Get_matrix(N, N, size_of_real, devPtrC, N, C, N)

  call cublas_Free(devPtrA)
  call cublas_Free(devPtrB)
  call cublas_Free(devPtrC)

  deallocate(A)
  deallocate(B)
  deallocate(C)

end program example_dgemm

Basically, I can test both CUBLAS and OPENACC using this simple code.
But, using CUBLAS needs some extra stuff. So what I need to do is copy
fortran.h, fortran_common.h and fortran.c from /usr/local/cuda/src.

Now, I do

 ./rungcc6.sh gcc -Wall -g -I/usr/local/cuda/include
-I/usr/local/cuda/src -DCUBLAS_GFORTRAN -c fortran.c

Finally, I do

./rungcc6.sh gfortran -Wall -g test.f90 fortran.o -fopenacc
-foffload=nvptx-none -foffload=-O3 -O3 -o gpu.x
-L/usr/local/cuda/lib64 -lcublas -lcudart

And finally nvprof ./gpu.x gives output

==9155== Profiling application: ./gpu.x
==9155== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 23.60%  13.952us         1  13.952us  13.952us  13.952us  MAIN__$_omp_fn$0
 23.55%  13.920us         1  13.920us  13.920us  13.920us  MAIN__$_omp_fn$1
 18.57%  10.976us        16     686ns     576ns  1.0880us  [CUDA memcpy HtoD]
 13.72%  8.1080us         2  4.0540us  2.0430us  6.0650us  [CUDA memcpy HtoH]
 12.07%  7.1360us         1  7.1360us  7.1360us  7.1360us  void
gemm_kernel2x2_core<double, bool=0, bool=0, bool=0, bool=0,
bool=0>(double*, double const *, double const *, int, int, int, int,
int, int, double*, double*, double, double, int)
  8.50%  5.0240us         3  1.6740us  1.5680us  1.8880us  [CUDA memcpy DtoH]

Clearly openacc loops run on GPU and DGEMM kernels also run on GPU.

So the only thing left now, is to put host_data so that the whole
process of cublas_alloc, cublas_set_matrix, cublas_get_matrix need not
be done.

If you can tell me how the patch is applied, I can test it with this
simple example.

Regards,
Vikram

On Mon, Aug 1, 2016 at 6:58 PM, Vikram Singh <vikramsingh001@gmail.com> wrote:
> I finally found some time to repeat my previous attempts, but with gcc 6.1.
> I'll try to be a little detailed so that people don't have to waste time in
> doing stuff I did.
>
> I used the same script as before
>
> https://github.com/olcf/OLCFHack15/blob/master/GCC5OffloadTest/auto-gcc5-offload-openacc-build-install.sh
>
> And modified it for the paths for gcc6.1. For some reason
>
> x86_64-pc-linux-gnu-accel-nvptx-none-gcc
>
> that builds in install/bin
>
> does not look for directories in $LD_LIBRARY_PATH. I found this using the
> -print-search-dirs option.
>
> x86_64-pc-linux-gnu-accel-nvptx-none-gcc -print-search-dirs
>
> So, I got a few errors for libraries not found for libraries that were
> present in install/lib64. I just copied those libraries to
> install/nvptx-none/lib/.
>
> With that I can now use gcc6.
>
> Onto the next task. I made a simple testcase. I'll copy paste the whole
> thing here. Lets call this test.f90
>
>
> program example_dgemm
>
>   use iso_c_binding
>   implicit none
>
>   integer                     :: N = 8
>   real(c_double), allocatable :: A(:,:), B(:, :), C(:, :)
>   integer                     :: size_of_real, i, j
>   integer*8 :: devPtrA, devPtrB, devPtrC
>
>   size_of_real = 16
>
>   allocate(A(N, N))
>   allocate(B(N, N))
>   allocate(C(N, N))
>
>   !$ACC PARALLEL COPY(A)
>   do i = 1, N
>       do j = 1, N
>           A(i, j) = i + j
>       end do
>   end do
>   !$ACC END PARALLEL
>   !$ACC PARALLEL COPY(B)
>   do i = 1, N
>       do j = 1, N
>           B(i, j) = j
>       end do
>   end do
>   !$ACC END PARALLEL
>
>   call cublas_Alloc(N*N, size_of_real, devPtrA)
>   call cublas_Alloc(N*N, size_of_real, devPtrB)
>   call cublas_Alloc(N*N, size_of_real, devPtrC)
>
>   ! Copy Fixed Data to the GPU
>   call cublas_Set_Matrix(N, N, size_of_real, A, N, devPtrA, N)
>   call cublas_Set_Matrix(N, N, size_of_real, C, N, devPtrC, N)
>
>   ! Copy data to the GPU
>   call cublas_Set_Matrix(N, N, size_of_real, B, N, devPtrB, N)
>
>   ! Do DGEMM on the GPU
>   call cublas_DGEMM('N', 'N', N, N, N, &
>        1.0_c_double, devPtrA, N, devPtrB, N, 0.0_c_double, devPtrC, N)
>
>   ! Copy data from the GPU
>   call cublas_Get_matrix(N, N, size_of_real, devPtrC, N, C, N)
>
>   call cublas_Free(devPtrA)
>   call cublas_Free(devPtrB)
>   call cublas_Free(devPtrC)
>
>   deallocate(A)
>   deallocate(B)
>   deallocate(C)
>
> end program example_dgemm
>
> Basically, I can test both CUBLAS and OPENACC using this simple code. But,
> using CUBLAS needs some extra stuff. So what I need to do is copy fortran.h,
> fortran_common.h and fortran.c from /usr/local/cuda/src.
>
> Now, I do
>
>  ./rungcc6.sh gcc -Wall -g -I/usr/local/cuda/include -I/usr/local/cuda/src
> -DCUBLAS_GFORTRAN -c fortran.c
>
> Finally, I do
>
> ./rungcc6.sh gfortran -Wall -g test.f90 fortran.o -fopenacc
> -foffload=nvptx-none -foffload=-O3 -O3 -o gpu.x -L/usr/local/cuda/lib64
> -lcublas -lcudart
>
> And finally nvprof ./gpu.x gives output
>
> ==9155== Profiling application: ./gpu.x
> ==9155== Profiling result:
> Time(%)      Time     Calls       Avg       Min       Max  Name
>  23.60%  13.952us         1  13.952us  13.952us  13.952us  MAIN__$_omp_fn$0
>  23.55%  13.920us         1  13.920us  13.920us  13.920us  MAIN__$_omp_fn$1
>  18.57%  10.976us        16     686ns     576ns  1.0880us  [CUDA memcpy
> HtoD]
>  13.72%  8.1080us         2  4.0540us  2.0430us  6.0650us  [CUDA memcpy
> HtoH]
>  12.07%  7.1360us         1  7.1360us  7.1360us  7.1360us  void
> gemm_kernel2x2_core<double, bool=0, bool=0, bool=0, bool=0, bool=0>(double*,
> double const *, double const *, int, int, int, int, int, int, double*,
> double*, double, double, int)
>   8.50%  5.0240us         3  1.6740us  1.5680us  1.8880us  [CUDA memcpy
> DtoH]
>
> Clearly openacc loops run on GPU and DGEMM kernels also run on GPU.
>
> So the only thing left now, is to put host_data so that the whole process of
> cublas_alloc, cublas_set_matrix, cublas_get_matrix need not be done.
>
> If you can tell me how the patch is applied, I can test it with this simple
> example.
>
> Regards,
> Vikram
>
>
> On Thu, May 12, 2016 at 6:34 PM, Thomas Schwinge <thomas@codesourcery.com>
> wrote:
>>
>> Hi!
>>
>> On Fri, 15 Apr 2016 13:59:53 +0300, Vikram Singh
>> <vikramsingh001@gmail.com> wrote:
>> > I checked libgomp.oacc-c-c++-common/context-*.c, and they seem to be
>> > exactly what I needed to start with.
>> >
>> > If I understand it correctly, I'll anyway be trying to implement
>> > something similar with PGI compilers, so I'll have something as a
>> > starting point.
>>
>> Any progress on that already?  I had a very quick look myself, but it's
>> not as easy as I thought...  A "courageous" use of "use cublas" in
>> Fortran code compiled with gfortran results in: "Fatal Error: Can't open
>> module file 'cublas.mod' for reading".  The problem is: Fortran
>> interfacing to C libraries (which cuBLAS is).
>>
>> <http://docs.nvidia.com/cuda/cublas/index.html#appendix-b-cublas-fortran-bindings>
>> has some instructions how to do it.  Anyone got that to work already?
>>
>> > But again, it will need the OpenACC host_data construct to be setup
>> > for gfortran to test.
>>
>> In
>>
>> <http://news.gmane.org/find-root.php?message_id=%3C2b4f59d5-be38-2814-27bb-73aa7ffb4b8f%40codesourcery.com%3E>,
>> Chung-Lin has now posted a patch (pending review) that should make the
>> OpenACC host_data construct usable in GCC Fortran.  (Problem discussed in
>>
>> <http://news.gmane.org/find-root.php?message_id=%3C878u0o6wwj.fsf%40kepler.schwinge.homeip.net%3E>
>> before.)
>>
>> For reference:
>>
>> > On Fri, Apr 15, 2016 at 11:57 AM, Thomas Schwinge
>> > <thomas@codesourcery.com> wrote:
>> > > On Fri, 15 Apr 2016 11:35:06 +0300, Vikram Singh
>> > > <vikramsingh001@gmail.com> wrote:
>> > >> Yes, I came to the conclusion that host_data
>> > >> would be the only way to do it in fortran.
>> > >>
>> > >> On the other hand, I though there were no plans to implement it in
>> > >> gfortran 6 either
>> > >
>> > > I still hope we'll get this (that is, <https://gcc.gnu.org/PR70598>)
>> > > fixed in time for the GCC 6.1 release.  I'll keep you posted.
>> > >
>> > >
>> > >> > I'm copying Jim, who is the author of this chapter in the
>> > >> > documentation
>> > >> > as well as the
>> > >> > libgomp/testsuite/libgomp.oacc-c-c++-common/context-*.c
>> > >> > test cases, and much of the relevant libgomp code, too, and who
>> > >> > should
>> > >> > please correct me if I'm wrong.  I'll make a note for later, that
>> > >> > we
>> > >> > should translate the libgomp.oacc-c-c++-common/context-*.c test
>> > >> > cases to
>> > >> > Fortran, and also replicate them using the OpenACC host_data
>> > >> > construct
>> > >> > (like in
>> > >> > libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c), and
>> > >> > the same for the documentation you referenced.  (Vikram, please
>> > >> > tell if
>> > >> > you're interested to work on these items.)
>> > >>
>> > >> I am not sure I understand what you want me to work on exactly. I am
>> > >> not really much of a C programmer, so I wouldn't be good at it. I
>> > >> would still like to help.
>> > >
>> > > Sorry for being unclear.  My idea/question has been whether you're
>> > > interested in helping by translating the documentation as well as the
>> > > libgomp.oacc-c-c++-common/context-*.c test cases from C to Fortran
>> > > (using
>> > > the OpenACC host_data construct instead of the acc_* functions).  If
>> > > yes,
>> > > then that's great, if not, then one of us will do it at some point.
>>
>>
>> Grüße
>>  Thomas
>
>

Re: [PATCH, Fortran, OpenACC] Fix PR70598, Fortran host_data ICE

[Jakub Jelinek]

On Fri, Jul 29, 2016 at 11:47:18PM +0800, Chung-Lin Tang wrote:
> On 2016/7/21 07:13 PM, Jakub Jelinek wrote:
> > Better put every && on a separate line if the whole if (...) doesn't fit
> > on a single line.
> > 
> >> > +			  && !n->sym->attr.cray_pointer
> >> > +			  && !n->sym->attr.cray_pointee)
> > This is too ugly.  I'd instead move the if after the cray pointer/pointee
> > tests, i.e.
> > if (n->sym->attr.cray_pointer)
> >   gfc_error (...);
> > else if (n->sym->attr.cray_pointee)
> >   gfc_error (...);
> > else if (n->sym->attr.flavor == FL_VARIABLE
> > 	 && !n->sym->as
> > 	 && !n->sym->attr.pointer)
> >   gfc_error (...);
> > 
> 
> Hi Jakub, I've adjusted the patch like you suggested.
> 
> Patch has been re-tested and applied to gomp-4_0-branch,
> okay for trunk as well?

Ok with suitable ChangeLog entry.

	Jakub

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Hi,

I tried the installation on yet another machine. It seems that the
installation goes along on Ubuntu, but for some reason not on Red Hat.
So I would request someone to please look at that.

Anyway, I installed gcc6.2 using the same steps. I believe host_data
has been implemented on gcc 6.2. So I tried the following code.

program example_dgemm

  use iso_c_binding

  implicit none

  integer, parameter        :: N = 8
  real(c_double)              :: A(N,N), B(N, N), C(N, N)
  integer                          ::  i, j

  !$acc enter data create(A, B)

  !$acc parallel present(A)
  do i = 1, N
      do j = 1, N
          A(i, j) = i + j
      end do
  end do
  !$acc end parallel
  !$acc parallel present(B)
  do i = 1, N
      do j = 1, N
          B(i, j) = j
      end do
  end do
  !$acc end parallel

  ! Do DGEMM on the GPU
  !$acc host_data use_device(A, B, C)
  call cublas_DGEMM('N', 'N', N, N, N, &
       1.0_c_double, A, N, B, N, 0.0_c_double, C, N)
   !$acc end host_data

end program example_dgemm

First, from the previous code, I had to change A, B, C from
allocatable to fixed size. Seems, gfortran host_data does not allow
allocatable arrays.

This code atleast compiles. Unfortunately on running it, I get the error,


Program received signal SIGSEGV: Segmentation fault - invalid memory reference.

Backtrace for this error:
#0  0x7f046e5bacaf in ???
#1  0x408b1d in cublas_dgemm_
at fortran.c:1461
#2  0x40477d in example_dgemm
at test.f90:37
#3  0x403e9c in main
at test.f90:42
Segmentation fault (core dumped)

It seems to me the data type is not being correctly translated to what
DGEMM requires.

Regards,
Vikram

Re: OpenACC-Library-Interoperability

[James Norris]

Vikram,

On 08/29/2016 08:58 AM, Vikram Singh wrote:
> Hi,
>
> I tried the installation on yet another machine. It seems that the
> installation goes along on Ubuntu, but for some reason not on Red Hat.
> So I would request someone to please look at that.

[snip snip]

What version of RedHat do you see the problem with?

Regards,
James Norris
Mentor Graphics

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Hi James,

Thanks for the response.

cat /proc/version

gives me

Linux version 2.6.32-573.8.1.el6.x86_64
(mockbuild@x86-033.build.eng.bos.redhat.com) (gcc version 4.4.7
20120313 (Red Hat 4.4.7-16) (GCC)

cat /etc/redhat-release

gives me

Red Hat Enterprise Linux Server release 6.7 (Santiago)

Vikram

On Mon, Aug 29, 2016 at 5:01 PM, James Norris <jnorris@codesourcery.com> wrote:
> Vikram,
>
> On 08/29/2016 08:58 AM, Vikram Singh wrote:
>>
>> Hi,
>>
>> I tried the installation on yet another machine. It seems that the
>> installation goes along on Ubuntu, but for some reason not on Red Hat.
>> So I would request someone to please look at that.
>
>
> [snip snip]
>
> What version of RedHat do you see the problem with?
>
> Regards,
> James Norris
> Mentor Graphics
>
>

Re: OpenACC-Library-Interoperability

[Cesar Philippidis]

On 08/29/2016 06:58 AM, Vikram Singh wrote:

> It seems to me the data type is not being correctly translated to what
> DGEMM requires.

Correct. Arrays in gfortran have different representations from those in
c. See
<https://gcc.gnu.org/onlinedocs/gfortran/Interoperability-with-C.html>
for instructions on how to call c functions from gfortran. I recently
posted a libgomp host_data test case which calls cublas here
<https://gcc.gnu.org/ml/gcc-patches/2016-08/msg00976.html>.

Cesar

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Hi Cesar,

Thanks for the pointers.

I tried to use your test case to make one for DGEMM.

program test

  use iso_c_binding

  implicit none

  integer(c_int), parameter :: N = 10
  integer(c_int) :: i, j
  real(c_double) :: x(N, N), y(N, N), z(N, N)
  character(kind=c_char)     :: flag

  interface
     subroutine cublasdgemm(transa, transb, m, n, k, alpha, A, lda, B, &
             ldb, beta, C, ldc) bind(c, name="cublasDgemm")
       use iso_c_binding
       character(kind=c_char)     :: transa, transb
       integer(kind=c_int), value :: m, n, k
       real(c_double), value      :: alpha
       type(*), dimension(*)      :: A
       integer(kind=c_int), value :: lda
       type(*), dimension(*)      :: B
       integer(kind=c_int), value :: ldb
       real(c_double), value      :: beta
       type(*), dimension(*)      :: C
       integer(kind=c_int), value :: ldc

     end subroutine cublasdgemm

  end interface

  do i = 1, N
     do j = 1, N
       x(i, j) = 4.0 * i
       y(i, j) = 3.0 + j
       z(i, j) = 0.0
     end do
  end do

  flag = 'N'
  !$acc data copyin (x, y) copy (z)
  !$acc host_data use_device (x, y, z)
  call cublasdgemm(flag, flag, n, n, n, 1.0_c_double, x, n, y, n,
0.0_c_double, z, n)
  !$acc end host_data

  !$acc update self(z)

  !$acc end data

  write(*, *) z

  call dgemm(flag, flag, n, n, n, 1.0_c_double, x, n, y, n, 0.0_c_double, z, n)

  write(*, *) z

end program test


z from the cublasdgemm call gives 0 everywhere, unlike the actual
dgemm call. The first line of output is

 ** On entry to DGEMM  parameter number 1 had an illegal value

In addition, I looked at your test case. You do not do

 !$acc update self(y)

Does that mean that the host variable is automatically updated in gfortran.

Regards,
Vikram

On Mon, Aug 29, 2016 at 5:34 PM, Cesar Philippidis
<cesar_philippidis@mentor.com> wrote:
> On 08/29/2016 06:58 AM, Vikram Singh wrote:
>
>> It seems to me the data type is not being correctly translated to what
>> DGEMM requires.
>
> Correct. Arrays in gfortran have different representations from those in
> c. See
> <https://gcc.gnu.org/onlinedocs/gfortran/Interoperability-with-C.html>
> for instructions on how to call c functions from gfortran. I recently
> posted a libgomp host_data test case which calls cublas here
> <https://gcc.gnu.org/ml/gcc-patches/2016-08/msg00976.html>.
>
> Cesar

Re: OpenACC-Library-Interoperability

[Cesar Philippidis]

On 08/29/2016 08:38 AM, Vikram Singh wrote:

> z from the cublasdgemm call gives 0 everywhere, unlike the actual
> dgemm call. The first line of output is
> 
>  ** On entry to DGEMM  parameter number 1 had an illegal value

Check the documentation for cublas. You may not be calling that function
with the proper arguments.

> In addition, I looked at your test case. You do not do
> 
>  !$acc update self(y)
> 
> Does that mean that the host variable is automatically updated in gfortran.

The host_data clause tells the compiler to use variables that have
already been mapped onto the accelerator via acc data or acc enter enter
data. You need to use the data directives to actually move the data as
appropriate. In host_data-2.f90 example, the value of 'y' was going to
be copyied out by the acc data region, so there was no need to use acc
update host.

You only need to use the update directive if you want to retain that
data on the accelerator.

Cesar

Re: OpenACC-Library-Interoperability

[Vikram Singh]

Thanks Cesar for the explanation.



On Mon, Aug 29, 2016 at 7:15 PM, Cesar Philippidis
<cesar_philippidis@mentor.com> wrote:
> On 08/29/2016 08:38 AM, Vikram Singh wrote:
>
>> z from the cublasdgemm call gives 0 everywhere, unlike the actual
>> dgemm call. The first line of output is
>>
>>  ** On entry to DGEMM  parameter number 1 had an illegal value
>
> Check the documentation for cublas. You may not be calling that function
> with the proper arguments.
>
>> In addition, I looked at your test case. You do not do
>>
>>  !$acc update self(y)
>>
>> Does that mean that the host variable is automatically updated in gfortran.
>
> The host_data clause tells the compiler to use variables that have
> already been mapped onto the accelerator via acc data or acc enter enter
> data. You need to use the data directives to actually move the data as
> appropriate. In host_data-2.f90 example, the value of 'y' was going to
> be copyied out by the acc data region, so there was no need to use acc
> update host.
>
> You only need to use the update directive if you want to retain that
> data on the accelerator.
>
> Cesar