[gomp4] Handle Fortran deviceptr clause.

Hi Thomas,

On 07/09/2015 03:29 AM, Thomas Schwinge wrote:
> Hi Jim!
>
> On Wed, 8 Jul 2015 13:00:16 -0500, James Norris <jnorris@codesourcery.com> wrote:
>> This patch adds handling of the deviceptr clause when
>> used within a Fortran program.
>
> Please motivate such non-obvious code changes by a test case.  At least
> to me, it's not at all obvious what's going on here...

Attached are two files which allow the testing of the fix.
I wasn't able to figure out how to get the testsuite to
compile a c source file and a fortran source file. So the
following should suffice until I figure something out.


/opt/codesourcery/trunk/bin/x86_64-none-linux-gnu-gcc -fopenacc -Wall -c 
dp-1.c
/opt/codesourcery/trunk/bin/x86_64-none-linux-gnu-gfortran -fopenacc 
-Wall -c dp-2.f90
/opt/codesourcery/trunk/bin/x86_64-none-linux-gnu-gcc -fopenacc -Wall -o 
dp dp-1.o dp-2.o


>
>> Committed to gomp-4_0-branch
>
>> +	* oacc-parallel.c (GOACC_parallel GOACC_data_start): Handle Fortran
>> +	deviceptr clause.
>
>> --- a/libgomp/oacc-parallel.c
>> +++ b/libgomp/oacc-parallel.c
>> @@ -211,6 +211,21 @@ GOACC_parallel (int device, void (*fn) (void *),
>>     thr = goacc_thread ();
>>     acc_dev = thr->dev;
>>
>> +  for (i = 0; i < mapnum; i++)
>> +    {
>> +      unsigned short kind1 = kinds[i] & 0xff;
>> +      unsigned short kind2 = kinds[i+1] & 0xff;
>> +
>> +      if ((kind1 == GOMP_MAP_FORCE_DEVICEPTR && kind2 == GOMP_MAP_POINTER)
>> +	   && (sizes[i + 1] == 0)
>> +	   && (hostaddrs[i] == *(void **)hostaddrs[i + 1]))
>> +	{
>> +	  kinds[i+1] = kinds[i];
>> +	  sizes[i+1] = sizeof (void *);
>> +	  hostaddrs[i] = NULL;
>> +	}
>> +    }
>
> Ugh.  That loop should be bounded by mapnum - 1 to avoid out-of-bounds
> array accesses.  And, such "voodoo" code constructs do need a comment,
> please

The 'Ugh' is fixed, comment added, and committed to gomp-4_0-branch.

> Why does this processing need to happen at run-time, in libgomp?
> Should something else be done during OMP lowering, for example?

The code in question is for the handling of the deviceptr clause
when presented via Fortran. As there is special code to handle
PSETs, which are Fortran specific, in the run-time, I felt that
there shouldn't be an issue with handling the deviceptr clause
there as well.

Thank you, thank you,
Jim
/* { dg-do run } */

#include <stdlib.h>
#include <openacc.h>
#include <stdio.h>
#include <string.h>

extern void subr1_ (int *);
extern void subr2_ (int *);
extern void subr3 (int *);
extern void subr4 (int *);

void
subr (int *a)
{
#pragma acc data deviceptr (a)
  {
#pragma acc parallel
    {
      int i;

      for (i = 0; i < 8; i++)
        a[i] = i + 1;
    }
  }
}

int
main (int argc, char **argv)
{
  int  N = 8;
  int  nbytes;
  int  *a, *b;
  int  i;

  nbytes = N * sizeof (int);

  a = (int *) acc_malloc (nbytes);
  b = (int *) malloc (nbytes);

  memset (&b[0], 0, nbytes);

  subr1_ (a);

  acc_memcpy_from_device (b, a, nbytes);

  for (i = 0; i < N; i++)
    {
      if (b[i] != i + 1)
	abort ();
    }

  memset (&b[0], 0, nbytes);

  subr2_ (a);

  acc_memcpy_from_device (b, a, nbytes);

  for (i = 0; i < N; i++)
    {
      if (b[i] != i + 1)
	abort ();
    }

  memset (&b[0], 0, nbytes);

  subr3 (a);

  acc_memcpy_from_device (b, a, nbytes);

  for (i = 0; i < N; i++)
    {
      if (b[i] != i + 1)
	abort ();
    }

  memset (&b[0], 0, nbytes);

  subr4 (a);

  acc_memcpy_from_device (b, a, nbytes);

  for (i = 0; i < N; i++)
    {
      if (b[i] != i + 1)
	abort ();
    }

  return 0;
}
subroutine subr1 (a)
  implicit none
  
  integer, dimension (8), intent (inout) :: a
  integer :: i

  i = 0

  !$acc parallel deviceptr (a)
    do i = 1, 8
      a(i) = i
    end do
  !$acc end parallel

end subroutine

subroutine subr2 (a)
  implicit none
  
  integer, dimension (8), intent (inout) :: a
  integer :: i

  i = 0

  !$acc data deviceptr (a)
  !$acc parallel
    do i = 1, 8
      a(i) = i
    end do
  !$acc end parallel
  !$acc end data

end subroutine

subroutine subr3 (a) bind (C, name = "subr3")
  use iso_c_binding
  implicit none
  
  integer (c_int32_t) , dimension (8), intent (inout) :: a
  integer :: i

  i = 0

  !$acc parallel deviceptr (a)
    do i = 1, 8
      a(i) = i
    end do
  !$acc end parallel

end subroutine

subroutine subr4 (a) bind (C, name = "subr4")
  use iso_c_binding
  implicit none
  
  integer (c_int32_t) , dimension (8), intent (inout) :: a
  integer :: i

  i = 0

  !$acc data deviceptr (a)
  !$acc parallel
    do i = 1, 8
      a(i) = i
    end do
  !$acc end parallel
  !$acc end data

end subroutine