AlexVlx wrote: > @efriedma-quic was kind enough to have a call where we discussed this a bit > more. I'll update tomorrow with a potential way forward, for the group's > consideration.
Following up, here's a possible approach to making progress, broken down in phases, (@efriedma-quic can correct me if I am misrepresenting any of these): 1. Have what is proposed here as an initial step, with the addition that we issue warnings on unguarded uses of builtins / ASM (similar to what `__builtin_available` / `@available` do), and we clean-up non-extern functions that become unreachable as a consequence of predicate expansion (i.e. `foo` can only be called from within this module, and it was only being called from a predicate guarded block, which was removed); 2. Add attribute based checking for predicate guarded areas: - Functions can be annotated either with the existing `target` attribute or with a new `target_can_invoke` (name up for bike-shedding) attribute; - Within a predicate guarded scope, if we encounter contradictions, e.g. we call a `target("gfx9000")` function, or a `target_can_invoke(builtin_only_on_gfx9000)`, within a `__builtin_amdgcn_processor_is("gfx8999")`, that is an error - This should reward users that go through the effort of annotating their functions, making it much harder to write bugs - I'm not entirely sure how to do this well yet (nested guarded regions, where to track the currently active guard etc.), and it probably needs a bit more design, hence why it's a different phase - It is a pre-requisite for any attempt at making these general, rather than target specific 3. In relation with generalisation, if we go in that direction (i.e. other targets are interested / we think there's merit into hoisting these into generic Clang builtins), we will have to look at whether or not we want a different IR representation (possibly / probably along the lines of what has been discussed here), for cases where a target must run some potentially disruptive optimisations before and cannot just do the expansion right after Clang. https://github.com/llvm/llvm-project/pull/134016 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits