From fb35fc36e6e57748802fe2d18bd962daaa5bbaa1 Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Wed, 30 Mar 2016 23:30:21 +0000 Subject: [CUDA] Make unattributed constexpr functions implicitly host+device. With this patch, by a constexpr function is implicitly host+device unless: a) it's a variadic function (variadic functions are not allowed on the device side), or b) it's preceeded by a __device__ overload in a system header. The restriction on overloading __host__ __device__ functions on the basis of their CUDA attributes remains in place, but we use (b) to allow us to define __device__ overloads for constexpr functions in cmath, which would otherwise be __host__ __device__ and thus not overloadable. You can disable this behavior with -fno-cuda-host-device-constexpr. Reviewers: tra, rnk, rsmith Subscribers: cfe-commits Differential Revision: http://reviews.llvm.org/D18380 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@264964 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Frontend/CompilerInvocation.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'lib/Frontend/CompilerInvocation.cpp') diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp index cc657ed5a3..5bb036ab26 100644 --- a/lib/Frontend/CompilerInvocation.cpp +++ b/lib/Frontend/CompilerInvocation.cpp @@ -1560,6 +1560,9 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, if (Args.hasArg(OPT_fcuda_allow_variadic_functions)) Opts.CUDAAllowVariadicFunctions = 1; + if (Args.hasArg(OPT_fno_cuda_host_device_constexpr)) + Opts.CUDAHostDeviceConstexpr = 0; + if (Opts.ObjC1) { if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) { StringRef value = arg->getValue(); -- cgit v1.2.1