r283637 - [CUDA] Do a better job at detecting wrong-side calls.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 7 18:07:11 PDT 2016
Author: jlebar
Date: Fri Oct 7 20:07:11 2016
New Revision: 283637
URL: http://llvm.org/viewvc/llvm-project?rev=283637&view=rev
Log:
[CUDA] Do a better job at detecting wrong-side calls.
Summary:
Move CheckCUDACall from ActOnCallExpr and BuildDeclRefExpr to
DiagnoseUseOfDecl. This lets us catch some edge cases we were missing,
specifically around class operators.
This necessitates a few other changes:
- Avoid emitting duplicate deferred diags in CheckCUDACall.
Previously we'd carefully placed our call to CheckCUDACall such that
it would only ever run once for a particular callsite. But now this
isn't the case.
- Emit deferred diagnostics from a template
specialization/instantiation's primary template, in addition to from
the specialization/instantiation itself. DiagnoseUseOfDecl ends up
putting the deferred diagnostics on the template, rather than the
specialization, so we need to check both.
Reviewers: rsmith
Subscribers: cfe-commits, tra
Differential Revision: https://reviews.llvm.org/D24573
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Sema/SemaCUDA.cpp
cfe/trunk/lib/Sema/SemaExpr.cpp
cfe/trunk/test/SemaCUDA/Inputs/cuda.h
cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=283637&r1=283636&r2=283637&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Fri Oct 7 20:07:11 2016
@@ -9267,16 +9267,27 @@ public:
void maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *FD,
const LookupResult &Previous);
+private:
+ /// Raw encodings of SourceLocations for which CheckCUDACall has emitted a
+ /// deferred "bad call" diagnostic. We use this to avoid emitting the same
+ /// deferred diag twice.
+ llvm::DenseSet<unsigned> LocsWithCUDACallDeferredDiags;
+
+public:
/// Check whether we're allowed to call Callee from the current context.
///
- /// If the call is never allowed in a semantically-correct program
- /// (CFP_Never), emits an error and returns false.
+ /// - If the call is never allowed in a semantically-correct program
+ /// (CFP_Never), emits an error and returns false.
+ ///
+ /// - If the call is allowed in semantically-correct programs, but only if
+ /// it's never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to
+ /// be emitted if and when the caller is codegen'ed, and returns true.
///
- /// If the call is allowed in semantically-correct programs, but only if it's
- /// never codegen'ed (CFP_WrongSide), creates a deferred diagnostic to be
- /// emitted if and when the caller is codegen'ed, and returns true.
+ /// Will only create deferred diagnostics for a given SourceLocation once,
+ /// so you can safely call this multiple times without generating duplicate
+ /// deferred errors.
///
- /// Otherwise, returns true without emitting any diagnostics.
+ /// - 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
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=283637&r1=283636&r2=283637&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Fri Oct 7 20:07:11 2016
@@ -2923,6 +2923,10 @@ void CodeGenModule::EmitGlobalFunctionDe
// non-error diags here, because order can be significant, e.g. with notes
// that follow errors.)
auto Diags = D->takeDeferredDiags();
+ if (auto *Templ = D->getPrimaryTemplate()) {
+ auto TemplDiags = Templ->getAsFunction()->takeDeferredDiags();
+ Diags.insert(Diags.end(), TemplDiags.begin(), TemplDiags.end());
+ }
bool HasError = llvm::any_of(Diags, [this](const PartialDiagnosticAt &PDAt) {
return getDiags().getDiagnosticLevel(PDAt.second.getDiagID(), PDAt.first) >=
DiagnosticsEngine::Error;
Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=283637&r1=283636&r2=283637&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Fri Oct 7 20:07:11 2016
@@ -495,7 +495,13 @@ bool Sema::CheckCUDACall(SourceLocation
Diag(Callee->getLocation(), diag::note_previous_decl) << Callee;
return false;
}
- if (Pref == Sema::CFP_WrongSide) {
+
+ // Insert into LocsWithCUDADeferredDiags to avoid emitting duplicate deferred
+ // diagnostics for the same location. Duplicate deferred diags are otherwise
+ // tricky to avoid, because, unlike with regular errors, sema checking
+ // proceeds unhindered when we omit a deferred diagnostic.
+ if (Pref == Sema::CFP_WrongSide &&
+ LocsWithCUDACallDeferredDiags.insert(Loc.getRawEncoding()).second) {
// 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()};
Modified: cfe/trunk/lib/Sema/SemaExpr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=283637&r1=283636&r2=283637&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExpr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExpr.cpp Fri Oct 7 20:07:11 2016
@@ -374,6 +374,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *
if (getLangOpts().CPlusPlus14 && FD->getReturnType()->isUndeducedType() &&
DeduceReturnType(FD, Loc))
return true;
+
+ if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
+ return true;
}
// [OpenMP 4.0], 2.15 declare reduction Directive, Restrictions
@@ -1743,11 +1746,6 @@ Sema::BuildDeclRefExpr(ValueDecl *D, Qua
const DeclarationNameInfo &NameInfo,
const CXXScopeSpec *SS, NamedDecl *FoundD,
const TemplateArgumentListInfo *TemplateArgs) {
- if (getLangOpts().CUDA)
- if (FunctionDecl *Callee = dyn_cast<FunctionDecl>(D))
- if (!CheckCUDACall(NameInfo.getLoc(), Callee))
- return ExprError();
-
bool RefersToCapturedVariable =
isa<VarDecl>(D) &&
NeedToCaptureVariable(cast<VarDecl>(D), NameInfo.getLoc());
@@ -5142,35 +5140,36 @@ static bool isNumberOfArgsValidForCall(S
return Callee->getMinRequiredArguments() <= NumArgs;
}
-static ExprResult ActOnCallExprImpl(Sema &S, Scope *Scope, Expr *Fn,
- SourceLocation LParenLoc,
- MultiExprArg ArgExprs,
- SourceLocation RParenLoc, Expr *ExecConfig,
- bool IsExecConfig) {
+/// ActOnCallExpr - Handle a call to Fn with the specified array of arguments.
+/// This provides the location of the left/right parens and a list of comma
+/// locations.
+ExprResult Sema::ActOnCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
+ MultiExprArg ArgExprs, SourceLocation RParenLoc,
+ Expr *ExecConfig, bool IsExecConfig) {
// Since this might be a postfix expression, get rid of ParenListExprs.
- ExprResult Result = S.MaybeConvertParenListExprToParenExpr(Scope, Fn);
+ ExprResult Result = MaybeConvertParenListExprToParenExpr(Scope, Fn);
if (Result.isInvalid()) return ExprError();
Fn = Result.get();
- if (checkArgsForPlaceholders(S, ArgExprs))
+ if (checkArgsForPlaceholders(*this, ArgExprs))
return ExprError();
- if (S.getLangOpts().CPlusPlus) {
+ if (getLangOpts().CPlusPlus) {
// If this is a pseudo-destructor expression, build the call immediately.
if (isa<CXXPseudoDestructorExpr>(Fn)) {
if (!ArgExprs.empty()) {
// Pseudo-destructor calls should not have any arguments.
- S.Diag(Fn->getLocStart(), diag::err_pseudo_dtor_call_with_args)
+ Diag(Fn->getLocStart(), diag::err_pseudo_dtor_call_with_args)
<< FixItHint::CreateRemoval(
SourceRange(ArgExprs.front()->getLocStart(),
ArgExprs.back()->getLocEnd()));
}
- return new (S.Context)
- CallExpr(S.Context, Fn, None, S.Context.VoidTy, VK_RValue, RParenLoc);
+ return new (Context)
+ CallExpr(Context, Fn, None, Context.VoidTy, VK_RValue, RParenLoc);
}
- if (Fn->getType() == S.Context.PseudoObjectTy) {
- ExprResult result = S.CheckPlaceholderExpr(Fn);
+ if (Fn->getType() == Context.PseudoObjectTy) {
+ ExprResult result = CheckPlaceholderExpr(Fn);
if (result.isInvalid()) return ExprError();
Fn = result.get();
}
@@ -5185,35 +5184,34 @@ static ExprResult ActOnCallExprImpl(Sema
if (Dependent) {
if (ExecConfig) {
- return new (S.Context) CUDAKernelCallExpr(
- S.Context, Fn, cast<CallExpr>(ExecConfig), ArgExprs,
- S.Context.DependentTy, VK_RValue, RParenLoc);
+ return new (Context) CUDAKernelCallExpr(
+ Context, Fn, cast<CallExpr>(ExecConfig), ArgExprs,
+ Context.DependentTy, VK_RValue, RParenLoc);
} else {
- return new (S.Context)
- CallExpr(S.Context, Fn, ArgExprs, S.Context.DependentTy, VK_RValue,
- RParenLoc);
+ return new (Context) CallExpr(
+ Context, Fn, ArgExprs, Context.DependentTy, VK_RValue, RParenLoc);
}
}
// Determine whether this is a call to an object (C++ [over.call.object]).
if (Fn->getType()->isRecordType())
- return S.BuildCallToObjectOfClassType(Scope, Fn, LParenLoc, ArgExprs,
- RParenLoc);
+ return BuildCallToObjectOfClassType(Scope, Fn, LParenLoc, ArgExprs,
+ RParenLoc);
- if (Fn->getType() == S.Context.UnknownAnyTy) {
- ExprResult result = rebuildUnknownAnyFunction(S, Fn);
+ if (Fn->getType() == Context.UnknownAnyTy) {
+ ExprResult result = rebuildUnknownAnyFunction(*this, Fn);
if (result.isInvalid()) return ExprError();
Fn = result.get();
}
- if (Fn->getType() == S.Context.BoundMemberTy) {
- return S.BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
- RParenLoc);
+ if (Fn->getType() == Context.BoundMemberTy) {
+ return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
+ RParenLoc);
}
}
// Check for overloaded calls. This can happen even in C due to extensions.
- if (Fn->getType() == S.Context.OverloadTy) {
+ if (Fn->getType() == Context.OverloadTy) {
OverloadExpr::FindResult find = OverloadExpr::find(Fn);
// We aren't supposed to apply this logic for if there'Scope an '&'
@@ -5221,17 +5219,17 @@ static ExprResult ActOnCallExprImpl(Sema
if (!find.HasFormOfMemberPointer) {
OverloadExpr *ovl = find.Expression;
if (UnresolvedLookupExpr *ULE = dyn_cast<UnresolvedLookupExpr>(ovl))
- return S.BuildOverloadedCallExpr(
+ return BuildOverloadedCallExpr(
Scope, Fn, ULE, LParenLoc, ArgExprs, RParenLoc, ExecConfig,
/*AllowTypoCorrection=*/true, find.IsAddressOfOperand);
- return S.BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
- RParenLoc);
+ return BuildCallToMemberFunction(Scope, Fn, LParenLoc, ArgExprs,
+ RParenLoc);
}
}
// If we're directly calling a function, get the appropriate declaration.
- if (Fn->getType() == S.Context.UnknownAnyTy) {
- ExprResult result = rebuildUnknownAnyFunction(S, Fn);
+ if (Fn->getType() == Context.UnknownAnyTy) {
+ ExprResult result = rebuildUnknownAnyFunction(*this, Fn);
if (result.isInvalid()) return ExprError();
Fn = result.get();
}
@@ -5256,10 +5254,10 @@ static ExprResult ActOnCallExprImpl(Sema
// with no explicit address space with the address space of the arguments
// in ArgExprs.
if ((FDecl =
- rewriteBuiltinFunctionDecl(&S, S.Context, FDecl, ArgExprs))) {
+ rewriteBuiltinFunctionDecl(this, Context, FDecl, ArgExprs))) {
NDecl = FDecl;
Fn = DeclRefExpr::Create(
- S.Context, FDecl->getQualifierLoc(), SourceLocation(), FDecl, false,
+ Context, FDecl->getQualifierLoc(), SourceLocation(), FDecl, false,
SourceLocation(), FDecl->getType(), Fn->getValueKind(), FDecl);
}
}
@@ -5268,8 +5266,8 @@ static ExprResult ActOnCallExprImpl(Sema
if (FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(NDecl)) {
if (CallingNDeclIndirectly &&
- !S.checkAddressOfFunctionIsAvailable(FD, /*Complain=*/true,
- Fn->getLocStart()))
+ !checkAddressOfFunctionIsAvailable(FD, /*Complain=*/true,
+ Fn->getLocStart()))
return ExprError();
// CheckEnableIf assumes that the we're passing in a sane number of args for
@@ -5279,42 +5277,22 @@ static ExprResult ActOnCallExprImpl(Sema
// number of args looks incorrect, don't do enable_if checks; we should've
// already emitted an error about the bad call.
if (FD->hasAttr<EnableIfAttr>() &&
- isNumberOfArgsValidForCall(S, FD, ArgExprs.size())) {
- if (const EnableIfAttr *Attr = S.CheckEnableIf(FD, ArgExprs, true)) {
- S.Diag(Fn->getLocStart(),
- isa<CXXMethodDecl>(FD)
- ? diag::err_ovl_no_viable_member_function_in_call
- : diag::err_ovl_no_viable_function_in_call)
+ isNumberOfArgsValidForCall(*this, FD, ArgExprs.size())) {
+ if (const EnableIfAttr *Attr = CheckEnableIf(FD, ArgExprs, true)) {
+ Diag(Fn->getLocStart(),
+ isa<CXXMethodDecl>(FD)
+ ? diag::err_ovl_no_viable_member_function_in_call
+ : diag::err_ovl_no_viable_function_in_call)
<< FD << FD->getSourceRange();
- S.Diag(FD->getLocation(),
- diag::note_ovl_candidate_disabled_by_enable_if_attr)
+ Diag(FD->getLocation(),
+ diag::note_ovl_candidate_disabled_by_enable_if_attr)
<< Attr->getCond()->getSourceRange() << Attr->getMessage();
}
}
}
- return S.BuildResolvedCallExpr(Fn, NDecl, LParenLoc, ArgExprs, RParenLoc,
- ExecConfig, IsExecConfig);
-}
-
-/// ActOnCallExpr - Handle a call to Fn with the specified array of arguments.
-/// This provides the location of the left/right parens and a list of comma
-/// locations.
-ExprResult Sema::ActOnCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc,
- MultiExprArg ArgExprs, SourceLocation RParenLoc,
- Expr *ExecConfig, bool IsExecConfig) {
- ExprResult Ret = ActOnCallExprImpl(*this, S, Fn, LParenLoc, ArgExprs,
- RParenLoc, ExecConfig, IsExecConfig);
-
- // If appropriate, check that this is a valid CUDA call (and emit an error if
- // the call is not allowed).
- if (getLangOpts().CUDA && Ret.isUsable())
- if (auto *Call = dyn_cast<CallExpr>(Ret.get()))
- if (auto *FD = Call->getDirectCallee())
- if (!CheckCUDACall(Call->getLocStart(), FD))
- return ExprError();
-
- return Ret;
+ return BuildResolvedCallExpr(Fn, NDecl, LParenLoc, ArgExprs, RParenLoc,
+ ExecConfig, IsExecConfig);
}
/// ActOnAsTypeExpr - create a new asType (bitcast) from the arguments.
Modified: cfe/trunk/test/SemaCUDA/Inputs/cuda.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/Inputs/cuda.h?rev=283637&r1=283636&r2=283637&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/Inputs/cuda.h (original)
+++ cfe/trunk/test/SemaCUDA/Inputs/cuda.h Fri Oct 7 20:07:11 2016
@@ -22,7 +22,9 @@ typedef struct cudaStream *cudaStream_t;
int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
cudaStream_t stream = 0);
-// Device-side placement new overloads.
+// Host- and device-side placement new overloads.
+void *operator new(__SIZE_TYPE__, void *p) { return p; }
+void *operator new[](__SIZE_TYPE__, void *p) { return p; }
__device__ void *operator new(__SIZE_TYPE__, void *p) { return p; }
__device__ void *operator new[](__SIZE_TYPE__, void *p) { return p; }
Modified: cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu?rev=283637&r1=283636&r2=283637&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu (original)
+++ cfe/trunk/test/SemaCUDA/call-host-fn-from-device.cu Fri Oct 7 20:07:11 2016
@@ -12,6 +12,9 @@ extern "C" void host_fn() {}
// expected-note at -4 {{'host_fn' declared here}}
// expected-note at -5 {{'host_fn' declared here}}
// expected-note at -6 {{'host_fn' declared here}}
+// expected-note at -7 {{'host_fn' declared here}}
+
+struct Dummy {};
struct S {
S() {}
@@ -34,6 +37,15 @@ struct T {
void h() {}
// expected-note at -1 {{'h' declared here}}
+
+ void operator+();
+ // expected-note at -1 {{'operator+' declared here}}
+
+ void operator-(const T&) {}
+ // expected-note at -1 {{'operator-' declared here}}
+
+ operator Dummy() { return Dummy(); }
+ // expected-note at -1 {{'operator Dummy' declared here}}
};
__host__ __device__ void T::hd3() {
@@ -92,3 +104,30 @@ template <typename T>
__host__ __device__ void fn_ptr_template() {
auto* ptr = &host_fn; // Not an error because the template isn't instantiated.
}
+
+__host__ __device__ void unaryOp() {
+ T t;
+ (void) +t; // expected-error {{reference to __host__ function 'operator+' in __host__ __device__ function}}
+}
+
+__host__ __device__ void binaryOp() {
+ T t;
+ (void) (t - t); // expected-error {{reference to __host__ function 'operator-' in __host__ __device__ function}}
+}
+
+__host__ __device__ void implicitConversion() {
+ T t;
+ Dummy d = t; // expected-error {{reference to __host__ function 'operator Dummy' in __host__ __device__ function}}
+}
+
+template <typename T>
+struct TmplStruct {
+ template <typename U> __host__ __device__ void fn() {}
+};
+
+template <>
+template <>
+__host__ __device__ void TmplStruct<int>::fn<int>() { host_fn(); }
+// expected-error at -1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}}
+
+__device__ void double_specialization() { TmplStruct<int>().fn<int>(); }
More information about the cfe-commits
mailing list