yaxunl marked an inline comment as done.
yaxunl added inline comments.
================
Comment at: clang/docs/LanguageExtensions.rst:275
+ // On amdgcn target
+ #if __has_target_feature("s-memtime-inst")
+ x = __builtin_amdgcn_s_memtime();
----------------
yaxunl wrote:
> yaxunl wrote:
> > aaron.ballman wrote:
> > > tra wrote:
> > > > 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.
> > > >
> > > >
> > > >
> > > >
> > > >
> > > > `__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.
> > >
> > > Why is that not a deficiency with `__has_builtin` that we should fix?
> > It is arguable whether this is a bug for `__has_builtin`. I tend to treat
> > it as a bug and think it should be fixed. However, fixing it may cause
> > regressions since there may be existing code expecting `__has_builtin` not
> > depending on availability of required target feature. This will limit the
> > usability of a fix for this issue.
> >
> > Even if we agree that this is a bug for `__has_builtin` which should be
> > fixed, this does conflict with adding `__has_target_feature`, as
> > `__has_target_feature` has more uses than determining whether a target
> > builtin is available, e.g. choosing different algorithms based on
> > availability of certain target feature, or using certain inline assembly
> > code.
> > It is arguable whether this is a bug for `__has_builtin`. I tend to treat
> > it as a bug and think it should be fixed. However, fixing it may cause
> > regressions since there may be existing code expecting `__has_builtin` not
> > depending on availability of required target feature. This will limit the
> > usability of a fix for this issue.
> >
> > Even if we agree that this is a bug for `__has_builtin` which should be
> > fixed, this does conflict with adding `__has_target_feature`, as
> > `__has_target_feature` has more uses than determining whether a target
> > builtin is available, e.g. choosing different algorithms based on
> > availability of certain target feature, or using certain inline assembly
> > code.
>
> sorry typo. this does not conflict
Opened https://reviews.llvm.org/D125829 to fix `__has_builtin`
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
