On Sun, Oct 19, 2025 at 2:27 PM H.J. Lu <[email protected]> wrote:
>
> On Fri, Oct 17, 2025 at 5:40 PM Josef Melcr <[email protected]> wrote:
> >
> > Hi,
> > I pushed the patched below to master. It is largely the same as the v5
> > (link: https://gcc.gnu.org/pipermail/gcc-patches/2025-October/697470.html),
> > but I took out all the oacc bits, as they did not apply (discussed in
> > https://gcc.gnu.org/pipermail/gcc-patches/2025-October/697527.html).
> >
> > ==========
> >
> > This patch enables constant propagation to outlined OpenMP kernels.
> > It does so using a new function attribute called ' callback' (note the
> > space).
> >
> > The attribute ' callback' captures the notion of a function calling one
> > of its arguments with some of its parameters as arguments. An OpenMP
> > example of such function is GOMP_parallel.
> > We implement the attribute with new callgraph edges called callback
> > edges. They are imaginary edges pointing from the caller of the function
> > with the attribute (e.g. caller of GOMP_parallel) to the body function
> > itself (e.g. the outlined OpenMP body). They share their call statement
> > with the edge from which they are derived (direct edge caller ->
> > GOMP_parallel
> > in this case). These edges allow passes such as ipa-cp to see the hidden
> > call site to the body function and optimize the function accordingly.
> >
> > To illustrate on an example, the body GOMP_parallel looks something
> > like this:
> >
> > void GOMP_parallel (void (*fn) (void *), void *data, /* ... */)
> > {
> > /* ... */
> > fn (data);
> > /* ... */
> > }
> >
> >
> > If we extend it with the attribute ' callback(1, 2)', we express that the
> > function calls its first argument and passes it its second argument.
> > This is represented in the call graph in this manner:
> >
> > direct indirect
> > caller -----------------> GOMP_parallel ---------------> fn
> > |
> > ----------------------> fn
> > callback
> >
> > The direct edge is then the callback-carrying edge, all new edges
> > are the derived callback edges.
> > While constant propagation is the main focus of this patch, callback
> > edges can be useful for different passes (for example, they improve icf
> > for OpenMP kernels), as they allow for address redirection.
> > If the outlined body function gets optimized and cloned, from body_fn to
> > body_fn.optimized, the callback edge allows us to replace the
> > address in the arguments list:
> >
> > GOMP_parallel (body_fn, &data_struct, /* ... */);
> >
> > becomes
> >
> > GOMP_parallel (body_fn.optimized, &data_struct, /* ... */);
> >
> > This redirection is possible for any function with the attribute.
> >
> > This callback attribute implementation is partially compatible with
> > clang's implementation. Its semantics, arguments and argument indexing
> > style are
> > the same, but we represent an unknown argument position with 0
> > (precedent set by attributes such as 'format'), while clang uses -1 or '?'.
> > We use the index 1 for the 'this' pointer in member functions, clang
> > uses 0. We also allow for multiple callback attributes on the same function,
> > while clang only allows one.
> >
> > The attribute is currently for GCC internal use only, thanks to the
> > space in its name. Originally, it was supposed to be called
> > 'callback' like its clang counterpart, but we cannot use this name, as
> > clang uses non-standard indexing style, leading to inconsistencies. The
> > attribute will be introduced into the public API as 'gnu::callback_only'
> > in a future patch.
> >
> > The attribute allows us to propagate constants into body functions of
> > OpenMP constructs. Currently, GCC won't propagate the value 'c' into the
> > OpenMP body in the following example:
> >
> > int a[100];
> > void test(int c) {
> > #pragma omp parallel for
> > for (int i = 0; i < c; i++) {
> > if (!__builtin_constant_p(c)) {
> > __builtin_abort();
> > }
> > a[i] = i;
> > }
> > }
> > int main() {
> > test(100);
> > return a[5] - 5;
> > }
> >
> > With this patch, the body function will get cloned and the constant 'c'
> > will get propagated.
> >
> > Some functions may utilize the attribute's infrastructure without being
> > declared with it, for example GOMP_task. These functions are special
> > cases and use the special case functions found in attr-callback.h. Special
> > cases use the attribute under certain circumstances, for example
> > GOMP_task uses it when the copy function is not being used required.
> >
> > gcc/ChangeLog:
> >
> > * Makefile.in: Add attr-callback.o to OBJS.
> > * builtin-attrs.def (ATTR_CALLBACK): Callback attr identifier.
> > (DEF_CALLBACK_ATTRIBUTE): Macro for callback attr creation.
> > (GOMP): Attr for libgomp functions.
> > (ATTR_CALLBACK_GOMP_LIST): ATTR_NOTHROW_LIST with GOMP callback
> > attr added.
> > * cgraph.cc (cgraph_add_edge_to_call_site_hash): Always hash the
> > callback-carrying edge.
> > (cgraph_node::get_edge): Always return the callback-carrying
> > edge.
> > (cgraph_edge::set_call_stmt): Add cascade for callback edges.
> > (symbol_table::create_edge): Allow callback edges to share call
> > stmts, initialize new flags.
> > (cgraph_edge::make_callback): New method, derives a new callback
> > edge.
> > (cgraph_edge::get_callback_carrying_edge): New method.
> > (cgraph_edge::first_callback_edge): Likewise.
> > (cgraph_edge::next_callback_edge): Likewise.
> > (cgraph_edge::purge_callback_edges): Likewise.
> > (cgraph_edge::redirect_callee): When redirecting a callback
> > edge, redirect its ref as well.
> > (cgraph_edge::redirect_call_stmt_to_callee): Add callback edge
> > redirection logic, set update_derived_edges to true hwne
> > redirecting the carrying edge.
> > (cgraph_node::remove_callers): Add cascade for callback edges.
> > (cgraph_edge::dump_edge_flags): Print callback flags.
> > (cgraph_node::verify_node): Add sanity checks for callback
> > edges.
> > * cgraph.h: Add new 1 bit flags and 16 bit callback_id to
> > cgraph_edge class.
> > * cgraphclones.cc (cgraph_edge::clone): Copy over callback data.
> > * cif-code.def (CALLBACK_EDGE): Add CIF_CALLBACK_EDGE code.
> > * ipa-cp.cc (purge_useless_callback_edges): New function,
> > deletes callback edges when necessary.
> > (ipcp_decision_stage): Call purge_useless_callback_edges.
> > * ipa-fnsummary.cc (ipa_call_summary_t::duplicate): Add
> > an exception for callback edges.
> > (analyze_function_body): Copy over summary from carrying to
> > callback edge.
> > * ipa-inline-analysis.cc (do_estimate_growth_1): Skip callback
> > edges when estimating growth.
> > * ipa-inline-transform.cc (inline_transform): Add redirection
> > cascade for callback edges.
> > * ipa-param-manipulation.cc
> > (drop_decl_attribute_if_params_changed_p): New function.
> > (ipa_param_adjustments::build_new_function_type): Add
> > args_modified out param.
> > (ipa_param_adjustments::adjust_decl): Drop callback attrs when
> > modifying args.
> > * ipa-param-manipulation.h: Adjust decl of
> > build_new_function_type.
> > * ipa-prop.cc (ipa_duplicate_jump_function): Add decl.
> > (init_callback_edge_summary): New function.
> > (ipa_compute_jump_functions_for_edge): Add callback edge
> > creation logic.
> > * lto-cgraph.cc (lto_output_edge): Stream out callback data.
> > (input_edge): Input callback data.
> > * omp-builtins.def (BUILT_IN_GOMP_PARALLEL_LOOP_STATIC): Use new
> > attr list.
> > (BUILT_IN_GOMP_PARALLEL_LOOP_GUIDED): Likewise.
> > (BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_DYNAMIC): Likewise.
> > (BUILT_IN_GOMP_PARALLEL_LOOP_NONMONOTONIC_RUNTIME): Likewise.
> > (BUILT_IN_GOMP_PARALLEL): Likewise.
> > (BUILT_IN_GOMP_PARALLEL_SECTIONS): Likewise.
> > (BUILT_IN_GOMP_TEAMS_REG): Likewise.
> > * tree-core.h (ECF_CB_1_2): New constant for callback(1,2).
> > * tree-inline.cc (copy_bb): Copy callback edges when copying the
> > carrying edge.
> > (redirect_all_calls): Redirect callback edges.
> > * tree.cc (set_call_expr_flags): Create callback attr according
> > to the ECF_CB flag.
> > * attr-callback.cc: New file.
> > * attr-callback.h: New file.
> >
> > gcc/c-family/ChangeLog:
> >
> > * c-attribs.cc: Define callback attr.
> >
> > gcc/fortran/ChangeLog:
> >
> > * f95-lang.cc (ATTR_CALLBACK_GOMP_LIST): New attr list
> > corresponding to the list in builtin-attrs.def.
> >
> > gcc/testsuite/ChangeLog:
> >
> > * gcc.dg/ipa/ipcp-cb-spec1.c: New test.
> > * gcc.dg/ipa/ipcp-cb-spec2.c: New test.
> > * gcc.dg/ipa/ipcp-cb1.c: New test.
>
> These run-time tests should be moved to libgomp to avoid
>
> xgcc: fatal error: cannot read spec file 'libgomp.spec': No such file
> or directory
> compilation terminated.
> compiler exited with status 1
> FAIL: gcc.dg/ipa/ipcp-cb1.c (test for excess errors)
No what Josef is doing is correct.
The testcases have already:
/* { dg-require-effective-target fopenmp } */
That is how gcc.dg/gomp/gomp.exp and others handles this already.
e.g. c-c++-common/gomp/assume-3.c c-c++-common/gomp/pr60823-2.c
gcc.dg/gomp/pr114075.c etc.
target-supports.exp has:
```
proc check_effective_target_fopenmp {} {
# nvptx/amdgcn can be built with the device-side bits of libgomp, but it
# does not make sense to test it as an openmp host.
if [istarget nvptx-*-*] { return 0 }
if [istarget amdgcn-*-*] { return 0 }
return [check_no_compiler_messages fopenmp object {
void foo (void) { }
} "-fopenmp"]
}
```
I am not sure why this is not working for you. Unless the other
testcases are failing in a similar way.
>
> --
> H.J.
