[PATCH] D12122: [CUDA] Add appropriate host/device attribute to target-specific builtins.

[Artem Belevich via cfe-commits]

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