diff mbox

[gomp4] Remove gomp_map_vars mem_map argument

Message ID 20141028165222.2f9a415a@octopus
State New
Headers show

Commit Message

Julian Brown Oct. 28, 2014, 4:52 p.m. UTC 
Hi,

This patch removes the now-redundant gomp_memory_mapping argument from
gomp_map_vars, introduced when OpenACC kept the structure in question
in a different place from OpenMP. Both now keep the memory map in the
gomp_device_descr structure, so there's no need to pass both that and
the memory map to the function explicitly.

OK for gomp4 branch?

Thanks,

Julian

ChangeLog

    libgomp/
    * target.c (gomp_map_vars): Remove MM argument.
    (GOMP_target, GOMP_target_data): Update calls to gomp_map_vars.
    * oacc-mem.c (acc_map_data, present_create_copy): Update calls to
    gomp_map_vars.
    * oacc-parallel.c (GOACC_parallel, GOACC_data_start): Likewise.
    * target.h (gomp_map_vars): Update prototype.

Comments

Julian Brown Oct. 28, 2014, 4:54 p.m. UTC | #1 
On Tue, 28 Oct 2014 16:52:22 +0000
Julian Brown <julian@codesourcery.com> wrote:

> Hi,
> 
> This patch removes the now-redundant gomp_memory_mapping argument from
> gomp_map_vars, introduced when OpenACC kept the structure in question
> in a different place from OpenMP. Both now keep the memory map in the
> gomp_device_descr structure, so there's no need to pass both that and
> the memory map to the function explicitly.
> 
> OK for gomp4 branch?

Forgot to say: this patch and the previous three have been tested with
no regressions (alongside a version of Bernd's PTX support patches) on
the gomp4 branch (libgomp tests).

Julian
Thomas Schwinge Oct. 28, 2014, 7:42 p.m. UTC | #2 
Hi Julian!

On Tue, 28 Oct 2014 16:52:22 +0000, Julian Brown <julian@codesourcery.com> wrote:
> This patch removes the now-redundant gomp_memory_mapping argument from
> gomp_map_vars, introduced when OpenACC kept the structure in question
> in a different place from OpenMP. Both now keep the memory map in the
> gomp_device_descr structure, so there's no need to pass both that and
> the memory map to the function explicitly.

Makes sense.

> OK for gomp4 branch?

Yes, thanks!  (But I have not reviewed this code/change in detail.)

>     libgomp/
>     * target.c (gomp_map_vars): Remove MM argument.
>     (GOMP_target, GOMP_target_data): Update calls to gomp_map_vars.
>     * oacc-mem.c (acc_map_data, present_create_copy): Update calls to
>     gomp_map_vars.
>     * oacc-parallel.c (GOACC_parallel, GOACC_data_start): Likewise.
>     * target.h (gomp_map_vars): Update prototype.


Grüße,
 Thomas
diff mbox

Patch

commit 3afc4e592a6d8a796ec0c44bb8dc808b1392fd29
Author: Julian Brown <julian@codesourcery.com>
Date:   Tue Oct 28 09:17:01 2014 -0700

    Remove gomp_map_vars mem_map argument

diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index d812f72..582a1e0 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -257,7 +257,7 @@  acc_map_data (void *h, void *d, size_t s)
       if (d != h)
         gomp_fatal ("cannot map data on shared-memory system");
 
-      tgt = gomp_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, false);
+      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, false);
     }
   else
     {
@@ -275,9 +275,8 @@  acc_map_data (void *h, void *d, size_t s)
 	gomp_fatal ("device address [%p, +%d] is already mapped", (void *)d,
 		    (int)s);
 
-      tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev,
-			   &acc_dev->mem_map, mapnum, &hostaddrs,
-			   &devaddrs, &sizes, &kinds, true, false);
+      tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, &devaddrs, &sizes,
+			   &kinds, true, false);
     }
 
   tgt->prev = acc_dev->openacc.data_environ;
@@ -383,9 +382,8 @@  present_create_copy (unsigned f, void *h, size_t s)
       else
         kinds = GOMP_MAP_ALLOC;
 
-      tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev,
-			   &acc_dev->mem_map, mapnum, &hostaddrs,
-			   NULL, &s, &kinds, true, false);
+      tgt = gomp_map_vars (acc_dev, mapnum, &hostaddrs, NULL, &s, &kinds, true,
+			   false);
 
       gomp_mutex_lock (&acc_dev->mem_map.lock);
 
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index b787df7..1639244 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -173,9 +173,8 @@  GOACC_parallel (int device, void (*fn) (void *), const void *openmp_target,
   else
     tgt_fn = (void (*)) fn;
 
-  tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev,
-		       &acc_dev->mem_map, mapnum, hostaddrs,
-		       NULL, sizes, kinds, true, false);
+  tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
+		       false);
 
   devaddrs = alloca (sizeof (void *) * mapnum);
   for (i = 0; i < mapnum; i++)
