diff --git a/libgomp/config/accel/target-indirect.c b/libgomp/config/accel/target-indirect.c index c60fd547cb674443edf9c1e7f401b406e3dffd2c..cfef1ddbc4950eb5dce47b4920cef9c41f3ed569 100644 --- a/libgomp/config/accel/target-indirect.c +++ b/libgomp/config/accel/target-indirect.c @@ -25,60 +25,73 @@ <http://www.gnu.org/licenses/>. */ #include <assert.h> +#include <string.h> #include "libgomp.h" -#define splay_tree_prefix indirect -#define splay_tree_c -#include "splay-tree.h" +struct indirect_map_t +{ + void *host_addr; + void *target_addr; +}; + +typedef struct indirect_map_t *hash_entry_type; + +static inline void * htab_alloc (size_t size) { return gomp_malloc (size); } +static inline void htab_free (void *ptr) { free (ptr); } + +#include "hashtab.h" + +static inline hashval_t +htab_hash (hash_entry_type element) +{ + return hash_pointer (element->host_addr); +} -volatile void **GOMP_INDIRECT_ADDR_MAP = NULL; +static inline bool +htab_eq (hash_entry_type x, hash_entry_type y) +{ + return x->host_addr == y->host_addr; +} -/* Use a splay tree to lookup the target address instead of using a - linear search. */ -#define USE_SPLAY_TREE_LOOKUP +void **GOMP_INDIRECT_ADDR_MAP = NULL; -#ifdef USE_SPLAY_TREE_LOOKUP +/* Use a hashtab to lookup the target address instead of using a linear + search. */ +#define USE_HASHTAB_LOOKUP -static struct indirect_splay_tree_s indirect_map; -static indirect_splay_tree_node indirect_array = NULL; +#ifdef USE_HASHTAB_LOOKUP -/* Build the splay tree used for host->target address lookups. */ +static htab_t indirect_htab = NULL; + +/* Build the hashtab used for host->target address lookups. */ void build_indirect_map (void) { size_t num_ind_funcs = 0; - volatile void **map_entry; - static int lock = 0; /* == gomp_mutex_t lock; gomp_mutex_init (&lock); */ + void **map_entry; if (!GOMP_INDIRECT_ADDR_MAP) return; - gomp_mutex_lock (&lock); - - if (!indirect_array) + if (!indirect_htab) { /* Count the number of entries in the NULL-terminated address map. */ for (map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry; map_entry += 2, num_ind_funcs++); - /* Build splay tree for address lookup. */ - indirect_array = gomp_malloc (num_ind_funcs * sizeof (*indirect_array)); - indirect_splay_tree_node array = indirect_array; + /* Build hashtab for address lookup. */ + indirect_htab = htab_create (num_ind_funcs); map_entry = GOMP_INDIRECT_ADDR_MAP; - for (int i = 0; i < num_ind_funcs; i++, array++) + for (int i = 0; i < num_ind_funcs; i++, map_entry += 2) { - indirect_splay_tree_key k = &array->key; - k->host_addr = (uint64_t) *map_entry++; - k->target_addr = (uint64_t) *map_entry++; - array->left = NULL; - array->right = NULL; - indirect_splay_tree_insert (&indirect_map, array); + struct indirect_map_t element = { *map_entry, NULL }; + hash_entry_type *slot = htab_find_slot (&indirect_htab, &element, + INSERT); + *slot = (hash_entry_type) map_entry; } } - - gomp_mutex_unlock (&lock); } void * @@ -88,15 +101,11 @@ GOMP_target_map_indirect_ptr (void *ptr) if (!ptr) return ptr; - assert (indirect_array); - - struct indirect_splay_tree_key_s k; - indirect_splay_tree_key node = NULL; - - k.host_addr = (uint64_t) ptr; - node = indirect_splay_tree_lookup (&indirect_map, &k); + assert (indirect_htab); - return node ? (void *) node->target_addr : ptr; + struct indirect_map_t element = { ptr, NULL }; + hash_entry_type entry = htab_find (indirect_htab, &element); + return entry ? entry->target_addr : ptr; } #else @@ -115,7 +124,7 @@ GOMP_target_map_indirect_ptr (void *ptr) assert (GOMP_INDIRECT_ADDR_MAP); - for (volatile void **map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry; + for (void **map_entry = GOMP_INDIRECT_ADDR_MAP; *map_entry; map_entry += 2) if (*map_entry == ptr) return (void *) *(map_entry + 1); diff --git a/libgomp/config/gcn/team.c b/libgomp/config/gcn/team.c index 61e9c616b678614e2207c00fc71e1fc6bfc8239a..bd3df448b5223aaa77330e0fc052b81b28a8febe 100644 --- a/libgomp/config/gcn/team.c +++ b/libgomp/config/gcn/team.c @@ -52,14 +52,15 @@ gomp_gcn_enter_kernel (void) { int threadid = __builtin_gcn_dim_pos (1); - /* Initialize indirect function support. */ - build_indirect_map (); - if (threadid == 0) { int numthreads = __builtin_gcn_dim_size (1); int teamid = __builtin_gcn_dim_pos(0); + /* Initialize indirect function support. */ + if (teamid == 0) + build_indirect_map (); + /* Set up the global state. Every team will do this, but that should be harmless. */ gomp_global_icv.nthreads_var = 16; diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c index 0cf5dad39cac899a86989ff38a34e87f112eb1f8..d5361917a24d267c4ff749f6908c6c4e6dd467f4 100644 --- a/libgomp/config/nvptx/team.c +++ b/libgomp/config/nvptx/team.c @@ -60,9 +60,6 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data) asm ("mov.u32 %0, %%tid.y;" : "=r" (tid)); asm ("mov.u32 %0, %%ntid.y;" : "=r" (ntids)); - /* Initialize indirect function support. */ - build_indirect_map (); - if (tid == 0) { gomp_global_icv.nthreads_var = ntids; @@ -74,6 +71,12 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data) nvptx_thrs = alloca (ntids * sizeof (*nvptx_thrs)); memset (nvptx_thrs, 0, ntids * sizeof (*nvptx_thrs)); + /* Initialize indirect function support. */ + unsigned int block_id; + asm ("mov.u32 %0, %%ctaid.x;" : "=r" (block_id)); + if (block_id == 0) + build_indirect_map (); + /* Find the low-latency heap details .... */ uint32_t *shared_pool; uint32_t shared_pool_size = 0; diff --git a/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c index 9fe190efce8fdbdcbbd06ccdd44b91f08dfd562e..545f1a9fcbffa3c45bcaa5ed0da47b39114a94ef 100644 --- a/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c +++ b/libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c @@ -17,17 +17,17 @@ int main (void) { switch (i % 3) { - case 0: fn_ptr[i] = &foo; - case 1: fn_ptr[i] = &bar; - case 2: fn_ptr[i] = &baz; + case 0: fn_ptr[i] = &foo; break; + case 1: fn_ptr[i] = &bar; break; + case 2: fn_ptr[i] = &baz; break; } expected += (*fn_ptr[i]) (); } -#pragma omp target teams distribute parallel for reduction(+: x) \ - map (to: fn_ptr) map (tofrom: x) - for (int i = 0; i < N; i++) - x += (*fn_ptr[i]) (); + #pragma omp target teams distribute parallel for \ + reduction (+: x) map (to: fn_ptr) map (tofrom: x) + for (int i = 0; i < N; i++) + x += (*fn_ptr[i]) (); return x - expected; } diff --git a/libgomp/testsuite/libgomp.fortran/declare-target-indirect-2.f90 b/libgomp/testsuite/libgomp.fortran/declare-target-indirect-2.f90 index 34dd27788931b6a52229541fcc21bf82032edd68..d3baa81dd079462a51dbaa6a820c4ee18ffa8135 100644 --- a/libgomp/testsuite/libgomp.fortran/declare-target-indirect-2.f90 +++ b/libgomp/testsuite/libgomp.fortran/declare-target-indirect-2.f90 @@ -1,5 +1,4 @@ ! { dg-do run } -! { dg-xfail-run-if "Requires libgomp bug fix pending review" { offload_device } } module m contains