Hi Paul-Antoine!
On 2026-07-02T15:22:34+0200, Paul-Antoine Arras <[email protected]> wrote:
> Add parameter to GOMP_{loop,distribute}_static_worksharing{,_start} carrying
> the
> total number of loop iterations. niter is unused for now, pending OMPT
> implementation.
Thanks! ..., but may I ask how you've tested this?
For NVIDIA GPU offloading, I'm seeing regressions all over the place in
libgomp/OpenMP testing. Are you're not seeing those? (Eitheer at
compile time as 'FAIL: [...] (test for excess errors)' if 'ptxas' is
available, otherwise at run time as 'FAIL: [...] execution test', as
'libgomp: cuLinkAddData (ptx_code) error: a PTX JIT compilation failed'.)
[...]
ptxas application ptx input, line 925; error : Type of argument does not
match formal parameter '%in_ar1'
ptxas application ptx input, line 925; error : Alignment of argument does
not match formal parameter '%in_ar1'
[...]
Etc.
Also:
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-1.c (test for excess
errors)
+FAIL: libgomp.c/../libgomp.c-c++-common/for-static-1.c
scan-amdgcn-amdhsa-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing \\(n" 2
+FAIL: libgomp.c/../libgomp.c-c++-common/for-static-1.c
scan-nvptx-none-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing \\(n" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-1.c
scan-amdgcn-amdhsa-offload-tree-dump-not optimized "omp_get_num_teams"
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-1.c
scan-nvptx-none-offload-tree-dump-not optimized "omp_get_num_teams"
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-1.c
scan-amdgcn-amdhsa-offload-tree-dump-not optimized "omp_get_team_num"
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-1.c
scan-nvptx-none-offload-tree-dump-not optimized "omp_get_team_num"
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c (test for excess
errors)
+FAIL: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-amdgcn-amdhsa-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_start \\(n" 2
+FAIL: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-nvptx-none-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_start \\(n" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-amdgcn-amdhsa-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_end \\(" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-nvptx-none-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_end \\(" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-amdgcn-amdhsa-offload-tree-dump-not optimized
"GOMP_distribute_static_worksharing \\("
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-nvptx-none-offload-tree-dump-not optimized
"GOMP_distribute_static_worksharing \\("
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-amdgcn-amdhsa-offload-tree-dump-not optimized
"GOMP_distribute_static_worksharing_dispatch"
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-2.c
scan-nvptx-none-offload-tree-dump-not optimized
"GOMP_distribute_static_worksharing_dispatch"
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-3.c (test for excess
errors)
+FAIL: libgomp.c/../libgomp.c-c++-common/for-static-3.c
scan-amdgcn-amdhsa-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_start \\(n" 2
+FAIL: libgomp.c/../libgomp.c-c++-common/for-static-3.c
scan-nvptx-none-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_start \\(n" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-3.c
scan-amdgcn-amdhsa-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_dispatch \\(" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-3.c
scan-nvptx-none-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_dispatch \\(" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-3.c
scan-amdgcn-amdhsa-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_end \\(" 2
+PASS: libgomp.c/../libgomp.c-c++-common/for-static-3.c
scan-nvptx-none-offload-tree-dump-times optimized
"GOMP_distribute_static_worksharing_end \\(" 2
Same for 'libgomp.c++'.
That's:
[...]
int n;
[...]
_18 = n_13;
_19 = __builtin_GOMP_distribute_static_worksharing (_18);
[...]
Etc.
That helpfully already points to the former issue; for example, for
'libgomp.c-c++-common/target-1.c', in PTX code, we declare:
.extern .func GOMP_distribute_static_worksharing (.param .u64 %in_ar0,
.param .u64 %in_ar1);
..., which is per:
DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_LONG, BT_COMPLEX_INT, BT_LONG)
That is '.param .u64 %in_ar1' is for the 'BT_LONG'.
But, we then call:
[...]
.reg .u32 %r40;
[...]
.param .u64 %out_arg1;
st.param.u64 [%out_arg1],%r112;
.param .u32 %out_arg2;
st.param.u32 [%out_arg2],%r40;
call GOMP_distribute_static_worksharing,(%out_arg1,%out_arg2);
[...]
..., that is with 'n' as '.u32' instead of '.u64', which 'ptxas' then
rejects as per the diagnostic cited above.
I assume, that's due to:
[...]
tree type, itype, [...];
[...]
itype = type = TREE_TYPE (fd->loop.v);
if (POINTER_TYPE_P (type))
itype = signed_type_for (type);
[...]
t = fold_convert (itype, t);
n = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE, true,
GSI_SAME_STMT);
[...]
tree packed = build_call_expr (decl, 1, n);
[...]
..., so 'n' isn't always 'long' ('.u64'), but dependant on the type of
'fd->loop.v'. Can we unconditionally cast 'n' to 'long' for purposes of
these calls, or what's a good approach, here? Actually, per my reading
of OpenMP/OMPT, we unconditionally (also for 32-bit configurations) have
'uint64_t count;', that is 'unsigned long long'.
Unsurprisingly, same issue for 'GOMP_loop_static_worksharing'.
I'm attaching a
"[WIP] openmp: Pass iteration count to static worksharing builtins for OMPT"
that indends to address these issues. Please verify, and if you agree,
merge into your commit, with:
Co-authored-by: Thomas Schwinge <[email protected]>
Minor: I personally try to avoid spurious formatting changes in code I'm
not actually modifying:
> - innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc),
> - OMP_CLAUSE__LOOPTEMP_);
> + innerc
> + = omp_find_clause (OMP_CLAUSE_CHAIN (innerc), OMP_CLAUSE__LOOPTEMP_);
> [...]
> - t = fold_build2 (TRUNC_DIV_EXPR, itype,
> - fold_build1 (NEGATE_EXPR, itype, t),
> + t = fold_build2 (TRUNC_DIV_EXPR, itype, fold_build1 (NEGATE_EXPR, itype,
> t),
> fold_build1 (NEGATE_EXPR, itype, step));
Especially the latter one confused me for a minute or two; I thought
you're replacing 't' with 'step' in the last expression. ;-)
Grüße
Thomas
> The call in omp-expand.cc is moved after the iteration count computation
> so that the iteration count is available to pass as the argument.
>
> gcc/ChangeLog:
>
> * builtin-types.def (BT_FN_COMPLEX_INT): Remove.
> (BT_FN_COMPLEX_INT_LONG): New.
> * omp-builtins.def (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING): Use
> BT_FN_COMPLEX_INT_LONG.
> (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START): Likewise.
> (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING): Likewise.
> (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START): Likewise.
> * omp-expand.cc (expand_omp_for_static_nochunk): Move worksharing
> builtin call after iteration count computation; pass n as niter.
> (expand_omp_for_static_chunk): Likewise.
>
> gcc/fortran/ChangeLog:
>
> * types.def (BT_FN_COMPLEX_INT): Remove.
> (BT_FN_COMPLEX_INT_LONG): New.
>
> libgomp/ChangeLog:
>
> * config/gcn/teams.c (GOMP_distribute_static_worksharing): Add
> long niter parameter (unused).
> (GOMP_distribute_static_worksharing_start): Likewise.
> * config/nvptx/teams.c (GOMP_distribute_static_worksharing):
> Likewise.
> (GOMP_distribute_static_worksharing_start): Likewise.
> * libgomp.texi: Update function signatures; fix real/imaginary
> description; document niter parameter.
> * libgomp_g.h (GOMP_loop_static_worksharing): Update prototype.
> (GOMP_distribute_static_worksharing): Likewise.
> * loop.c (GOMP_loop_static_worksharing): Add long niter parameter
> (unused).
> (GOMP_loop_static_worksharing_start): Likewise.
> * teams.c (GOMP_distribute_static_worksharing): Likewise.
> (GOMP_distribute_static_worksharing_start): Likewise.
> * testsuite/libgomp.c-c++-common/for-static-1.c: Update scan
> pattern to match niter argument.
> * testsuite/libgomp.c-c++-common/for-static-2.c: Likewise.
> * testsuite/libgomp.c-c++-common/for-static-3.c: Likewise.
>
> gcc/testsuite/ChangeLog:
>
> * c-c++-common/gomp/for-static-1.c: Update scan pattern to match
> niter argument.
> * c-c++-common/gomp/for-static-2.c: Likewise.
> * c-c++-common/gomp/for-static-3.c: Likewise.
> ---
> gcc/builtin-types.def | 2 +-
> gcc/fortran/types.def | 2 +-
> gcc/omp-builtins.def | 9 +-
> gcc/omp-expand.cc | 168 +++++++++---------
> .../c-c++-common/gomp/for-static-1.c | 4 +-
> .../c-c++-common/gomp/for-static-2.c | 8 +-
> .../c-c++-common/gomp/for-static-3.c | 4 +-
> libgomp/config/gcn/teams.c | 7 +-
> libgomp/config/nvptx/teams.c | 7 +-
> libgomp/libgomp.texi | 9 +-
> libgomp/libgomp_g.h | 4 +-
> libgomp/loop.c | 7 +-
> libgomp/teams.c | 7 +-
> .../libgomp.c-c++-common/for-static-1.c | 2 +-
> .../libgomp.c-c++-common/for-static-2.c | 2 +-
> .../libgomp.c-c++-common/for-static-3.c | 2 +-
> 16 files changed, 127 insertions(+), 117 deletions(-)
>
> diff --git gcc/builtin-types.def gcc/builtin-types.def
> index 55ed6f0cae1..25b0aa6ffc3 100644
> --- gcc/builtin-types.def
> +++ gcc/builtin-types.def
> @@ -236,7 +236,6 @@ DEF_FUNCTION_TYPE_0 (BT_FN_CONST_STRING, BT_CONST_STRING)
> DEF_FUNCTION_TYPE_0 (BT_FN_PID, BT_PID)
> DEF_FUNCTION_TYPE_0 (BT_FN_INT, BT_INT)
> DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT)
> -DEF_FUNCTION_TYPE_0 (BT_FN_COMPLEX_INT, BT_COMPLEX_INT)
> DEF_FUNCTION_TYPE_0 (BT_FN_ULONG, BT_ULONG)
> DEF_FUNCTION_TYPE_0 (BT_FN_ULONGLONG, BT_ULONGLONG)
> DEF_FUNCTION_TYPE_0 (BT_FN_FLOAT, BT_FLOAT)
> @@ -429,6 +428,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_UINT32, BT_UINT16,
> BT_UINT32)
> DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_UINT16, BT_UINT32, BT_UINT16)
> DEF_FUNCTION_TYPE_1 (BT_FN_INT_FENV_T_PTR, BT_INT, BT_FENV_T_PTR)
> DEF_FUNCTION_TYPE_1 (BT_FN_INT_CONST_FENV_T_PTR, BT_INT, BT_CONST_FENV_T_PTR)
> +DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_LONG, BT_COMPLEX_INT, BT_LONG)
>
> DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
>
> diff --git gcc/fortran/types.def gcc/fortran/types.def
> index 6f98e3e48ff..1b362bc1d6f 100644
> --- gcc/fortran/types.def
> +++ gcc/fortran/types.def
> @@ -77,7 +77,6 @@ DEF_POINTER_TYPE (BT_PTR_PTR, BT_PTR)
> DEF_FUNCTION_TYPE_0 (BT_FN_BOOL, BT_BOOL)
> DEF_FUNCTION_TYPE_0 (BT_FN_PTR, BT_PTR)
> DEF_FUNCTION_TYPE_0 (BT_FN_INT, BT_INT)
> -DEF_FUNCTION_TYPE_0 (BT_FN_COMPLEX_INT, BT_COMPLEX_INT)
> DEF_FUNCTION_TYPE_0 (BT_FN_UINT, BT_UINT)
> DEF_FUNCTION_TYPE_0 (BT_FN_VOID, BT_VOID)
> DEF_FUNCTION_TYPE_0 (BT_FN_ULONGLONG, BT_ULONGLONG)
> @@ -91,6 +90,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
> DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
> DEF_FUNCTION_TYPE_1 (BT_FN_VOID_BOOL, BT_VOID, BT_BOOL)
> DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_INT, BT_BOOL, BT_INT)
> +DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_LONG, BT_COMPLEX_INT, BT_LONG)
>
> DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
>
> diff --git gcc/omp-builtins.def gcc/omp-builtins.def
> index 465aeebefdb..d34a0ff8ee8 100644
> --- gcc/omp-builtins.def
> +++ gcc/omp-builtins.def
> @@ -510,11 +510,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_HAS_MASKED_THREAD_NUM,
> ATTR_CONST_NOTHROW_LEAF_LIST)
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_MASKED_END, "GOMP_masked_end", BT_FN_VOID,
> ATTR_NOTHROW_LEAF_LIST)
> +/* TODO review these function attributes once OMPT implementation done */
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING,
> - "GOMP_loop_static_worksharing", BT_FN_COMPLEX_INT,
> + "GOMP_loop_static_worksharing", BT_FN_COMPLEX_INT_LONG,
> ATTR_CONST_NOTHROW_LEAF_LIST)
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START,
> - "GOMP_loop_static_worksharing_start", BT_FN_COMPLEX_INT,
> + "GOMP_loop_static_worksharing_start", BT_FN_COMPLEX_INT_LONG,
> ATTR_CONST_NOTHROW_LEAF_LIST)
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_END,
> "GOMP_loop_static_worksharing_end", BT_FN_VOID,
> @@ -523,10 +524,10 @@ DEF_GOMP_BUILTIN
> (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_DISPATCH,
> "GOMP_loop_static_worksharing_dispatch", BT_FN_VOID,
> ATTR_NOTHROW_LEAF_LIST)
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING,
> - "GOMP_distribute_static_worksharing", BT_FN_COMPLEX_INT,
> + "GOMP_distribute_static_worksharing", BT_FN_COMPLEX_INT_LONG,
> ATTR_CONST_NOTHROW_LEAF_LIST)
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START,
> - "GOMP_distribute_static_worksharing_start", BT_FN_COMPLEX_INT,
> + "GOMP_distribute_static_worksharing_start",
> BT_FN_COMPLEX_INT_LONG,
> ATTR_CONST_NOTHROW_LEAF_LIST)
> DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_END,
> "GOMP_distribute_static_worksharing_end", BT_FN_VOID,
> diff --git gcc/omp-expand.cc gcc/omp-expand.cc
> index ad380b27d06..c4d84ced5a2 100644
> --- gcc/omp-expand.cc
> +++ gcc/omp-expand.cc
> @@ -5190,42 +5190,6 @@ expand_omp_for_static_nochunk (struct omp_region
> *region,
> release_ssa_name (gimple_assign_lhs (g));
> }
> }
> - /* Fetch the thread/team id and the number of threads/teams in a single
> - call to GOMP_loop_static_worksharing or
> GOMP_distribute_static_worksharing.
> - The helper returns both values packed into one complex int, with
> - the id as the imaginary part and the count as the real part. Returning
> - (rather than writing through pointers) keeps both values as plain SSA
> - names, which lets later passes - notably IPA-CP propagating constants
> - into the outlined kernel - reason about them. */
> - tree decl;
> - switch (gimple_omp_for_kind (fd->for_stmt))
> - {
> - case GF_OMP_FOR_KIND_FOR:
> - decl = builtin_decl_explicit (
> - flag_openmp_ompt ? BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START
> - : BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING);
> - break;
> - case GF_OMP_FOR_KIND_DISTRIBUTE:
> - decl = builtin_decl_explicit (
> - flag_openmp_ompt ? BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START
> - : BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING);
> - break;
> - default:
> - gcc_unreachable ();
> - }
> - {
> - tree packed = build_call_expr (decl, 0);
> - packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE,
> - true, GSI_SAME_STMT);
> - threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed);
> - threadid = fold_convert (itype, threadid);
> - threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
> - true, GSI_SAME_STMT);
> - nthreads = fold_build1 (REALPART_EXPR, integer_type_node, packed);
> - nthreads = fold_convert (itype, nthreads);
> - nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
> - true, GSI_SAME_STMT);
> - }
>
> n1 = fd->loop.n1;
> n2 = fd->loop.n2;
> @@ -5236,31 +5200,69 @@ expand_omp_for_static_nochunk (struct omp_region
> *region,
> OMP_CLAUSE__LOOPTEMP_);
> gcc_assert (innerc);
> n1 = OMP_CLAUSE_DECL (innerc);
> - innerc = omp_find_clause (OMP_CLAUSE_CHAIN (innerc),
> - OMP_CLAUSE__LOOPTEMP_);
> + innerc
> + = omp_find_clause (OMP_CLAUSE_CHAIN (innerc), OMP_CLAUSE__LOOPTEMP_);
> gcc_assert (innerc);
> n2 = OMP_CLAUSE_DECL (innerc);
> }
> - n1 = force_gimple_operand_gsi (&gsi, fold_convert (type, n1),
> - true, NULL_TREE, true, GSI_SAME_STMT);
> - n2 = force_gimple_operand_gsi (&gsi, fold_convert (itype, n2),
> - true, NULL_TREE, true, GSI_SAME_STMT);
> - step = force_gimple_operand_gsi (&gsi, fold_convert (itype, step),
> - true, NULL_TREE, true, GSI_SAME_STMT);
> + n1 = force_gimple_operand_gsi (&gsi, fold_convert (type, n1), true,
> NULL_TREE,
> + true, GSI_SAME_STMT);
> + n2 = force_gimple_operand_gsi (&gsi, fold_convert (itype, n2), true,
> + NULL_TREE, true, GSI_SAME_STMT);
> + step = force_gimple_operand_gsi (&gsi, fold_convert (itype, step), true,
> + NULL_TREE, true, GSI_SAME_STMT);
>
> t = build_int_cst (itype, (fd->loop.cond_code == LT_EXPR ? -1 : 1));
> t = fold_build2 (PLUS_EXPR, itype, step, t);
> t = fold_build2 (PLUS_EXPR, itype, t, n2);
> t = fold_build2 (MINUS_EXPR, itype, t, fold_convert (itype, n1));
> if (TYPE_UNSIGNED (itype) && fd->loop.cond_code == GT_EXPR)
> - t = fold_build2 (TRUNC_DIV_EXPR, itype,
> - fold_build1 (NEGATE_EXPR, itype, t),
> + t = fold_build2 (TRUNC_DIV_EXPR, itype, fold_build1 (NEGATE_EXPR, itype,
> t),
> fold_build1 (NEGATE_EXPR, itype, step));
> else
> t = fold_build2 (TRUNC_DIV_EXPR, itype, t, step);
> t = fold_convert (itype, t);
> n = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE, true,
> GSI_SAME_STMT);
>
> + {
> + /* Fetch the thread/team id and the number of threads/teams in a single
> + call to GOMP_loop_static_worksharing or
> + GOMP_distribute_static_worksharing. The helper returns both values
> packed
> + into one complex int, with the id as the imaginary part and the count
> as
> + the real part. Returning (rather than writing through pointers) keeps
> + both values as plain SSA names, which lets later passes - notably
> IPA-CP
> + propagating constants into the outlined kernel - reason about them.
> + Also pass the total number of iterations for OMPT. */
> + tree decl;
> + switch (gimple_omp_for_kind (fd->for_stmt))
> + {
> + case GF_OMP_FOR_KIND_FOR:
> + decl = builtin_decl_explicit (
> + flag_openmp_ompt ? BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START
> + : BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING);
> + break;
> + case GF_OMP_FOR_KIND_DISTRIBUTE:
> + decl = builtin_decl_explicit (
> + flag_openmp_ompt ? BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START
> + : BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING);
> + break;
> + default:
> + gcc_unreachable ();
> + }
> +
> + tree packed = build_call_expr (decl, 1, n);
> + packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE,
> + true, GSI_SAME_STMT);
> + threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed);
> + threadid = fold_convert (itype, threadid);
> + threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
> + true, GSI_SAME_STMT);
> + nthreads = fold_build1 (REALPART_EXPR, integer_type_node, packed);
> + nthreads = fold_convert (itype, nthreads);
> + nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
> + true, GSI_SAME_STMT);
> + }
> +
> q = create_tmp_reg (itype, "q");
> t = fold_build2 (TRUNC_DIV_EXPR, itype, n, nthreads);
> t = force_gimple_operand_gsi (&gsi, t, false, NULL_TREE, true,
> GSI_SAME_STMT);
> @@ -6008,42 +6010,6 @@ expand_omp_for_static_chunk (struct omp_region *region,
> release_ssa_name (gimple_assign_lhs (g));
> }
> }
> - /* Fetch the thread/team id and the number of threads/teams in a single
> - call to GOMP_loop_static_worksharing or
> GOMP_distribute_static_worksharing.
> - The helper returns both values packed into one complex int, with
> - the id as the imaginary part and the count as the real part. Returning
> - (rather than writing through pointers) keeps both values as plain SSA
> - names, which lets later passes - notably IPA-CP propagating constants
> - into the outlined kernel - reason about them. */
> - tree decl;
> - switch (gimple_omp_for_kind (fd->for_stmt))
> - {
> - case GF_OMP_FOR_KIND_FOR:
> - decl = builtin_decl_explicit (
> - flag_openmp_ompt ? BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START
> - : BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING);
> - break;
> - case GF_OMP_FOR_KIND_DISTRIBUTE:
> - decl = builtin_decl_explicit (
> - flag_openmp_ompt ? BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START
> - : BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING);
> - break;
> - default:
> - gcc_unreachable ();
> - }
> - {
> - tree packed = build_call_expr (decl, 0);
> - packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE,
> - true, GSI_SAME_STMT);
> - threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed);
> - threadid = fold_convert (itype, threadid);
> - threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
> - true, GSI_SAME_STMT);
> - nthreads = fold_build1 (REALPART_EXPR, integer_type_node, packed);
> - nthreads = fold_convert (itype, nthreads);
> - nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
> - true, GSI_SAME_STMT);
> - }
>
> n1 = fd->loop.n1;
> n2 = fd->loop.n2;
> @@ -6086,6 +6052,44 @@ expand_omp_for_static_chunk (struct omp_region *region,
> n = force_gimple_operand_gsi (&gsi, t, true, NULL_TREE,
> true, GSI_SAME_STMT);
>
> + {
> + /* Fetch the thread/team id and the number of threads/teams in a single
> + call to GOMP_loop_static_worksharing or
> + GOMP_distribute_static_worksharing. The helper returns both values
> packed
> + into one complex int, with the id as the imaginary part and the count
> as
> + the real part. Returning (rather than writing through pointers) keeps
> + both values as plain SSA names, which lets later passes - notably
> IPA-CP
> + propagating constants into the outlined kernel - reason about them.
> + Also pass the total number of iterations for OMPT. */
> + tree decl;
> + switch (gimple_omp_for_kind (fd->for_stmt))
> + {
> + case GF_OMP_FOR_KIND_FOR:
> + decl = builtin_decl_explicit (
> + flag_openmp_ompt ? BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START
> + : BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING);
> + break;
> + case GF_OMP_FOR_KIND_DISTRIBUTE:
> + decl = builtin_decl_explicit (
> + flag_openmp_ompt ? BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START
> + : BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING);
> + break;
> + default:
> + gcc_unreachable ();
> + }
> + tree packed = build_call_expr (decl, 1, n);
> + packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE, true,
> + GSI_SAME_STMT);
> + threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed);
> + threadid = fold_convert (itype, threadid);
> + threadid = force_gimple_operand_gsi (&gsi, threadid, true, NULL_TREE,
> true,
> + GSI_SAME_STMT);
> + nthreads = fold_build1 (REALPART_EXPR, integer_type_node, packed);
> + nthreads = fold_convert (itype, nthreads);
> + nthreads = force_gimple_operand_gsi (&gsi, nthreads, true, NULL_TREE,
> true,
> + GSI_SAME_STMT);
> + }
> +
> trip_var = create_tmp_reg (itype, ".trip");
> if (gimple_in_ssa_p (cfun))
> {
> diff --git gcc/testsuite/c-c++-common/gomp/for-static-1.c
> gcc/testsuite/c-c++-common/gomp/for-static-1.c
> index 50d4e176f3d..4a5cee2a5cd 100644
> --- gcc/testsuite/c-c++-common/gomp/for-static-1.c
> +++ gcc/testsuite/c-c++-common/gomp/for-static-1.c
> @@ -9,8 +9,8 @@
>
> #include "for-static.h"
>
> -/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing \\(" 2
> "ompexp" } } */
> -/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing
> \\(" 2 "ompexp" } } */
> +/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing \\(n" 2
> "ompexp" } } */
> +/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing
> \\(n" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-not "omp_get_num_threads" "ompexp" } } */
> /* { dg-final { scan-tree-dump-not "omp_get_thread_num" "ompexp" } } */
> /* { dg-final { scan-tree-dump-not "omp_get_num_teams" "ompexp" } } */
> diff --git gcc/testsuite/c-c++-common/gomp/for-static-2.c
> gcc/testsuite/c-c++-common/gomp/for-static-2.c
> index c3d25b97bed..c42413beb4a 100644
> --- gcc/testsuite/c-c++-common/gomp/for-static-2.c
> +++ gcc/testsuite/c-c++-common/gomp/for-static-2.c
> @@ -5,11 +5,11 @@
>
> #include "for-static.h"
>
> -/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start
> \\(" 2 "ompexp" } } */
> +/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start
> \\(n" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_end \\("
> 2 "ompexp" } } */
> -/* { dg-final { scan-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(" 2 "ompexp" } } */
> +/* { dg-final { scan-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(n" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_end
> \\(" 2 "ompexp" } } */
> -/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing \\("
> "ompexp" } } */
> +/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing \\(n"
> "ompexp" } } */
> /* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing_dispatch"
> "ompexp" } } */
> -/* { dg-final { scan-tree-dump-not "GOMP_distribute_static_worksharing \\("
> "ompexp" } } */
> +/* { dg-final { scan-tree-dump-not "GOMP_distribute_static_worksharing \\(n"
> "ompexp" } } */
> /* { dg-final { scan-tree-dump-not
> "GOMP_distribute_static_worksharing_dispatch" "ompexp" } } */
> diff --git gcc/testsuite/c-c++-common/gomp/for-static-3.c
> gcc/testsuite/c-c++-common/gomp/for-static-3.c
> index 6fd1da6c407..fe344830f30 100644
> --- gcc/testsuite/c-c++-common/gomp/for-static-3.c
> +++ gcc/testsuite/c-c++-common/gomp/for-static-3.c
> @@ -6,9 +6,9 @@
>
> #include "for-static.h"
>
> -/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start
> \\(" 2 "ompexp" } } */
> +/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start
> \\(n" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_dispatch
> \\(" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_end \\("
> 2 "ompexp" } } */
> -/* { dg-final { scan-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(" 2 "ompexp" } } */
> +/* { dg-final { scan-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(n" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-times
> "GOMP_distribute_static_worksharing_dispatch \\(" 2 "ompexp" } } */
> /* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_end
> \\(" 2 "ompexp" } } */
> diff --git libgomp/config/gcn/teams.c libgomp/config/gcn/teams.c
> index 99f807ac876..8e2a56ec9cb 100644
> --- libgomp/config/gcn/teams.c
> +++ libgomp/config/gcn/teams.c
> @@ -39,10 +39,11 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned
> int num_teams,
> }
>
> /* For a distribute construct with static schedule, return the team ID and
> - number of teams packed into a single complexvalue. */
> + number of teams packed into a single complex value. NITER is the total
> + number of iterations. */
>
> _Complex int
> -GOMP_distribute_static_worksharing (void)
> +GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
> {
> int __lds *gomp_team_num = (int __lds *) GOMP_TEAM_NUM;
> unsigned tid = *gomp_team_num;
> @@ -53,7 +54,7 @@ GOMP_distribute_static_worksharing (void)
> /* OMPT variant enabled by -fopenmp-ompt. */
>
> _Complex int
> -GOMP_distribute_static_worksharing_start (void)
> +GOMP_distribute_static_worksharing_start (long niter __attribute__
> ((unused)))
> {
> int __lds *gomp_team_num = (int __lds *) GOMP_TEAM_NUM;
> unsigned tid = *gomp_team_num;
> diff --git libgomp/config/nvptx/teams.c libgomp/config/nvptx/teams.c
> index 3e6ea287276..927a3a14366 100644
> --- libgomp/config/nvptx/teams.c
> +++ libgomp/config/nvptx/teams.c
> @@ -42,10 +42,11 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned
> int num_teams,
> }
>
> /* For a distribute construct with static schedule, return the team ID and
> - number of teams packed into a single complexvalue. */
> + number of teams packed into a single complex value. NITER is the total
> + number of iterations. */
>
> _Complex int
> -GOMP_distribute_static_worksharing (void)
> +GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
> {
> unsigned tid = __gomp_team_num;
> unsigned nteams = gomp_num_teams_var + 1;
> @@ -55,7 +56,7 @@ GOMP_distribute_static_worksharing (void)
> /* OMPT variant enabled by -fopenmp-ompt. */
>
> _Complex int
> -GOMP_distribute_static_worksharing_start (void)
> +GOMP_distribute_static_worksharing_start (long niter __attribute__
> ((unused)))
> {
> unsigned tid = __gomp_team_num;
> unsigned nteams = gomp_num_teams_var + 1;
> diff --git libgomp/libgomp.texi libgomp/libgomp.texi
> index 83b90e6dc53..61a0481078f 100644
> --- libgomp/libgomp.texi
> +++ libgomp/libgomp.texi
> @@ -8318,17 +8318,18 @@ but since its supposed to remain unchanged, we can
> also not if we like.
> For @code{schedule(static)} without @code{ordered}, the compiler uses:
>
> @smallexample
> - _Complex int GOMP_loop_static_worksharing (void)
> + _Complex int GOMP_loop_static_worksharing (long)
> @end smallexample
>
> -which returns the thread id in the real part and the thread count in the
> imaginary
> -part as a single value, avoiding separate @code{omp_get_thread_num} /
> +which returns the thread id in the imaginary part and the thread count in the
> +real part as a single value, avoiding separate @code{omp_get_thread_num} /
> @code{omp_get_num_threads} calls that OMPT cannot distinguish from user
> queries.
> +For OMPT, the total number of loop iterations is passed as argument.
>
> Similarly, for @code{distribute} constructs:
>
> @smallexample
> - _Complex int GOMP_distribute_static_worksharing (void)
> + _Complex int GOMP_distribute_static_worksharing (long)
> @end smallexample
>
> returns the team id and team count.
> diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h
> index a3097d5b093..49f3a06e579 100644
> --- libgomp/libgomp_g.h
> +++ libgomp/libgomp_g.h
> @@ -298,7 +298,7 @@ extern unsigned GOMP_parallel_reductions (void (*) (void
> *), void *, unsigned,
> extern bool GOMP_cancel (int, bool);
> extern bool GOMP_cancellation_point (int);
> extern bool GOMP_has_masked_thread_num (int);
> -extern __complex__ int GOMP_loop_static_worksharing (void);
> +extern __complex__ int GOMP_loop_static_worksharing (long);
>
> /* task.c */
>
> @@ -379,7 +379,7 @@ extern void *GOMP_target_map_indirect_ptr (void *);
>
> extern void GOMP_teams_reg (void (*) (void *), void *, unsigned, unsigned,
> unsigned);
> -extern __complex__ int GOMP_distribute_static_worksharing (void);
> +extern __complex__ int GOMP_distribute_static_worksharing (long);
>
> /* allocator.c */
>
> diff --git libgomp/loop.c libgomp/loop.c
> index d53d994ee34..889ac282e89 100644
> --- libgomp/loop.c
> +++ libgomp/loop.c
> @@ -1188,10 +1188,11 @@ GOMP_loop_ordered_guided_next (long *istart, long
> *iend)
> #endif
>
> /* For a worksharing-loop construct with static schedule, return the thread
> ID
> - and number of threads packed into a single complex value. */
> + and number of threads packed into a single complex value. NITER is the
> total
> + number of iterations. */
>
> _Complex int
> -GOMP_loop_static_worksharing (void)
> +GOMP_loop_static_worksharing (long niter __attribute__ ((unused)))
> {
> struct gomp_team *team = gomp_thread ()->ts.team;
> unsigned tid = gomp_thread ()->ts.team_id;
> @@ -1202,7 +1203,7 @@ GOMP_loop_static_worksharing (void)
> /* OMPT variant enabled by -fopenmp-ompt. */
>
> _Complex int
> -GOMP_loop_static_worksharing_start (void)
> +GOMP_loop_static_worksharing_start (long niter __attribute__ ((unused)))
> {
> struct gomp_team *team = gomp_thread ()->ts.team;
> unsigned tid = gomp_thread ()->ts.team_id;
> diff --git libgomp/teams.c libgomp/teams.c
> index f501123d647..319cfee0db5 100644
> --- libgomp/teams.c
> +++ libgomp/teams.c
> @@ -59,10 +59,11 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned
> int num_teams,
> }
>
> /* For a distribute construct with static schedule, return the team ID and
> - number of teams packed into a single complex value. */
> + number of teams packed into a single complex value. NITER is the total
> + number of iterations. */
>
> _Complex int
> -GOMP_distribute_static_worksharing (void)
> +GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
> {
> struct gomp_thread *thr = gomp_thread ();
> unsigned tid = thr->team_num;
> @@ -73,7 +74,7 @@ GOMP_distribute_static_worksharing (void)
> /* OMPT variant enabled by -fopenmp-ompt. */
>
> _Complex int
> -GOMP_distribute_static_worksharing_start (void)
> +GOMP_distribute_static_worksharing_start (long niter __attribute__
> ((unused)))
> {
> struct gomp_thread *thr = gomp_thread ();
> unsigned tid = thr->team_num;
> diff --git libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
> libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
> index b304c8c10ec..3715a5efda4 100644
> --- libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
> +++ libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
> @@ -8,6 +8,6 @@
>
> #include "for-static.h"
>
> -/* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing \\(" 2 "optimized" } } */
> +/* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing \\(n" 2 "optimized" } } */
> /* { dg-final { scan-offload-tree-dump-not "omp_get_num_teams" "optimized" }
> } */
> /* { dg-final { scan-offload-tree-dump-not "omp_get_team_num" "optimized" }
> } */
> diff --git libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
> libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
> index 83e577a3298..a0b15753155 100644
> --- libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
> +++ libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
> @@ -5,7 +5,7 @@
>
> #include "for-static.h"
>
> -/* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(" 2 "optimized" } } */
> +/* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(n" 2 "optimized" } } */
> /* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_end \\(" 2 "optimized" } } */
> /* { dg-final { scan-offload-tree-dump-not
> "GOMP_distribute_static_worksharing \\(" "optimized" } } */
> /* { dg-final { scan-offload-tree-dump-not
> "GOMP_distribute_static_worksharing_dispatch" "optimized" } } */
> diff --git libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
> libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
> index e7b904eedc3..afbac717f62 100644
> --- libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
> +++ libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
> @@ -6,6 +6,6 @@
>
> #include "for-static.h"
>
> -/* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(" 2 "optimized" } } */
> +/* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_start \\(n" 2 "optimized" } } */
> /* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_dispatch \\(" 2 "optimized" } } */
> /* { dg-final { scan-offload-tree-dump-times
> "GOMP_distribute_static_worksharing_end \\(" 2 "optimized" } } */
> --
> 2.39.5
>From a3bf195b235b0880157f057d13be969e9768fecf Mon Sep 17 00:00:00 2001
From: Thomas Schwinge <[email protected]>
Date: Tue, 7 Jul 2026 09:11:05 +0200
Subject: [PATCH] [WIP] openmp: Pass iteration count to static worksharing
builtins for OMPT
---
gcc/builtin-types.def | 2 +-
gcc/fortran/types.def | 2 +-
gcc/omp-builtins.def | 12 ++++++++----
gcc/omp-expand.cc | 7 ++++---
gcc/testsuite/c-c++-common/gomp/for-static-1.c | 7 +++++--
gcc/testsuite/c-c++-common/gomp/for-static-2.c | 16 ++++++++++------
gcc/testsuite/c-c++-common/gomp/for-static-3.c | 18 ++++++++++++------
libgomp/config/gcn/teams.c | 6 ++++--
libgomp/config/nvptx/teams.c | 6 ++++--
libgomp/libgomp.texi | 4 ++--
libgomp/libgomp_g.h | 15 ++++++++++++---
libgomp/loop.c | 6 ++++--
libgomp/teams.c | 6 ++++--
.../libgomp.c-c++-common/for-static-1.c | 2 +-
.../libgomp.c-c++-common/for-static-2.c | 4 ++--
.../libgomp.c-c++-common/for-static-3.c | 6 +++---
16 files changed, 77 insertions(+), 42 deletions(-)
diff --git a/gcc/builtin-types.def b/gcc/builtin-types.def
index 4f63da292b4..c17ded784c0 100644
--- a/gcc/builtin-types.def
+++ b/gcc/builtin-types.def
@@ -428,7 +428,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_UINT16_UINT32, BT_UINT16, BT_UINT32)
DEF_FUNCTION_TYPE_1 (BT_FN_UINT32_UINT16, BT_UINT32, BT_UINT16)
DEF_FUNCTION_TYPE_1 (BT_FN_INT_FENV_T_PTR, BT_INT, BT_FENV_T_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_INT_CONST_FENV_T_PTR, BT_INT, BT_CONST_FENV_T_PTR)
-DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_LONG, BT_COMPLEX_INT, BT_LONG)
+DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_ULONGLONG, BT_COMPLEX_INT, BT_ULONGLONG)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
diff --git a/gcc/fortran/types.def b/gcc/fortran/types.def
index 9e4166bce63..e865fc9b08d 100644
--- a/gcc/fortran/types.def
+++ b/gcc/fortran/types.def
@@ -90,7 +90,7 @@ DEF_FUNCTION_TYPE_1 (BT_FN_PTR_PTR, BT_PTR, BT_PTR)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_INT, BT_VOID, BT_INT)
DEF_FUNCTION_TYPE_1 (BT_FN_VOID_BOOL, BT_VOID, BT_BOOL)
DEF_FUNCTION_TYPE_1 (BT_FN_BOOL_INT, BT_BOOL, BT_INT)
-DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_LONG, BT_COMPLEX_INT, BT_LONG)
+DEF_FUNCTION_TYPE_1 (BT_FN_COMPLEX_INT_ULONGLONG, BT_COMPLEX_INT, BT_ULONGLONG)
DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR)
diff --git a/gcc/omp-builtins.def b/gcc/omp-builtins.def
index 0dd6e679572..efa3f9804c8 100644
--- a/gcc/omp-builtins.def
+++ b/gcc/omp-builtins.def
@@ -522,10 +522,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_MASKED_END, "GOMP_masked_end", BT_FN_VOID,
ATTR_NOTHROW_LEAF_LIST)
/* TODO review these function attributes once OMPT implementation done */
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING,
- "GOMP_loop_static_worksharing", BT_FN_COMPLEX_INT_LONG,
+ "GOMP_loop_static_worksharing",
+ BT_FN_COMPLEX_INT_ULONGLONG,
ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_START,
- "GOMP_loop_static_worksharing_start", BT_FN_COMPLEX_INT_LONG,
+ "GOMP_loop_static_worksharing_start",
+ BT_FN_COMPLEX_INT_ULONGLONG,
ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_END,
"GOMP_loop_static_worksharing_end", BT_FN_VOID,
@@ -534,10 +536,12 @@ DEF_GOMP_BUILTIN (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING_DISPATCH,
"GOMP_loop_static_worksharing_dispatch", BT_FN_VOID,
ATTR_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING,
- "GOMP_distribute_static_worksharing", BT_FN_COMPLEX_INT_LONG,
+ "GOMP_distribute_static_worksharing",
+ BT_FN_COMPLEX_INT_ULONGLONG,
ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_START,
- "GOMP_distribute_static_worksharing_start", BT_FN_COMPLEX_INT_LONG,
+ "GOMP_distribute_static_worksharing_start",
+ BT_FN_COMPLEX_INT_ULONGLONG,
ATTR_CONST_NOTHROW_LEAF_LIST)
DEF_GOMP_BUILTIN (BUILT_IN_GOMP_DISTRIBUTE_STATIC_WORKSHARING_END,
"GOMP_distribute_static_worksharing_end", BT_FN_VOID,
diff --git a/gcc/omp-expand.cc b/gcc/omp-expand.cc
index c4d84ced5a2..bce930fc108 100644
--- a/gcc/omp-expand.cc
+++ b/gcc/omp-expand.cc
@@ -5249,8 +5249,8 @@ expand_omp_for_static_nochunk (struct omp_region *region,
default:
gcc_unreachable ();
}
-
- tree packed = build_call_expr (decl, 1, n);
+ tree n_ull = fold_convert (long_long_unsigned_type_node, n);
+ tree packed = build_call_expr (decl, 1, n_ull);
packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE,
true, GSI_SAME_STMT);
threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed);
@@ -6077,7 +6077,8 @@ expand_omp_for_static_chunk (struct omp_region *region,
default:
gcc_unreachable ();
}
- tree packed = build_call_expr (decl, 1, n);
+ tree n_ull = fold_convert (long_long_unsigned_type_node, n);
+ tree packed = build_call_expr (decl, 1, n_ull);
packed = force_gimple_operand_gsi (&gsi, packed, true, NULL_TREE, true,
GSI_SAME_STMT);
threadid = fold_build1 (IMAGPART_EXPR, integer_type_node, packed);
diff --git a/gcc/testsuite/c-c++-common/gomp/for-static-1.c b/gcc/testsuite/c-c++-common/gomp/for-static-1.c
index 4a5cee2a5cd..09e533c931d 100644
--- a/gcc/testsuite/c-c++-common/gomp/for-static-1.c
+++ b/gcc/testsuite/c-c++-common/gomp/for-static-1.c
@@ -9,8 +9,11 @@
#include "for-static.h"
-/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing \\(n" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing \\(n" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing \\(\[^\\)\]" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing_start" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing_dispatch" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing_end" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing \\(\[^\\)\]" 2 "ompexp" } } */
/* { dg-final { scan-tree-dump-not "omp_get_num_threads" "ompexp" } } */
/* { dg-final { scan-tree-dump-not "omp_get_thread_num" "ompexp" } } */
/* { dg-final { scan-tree-dump-not "omp_get_num_teams" "ompexp" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/for-static-2.c b/gcc/testsuite/c-c++-common/gomp/for-static-2.c
index c42413beb4a..d99f59d7471 100644
--- a/gcc/testsuite/c-c++-common/gomp/for-static-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/for-static-2.c
@@ -5,11 +5,15 @@
#include "for-static.h"
-/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start \\(n" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_end \\(" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_start \\(n" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_end \\(" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing \\(n" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing \\(" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start \\(\[^\\)\]" 2 "ompexp" } } */
/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing_dispatch" "ompexp" } } */
-/* { dg-final { scan-tree-dump-not "GOMP_distribute_static_worksharing \\(n" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_end \\(\\)" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_distribute_static_worksharing \\(" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_start \\(\[^\\)\]" 2 "ompexp" } } */
/* { dg-final { scan-tree-dump-not "GOMP_distribute_static_worksharing_dispatch" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_end \\(\\)" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_num_threads" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_thread_num" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_num_teams" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_team_num" "ompexp" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/for-static-3.c b/gcc/testsuite/c-c++-common/gomp/for-static-3.c
index fe344830f30..8e45223c353 100644
--- a/gcc/testsuite/c-c++-common/gomp/for-static-3.c
+++ b/gcc/testsuite/c-c++-common/gomp/for-static-3.c
@@ -6,9 +6,15 @@
#include "for-static.h"
-/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start \\(n" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_dispatch \\(" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_end \\(" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_start \\(n" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_dispatch \\(" 2 "ompexp" } } */
-/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_end \\(" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_loop_static_worksharing \\(" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_start \\(\[^\\)\]" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_dispatch \\(\\)" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_loop_static_worksharing_end \\(\\)" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "GOMP_distribute_static_worksharing \\(" "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_start \\(\[^\\)\]" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_dispatch \\(\\)" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing_end \\(\\)" 2 "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_num_threads" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_thread_num" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_num_teams" "ompexp" } } */
+/* { dg-final { scan-tree-dump-not "omp_get_team_num" "ompexp" } } */
diff --git a/libgomp/config/gcn/teams.c b/libgomp/config/gcn/teams.c
index 8e2a56ec9cb..4354131f9b7 100644
--- a/libgomp/config/gcn/teams.c
+++ b/libgomp/config/gcn/teams.c
@@ -43,7 +43,8 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned int num_teams,
number of iterations. */
_Complex int
-GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
+GOMP_distribute_static_worksharing (unsigned long long niter
+ __attribute__ ((unused)))
{
int __lds *gomp_team_num = (int __lds *) GOMP_TEAM_NUM;
unsigned tid = *gomp_team_num;
@@ -54,7 +55,8 @@ GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
/* OMPT variant enabled by -fopenmp-ompt. */
_Complex int
-GOMP_distribute_static_worksharing_start (long niter __attribute__ ((unused)))
+GOMP_distribute_static_worksharing_start (unsigned long long niter
+ __attribute__ ((unused)))
{
int __lds *gomp_team_num = (int __lds *) GOMP_TEAM_NUM;
unsigned tid = *gomp_team_num;
diff --git a/libgomp/config/nvptx/teams.c b/libgomp/config/nvptx/teams.c
index 927a3a14366..e04f8bbbdb1 100644
--- a/libgomp/config/nvptx/teams.c
+++ b/libgomp/config/nvptx/teams.c
@@ -46,7 +46,8 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned int num_teams,
number of iterations. */
_Complex int
-GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
+GOMP_distribute_static_worksharing (unsigned long long niter
+ __attribute__ ((unused)))
{
unsigned tid = __gomp_team_num;
unsigned nteams = gomp_num_teams_var + 1;
@@ -56,7 +57,8 @@ GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
/* OMPT variant enabled by -fopenmp-ompt. */
_Complex int
-GOMP_distribute_static_worksharing_start (long niter __attribute__ ((unused)))
+GOMP_distribute_static_worksharing_start (unsigned long long niter
+ __attribute__ ((unused)))
{
unsigned tid = __gomp_team_num;
unsigned nteams = gomp_num_teams_var + 1;
diff --git a/libgomp/libgomp.texi b/libgomp/libgomp.texi
index 61a0481078f..5d29da0984d 100644
--- a/libgomp/libgomp.texi
+++ b/libgomp/libgomp.texi
@@ -8318,7 +8318,7 @@ but since its supposed to remain unchanged, we can also not if we like.
For @code{schedule(static)} without @code{ordered}, the compiler uses:
@smallexample
- _Complex int GOMP_loop_static_worksharing (long)
+ _Complex int GOMP_loop_static_worksharing (unsigned long long)
@end smallexample
which returns the thread id in the imaginary part and the thread count in the
@@ -8329,7 +8329,7 @@ For OMPT, the total number of loop iterations is passed as argument.
Similarly, for @code{distribute} constructs:
@smallexample
- _Complex int GOMP_distribute_static_worksharing (long)
+ _Complex int GOMP_distribute_static_worksharing (unsigned long long)
@end smallexample
returns the team id and team count.
diff --git a/libgomp/libgomp_g.h b/libgomp/libgomp_g.h
index b2329373795..392b945dbf2 100644
--- a/libgomp/libgomp_g.h
+++ b/libgomp/libgomp_g.h
@@ -144,6 +144,14 @@ extern void GOMP_loop_end (void);
extern void GOMP_loop_end_nowait (void);
extern bool GOMP_loop_end_cancel (void);
+extern _Complex int GOMP_loop_static_worksharing (unsigned long long);
+extern _Complex int GOMP_loop_static_worksharing_start (unsigned long long);
+extern void GOMP_loop_static_worksharing_dispatch (void);
+extern void GOMP_loop_static_worksharing_end (void);
+
+extern bool GOMP_has_masked_thread_num (int);
+extern void GOMP_masked_end (void);
+
/* loop_ull.c */
extern bool GOMP_loop_ull_static_start (bool, unsigned long long,
@@ -297,8 +305,6 @@ extern unsigned GOMP_parallel_reductions (void (*) (void *), void *, unsigned,
unsigned);
extern bool GOMP_cancel (int, bool);
extern bool GOMP_cancellation_point (int);
-extern bool GOMP_has_masked_thread_num (int);
-extern __complex__ int GOMP_loop_static_worksharing (long);
/* task.c */
@@ -379,7 +385,10 @@ extern void *GOMP_target_map_indirect_ptr (void *);
extern void GOMP_teams_reg (void (*) (void *), void *, unsigned, unsigned,
unsigned);
-extern __complex__ int GOMP_distribute_static_worksharing (long);
+extern _Complex int GOMP_distribute_static_worksharing (unsigned long long);
+extern _Complex int GOMP_distribute_static_worksharing_start (unsigned long long);
+extern void GOMP_distribute_static_worksharing_dispatch (void);
+extern void GOMP_distribute_static_worksharing_end (void);
/* allocator.c */
diff --git a/libgomp/loop.c b/libgomp/loop.c
index f3cc9f698c9..f5bb6795c3d 100644
--- a/libgomp/loop.c
+++ b/libgomp/loop.c
@@ -1192,7 +1192,8 @@ GOMP_loop_ordered_guided_next (long *istart, long *iend)
number of iterations. */
_Complex int
-GOMP_loop_static_worksharing (long niter __attribute__ ((unused)))
+GOMP_loop_static_worksharing (unsigned long long niter
+ __attribute__ ((unused)))
{
struct gomp_team *team = gomp_thread ()->ts.team;
unsigned tid = gomp_thread ()->ts.team_id;
@@ -1203,7 +1204,8 @@ GOMP_loop_static_worksharing (long niter __attribute__ ((unused)))
/* OMPT variant enabled by -fopenmp-ompt. */
_Complex int
-GOMP_loop_static_worksharing_start (long niter __attribute__ ((unused)))
+GOMP_loop_static_worksharing_start (unsigned long long niter
+ __attribute__ ((unused)))
{
struct gomp_team *team = gomp_thread ()->ts.team;
unsigned tid = gomp_thread ()->ts.team_id;
diff --git a/libgomp/teams.c b/libgomp/teams.c
index 319cfee0db5..71c40ef5b46 100644
--- a/libgomp/teams.c
+++ b/libgomp/teams.c
@@ -63,7 +63,8 @@ GOMP_teams_reg (void (*fn) (void *), void *data, unsigned int num_teams,
number of iterations. */
_Complex int
-GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
+GOMP_distribute_static_worksharing (unsigned long long
+ niter __attribute__ ((unused)))
{
struct gomp_thread *thr = gomp_thread ();
unsigned tid = thr->team_num;
@@ -74,7 +75,8 @@ GOMP_distribute_static_worksharing (long niter __attribute__ ((unused)))
/* OMPT variant enabled by -fopenmp-ompt. */
_Complex int
-GOMP_distribute_static_worksharing_start (long niter __attribute__ ((unused)))
+GOMP_distribute_static_worksharing_start (unsigned long long
+ niter __attribute__ ((unused)))
{
struct gomp_thread *thr = gomp_thread ();
unsigned tid = thr->team_num;
diff --git a/libgomp/testsuite/libgomp.c-c++-common/for-static-1.c b/libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
index 3715a5efda4..2eb63e3d1ad 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/for-static-1.c
@@ -8,6 +8,6 @@
#include "for-static.h"
-/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing \\(n" 2 "optimized" } } */
+/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing \\(\[^\\)\]" 2 "optimized" } } */
/* { dg-final { scan-offload-tree-dump-not "omp_get_num_teams" "optimized" } } */
/* { dg-final { scan-offload-tree-dump-not "omp_get_team_num" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/for-static-2.c b/libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
index a0b15753155..ff756a38489 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/for-static-2.c
@@ -5,7 +5,7 @@
#include "for-static.h"
-/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_start \\(n" 2 "optimized" } } */
-/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_end \\(" 2 "optimized" } } */
+/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_start \\(\[^\\)\]" 2 "optimized" } } */
+/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_end \\(\\)" 2 "optimized" } } */
/* { dg-final { scan-offload-tree-dump-not "GOMP_distribute_static_worksharing \\(" "optimized" } } */
/* { dg-final { scan-offload-tree-dump-not "GOMP_distribute_static_worksharing_dispatch" "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/for-static-3.c b/libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
index afbac717f62..d783b2c7c6b 100644
--- a/libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
+++ b/libgomp/testsuite/libgomp.c-c++-common/for-static-3.c
@@ -6,6 +6,6 @@
#include "for-static.h"
-/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_start \\(n" 2 "optimized" } } */
-/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_dispatch \\(" 2 "optimized" } } */
-/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_end \\(" 2 "optimized" } } */
+/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_start \\(\[^\\)\]" 2 "optimized" } } */
+/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_dispatch \\(\\)" 2 "optimized" } } */
+/* { dg-final { scan-offload-tree-dump-times "GOMP_distribute_static_worksharing_end \\(\\)" 2 "optimized" } } */
--
2.53.0
