[PATCH] D125555: [clang] Add __has_target_feature

[Artem Belevich via Phabricator via cfe-commits]

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