tra added a comment.
Oops. forgot to hit 'submit' last week, so there's some overlap with Aaron's
question.
================
Comment at: clang/docs/LanguageExtensions.rst:275
+ // On amdgcn target
+ #if __has_target_feature("s-memtime-inst")
+ x = __builtin_amdgcn_s_memtime();
----------------
yaxunl wrote:
> tra wrote:
> > Do you have a better example? This particular case could've been handled by
> > existing `__has_builtin()`.
> >
> > While I could see usefulness of checking features (e.g. for CUDA/NVPTX it
> > could be used to chose inline assembly supported only by newer PTX
> > versions, but even then that choice could be made using existing methods,
> > even if they are not always direct (e.g. by using CUDA_VERSION as a proxy
> > for "new enough PTX version").
> >
> `__has_builtin` returns 1 as long as the builtin is known to clang, even if
> the builtin is not supported by the target CPU. This is because the required
> target feature for a builtin is in ASTContext, whereas `__has_builtin` is
> evaluated in preprocessor, where the information is not known.
I'm missing something. `__has_target_feature("s-memtime-inst")` is also
evaluated by preprocessor, right next to where `__has_target_builtin` is
processed.
I guess what you're saying is that preprocessor does not pay attention to the
target feature conditions attached to those builtins.
Those are evaluted in `CodeGenFunction::checkTargetFeatures`.
This means that in order to use `__has_feature()` to detect availability of
target builtins would effectively require the user to specify exactly the same
set of conditions, as applied to the builtin. That will work for builtins that
map to features 1:1, but it will be prone to breaking for NVPTX builtins that
carry fairly large set of feature requirements and that set changes every time
we add a new PTX or GPU variant, which happens fairly often.
I wonder if it may be better to generalize target builtin feature evaluation
and use it from both preprocessor and the codegen and just stick with
`__has_builtin`, only now working correctly to indicate availability of the
target builtins.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D125555/new/
https://reviews.llvm.org/D125555
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
