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);
