[Patch] OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]

[Patch] OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]

[Tobias Burnus]

[-- Attachment #1.1: Type: text/plain, Size: 946 bytes --]

Using dummy procedures in a target region with 'defaultmap(none)' leads to:

   Error: 'g' not specified in enclosing 'target'

and this cannot be fixed by using 'firstprivate' as non-pointer dummy routines
are rejected as "Error: Object 'g' is not a variable".

Fixed by doing the same for mapping as for data sharing: using predetermined
firstprivate.

BTW: Only since GCC 14, 'declare target indirect' makes it possible to
simply use dummy procedures and procedures pointers in a target region.

Comments? Suggestions?

Tobias

PS: Procedure pointers aren't variables either, but they act even more like
variables as they permit changing pointer association such that '(first)private'
vs. 'shared'/'map' can both make sense. — GCC accepts those in (nearly) all clauses,
ifort only in (first)private while flang not at all. The spec is somewhat silent
about it. This is tracked in the same PR (PR114283) and in the specification
issue #3823.

[-- Attachment #2: dummyproc.diff --]
[-- Type: text/x-patch, Size: 2504 bytes --]

OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]

Dummy procedures look similar to variables but aren't - neither in Fortran
nor in OpenMP. As the middle end sees PARM_DECLs, mark them as predetermined
firstprivate for mapping (as already done in gfc_omp_predetermined_sharing).

This does not address the isses related to procedure pointers, which are
still discussed on spec level [see PR].

	PR fortran/114283

gcc/fortran/ChangeLog:

	* trans-openmp.cc (gfc_omp_predetermined_mapping): Map dummy
	procedures as firstprivate.

gcc/testsuite/ChangeLog:

	* gfortran.dg/gomp/target4.f90: New test.

 gcc/fortran/trans-openmp.cc                |  9 +++++++++
 gcc/testsuite/gfortran.dg/gomp/target4.f90 | 18 ++++++++++++++++++
 2 files changed, 27 insertions(+)

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index a2bf15665b3..1dba47126ed 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -343,6 +343,15 @@ gfc_omp_predetermined_mapping (tree decl)
 	    && GFC_DECL_SAVED_DESCRIPTOR (decl)))
     return OMP_CLAUSE_DEFAULTMAP_TO;
 
+  /* Dummy procedures aren't considered variables by OpenMP, thus are
+     disallowed in OpenMP clauses.  They are represented as PARM_DECLs
+     in the middle-end, so return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE here
+     to avoid complaining about their uses with defaultmap(none).  */
+  if (TREE_CODE (decl) == PARM_DECL
+      && TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+      && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == FUNCTION_TYPE)
+    return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE;
+
   /* These are either array or derived parameters, or vtables.  */
   if (VAR_P (decl) && TREE_READONLY (decl)
       && (TREE_STATIC (decl) || DECL_EXTERNAL (decl)))
diff --git a/gcc/testsuite/gfortran.dg/gomp/target4.f90 b/gcc/testsuite/gfortran.dg/gomp/target4.f90
new file mode 100644
index 00000000000..09364e707f1
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/gomp/target4.f90
@@ -0,0 +1,18 @@
+! { dg-additional-options "-fdump-tree-gimple" }
+
+! PR fortran/114283
+
+! { dg-final { scan-tree-dump "#pragma omp parallel default\\(none\\) firstprivate\\(g\\)" "gimple" } }
+! { dg-final { scan-tree-dump "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) defaultmap\\(none\\) firstprivate\\(g\\)" "gimple" } }
+
+subroutine f(g)
+procedure() :: g
+
+!$omp parallel default(none)
+  call g
+!$omp end parallel
+
+!$omp target defaultmap(none)
+  call g
+!$omp end target
+end

Re: [Patch] OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]

[Jakub Jelinek]

On Mon, Mar 11, 2024 at 11:07:46AM +0100, Tobias Burnus wrote:
> Using dummy procedures in a target region with 'defaultmap(none)' leads to:
> 
>   Error: 'g' not specified in enclosing 'target'
> 
> and this cannot be fixed by using 'firstprivate' as non-pointer dummy routines
> are rejected as "Error: Object 'g' is not a variable".
> 
> Fixed by doing the same for mapping as for data sharing: using predetermined
> firstprivate.
> 
> BTW: Only since GCC 14, 'declare target indirect' makes it possible to
> simply use dummy procedures and procedures pointers in a target region.

So firstprivate clause handling remaps them then if declare target indirect
is used?
If so, the patch looks reasonable to me.

	Jakub

Re: [Patch] OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]

[Tobias Burnus]

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

Jakub Jelinek wrote:

> So firstprivate clause handling remaps them then if declare target indirect
> is used? If so, the patch looks reasonable to me.

