From ba122ab42fe54aee3427dc61b765cc8a9dad9d85 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 llvm-svn: 264964 --- clang/lib/Frontend/CompilerInvocation.cpp | 3 +++ 1 file changed, 3 insertions(+) (limited to 'clang/lib/Frontend/CompilerInvocation.cpp') diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index cc657ed..5bb036a 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/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.1