diff mbox series

[OpenACC] Elaborate/simplify 'exit data' 'finalize' handling (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior)

Message ID 87y2vgz526.fsf@euler.schwinge.homeip.net
State New
Headers show
Series [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling (was: [OpenACC] Update OpenACC data clause semantics to the 2.5 behavior) | expand

Commit Message

Thomas Schwinge Dec. 13, 2019, 2:13 p.m. UTC 
Hi!

Julian, Tobias, regarding the following OpenACC 'exit data' 'finalize'
handling:

On 2018-05-25T13:01:58-0700, Cesar Philippidis <cesar@codesourcery.com> wrote:
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c

> @@ -10859,6 +10849,53 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)

> +  else if (TREE_CODE (expr) == OACC_EXIT_DATA
> +	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
> +			       OMP_CLAUSE_FINALIZE))
> +    {
> +      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
> +	 semantics apply to all mappings of this OpenACC directive.  */
> +      bool finalize_marked = false;
> +      for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
> +	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> +	  switch (OMP_CLAUSE_MAP_KIND (c))
> +	    {
> +	    case GOMP_MAP_FROM:
> +	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
> +	      finalize_marked = true;
> +	      break;
> +	    case GOMP_MAP_RELEASE:
> +	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
> +	      finalize_marked = true;
> +	      break;
> +	    default:
> +	      /* Check consistency: libgomp relies on the very first data
> +		 mapping clause being marked, so make sure we did that before
> +		 any other mapping clauses.  */
> +	      gcc_assert (finalize_marked);
> +	      break;
> +	    }
> +    }

> --- a/libgomp/oacc-parallel.c
> +++ b/libgomp/oacc-parallel.c

> @@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,

> +  /* Determine whether "finalize" semantics apply to all mappings of this
> +     OpenACC directive.  */
> +  bool finalize = false;
> +  if (mapnum > 0)
> +    {
> +      unsigned char kind = kinds[0] & 0xff;
> +      if (kind == GOMP_MAP_DELETE
> +	  || kind == GOMP_MAP_FORCE_FROM)
> +	finalize = true;
> +    }
> +

> @@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t mapnum,

>  	    switch (kind)
>  	      {
> -	      case GOMP_MAP_POINTER:
> -		gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> -					 == GOMP_MAP_FORCE_FROM,
> -					 async, 1);
> -		break;
> +	      case GOMP_MAP_RELEASE:
>  	      case GOMP_MAP_DELETE:
> -		acc_delete (hostaddrs[i], sizes[i]);
> +		if (acc_is_present (hostaddrs[i], sizes[i]))
> +		  {
> +		    if (finalize)
> +		      acc_delete_finalize (hostaddrs[i], sizes[i]);
> +		    else
> +		      acc_delete (hostaddrs[i], sizes[i]);
> +		  }
>  		break;
> +	      case GOMP_MAP_FROM:
>  	      case GOMP_MAP_FORCE_FROM:
> -		acc_copyout (hostaddrs[i], sizes[i]);
> +		if (finalize)
> +		  acc_copyout_finalize (hostaddrs[i], sizes[i]);
> +		else
> +		  acc_copyout (hostaddrs[i], sizes[i]);
>  		break;
>  	      default:
>  		gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
> @@ -385,10 +489,12 @@ GOACC_enter_exit_data (int device, size_t mapnum,
>  	  }
>  	else
>  	  {
> -	    gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] & 0xff)
> -				     == GOMP_MAP_FORCE_FROM, async, 3);
> +[...]
> +	    gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
> +				     finalize, pointer);

... does the attached patch "[OpenACC] Elaborate/simplify 'exit data'
'finalize' handling" (with "No functional changes") match your
understanding of what's going on?  If approving this patch, please
respond with "Reviewed-by: NAME <EMAIL>" so that your effort will be
recorded in the commit log, see <https://gcc.gnu.org/wiki/Reviewed-by>.