[I have now updated the patch to turn the testcase to ensure
that is also keeps works at runtime.]

OpenMP leaves it a bit open when the remapping has to happen,
but one can construct cases – in particular with unified-shared memory –
where it is not possible to do this upon entry to a target region.

Thus, it has to be done when the function is invoked, e.g.

i = (*g) ();

is turned (in the target region but only on the device side) into

i = (*GOMP_target_map_indirect_ptr (g)) ();

Thus, as long as the host pointer value is transferred to the device,
it works – as the lookup is done on the device side. Directly using a
device address (remap when mapping to the target) will also not shorten
the lookup, i.e. there is no need for it.

Does it still look reasonable to you?

Tobias

PS: The current OpenMP specification, it is listed mainly described via
the glossary (newest change is the addition of dummy procedure):

"indirect device invocation – An indirect call to the _device_ version of 
a _procedure_ on a _device_ other than the _host-device_, through a 
function pointer (C/C++), a pointer to a member function (C++), a dummy 
procedure (Fortran), or a procedure pointer (Fortran) that refers to the 
host version of the _procedure_."

[-- Attachment #2: dummyproc-v2.diff --]
[-- Type: text/x-patch, Size: 3176 bytes --]

OpenMP/Fortran: Fix defaultmap(none) issue with dummy procedures [PR114283]

Dummy procedures look similar to variables but aren't - neither in Fortran
nor in OpenMP. As the middle end sees PARM_DECLs, mark them as predetermined
firstprivate for mapping (as already done in gfc_omp_predetermined_sharing).

This does not address the isses related to procedure pointers, which are
still discussed on spec level [see PR].

	PR fortran/114283

gcc/fortran/ChangeLog:

	* trans-openmp.cc (gfc_omp_predetermined_mapping): Map dummy
	procedures as firstprivate.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/declare-target-indirect-4.f90: New test.

 gcc/fortran/trans-openmp.cc                        |  9 +++++
 .../libgomp.fortran/declare-target-indirect-4.f90  | 43 ++++++++++++++++++++++
 2 files changed, 52 insertions(+)

diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc
index a2bf15665b3..1dba47126ed 100644
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -343,6 +343,15 @@ gfc_omp_predetermined_mapping (tree decl)
 	    && GFC_DECL_SAVED_DESCRIPTOR (decl)))
     return OMP_CLAUSE_DEFAULTMAP_TO;
 
+  /* Dummy procedures aren't considered variables by OpenMP, thus are
+     disallowed in OpenMP clauses.  They are represented as PARM_DECLs
+     in the middle-end, so return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE here
+     to avoid complaining about their uses with defaultmap(none).  */
+  if (TREE_CODE (decl) == PARM_DECL
+      && TREE_CODE (TREE_TYPE (decl)) == POINTER_TYPE
+      && TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) == FUNCTION_TYPE)
+    return OMP_CLAUSE_DEFAULTMAP_FIRSTPRIVATE;
+
   /* These are either array or derived parameters, or vtables.  */
   if (VAR_P (decl) && TREE_READONLY (decl)
       && (TREE_STATIC (decl) || DECL_EXTERNAL (decl)))
diff --git a/libgomp/testsuite/libgomp.fortran/declare-target-indirect-4.f90 b/libgomp/testsuite/libgomp.fortran/declare-target-indirect-4.f90
new file mode 100644
index 00000000000..43f4295494c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.fortran/declare-target-indirect-4.f90
@@ -0,0 +1,43 @@
+! { dg-additional-options "-fdump-tree-gimple" }
+
+! PR fortran/114283
+
+! { dg-final { scan-tree-dump "#pragma omp parallel shared\\(i\\) if\\(0\\) default\\(none\\) firstprivate\\(g\\)" "gimple" } }
+! { dg-final { scan-tree-dump "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) firstprivate\\(h\\) map\\(from:j \\\[len: 4\\\]\\) defaultmap\\(none\\)" "gimple" } }
+
+
+module m
+  implicit none (type, external)
+  !$omp declare target indirect enter(f1, f2)
+contains
+  integer function f1 ()
+    f1 = 99
+  end
+  integer function f2 ()
+    f2 = 89
+  end
+end module m
+
+use m
+implicit none (type, external)
+call sub1(f1)
+call sub2(f2)
+contains
+  subroutine sub1(g)
+    procedure(integer) :: g
+    integer :: i
+    !$omp parallel default(none) if(.false.) shared(i)
+      i = g ()
+    !$omp end parallel
+    if (i /= 99) stop 1
+  end
+
+  subroutine sub2(h)
+    procedure(integer) :: h
+    integer :: j
+    !$omp target defaultmap(none) map(from:j)
+      j = h ()
+    !$omp end target
+    if (j /= 89) stop 1
+  end
+end