[gcc r17-2448] OpenMP: uses_allocators - fix Fortran memspace use + multiple clauses [PR126279]
Tobias Burnus via Gcc-cvs <[email protected]>
| Newsgroups | gmane.comp.gcc.cvs |
|---|---|
| Message-ID | <[email protected]> |
https://gcc.gnu.org/g:9e2e40f4ed8f0797963a0166d2568d23c2d25024 commit r17-2448-g9e2e40f4ed8f0797963a0166d2568d23c2d25024 Author: Tobias Burnus <[email protected]> Date: Thu Jul 16 12:11:38 2026 +0200 OpenMP: uses_allocators - fix Fortran memspace use + multiple clauses [PR126279] The memspace modifier to uses_allocators was converted in Fortran not to an integer number but to an external variable - which failed then at link time. Additionally, it turned out that - at least with host fallback - multiple uses_allocator clauses aren't supported as newer ones overrode previous ones; shifting it by gomp_omp_allocator_data_size data helped. OpenMP requires 'a constant array, have constant values' for C/C++ and for Fortran a 'named constant of rank one'; it turned out that the FEs already generate, e.g. static struct omp_alloctrait_t traits[1] = {{.key=2, .value=1024}}; which can be directly passed to the runtime library - no need to expand the initializer values into a new variable and then take the address of that one. For C/C++, this might be new since commit r17-2414-g9e9e22587f7954 that set calls 'mark_exp_read' to avoid an 'unused but set' warning for the traits variable [not checked] - but in any case it is now consistent. PR fortran/126279 gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_trans_omp_clauses): Convert memspace constant to an integer not to an external variable. gcc/ChangeLog: * omp-low.cc (lower_omp_target): Avoid creating a second static const variable for the trait array used by uses_allocators. libgomp/ChangeLog: * target.c (copy_firstprivate_data): After storing users_allocators' allocator, shift the offset by its storage size to avoid overriding. * testsuite/libgomp.fortran/uses_allocators-1.f90: Use kind c_size_t instead of kind 8 for 32bit compatibility. * testsuite/libgomp.fortran/uses_allocators-4.f90: Update expected tree dump. * testsuite/libgomp.fortran/uses_allocators-7.f90: Extend testcase now that user-defined allocators are supported by uses_allocators. Diff: --- gcc/fortran/trans-openmp.cc | 8 ++++++-- gcc/omp-low.cc | 12 ------------ libgomp/target.c | 1 + .../libgomp.fortran/uses_allocators-1.f90 | 2 +- .../libgomp.fortran/uses_allocators-4.f90 | 8 ++++---- .../libgomp.fortran/uses_allocators-7.f90 | 22 +++++++++++++++++++--- 6 files changed, 31 insertions(+), 22 deletions(-) diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 7243301cae0b..aa4e8ebfa239 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -5494,9 +5494,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_USES_ALLOCATORS_ALLOCATOR(node) = t; if (n->u.memspace_sym) { + gcc_checking_assert (n->u.memspace_sym->attr.flavor + == FL_PARAMETER + && !n->u.memspace_sym->attr.dimension); n->u.memspace_sym->attr.referenced = true; - OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (node) - = gfc_get_symbol_decl (n->u.memspace_sym); + gfc_init_se (&se, NULL); + gfc_conv_expr (&se, n->u.memspace_sym->value); + OMP_CLAUSE_USES_ALLOCATORS_MEMSPACE (node) = se.expr; } if (n->u2.traits_sym) { diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index 33082fba1a36..4bab2c56a23d 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -14051,21 +14051,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) ntraits = integer_zero_node; traits_var = null_pointer_node; } - else if (DECL_INITIAL (traits)) - { - location_t loc = OMP_CLAUSE_LOCATION (c); - ntraits = array_type_nelts_top (TREE_TYPE (traits)); - tree t = DECL_INITIAL (traits); - t = get_initialized_tmp_var (t, &ilist, NULL); - traits_var = build_fold_addr_expr_loc (loc, t); - } else { - /* This happens for VLAs, which probably aren't useful - because they can't be const initialized in the same - scope.... is there something else? */ location_t loc = OMP_CLAUSE_LOCATION (c); - gcc_assert (TREE_CODE (TREE_TYPE (traits)) == ARRAY_TYPE); ntraits = array_type_nelts_top (TREE_TYPE (traits)); traits_var = build_fold_addr_expr_loc (loc, traits); } diff --git a/libgomp/target.c b/libgomp/target.c index e49eacc00713..0f5091521200 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -3282,6 +3282,7 @@ copy_firstprivate_data (char *tgt, size_t mapnum, void **hostaddrs, tgt + tgt_size, hostaddrs[i]); hostaddrs[i] = (void *) allocator; + tgt_size += gomp_omp_allocator_data_size; } } diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 index cdabbed04e4f..329723a065f5 100644 --- a/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-1.f90 @@ -10,7 +10,7 @@ integer(omp_allocator_handle_kind) :: var !$omp target uses_allocators(memspace(omp_default_mem_space), traits(mem2) : var) defaultmap(none) block; type(c_ptr) ::c -c = omp_alloc(omp_default_mem_space, 20_8) +c = omp_alloc(omp_default_mem_space, 20_c_size_t) end block !$omp target uses_allocators(omp_default_mem_alloc, var(mem2)) ! { dg-warning "The specification of arguments to 'uses_allocators' at \\(1\\) where each item is of the form 'allocator\\(traits\\)' is deprecated since OpenMP 5.2; instead use 'uses_allocators\\(traits\\(mem2\\): var\\)' \\\[-Wdeprecated-openmp\\\]" } block; end block diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 index 0efdc9ef54b0..8338bb2d0306 100644 --- a/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-4.f90 @@ -48,12 +48,12 @@ end program main ! { dg-final { scan-tree-dump "pragma omp target" "original" } } ! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : foo\\)" "original" } } ! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\)" "original" } } -! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(omp_low_lat_mem_space\\), traits\\(\\) : bar\\)" "original" } } -! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\) : bar\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(4\\), traits\\(\\) : bar\\)" "original" } } +! { dg-final { scan-tree-dump "pragma omp target uses_allocators\\(memspace\\(3\\), traits\\(traits_array\\) : bar\\)" "original" } } ! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) allocate\\(allocator\\(bar\\):arr\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : bar\\)" "gimple" } } ! { dg-final { scan-tree-dump "pragma omp target" "gimple" } } ! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : foo\\)" "gimple" } } ! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(\\), traits\\(traits_array\\) : bar\\)" "gimple" } } -! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(omp_low_lat_mem_space\\), traits\\(\\) : bar\\)" "gimple" } } -! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(omp_high_bw_mem_space\\), traits\\(traits_array\\) : bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(4\\), traits\\(\\) : bar\\)" "gimple" } } +! { dg-final { scan-tree-dump "pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(3\\), traits\\(traits_array\\) : bar\\)" "gimple" } } diff --git a/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 b/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 index d0a603b5ced1..7605978c2b3d 100644 --- a/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 +++ b/libgomp/testsuite/libgomp.fortran/uses_allocators-7.f90 @@ -9,13 +9,27 @@ program main integer(omp_allocator_handle_kind) :: my_alloc, my, my2, my3, my4 type(omp_alloctrait), parameter :: trait(*) = [omp_alloctrait(omp_atk_alignment, 128)] type(omp_alloctrait), parameter :: t(*) = [omp_alloctrait:: ] - type(omp_alloctrait), parameter :: t2(*) = [omp_alloctrait:: ] + type(omp_alloctrait), parameter :: t2(*) = [omp_alloctrait:: omp_alloctrait(omp_atk_alignment, 1024)] - ! FIXME - improve check that that ';' is handled + my = -1; my2 = -2; my3 = -3; my4 = -4 !$omp target uses_allocators(traits(t), memspace(omp_high_bw_mem_space) : my; omp_default_mem_alloc, omp_null_allocator; my2; traits(t2) : my3; memspace(omp_large_cap_mem_space) : my4) block + integer, pointer :: ip3(:) + type(c_ptr) :: c1, c2, c3, c4 + c1 = omp_alloc (c_sizeof(x) * 1, my) + c2 = omp_alloc (c_sizeof(x) * 2, my2) + c3 = omp_alloc (c_sizeof(x) * 3, my3) + c4 = omp_alloc (c_sizeof(x) * 4, my4) + call c_f_pointer(c3, ip3, shape=[3]) + if (mod (TRANSFER (loc(ip3), iptr), 1024) /= 0) & + stop 1 + call omp_free (c1, omp_null_allocator) + call omp_free (c2, omp_null_allocator) + call omp_free (c3, omp_null_allocator) + call omp_free (c4, omp_null_allocator) end block + x = 1; xbuf = 0 !$omp target uses_allocators(omp_low_lat_mem_alloc) map(tofrom: x, xbuf) defaultmap(none) !$omp parallel allocate(allocator(omp_low_lat_mem_alloc), align(128): x, xbuf) if(.false.) firstprivate(x, xbuf) if (mod (TRANSFER (loc(x), iptr), 128) /= 0) & @@ -55,4 +69,6 @@ end ! { dg-final { scan-tree-dump-times "#pragma omp target .*uses_allocators\\(memspace\\(\\), traits\\(trait\\) : my_alloc\\)" 1 "gimple" } } ! { dg-final { scan-tree-dump-times "#pragma omp target .*uses_allocators\\(memspace\\(\\), traits\\(\\) : my_alloc\\)" 1 "gimple" } } -! { dg-final { scan-tree-dump-times "#pragma omp target .*uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\)" 1 "original" { xfail *-*-* } } } + +! { dg-final { scan-tree-dump-times "#pragma omp target uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : 0B\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target num_teams\\(-2\\) thread_limit\\(0\\) uses_allocators\\(memspace\\(3\\), traits\\(t\\) : my\\) uses_allocators\\(memspace\\(\\), traits\\(\\) : my2\\) uses_allocators\\(memspace\\(\\), traits\\(t2\\) : my3\\) uses_allocators\\(memspace\\(1\\), traits\\(\\) : my4\\)" 1 "gimple" } }