(It will be a separate discussion to change the 'GOMP_MAP_POINTER',
'GOMP_MAP_TO_PSET' stuff later on -- thinking about the changes from
Julian's big "OpenACC reference count overhaul" as well as
<https://gcc.gnu.org/PR92929> "OpenACC/OpenMP 'target' 'exit
data'/'update' optimizations".  That patch here is just meant to document
what's going at present, and simplify things as a preparation for other
changes.)


Grüße
 Thomas

Comments

Julian Brown Dec. 13, 2019, 11:34 p.m. UTC | #1 
On Fri, 13 Dec 2019 15:13:53 +0100
Thomas Schwinge <thomas@codesourcery.com> wrote:

> Hi!
> 
> Julian, Tobias, regarding the following OpenACC 'exit data' 'finalize'
> handling:
> 
> On 2018-05-25T13:01:58-0700, Cesar Philippidis
> <cesar@codesourcery.com> wrote:
> > --- a/gcc/gimplify.c
> > +++ b/gcc/gimplify.c  
> 
> > @@ -10859,6 +10849,53 @@ gimplify_omp_target_update (tree *expr_p,
> > gimple_seq *pre_p)  
> 
> > +  else if (TREE_CODE (expr) == OACC_EXIT_DATA
> > +	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
> > +			       OMP_CLAUSE_FINALIZE))
> > +    {
> > +      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that
> > "finalize"
> > +	 semantics apply to all mappings of this OpenACC
> > directive.  */
> > +      bool finalize_marked = false;
> > +      for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c =
> > OMP_CLAUSE_CHAIN (c))
> > +	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
> > +	  switch (OMP_CLAUSE_MAP_KIND (c))
> > +	    {
> > +	    case GOMP_MAP_FROM:
> > +	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
> > +	      finalize_marked = true;
> > +	      break;
> > +	    case GOMP_MAP_RELEASE:
> > +	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
> > +	      finalize_marked = true;
> > +	      break;
> > +	    default:
> > +	      /* Check consistency: libgomp relies on the very
> > first data
> > +		 mapping clause being marked, so make sure we did
> > that before
> > +		 any other mapping clauses.  */
> > +	      gcc_assert (finalize_marked);
> > +	      break;
> > +	    }
> > +    }  
> 
> > --- a/libgomp/oacc-parallel.c
> > +++ b/libgomp/oacc-parallel.c  
> 
> > @@ -286,6 +360,17 @@ GOACC_enter_exit_data (int device, size_t
> > mapnum,  
> 
> > +  /* Determine whether "finalize" semantics apply to all mappings
> > of this
> > +     OpenACC directive.  */
> > +  bool finalize = false;
> > +  if (mapnum > 0)
> > +    {
> > +      unsigned char kind = kinds[0] & 0xff;
> > +      if (kind == GOMP_MAP_DELETE
> > +	  || kind == GOMP_MAP_FORCE_FROM)
> > +	finalize = true;
> > +    }
> > +  
> 
> > @@ -360,22 +458,28 @@ GOACC_enter_exit_data (int device, size_t
> > mapnum,  
> 
> >  	    switch (kind)
> >  	      {
> > -	      case GOMP_MAP_POINTER:
> > -		gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] &
> > 0xff)
> > -					 == GOMP_MAP_FORCE_FROM,
> > -					 async, 1);
> > -		break;
> > +	      case GOMP_MAP_RELEASE:
> >  	      case GOMP_MAP_DELETE:
> > -		acc_delete (hostaddrs[i], sizes[i]);
> > +		if (acc_is_present (hostaddrs[i], sizes[i]))
> > +		  {
> > +		    if (finalize)
> > +		      acc_delete_finalize (hostaddrs[i], sizes[i]);
> > +		    else
> > +		      acc_delete (hostaddrs[i], sizes[i]);
> > +		  }
> >  		break;
> > +	      case GOMP_MAP_FROM:
> >  	      case GOMP_MAP_FORCE_FROM:
> > -		acc_copyout (hostaddrs[i], sizes[i]);
> > +		if (finalize)
> > +		  acc_copyout_finalize (hostaddrs[i], sizes[i]);
> > +		else
> > +		  acc_copyout (hostaddrs[i], sizes[i]);
> >  		break;
> >  	      default:
> >  		gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED
> > kind 0x%.2x", @@ -385,10 +489,12 @@ GOACC_enter_exit_data (int
> > device, size_t mapnum, }
> >  	else
> >  	  {
> > -	    gomp_acc_remove_pointer (hostaddrs[i], (kinds[i] &
> > 0xff)
> > -				     == GOMP_MAP_FORCE_FROM,
> > async, 3); +[...]
> > +	    gomp_acc_remove_pointer (hostaddrs[i], sizes[i],
> > copyfrom, async,
> > +				     finalize, pointer);  
> 
> ... does the attached patch "[OpenACC] Elaborate/simplify 'exit data'
> 'finalize' handling" (with "No functional changes") match your
> understanding of what's going on?

Your patch looks OK to me, FWIW. As you mentioned at some point though,
it might be good to get rid of this style of finalize handling,
replacing it with a flag passed to GOACC_exit_data -- presuming that at
the same time, we separate out the needlessly-dual-purpose
GOACC_enter_exit_data API entry point into "enter" and "exit" halves.

Thanks,

Julian
Thomas Schwinge Dec. 18, 2019, 5:09 p.m. UTC | #2 
Hi!

On 2019-12-13T23:34:15+0000, Julian Brown <julian@codesourcery.com> wrote:
> On Fri, 13 Dec 2019 15:13:53 +0100
> Thomas Schwinge <thomas@codesourcery.com> wrote:
>> Julian, Tobias, regarding the following OpenACC 'exit data' 'finalize'
>> handling:
>> 
>> On 2018-05-25T13:01:58-0700, Cesar Philippidis
>> <cesar@codesourcery.com> wrote:
>> > [...]
>> 
>> ... does the attached patch "[OpenACC] Elaborate/simplify 'exit data'
>> 'finalize' handling" (with "No functional changes") match your
>> understanding of what's going on?
>
> Your patch looks OK to me, FWIW.

Thanks for the review.

See attached "[OpenACC] Elaborate/simplify 'exit data' 'finalize'
handling"; committed to trunk in r279531.


> As you mentioned at some point though,
> it might be good to get rid of this style of finalize handling,
> replacing it with a flag passed to GOACC_exit_data

Actually, as recently discussed in a different context, I'm now steering
into the opposite direction: make all that explicit in the mapping kinds.
The reason is: while it's true that currently the OpenACC 'finalize'
clause applies to all data clauses on the directive, that's not the only
way: it might at some point become a flag for each individual clause
('copyout(finalize: [...])', or something like that) -- like OpenMP
already does, as far as I remember, so we need to support that
per-mapping kind anyway.


> -- presuming that at
> the same time, we separate out the needlessly-dual-purpose
> GOACC_enter_exit_data API entry point into "enter" and "exit" halves.

Indeed -- while that one's not a problem, it's still a bit "uh".  But,
for the sake of backwards compatibility..., it'll stay this way until we
do any other breaking changes.


Grüße
 Thomas
diff mbox series

Patch

From 283577c63b374c3e368e3c0b68b90e19085f193c Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <thomas@codesourcery.com>
Date: Fri, 13 Dec 2019 13:56:51 +0100
Subject: [PATCH] [OpenACC] Elaborate/simplify 'exit data' 'finalize' handling

No functional changes.
---
 gcc/gimplify.c                                | 23 +++++++++++--------
 gcc/testsuite/c-c++-common/goacc/finalize-1.c | 11 ++++++++-
 gcc/testsuite/gfortran.dg/goacc/finalize-1.f  | 10 ++++++++
 libgomp/oacc-mem.c                            | 14 +++--------
 4 files changed, 36 insertions(+), 22 deletions(-)

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 9073680cb31..60a80cb8098 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -12738,27 +12738,30 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
 			       OMP_CLAUSE_FINALIZE))
     {
-      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
-	 semantics apply to all mappings of this OpenACC directive.  */
-      bool finalize_marked = false;
+      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote "finalize"
+	 semantics.  */
+      bool have_clause = false;
       for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
 	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
 	  switch (OMP_CLAUSE_MAP_KIND (c))
 	    {
 	    case GOMP_MAP_FROM:
 	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_FROM);
-	      finalize_marked = true;
+	      have_clause = true;
 	      break;
 	    case GOMP_MAP_RELEASE:
 	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
-	      finalize_marked = true;
+	      have_clause = true;
 	      break;
-	    default:
-	      /* Check consistency: libgomp relies on the very first data
-		 mapping clause being marked, so make sure we did that before
-		 any other mapping clauses.  */
-	      gcc_assert (finalize_marked);
+	    case GOMP_MAP_POINTER:
+	    case GOMP_MAP_TO_PSET:
+	      /* TODO PR92929: we may see these here, but they'll always follow
+		 one of the clauses above, and will be handled by libgomp as
+		 one group, so no handling required here.  */
+	      gcc_assert (have_clause);
 	      break;
+	    default:
+	      gcc_unreachable ();
 	    }
     }
   stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