@@ -217,7 +216,7 @@  GOACC_data_start (int device, const void *openmp_target, size_t mapnum,
   if ((acc_dev->capabilities & TARGET_CAP_SHARED_MEM)
       || !if_clause_condition_value)
     {
-      tgt = gomp_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, true, false);
+      tgt = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, true, false);
       tgt->prev = thr->mapped_data;
       thr->mapped_data = tgt;
 
@@ -225,9 +224,8 @@  GOACC_data_start (int device, const void *openmp_target, size_t mapnum,
     }
 
   gomp_notify ("  %s: prepare mappings\n", __FUNCTION__);
-  tgt = gomp_map_vars ((struct gomp_device_descr *) acc_dev,
-		       &acc_dev->mem_map, mapnum, hostaddrs,
-		       NULL, sizes, kinds, true, false);
+  tgt = gomp_map_vars (acc_dev, mapnum, hostaddrs, NULL, sizes, kinds, true,
+		       false);
   gomp_notify ("  %s: mappings prepared\n", __FUNCTION__);
   tgt->prev = thr->mapped_data;
   thr->mapped_data = tgt;
diff --git a/libgomp/target.c b/libgomp/target.c
index 615ba6b..507488e 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -134,14 +134,14 @@  get_kind (bool is_openacc, void *kinds, int idx)
 }
 
 attribute_hidden struct target_mem_desc *
-gomp_map_vars (struct gomp_device_descr *devicep,
-	       struct gomp_memory_mapping *mm, size_t mapnum,
-	       void **hostaddrs, void **devaddrs, size_t *sizes,
-	       void *kinds, bool is_openacc, bool is_target)
+gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
+	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
+	       bool is_openacc, bool is_target)
 {
   size_t i, tgt_align, tgt_size, not_found_cnt = 0;
   const int rshift = is_openacc ? 8 : 3;
   const int typemask = is_openacc ? 0xff : 0x7;
+  struct gomp_memory_mapping *mm = &devicep->mem_map;
   struct splay_tree_key_s cur_node;
   struct target_mem_desc *tgt
     = gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
@@ -861,8 +861,8 @@  GOMP_target (int device, void (*fn) (void *), const void *openmp_target,
   gomp_mutex_unlock (&mm->lock);
 
   struct target_mem_desc *tgt_vars
-    = gomp_map_vars (devicep, &devicep->mem_map, mapnum, hostaddrs, NULL,
-		     sizes, kinds, false, true);
+    = gomp_map_vars (devicep, mapnum, hostaddrs, NULL, sizes, kinds, false,
+		     true);
   struct gomp_thread old_thr, *thr = gomp_thread ();
   old_thr = *thr;
   memset (thr, '\0', sizeof (*thr));
@@ -901,17 +901,15 @@  GOMP_target_data (int device, const void *openmp_target, size_t mapnum,
 	     new #pragma omp target data, otherwise GOMP_target_end_data
 	     would get out of sync.  */
 	  struct target_mem_desc *tgt
-	    = gomp_map_vars (NULL, NULL, 0, NULL, NULL, NULL, NULL, false,
-			     false);
+	    = gomp_map_vars (NULL, 0, NULL, NULL, NULL, NULL, false, false);
 	  tgt->prev = icv->target_data;
 	  icv->target_data = tgt;
 	}
       return;
     }
 
-  struct target_mem_desc *tgt
-    = gomp_map_vars (devicep, &devicep->mem_map, mapnum, hostaddrs, NULL, sizes,
-		     kinds, false, false);
+  struct target_mem_desc *tgt = gomp_map_vars (devicep, mapnum, hostaddrs, NULL,
+					       sizes, kinds, false, false);
   struct gomp_task_icv *icv = gomp_icv (true);
   tgt->prev = icv->target_data;
   icv->target_data = tgt;
diff --git a/libgomp/target.h b/libgomp/target.h
index b5773e2..d4c1120 100644
--- a/libgomp/target.h
+++ b/libgomp/target.h
@@ -199,10 +199,9 @@  struct gomp_device_descr
 };
 
 extern struct target_mem_desc *
-gomp_map_vars (struct gomp_device_descr *devicep,
-	       struct gomp_memory_mapping *mm, size_t mapnum,
-	       void **hostaddrs, void **devaddrs, size_t *sizes,
-	       void *kinds, bool is_openacc, bool is_target);
+gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
+	       void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
+	       bool is_openacc, bool is_target);
 
 extern void
 gomp_copy_from_async (struct target_mem_desc *tgt);