[PATCH] D42513: [CUDA] Added partial support for CUDA-9.1

[Justin Lebar via Phabricator via cfe-commits]

jlebar added inline comments.


================
Comment at: clang/lib/Headers/__clang_cuda_device_functions.h:32
+
+#define __DEVICE__ static __device__ __forceinline__
+// There are number of functions that became compiler builtins in CUDA-9 and are
----------------
tra wrote:
> jlebar wrote:
> > I don't think we should need `__forceinline__`?  I'm sure they had that in their headers, but clang doesn't rely on its inliner for correctness, unlike some *other* compilers.  :)
> > 
> > I'm also not sure about `static`.  I think we mark all CUDA code as internal, so it shouldn't be necessary?  If they're not static, they should be marked as `inline`, though.
> The idea is to make those wrappers always inlined, so they are effectively replaced with a call to __nv_* libdevice function. Whether the callee itself get inlined is controlled by the attributes in the libdevice bitcode.
> 
> `__attribute__((always_inline))` which `__forceinline__` expands to is handled differently from regular inline. We have Transforms/IPO/AlwaysInliner.cpp which handles them with -O1. Higher opt levels will presumably inline these functions, too. Without always_inline these wrappers will remain uninlined at -O1. Replacing it with just `inline` will work OK (only affects -O1 in a minor way), but `__forceinline__` is a bit closer to what I want from these wrappers.
> 
> As for the `static`, unused non-static functions will be emitted by clang -- as far as clang is concerned they are externally visible. I've checked and we do end up emitting them in PTX, even though they don't have .visible directive (I think that's what you referred to when you mentioned 'mark as internal'). I don't think we want that.
> 
> I'll keep this combination for now. We can adjust it later, if necessary.
Can you help me understand what makes you want this particular inlining behavior (i.e., always inline, even at -O0) for these wrappers?

Does e.g. libc++ do the same thing for the functions it has that are thin wrappers around e.g. libc functions?

> As for the static, unused non-static functions will be emitted by clang -- as far as clang is concerned they are externally visible. I've checked and we do end up emitting them in PTX, even though they don't have .visible directive (I think that's what you referred to when you mentioned 'mark as internal'). I don't think we want that.

Agree.  This is kind of unfortunate, though -- it's not How You Write Headers.  Could we add a comment?


================
Comment at: clang/lib/Headers/__clang_cuda_runtime_wrapper.h:144
+// Declare or define device-side functions that particular CUDA version relies
+// on but does (no longer) declare in its headers. E.g. some of the device-side
+// functions that used to be implemented in a header in CUDA-8, became NVCC's
----------------
tra wrote:
> jlebar wrote:
> > Sentence doesn't read right if you skip the parens.
> > 
> > But also: Our header provides *everything* from cuda/device_functions_decls.h, not just the things that were once in that header and, in newer CUDA versions, are no longer there, right?  I'm actually even more confused now that I read the big header, because here we say the header "defines or declares" these functions, but that suggests that for every function, it decides whether to define or to declare it, but that's not the case...
> I should rephrase that then.
> 
> `__clang_cuda_device_functions.h` does two jobs. It always includes `__clang_cuda_libdevice_declares.h` which provides **declarations** for the libdevice functions. Those are gone in CUDA-9.1.
> 
> It also provides **definitions** for the standard library and __* device-side functions that call libdevice.
> 
> I can explicitly include _clang_cuda_libdevice_declares.h  here and update the comment to make more sense.
I think that would help clarify things, for sure.


https://reviews.llvm.org/D42513