tra marked an inline comment as done.
================
Comment at: lib/Sema/SemaDecl.cpp:11166
@@ +11165,3 @@
+ !FD->hasAttr<CUDADeviceAttr>() && !FD->hasAttr<CUDAHostAttr>()) {
+ if (getLangOpts().CUDAIsDevice)
+ FD->addAttr(CUDADeviceAttr::CreateImplicit(Context,
FD->getLocation()));
----------------
eliben wrote:
> It is not immediately clear why you would mark target-specific builtins as
> __host__ in host compilation mode. So for example __builtin_ptx_read_tid_x
> would be callable from a host function when compiling in host mode? Can you
> clarify this with a comment here, and also add a relevant test?
Target triple used for particular compilation pass is assumed to be valid for
this particular compilation mode. Builtins provided by the target are therefore
assumed to be intended to be used in that compilation mode. builtins.cu already
includes test cases for using target-specific builtins from different targets
and and non-target-specific builtins in host and device functions in host and
device compilation.
Your example scenario (__builtin_ptx_read_tid_x getting host attribute) is
possible only if you compile in host mode *and* use --triple
nvptx-unknown-cuda. IMO, __host__ would be an appropriate attribute to assign
to builtins in this scenario, even though I don't think we can currently do
anything useful with NVPTX as the host at the moment. If you attempt to use
__builtin_ptx_read_tid_x() during host compilation using non-NVPTX target
(which is a typical scenario), then compilation will fail because that
particular builtin would not be available at all.
That said, another corner case would be compiling CUDA for device with the same
architecture as host. Again, it's not a practical scenario at the moment.
If/when we get to the point when we may want to do that, it should be easy to
check for it and treat builtins as __host__ __device__ which would reflect
their intended use domain.
http://reviews.llvm.org/D12122
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