index 94820290b94..3d64b2e7cb3 100644
--- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -4,8 +4,10 @@ 
 
 extern int del_r;
 extern float del_f[3];
+extern char *del_f_p;
 extern double cpo_r[8];
 extern long cpo_f;
+extern char *cpo_f_p;
 
 void f ()
 {
@@ -17,6 +19,10 @@  void f ()
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
    { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
 
+#pragma acc exit data finalize delete (del_f_p[2:5])
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
+   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
+
 #pragma acc exit data copyout (cpo_r)
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
    { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
@@ -24,5 +30,8 @@  void f ()
 #pragma acc exit data copyout (cpo_f) finalize
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } }
    { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
-}
 
+#pragma acc exit data copyout (cpo_f_p[4:10]) finalize
+/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
+   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index 5c7a921a2e3..ca642156e9f 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -6,8 +6,10 @@ 
       IMPLICIT NONE
       INTEGER :: del_r
       REAL, DIMENSION (3) :: del_f
+      INTEGER (1), DIMENSION (:), ALLOCATABLE :: del_f_p
       DOUBLE PRECISION, DIMENSION (8) :: cpo_r
       LOGICAL :: cpo_f
+      INTEGER (1), DIMENSION (:), ALLOCATABLE :: cpo_f_p
 
 !$ACC EXIT DATA DELETE (del_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_r\\);$" 1 "original" } }
@@ -17,6 +19,10 @@ 
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
+!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+
 !$ACC EXIT DATA COPYOUT (cpo_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
@@ -24,4 +30,8 @@ 
 !$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
+
+!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 311f9585f77..291ef9192b9 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -1058,17 +1058,6 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
   thr = goacc_thread ();
   acc_dev = thr->dev;
 
-  /* Determine whether "finalize" semantics apply to all mappings of this
-     OpenACC directive.  */
-  bool finalize = false;
-  if (mapnum > 0)
-    {
-      unsigned char kind = kinds[0] & 0xff;
-      if (kind == GOMP_MAP_DELETE
-	  || kind == GOMP_MAP_FORCE_FROM)
-	finalize = true;
-    }
-
   /* Determine if this is an "acc enter data".  */
   for (i = 0; i < mapnum; ++i)
     {
@@ -1221,6 +1210,9 @@  GOACC_enter_exit_data (int flags_m, size_t mapnum, void **hostaddrs,
       {
 	unsigned char kind = kinds[i] & 0xff;
 
+	bool finalize = (kind == GOMP_MAP_DELETE
+			 || kind == GOMP_MAP_FORCE_FROM);
+
 	int pointer = find_pointer (i, mapnum, kinds);
 
 	if (!pointer)
-- 
2.17.1