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.
commit 3afc4e592a6d8a796ec0c44bb8dc808b1392fd29
Author: Julian Brown <jul...@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);