Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -6710,7 +6710,7 @@ "conflicting __device__ function declared here">; def err_cuda_device_exceptions : Error< "cannot use '%0' in " - "%select{__device__|__global__|__host__|__host__ __device__}1 function %2">; + "%select{__device__|__global__|__host__|__host__ __device__}1 function">; def err_dynamic_var_init : Error< "dynamic initialization is not supported for " "__device__, __constant__, and __shared__ variables.">; Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -9189,6 +9189,115 @@ QualType FieldTy, bool IsMsStruct, Expr *BitWidth, bool *ZeroWidth = nullptr); + /// Diagnostic builder for CUDA errors which may or may not be deferred. + /// + /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch) + /// which are not allowed to appear inside __device__ functions and are + /// allowed to appear in __host__ __device__ functions only if the host+device + /// function is never codegen'ed. + /// + /// To handle this, we use the notion of "deferred diagnostics", where we + /// attach a diagnostic to a FunctionDecl that's emitted iff it's codegen'ed. + /// + /// This class lets you emit either a regular diagnostic, a deferred + /// diagnostic, or no diagnostic at all, according to an argument you pass to + /// its constructor, thus simplifying the process of creating these "maybe + /// deferred" diagnostics. + class CUDADiagBuilder { + public: + enum Kind { + /// Emit no diagnostics. + NOP, + /// Emit the diagnostic immediately (i.e., behave like Sema::Diag()). + IMMEDIATE, + /// Create a deferred diagnostic, which is emitted only if the function + /// it's attached to is codegen'ed. + DEFERRED + }; + + CUDADiagBuilder(Kind K, SourceLocation Loc, unsigned DiagID, + FunctionDecl *Fn, Sema &S) { + switch (K) { + case NOP: + break; + case IMMEDIATE: + ImmediateDiagBuilder.emplace(S.Diag(Loc, DiagID)); + break; + case DEFERRED: + assert(Fn && "Must have a function to attach the deferred diag to."); + PartialDiagInfo.emplace(Loc, DiagID, Fn); + break; + } + } + + /// Convertible to bool: True if we immediately emitted an error, false if + /// we didn't emit an error or we created a deferred error. + /// + /// Example usage: + /// + /// if (CUDADiagBuilder(...) << foo << bar) + /// return ExprError(); + /// + operator bool() const { return ImmediateDiagBuilder.hasValue(); } + + template + friend const CUDADiagBuilder &operator<<(const CUDADiagBuilder &Diag, + const T &Value) { + if (Diag.ImmediateDiagBuilder.hasValue()) + *Diag.ImmediateDiagBuilder << Value; + else if (Diag.PartialDiagInfo.hasValue()) + Diag.PartialDiagInfo->PD << Value; + return Diag; + } + + private: + struct PartialDiagnosticInfo { + PartialDiagnosticInfo(SourceLocation Loc, unsigned DiagID, + FunctionDecl *Fn) + : Loc(Loc), PD(PartialDiagnostic::NullDiagnostic()), Fn(Fn) { + // We have to do this odd dance to create our PartialDiagnostic (first + // creating a NullDiagnostic(), then calling Reset()) because we want + // its storage to be allocated with operator new, not in an arena. + PD.Reset(DiagID); + } + + ~PartialDiagnosticInfo() { Fn->addDeferredDiag({Loc, std::move(PD)}); } + + SourceLocation Loc; + PartialDiagnostic PD; + FunctionDecl *Fn; + }; + + // Invariant: At most one of these Optionals has a value. + // FIXME: Switch these to a Variant once that exists. + llvm::Optional ImmediateDiagBuilder; + llvm::Optional PartialDiagInfo; + }; + + /// Creates a CUDADiagBuilder that emits the diagnostic if the current context + /// is "used as device code". + /// + /// - If CurContext is a __host__ function, does not emit any diagnostics. + /// - If CurContext is a __device__ or __global__ function, emits the + /// diagnostics immediately. + /// - If CurContext is a __host__ __device__ function and we are compiling for + /// the device, creates a deferred diagnostic which is emitted if and when + /// the function is codegen'ed. + /// + /// Example usage: + /// + /// // Variable-length arrays are not allowed in CUDA device code. + /// if (CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget()) + /// return ExprError(); + /// // Otherwise, continue parsing as normal. + CUDADiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID); + + /// Creates a CUDADiagBuilder that emits the diagnostic if the current context + /// is "used as host code". + /// + /// Same as CUDADiagIfDeviceCode, with "host" and "device" switched. + CUDADiagBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID); + enum CUDAFunctionTarget { CFT_Device, CFT_Global, @@ -9197,8 +9306,18 @@ CFT_InvalidTarget }; + /// Determines whether the given function is a CUDA device/host/kernel/etc. + /// function. + /// + /// Use this rather than examining the function's attributes yourself -- you + /// will get it wrong. Returns CFT_Host if D is null. CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D); + /// Gets the CUDA target for the current context. + CUDAFunctionTarget CurrentCUDATarget() { + return IdentifyCUDATarget(dyn_cast(CurContext)); + } + // CUDA function call preference. Must be ordered numerically from // worst to best. enum CUDAFunctionPreference { @@ -9249,21 +9368,6 @@ /// Otherwise, returns true without emitting any diagnostics. bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee); - /// Check whether a 'try' or 'throw' expression is allowed within the current - /// context, and raise an error or create a deferred error, as appropriate. - /// - /// 'try' and 'throw' are never allowed in CUDA __device__ functions, and are - /// allowed in __host__ __device__ functions only if those functions are never - /// codegen'ed for the device. - /// - /// ExprTy should be the string "try" or "throw", as appropriate. - bool CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy); - - /// Check whether it's legal for us to create a variable-length array in the - /// current context. Returns true if the VLA is OK; returns false and emits - /// an error otherwise. - bool CheckCUDAVLA(SourceLocation Loc); - /// Set __device__ or __host__ __device__ attributes on the given lambda /// operator() method. /// Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -42,6 +42,10 @@ /// IdentifyCUDATarget - Determine the CUDA compilation target for this function Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { + // Code that lives outside a function is run on the host. + if (D == nullptr) + return CFT_Host; + if (D->hasAttr()) return CFT_InvalidTarget; @@ -95,9 +99,8 @@ Sema::IdentifyCUDAPreference(const FunctionDecl *Caller, const FunctionDecl *Callee) { assert(Callee && "Callee must be valid."); + CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller); CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee); - CUDAFunctionTarget CallerTarget = - (Caller != nullptr) ? IdentifyCUDATarget(Caller) : Sema::CFT_Host; // If one of the targets is invalid, the check always fails, no matter what // the other target is. @@ -481,6 +484,45 @@ NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context)); } +Sema::CUDADiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc, + unsigned DiagID) { + assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); + CUDADiagBuilder::Kind DiagKind; + switch (CurrentCUDATarget()) { + case CFT_Global: + case CFT_Device: + DiagKind = CUDADiagBuilder::IMMEDIATE; + break; + case CFT_HostDevice: + DiagKind = getLangOpts().CUDAIsDevice ? CUDADiagBuilder::DEFERRED + : CUDADiagBuilder::NOP; + break; + default: + DiagKind = CUDADiagBuilder::NOP; + } + return CUDADiagBuilder(DiagKind, Loc, DiagID, + dyn_cast(CurContext), *this); +} + +Sema::CUDADiagBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc, + unsigned DiagID) { + assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); + CUDADiagBuilder::Kind DiagKind; + switch (CurrentCUDATarget()) { + case CFT_Host: + DiagKind = CUDADiagBuilder::IMMEDIATE; + break; + case CFT_HostDevice: + DiagKind = getLangOpts().CUDAIsDevice ? CUDADiagBuilder::NOP + : CUDADiagBuilder::DEFERRED; + break; + default: + DiagKind = CUDADiagBuilder::NOP; + } + return CUDADiagBuilder(DiagKind, Loc, DiagID, + dyn_cast(CurContext), *this); +} + bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) { assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); assert(Callee && "Callee may not be null."); @@ -488,75 +530,26 @@ if (!Caller) return true; - Sema::CUDAFunctionPreference Pref = IdentifyCUDAPreference(Caller, Callee); - if (Pref == Sema::CFP_Never) { - Diag(Loc, diag::err_ref_bad_target) << IdentifyCUDATarget(Callee) << Callee - << IdentifyCUDATarget(Caller); - Diag(Callee->getLocation(), diag::note_previous_decl) << Callee; - return false; + CUDADiagBuilder::Kind DiagKind; + switch (IdentifyCUDAPreference(Caller, Callee)) { + case CFP_Never: + DiagKind = CUDADiagBuilder::IMMEDIATE; + break; + case CFP_WrongSide: + assert(Caller && "WrongSide calls require a non-null caller"); + DiagKind = CUDADiagBuilder::DEFERRED; + break; + default: + DiagKind = CUDADiagBuilder::NOP; } - if (Pref == Sema::CFP_WrongSide) { - // We have to do this odd dance to create our PartialDiagnostic because we - // want its storage to be allocated with operator new, not in an arena. - PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()}; - ErrPD.Reset(diag::err_ref_bad_target); - ErrPD << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller); - Caller->addDeferredDiag({Loc, std::move(ErrPD)}); - PartialDiagnostic NotePD{PartialDiagnostic::NullDiagnostic()}; - NotePD.Reset(diag::note_previous_decl); - NotePD << Callee; - Caller->addDeferredDiag({Callee->getLocation(), std::move(NotePD)}); - - // This is not immediately an error, so return true. The deferred errors - // will be emitted if and when Caller is codegen'ed. - return true; - } - return true; -} - -bool Sema::CheckCUDAExceptionExpr(SourceLocation Loc, StringRef ExprTy) { - assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); - FunctionDecl *CurFn = dyn_cast(CurContext); - if (!CurFn) - return true; - CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn); - - // Raise an error immediately if this is a __global__ or __device__ function. - // If it's a __host__ __device__ function, enqueue a deferred error which will - // be emitted if the function is codegen'ed for device. - if (Target == CFT_Global || Target == CFT_Device) { - Diag(Loc, diag::err_cuda_device_exceptions) << ExprTy << Target << CurFn; - return false; - } - if (Target == CFT_HostDevice && getLangOpts().CUDAIsDevice) { - PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()}; - ErrPD.Reset(diag::err_cuda_device_exceptions); - ErrPD << ExprTy << Target << CurFn; - CurFn->addDeferredDiag({Loc, std::move(ErrPD)}); - return false; - } - return true; -} - -bool Sema::CheckCUDAVLA(SourceLocation Loc) { - assert(getLangOpts().CUDA && "Should only be called during CUDA compilation"); - FunctionDecl *CurFn = dyn_cast(CurContext); - if (!CurFn) - return true; - CUDAFunctionTarget Target = IdentifyCUDATarget(CurFn); - if (Target == CFT_Global || Target == CFT_Device) { - Diag(Loc, diag::err_cuda_vla) << Target; - return false; - } - if (Target == CFT_HostDevice && getLangOpts().CUDAIsDevice) { - PartialDiagnostic ErrPD{PartialDiagnostic::NullDiagnostic()}; - ErrPD.Reset(diag::err_cuda_vla); - ErrPD << Target; - CurFn->addDeferredDiag({Loc, std::move(ErrPD)}); - return false; - } - return true; + bool IsImmediateErr = + CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this) + << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller); + CUDADiagBuilder(DiagKind, Callee->getLocation(), diag::note_previous_decl, + Caller, *this) + << Callee; + return !IsImmediateErr; } void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) { Index: clang/lib/Sema/SemaExprCXX.cpp =================================================================== --- clang/lib/Sema/SemaExprCXX.cpp +++ clang/lib/Sema/SemaExprCXX.cpp @@ -685,7 +685,8 @@ // Exceptions aren't allowed in CUDA device code. if (getLangOpts().CUDA) - CheckCUDAExceptionExpr(OpLoc, "throw"); + CUDADiagIfDeviceCode(OpLoc, diag::err_cuda_device_exceptions) + << "throw" << CurrentCUDATarget(); if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope()) Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw"; Index: clang/lib/Sema/SemaStmt.cpp =================================================================== --- clang/lib/Sema/SemaStmt.cpp +++ clang/lib/Sema/SemaStmt.cpp @@ -3646,7 +3646,8 @@ // Exceptions aren't allowed in CUDA device code. if (getLangOpts().CUDA) - CheckCUDAExceptionExpr(TryLoc, "try"); + CUDADiagIfDeviceCode(TryLoc, diag::err_cuda_device_exceptions) + << "try" << CurrentCUDATarget(); if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope()) Diag(TryLoc, diag::err_omp_simd_region_cannot_use_stmt) << "try"; Index: clang/lib/Sema/SemaType.cpp =================================================================== --- clang/lib/Sema/SemaType.cpp +++ clang/lib/Sema/SemaType.cpp @@ -2249,8 +2249,8 @@ return QualType(); } // CUDA device code doesn't support VLAs. - if (getLangOpts().CUDA && T->isVariableArrayType() && !CheckCUDAVLA(Loc)) - return QualType(); + if (getLangOpts().CUDA && T->isVariableArrayType()) + CUDADiagIfDeviceCode(Loc, diag::err_cuda_vla) << CurrentCUDATarget(); // If this is not C99, extwarn about VLA's and C99 array size modifiers. if (!getLangOpts().C99) { Index: clang/test/SemaCUDA/exceptions-host-device.cu =================================================================== --- clang/test/SemaCUDA/exceptions-host-device.cu +++ clang/test/SemaCUDA/exceptions-host-device.cu @@ -14,8 +14,8 @@ throw NULL; try {} catch(void*) {} #ifndef HOST - // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function 'hd1'}} - // expected-error@-3 {{cannot use 'try' in __host__ __device__ function 'hd1'}} + // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}} + // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}} #endif } @@ -31,8 +31,8 @@ throw NULL; try {} catch(void*) {} #ifndef HOST - // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function 'hd3'}} - // expected-error@-3 {{cannot use 'try' in __host__ __device__ function 'hd3'}} + // expected-error@-3 {{cannot use 'throw' in __host__ __device__ function}} + // expected-error@-3 {{cannot use 'try' in __host__ __device__ function}} #endif } __device__ void call_hd3() { hd3(); } Index: clang/test/SemaCUDA/exceptions.cu =================================================================== --- clang/test/SemaCUDA/exceptions.cu +++ clang/test/SemaCUDA/exceptions.cu @@ -9,13 +9,13 @@ } __device__ void device() { throw NULL; - // expected-error@-1 {{cannot use 'throw' in __device__ function 'device'}} + // expected-error@-1 {{cannot use 'throw' in __device__ function}} try {} catch(void*) {} - // expected-error@-1 {{cannot use 'try' in __device__ function 'device'}} + // expected-error@-1 {{cannot use 'try' in __device__ function}} } __global__ void kernel() { throw NULL; - // expected-error@-1 {{cannot use 'throw' in __global__ function 'kernel'}} + // expected-error@-1 {{cannot use 'throw' in __global__ function}} try {} catch(void*) {} - // expected-error@-1 {{cannot use 'try' in __global__ function 'kernel'}} + // expected-error@-1 {{cannot use 'try' in __global__ function}} }