[PATCH] D18328: [CUDA] Add option to mark most functions inside <complex> as host+device.

Richard Smith via cfe-commits cfe-commits at lists.llvm.org
Mon Mar 21 16:32:54 PDT 2016


rsmith added inline comments.

================
Comment at: lib/Sema/SemaCUDA.cpp:464-465
@@ +463,4 @@
+// <complex> without passing -fcuda-allow-std-complex.
+// TODO: Output a nvcc-compat warning if you try to use a non-constexpr function
+// from <complex> -- nvcc only lets you use constexpr functions.
+bool Sema::declShouldBeCUDAHostDevice(const FunctionDecl &FD) {
----------------
jlebar wrote:
> rsmith wrote:
> > Does nvcc do this "`constexpr` implies `__host__ __device__`" thing only for functions declared within <complex>, or for all functions?
> > 
> > Another alternative strategy: a wrapper `<complex>` header that does this:
> > 
> >     #include // ... union of includes from libc++ and libstdc++ <complex>
> >     #define constexpr __host__ __device__ constexpr
> >     #include_next <complex>
> >     #undef constexpr
> > Does nvcc do this "constexpr implies __host__ __device__" thing only for functions declared within <complex>, or for all functions?
> 
> All functions.  Although std::complex is the main use I've observed.
> 
> > Another alternative strategy: a wrapper <complex> header that does this:
> 
> That one is quite clever, although I'm not sure about enumerating all of the includes from the headers.  I guess that should be reasonably stable...
> 
> I think I would like to get full complex support, though, if we can agree on a path towards that.  The current limitation is silly, it seems clear that people want this, and the constexpr thing gives you but a shadow of the actual library.
Supporting a "`constexpr` implies `__host__ __device__`" feature for all functions seems a lot cleaner than the approach taken by this patch, and will presumably improve NVCC compatibility in other cases too (though perhaps they're quite rare). This seems like a very odd pair of features to link in this way, but if we're going to have something weird like this to support existing NVCC-targeting code, using the same approach may be better. This would also mean we would not be further extending NVCC's extension.


http://reviews.llvm.org/D18328





More information about the cfe-commits mailing list