r264964 - [CUDA] Make unattributed constexpr functions implicitly host+device.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 30 16:30:22 PDT 2016
Author: jlebar
Date: Wed Mar 30 18:30:21 2016
New Revision: 264964
URL: http://llvm.org/viewvc/llvm-project?rev=264964&view=rev
Log:
[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
Added:
cfe/trunk/test/SemaCUDA/Inputs/overload.h
cfe/trunk/test/SemaCUDA/host-device-constexpr.cu
cfe/trunk/test/SemaCUDA/no-host-device-constexpr.cu
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/include/clang/Basic/LangOptions.def
cfe/trunk/include/clang/Driver/CC1Options.td
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/Frontend/CompilerInvocation.cpp
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaDecl.cpp
cfe/trunk/lib/Sema/SemaOverload.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Wed Mar 30 18:30:21 2016
@@ -6491,6 +6491,12 @@ def err_variadic_device_fn : Error<
def err_va_arg_in_device : Error<
"CUDA device code does not support va_arg">;
def err_alias_not_supported_on_nvptx : Error<"CUDA does not support aliases">;
+def err_cuda_unattributed_constexpr_cannot_overload_device : Error<
+ "constexpr function '%0' without __host__ or __device__ attributes cannot "
+ "overload __device__ function with same signature. Add a __host__ "
+ "attribute, or build with -fno-cuda-host-device-constexpr.">;
+def note_cuda_conflicting_device_function_declared_here : Note<
+ "conflicting __device__ function declared here">;
def err_dynamic_var_init : Error<
"dynamic initialization is not supported for "
"__device__, __constant__, and __shared__ variables.">;
Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Wed Mar 30 18:30:21 2016
@@ -172,6 +172,7 @@ LANGOPT(OpenMPIsDevice , 1, 0, "Gener
LANGOPT(CUDAIsDevice , 1, 0, "compiling for CUDA device")
LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "allowing variadic functions in CUDA device code")
+LANGOPT(CUDAHostDeviceConstexpr, 1, 1, "treating unattributed constexpr functions as __host__ __device__")
LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Wed Mar 30 18:30:21 2016
@@ -691,6 +691,8 @@ def fcuda_include_gpubinary : Separate<[
HelpText<"Incorporate CUDA device-side binary into host object file.">;
def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">,
HelpText<"Allow variadic functions in CUDA device code.">;
+def fno_cuda_host_device_constexpr : Flag<["-"], "fno-cuda-host-device-constexpr">,
+ HelpText<"Don't treat unattributed constexpr functions as __host__ __device__.">;
//===----------------------------------------------------------------------===//
// OpenMP Options
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Wed Mar 30 18:30:21 2016
@@ -2192,7 +2192,8 @@ public:
const LookupResult &OldDecls,
NamedDecl *&OldDecl,
bool IsForUsingDecl);
- bool IsOverload(FunctionDecl *New, FunctionDecl *Old, bool IsForUsingDecl);
+ bool IsOverload(FunctionDecl *New, FunctionDecl *Old, bool IsForUsingDecl,
+ bool ConsiderCudaAttrs = true);
/// \brief Checks availability of the function depending on the current
/// function context.Inside an unavailable function,unavailability is ignored.
@@ -8904,6 +8905,11 @@ public:
return IdentifyCUDAPreference(Caller, Callee) == CFP_Never;
}
+ /// May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD,
+ /// depending on FD and the current compilation settings.
+ void maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *FD,
+ const LookupResult &Previous);
+
/// Finds a function in \p Matches with highest calling priority
/// from \p Caller context and erases all functions with lower
/// calling priority.
Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Wed Mar 30 18:30:21 2016
@@ -1560,6 +1560,9 @@ static void ParseLangArgs(LangOptions &O
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();
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Wed Mar 30 18:30:21 2016
@@ -11,12 +11,14 @@
///
//===----------------------------------------------------------------------===//
-#include "clang/Sema/Sema.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/Decl.h"
#include "clang/AST/ExprCXX.h"
#include "clang/Lex/Preprocessor.h"
+#include "clang/Sema/Lookup.h"
+#include "clang/Sema/Sema.h"
#include "clang/Sema/SemaDiagnostic.h"
+#include "clang/Sema/Template.h"
#include "llvm/ADT/Optional.h"
#include "llvm/ADT/SmallVector.h"
using namespace clang;
@@ -381,3 +383,50 @@ bool Sema::isEmptyCudaConstructor(Source
return true;
}
+
+// With -fcuda-host-device-constexpr, an unattributed constexpr function is
+// treated as implicitly __host__ __device__, unless:
+// * it is a variadic function (device-side variadic functions are not
+// allowed), or
+// * a __device__ function with this signature was already declared, in which
+// case in which case we output an error, unless the __device__ decl is in a
+// system header, in which case we leave the constexpr function unattributed.
+void Sema::maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *NewD,
+ const LookupResult &Previous) {
+ assert(getLangOpts().CUDA && "May be called only for CUDA compilations.");
+ if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() ||
+ NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() ||
+ NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>())
+ return;
+
+ // Is D a __device__ function with the same signature as NewD, ignoring CUDA
+ // attributes?
+ auto IsMatchingDeviceFn = [&](NamedDecl *D) {
+ if (UsingShadowDecl *Using = dyn_cast<UsingShadowDecl>(D))
+ D = Using->getTargetDecl();
+ FunctionDecl *OldD = D->getAsFunction();
+ return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
+ !OldD->hasAttr<CUDAHostAttr>() &&
+ !IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
+ /* ConsiderCudaAttrs = */ false);
+ };
+ auto It = llvm::find_if(Previous, IsMatchingDeviceFn);
+ if (It != Previous.end()) {
+ // We found a __device__ function with the same name and signature as NewD
+ // (ignoring CUDA attrs). This is an error unless that function is defined
+ // in a system header, in which case we simply return without making NewD
+ // host+device.
+ NamedDecl *Match = *It;
+ if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
+ Diag(NewD->getLocation(),
+ diag::err_cuda_unattributed_constexpr_cannot_overload_device)
+ << NewD->getName();
+ Diag(Match->getLocation(),
+ diag::note_cuda_conflicting_device_function_declared_here);
+ }
+ return;
+ }
+
+ NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
+ NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+}
Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Wed Mar 30 18:30:21 2016
@@ -8009,6 +8009,9 @@ Sema::ActOnFunctionDeclarator(Scope *S,
// Handle attributes.
ProcessDeclAttributes(S, NewFD, D);
+ if (getLangOpts().CUDA)
+ maybeAddCUDAHostDeviceAttrs(S, NewFD, Previous);
+
if (getLangOpts().OpenCL) {
// OpenCL v1.1 s6.5: Using an address space qualifier in a function return
// type declaration will generate a compilation error.
Modified: cfe/trunk/lib/Sema/SemaOverload.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOverload.cpp?rev=264964&r1=264963&r2=264964&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOverload.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOverload.cpp Wed Mar 30 18:30:21 2016
@@ -992,7 +992,7 @@ Sema::CheckOverload(Scope *S, FunctionDe
}
bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old,
- bool UseMemberUsingDeclRules) {
+ bool UseMemberUsingDeclRules, bool ConsiderCudaAttrs) {
// C++ [basic.start.main]p2: This function shall not be overloaded.
if (New->isMain())
return false;
@@ -1125,7 +1125,7 @@ bool Sema::IsOverload(FunctionDecl *New,
return true;
}
- if (getLangOpts().CUDA) {
+ if (getLangOpts().CUDA && ConsiderCudaAttrs) {
CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New),
OldTarget = IdentifyCUDATarget(Old);
if (NewTarget == CFT_InvalidTarget || NewTarget == CFT_Global)
Added: cfe/trunk/test/SemaCUDA/Inputs/overload.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/Inputs/overload.h?rev=264964&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/Inputs/overload.h (added)
+++ cfe/trunk/test/SemaCUDA/Inputs/overload.h Wed Mar 30 18:30:21 2016
@@ -0,0 +1,8 @@
+// This header is used by tests which are interested in __device__ functions
+// which appear in a system header.
+
+__device__ int OverloadMe();
+
+namespace ns {
+using ::OverloadMe;
+}
Added: cfe/trunk/test/SemaCUDA/host-device-constexpr.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/host-device-constexpr.cu?rev=264964&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/host-device-constexpr.cu (added)
+++ cfe/trunk/test/SemaCUDA/host-device-constexpr.cu Wed Mar 30 18:30:21 2016
@@ -0,0 +1,69 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -isystem %S/Inputs %s -fcuda-is-device
+
+#include "Inputs/cuda.h"
+
+// Declares one function and pulls it into namespace ns:
+//
+// __device__ int OverloadMe();
+// namespace ns { using ::OverloadMe; }
+//
+// Clang cares that this is done in a system header.
+#include <overload.h>
+
+// Opaque type used to determine which overload we're invoking.
+struct HostReturnTy {};
+
+// These shouldn't become host+device because they already have attributes.
+__host__ constexpr int HostOnly() { return 0; }
+// expected-note at -1 0+ {{not viable}}
+__device__ constexpr int DeviceOnly() { return 0; }
+// expected-note at -1 0+ {{not viable}}
+
+constexpr int HostDevice() { return 0; }
+
+// This should be a host-only function, because there's a previous __device__
+// overload in <overload.h>.
+constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
+
+namespace ns {
+// The "using" statement in overload.h should prevent OverloadMe from being
+// implicitly host+device.
+constexpr HostReturnTy OverloadMe() { return HostReturnTy(); }
+} // namespace ns
+
+// This is an error, because NonSysHdrOverload was not defined in a system
+// header.
+__device__ int NonSysHdrOverload() { return 0; }
+// expected-note at -1 {{conflicting __device__ function declared here}}
+constexpr int NonSysHdrOverload() { return 0; }
+// expected-error at -1 {{constexpr function 'NonSysHdrOverload' without __host__ or __device__ attributes}}
+
+// Variadic device functions are not allowed, so this is just treated as
+// host-only.
+constexpr void Variadic(const char*, ...);
+// expected-note at -1 {{call to __host__ function from __device__ function}}
+
+__host__ void HostFn() {
+ HostOnly();
+ DeviceOnly(); // expected-error {{no matching function}}
+ HostReturnTy x = OverloadMe();
+ HostReturnTy y = ns::OverloadMe();
+ Variadic("abc", 42);
+}
+
+__device__ void DeviceFn() {
+ HostOnly(); // expected-error {{no matching function}}
+ DeviceOnly();
+ int x = OverloadMe();
+ int y = ns::OverloadMe();
+ Variadic("abc", 42); // expected-error {{no matching function}}
+}
+
+__host__ __device__ void HostDeviceFn() {
+#ifdef __CUDA_ARCH__
+ int y = OverloadMe();
+#else
+ constexpr HostReturnTy y = OverloadMe();
+#endif
+}
Added: cfe/trunk/test/SemaCUDA/no-host-device-constexpr.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/no-host-device-constexpr.cu?rev=264964&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/no-host-device-constexpr.cu (added)
+++ cfe/trunk/test/SemaCUDA/no-host-device-constexpr.cu Wed Mar 30 18:30:21 2016
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fno-cuda-host-device-constexpr -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fno-cuda-host-device-constexpr -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+// Check that, with -fno-cuda-host-device-constexpr, constexpr functions are
+// host-only, and __device__ constexpr functions are still device-only.
+
+constexpr int f() { return 0; } // expected-note {{not viable}}
+__device__ constexpr int g() { return 0; } // expected-note {{not viable}}
+
+void __device__ foo() {
+ f(); // expected-error {{no matching function}}
+ g();
+}
+
+void __host__ foo() {
+ f();
+ g(); // expected-error {{no matching function}}
+}
More information about the cfe-commits
mailing list