Hi Thomas, On 07/07/2026 09:19, Thomas Schwinge wrote:
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]>
I had run the testsuite on amdgcn and fixed the loop-counter-variable-type issue, but it seems to have gotten lost in a rebase or something :sigh: Anyway, thanks for the fixup patch -- I re-ran the testsuite with it and it now works.
The attached patch has it squashed in, the trailer added and the ChangeLog updated.
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. ;-)
I agree; the formatter being greedy...
Grüße ThomasThe 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.defindex 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.defindex 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 smallexamplereturns 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.cindex 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
-- PA
commit cfdd816add4fd5ca97ffc06d5f6e2a8f6b643a39 Author: Paul-Antoine Arras <[email protected]> Date: Wed Jul 1 18:51:07 2026 +0200 openmp: Pass iteration count to static worksharing builtins for OMPT Add parameter to GOMP_{loop,distribute}_static_worksharing{,_start} carrying the total number of loop iterations. niter is unused for now, pending OMPT implementation. 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_ULONGLONG): New. * omp-builtins.def (BUILT_IN_GOMP_LOOP_STATIC_WORKSHARING): Use BT_FN_COMPLEX_INT_ULONGLONG. (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_ULONGLONG): 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. Co-authored-by: Thomas Schwinge <[email protected]> diff --git gcc/builtin-types.def gcc/builtin-types.def index 55ed6f0cae1..d11e6e736a2 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_ULONGLONG, BT_COMPLEX_INT, BT_ULONGLONG) DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR) diff --git gcc/fortran/types.def gcc/fortran/types.def index 6f98e3e48ff..b694e689360 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_ULONGLONG, BT_COMPLEX_INT, BT_ULONGLONG) DEF_POINTER_TYPE (BT_PTR_FN_VOID_PTR, BT_FN_VOID_PTR) diff --git gcc/omp-builtins.def gcc/omp-builtins.def index 465aeebefdb..3b0d5572043 100644 --- gcc/omp-builtins.def +++ gcc/omp-builtins.def @@ -510,11 +510,14 @@ 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_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, + "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, @@ -523,10 +526,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, + "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, + "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 gcc/omp-expand.cc gcc/omp-expand.cc index ad380b27d06..bce930fc108 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_); 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), 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 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); + 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,45 @@ 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 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); + 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..09e533c931d 100644 --- gcc/testsuite/c-c++-common/gomp/for-static-1.c +++ 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 \\(" 2 "ompexp" } } */ -/* { dg-final { scan-tree-dump-times "GOMP_distribute_static_worksharing \\(" 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 gcc/testsuite/c-c++-common/gomp/for-static-2.c gcc/testsuite/c-c++-common/gomp/for-static-2.c index c3d25b97bed..d99f59d7471 100644 --- gcc/testsuite/c-c++-common/gomp/for-static-2.c +++ 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 \\(" 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_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-not "GOMP_loop_static_worksharing_dispatch" "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 gcc/testsuite/c-c++-common/gomp/for-static-3.c gcc/testsuite/c-c++-common/gomp/for-static-3.c index 6fd1da6c407..8e45223c353 100644 --- gcc/testsuite/c-c++-common/gomp/for-static-3.c +++ 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 \\(" 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_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 libgomp/config/gcn/teams.c libgomp/config/gcn/teams.c index 99f807ac876..4354131f9b7 100644 --- libgomp/config/gcn/teams.c +++ libgomp/config/gcn/teams.c @@ -39,10 +39,12 @@ 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 (unsigned long long niter + __attribute__ ((unused))) { int __lds *gomp_team_num = (int __lds *) GOMP_TEAM_NUM; unsigned tid = *gomp_team_num; @@ -53,7 +55,8 @@ GOMP_distribute_static_worksharing (void) /* OMPT variant enabled by -fopenmp-ompt. */ _Complex int -GOMP_distribute_static_worksharing_start (void) +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 libgomp/config/nvptx/teams.c libgomp/config/nvptx/teams.c index 3e6ea287276..e04f8bbbdb1 100644 --- libgomp/config/nvptx/teams.c +++ libgomp/config/nvptx/teams.c @@ -42,10 +42,12 @@ 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 (unsigned long long niter + __attribute__ ((unused))) { unsigned tid = __gomp_team_num; unsigned nteams = gomp_num_teams_var + 1; @@ -55,7 +57,8 @@ GOMP_distribute_static_worksharing (void) /* OMPT variant enabled by -fopenmp-ompt. */ _Complex int -GOMP_distribute_static_worksharing_start (void) +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 libgomp/libgomp.texi libgomp/libgomp.texi index 83b90e6dc53..5d29da0984d 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 (unsigned long 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 (unsigned long long) @end smallexample returns the team id and team count. diff --git libgomp/libgomp_g.h libgomp/libgomp_g.h index a3097d5b093..3c7787358ad 100644 --- libgomp/libgomp_g.h +++ 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 (void); /* 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 (void); +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 libgomp/loop.c libgomp/loop.c index d53d994ee34..2dd879e1cce 100644 --- libgomp/loop.c +++ libgomp/loop.c @@ -1188,10 +1188,12 @@ 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 (unsigned long long niter + __attribute__ ((unused))) { struct gomp_team *team = gomp_thread ()->ts.team; unsigned tid = gomp_thread ()->ts.team_id; @@ -1202,7 +1204,8 @@ GOMP_loop_static_worksharing (void) /* OMPT variant enabled by -fopenmp-ompt. */ _Complex int -GOMP_loop_static_worksharing_start (void) +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 libgomp/teams.c libgomp/teams.c index f501123d647..71c40ef5b46 100644 --- libgomp/teams.c +++ libgomp/teams.c @@ -59,10 +59,12 @@ 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 (unsigned long long + niter __attribute__ ((unused))) { struct gomp_thread *thr = gomp_thread (); unsigned tid = thr->team_num; @@ -73,7 +75,8 @@ GOMP_distribute_static_worksharing (void) /* OMPT variant enabled by -fopenmp-ompt. */ _Complex int -GOMP_distribute_static_worksharing_start (void) +GOMP_distribute_static_worksharing_start (unsigned long 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..2eb63e3d1ad 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 \\(\[^\\)\]" 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..ff756a38489 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_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 libgomp/testsuite/libgomp.c-c++-common/for-static-3.c libgomp/testsuite/libgomp.c-c++-common/for-static-3.c index e7b904eedc3..d783b2c7c6b 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_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" } } */
