diff options
| author | Kwok Cheung Yeung <kcyeung@baylibre.com> | 2024-03-22 18:07:43 +0000 |
|---|---|---|
| committer | Kwok Cheung Yeung <kcyeung@baylibre.com> | 2024-03-22 18:09:40 +0000 |
| commit | 637e76b90e8b045c5e25206a41e3be55deace8d5 (patch) | |
| tree | a33446e0aea1cf91b0628c4d3f0abce0987f30b6 /libgomp/testsuite | |
| parent | 65107faad79354a75844d8dba053be6509200504 (diff) | |
| download | gcc-637e76b90e8b045c5e25206a41e3be55deace8d5.zip gcc-637e76b90e8b045c5e25206a41e3be55deace8d5.tar.gz gcc-637e76b90e8b045c5e25206a41e3be55deace8d5.tar.bz2 | |
openmp: Change to using a hashtab to lookup offload target addresses for indirect function calls
A splay-tree was previously used to lookup equivalent target addresses
for a given host address on offload targets. However, as splay-trees can
modify their structure on lookup, they are not suitable for concurrent
access from separate teams/threads without some form of locking. This
patch changes the lookup data structure to a hashtab instead, which does
not have these issues.
The call to build_indirect_map to initialize the data structure is now
called from just the first thread of the first team to avoid redundant
calls to this function.
2024-03-22 Kwok Cheung Yeung <kcyeung@baylibre.com>
libgomp/
* config/accel/target-indirect.c: Include string.h and hashtab.h.
Remove include of splay-tree.h. Update comments.
(splay_tree_prefix, splay_tree_c): Delete.
(struct indirect_map_t): New.
(hash_entry_type, htab_alloc, htab_free, htab_hash, htab_eq): New.
(GOMP_INDIRECT_ADD_MAP): Remove volatile qualifier.
(USE_SPLAY_TREE_LOOKUP): Rename to...
(USE_HASHTAB_LOOKUP): ..this.
(indirect_map, indirect_array): Delete.
(indirect_htab): New.
(build_indirect_map): Remove locking. Build indirect map using
hashtab.
(GOMP_target_map_indirect_ptr): Use indirect_htab to lookup target
address.
(GOMP_target_map_indirect_ptr): Remove volatile qualifier.
* config/gcn/team.c (gomp_gcn_enter_kernel): Call build_indirect_map
from first thread of first team only.
* config/nvptx/team.c (gomp_nvptx_main): Likewise.
* testsuite/libgomp.c-c++-common/declare-target-indirect-2.c (main):
Add missing break statements.
* testsuite/libgomp.fortran/declare-target-indirect-2.f90: Remove
xfail.
Diffstat (limited to 'libgomp/testsuite')
| -rw-r--r-- | libgomp/testsuite/libgomp.c-c++-common/declare-target-indirect-2.c | 14 | ||||
| -rw-r--r-- | libgomp/testsuite/libgomp.fortran/declare-target-indirect-2.f90 | 1 |
2 files changed, 7 insertions, 8 deletions
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 9fe190e..545f1a9 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 34dd277..d3baa81 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 |