jlebar created this revision.
jlebar added a reviewer: rnk.
jlebar added subscribers: tra, cfe-commits.
Together these let you easily create diagnostics that
- are never emitted for host code
- are always emitted for __device__ and __global__ functions, and
- are emitted for __host__ __device__ functions iff these functions are
codegen'ed.
At the moment there are only three diagnostics that need this treatment,
but I have more to add, and it's not sustainable to write code for emitting
every such diagnostic twice, and from a special wrapper in SemaCUDA.cpp.
While we're at it, don't emit the function name in
err_cuda_device_exceptions: It's not necessary to print it, and making
this work in the new framework in the face of a null value for
dyn_cast<FunctionDecl>(CurContext) isn't worth the effort.
https://reviews.llvm.org/D25139
Files:
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/include/clang/Sema/Sema.h
clang/lib/Sema/SemaCUDA.cpp
clang/lib/Sema/SemaExprCXX.cpp
clang/lib/Sema/SemaStmt.cpp
clang/lib/Sema/SemaType.cpp
clang/test/SemaCUDA/exceptions-host-device.cu
clang/test/SemaCUDA/exceptions.cu
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}}
}
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/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/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/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/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<CUDAInvalidTargetAttr>())
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,83 +484,53 @@
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<FunctionDecl>(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.");
FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
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<FunctionDecl>(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<FunctionDecl>(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;
+ CUDADiagBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
+ << IdentifyCUDATarget(Callee) << Callee << IdentifyCUDATarget(Caller);
+ return (CUDADiagBuilder(DiagKind, Callee->getLocation(),
+ diag::note_previous_decl, Caller, *this)
+ << Callee)
+ .IsDeferredOrNop();
}
void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -9189,16 +9189,123 @@
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;
+ }
+ }
+
+ /// Returns true if our Kind is DEFERRED or NOP.
+ bool IsDeferredOrNop() const { return !ImmediateDiagBuilder.hasValue(); }
+
+ template <typename T>
+ 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<Sema::SemaDiagnosticBuilder> ImmediateDiagBuilder;
+ llvm::Optional<PartialDiagnosticInfo> 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()).IsDeferredOrNop())
+ /// return ExprError();
+ /// // Otherwise, continue parsing as normal.
+ CUDADiagBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+
enum CUDAFunctionTarget {
CFT_Device,
CFT_Global,
CFT_Host,
CFT_HostDevice,
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<FunctionDecl>(CurContext));
+ }
+
// CUDA function call preference. Must be ordered numerically from
// worst to best.
enum CUDAFunctionPreference {
@@ -9249,21 +9356,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/include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -6711,7 +6711,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.">;
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
