yaxunl updated this revision to Diff 180848.
yaxunl added a comment.
Add test for `__host__ __device__`.
Removing the flag IsParsingTemplateArgument in Sema. Instead, check
ExprEvalContexts
for disabling checking device/host consistency.
I did not use ExprEvalContext Unevaluated to condition the check because
the issue happens with ExprEvalContext ConstantEvaluated. Also we do not want to
totally remove the check, we just want to defer the check until the arg
evaluation is done.
When the deferred check is performed, ExprEvalContext is still in
ConstantEvaluated but
its kind is no longer EK_TemplateArgument, therefore we can use the expr kind
to condition
the check.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D56411/new/
https://reviews.llvm.org/D56411
Files:
include/clang/Sema/Sema.h
lib/Sema/SemaCUDA.cpp
lib/Sema/SemaExpr.cpp
lib/Sema/SemaTemplate.cpp
Index: lib/Sema/SemaTemplate.cpp
===================================================================
--- lib/Sema/SemaTemplate.cpp
+++ lib/Sema/SemaTemplate.cpp
@@ -4784,8 +4784,8 @@
TemplateArgument Result;
unsigned CurSFINAEErrors = NumSFINAEErrors;
ExprResult Res =
- CheckTemplateArgument(NTTP, NTTPType, Arg.getArgument().getAsExpr(),
- Result, CTAK);
+ CheckTemplateArgument(NTTP, NTTPType, Arg.getArgument().getAsExpr(),
+ Result, CTAK, dyn_cast<TemplateDecl>(Template));
if (Res.isInvalid())
return true;
// If the current template argument causes an error, give up now.
@@ -6154,6 +6154,27 @@
return true;
}
+namespace {
+bool CheckCUDATemplateArgument(Sema &S, Expr *Arg, TemplateDecl *Template) {
+ if (Template) {
+ Expr *E = Arg;
+ if (UnaryOperator *UO = dyn_cast<UnaryOperator>(E)) {
+ E = UO ? UO->getSubExpr() : nullptr;
+ }
+ if (DeclRefExpr *DRE = dyn_cast_or_null<DeclRefExpr>(E)) {
+ ValueDecl *Entity = DRE ? DRE->getDecl() : nullptr;
+ if (Entity) {
+ if (auto Callee = dyn_cast<FunctionDecl>(Entity))
+ if (auto Caller =
+ dyn_cast<FunctionDecl>(Template->getTemplatedDecl()))
+ if (!S.CheckCUDACall(Arg->getBeginLoc(), Callee, Caller))
+ return false;
+ }
+ }
+ }
+ return true;
+}
+} // namespace
/// Check a template argument against its corresponding
/// non-type template parameter.
///
@@ -6164,7 +6185,8 @@
ExprResult Sema::CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
QualType ParamType, Expr *Arg,
TemplateArgument &Converted,
- CheckTemplateArgumentKind CTAK) {
+ CheckTemplateArgumentKind CTAK,
+ TemplateDecl *Template) {
SourceLocation StartLoc = Arg->getBeginLoc();
// If the parameter type somehow involves auto, deduce the type now.
@@ -6561,7 +6583,11 @@
if (FunctionDecl *Fn = ResolveAddressOfOverloadedFunction(Arg, ParamType,
true,
FoundResult)) {
- if (DiagnoseUseOfDecl(Fn, Arg->getBeginLoc()))
+ if (DiagnoseUseOfDecl(Fn, Arg->getBeginLoc(),
+ /*UnknownObjCClass=*/nullptr,
+ /*ObjCPropertyAccess=*/false,
+ /*AvoidPartialAvailabilityChecks=*/false,
+ /*ClassReciever=*/nullptr, Template))
return ExprError();
Arg = FixOverloadedFunctionReference(Arg, FoundResult, Fn);
@@ -6570,6 +6596,9 @@
return ExprError();
}
+ if (!CheckCUDATemplateArgument(*this, Arg, Template))
+ return ExprError();
+
if (!ParamType->isMemberPointerType()) {
if (CheckTemplateArgumentAddressOfObjectOrFunction(*this, Param,
ParamType,
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -213,7 +213,8 @@
const ObjCInterfaceDecl *UnknownObjCClass,
bool ObjCPropertyAccess,
bool AvoidPartialAvailabilityChecks,
- ObjCInterfaceDecl *ClassReceiver) {
+ ObjCInterfaceDecl *ClassReceiver,
+ TemplateDecl *Template) {
SourceLocation Loc = Locs.front();
if (getLangOpts().CPlusPlus && isa<FunctionDecl>(D)) {
// If there were any diagnostics suppressed by template argument deduction,
@@ -270,7 +271,11 @@
DeduceReturnType(FD, Loc))
return true;
- if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
+ if (getLangOpts().CUDA &&
+ !CheckCUDACall(
+ Loc, FD,
+ Template ? dyn_cast<FunctionDecl>(Template->getTemplatedDecl())
+ : nullptr))
return true;
}
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -833,12 +833,19 @@
}
}
-bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
+bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee,
+ FunctionDecl *Caller) {
assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
assert(Callee && "Callee may not be null.");
+
+ if (ExprEvalContexts.back().ExprContext ==
+ ExpressionEvaluationContextRecord::ExpressionKind::EK_TemplateArgument)
+ return true;
+
// FIXME: Is bailing out early correct here? Should we instead assume that
// the caller is a global initializer?
- FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
+ if (!Caller)
+ Caller = dyn_cast<FunctionDecl>(CurContext);
if (!Caller)
return true;
Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -4010,11 +4010,15 @@
// Expression Parsing Callbacks: SemaExpr.cpp.
bool CanUseDecl(NamedDecl *D, bool TreatUnavailableAsInvalid);
+
+ // \param TemplateReceiver is the receiving instantiated template declaration
+ // if it is not a null pointer.
bool DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
const ObjCInterfaceDecl *UnknownObjCClass = nullptr,
bool ObjCPropertyAccess = false,
bool AvoidPartialAvailabilityChecks = false,
- ObjCInterfaceDecl *ClassReciever = nullptr);
+ ObjCInterfaceDecl *ClassReciever = nullptr,
+ TemplateDecl *TemplateReceiver = nullptr);
void NoteDeletedFunction(FunctionDecl *FD);
void NoteDeletedInheritingConstructor(CXXConstructorDecl *CD);
std::string getDeletedOrUnavailableSuffix(const FunctionDecl *FD);
@@ -6453,10 +6457,12 @@
bool CheckTemplateArgument(TemplateTypeParmDecl *Param,
TypeSourceInfo *Arg);
- ExprResult CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
- QualType InstantiatedParamType, Expr *Arg,
- TemplateArgument &Converted,
- CheckTemplateArgumentKind CTAK = CTAK_Specified);
+ ExprResult
+ CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
+ QualType InstantiatedParamType, Expr *Arg,
+ TemplateArgument &Converted,
+ CheckTemplateArgumentKind CTAK = CTAK_Specified,
+ TemplateDecl *Template = nullptr);
bool CheckTemplateTemplateArgument(TemplateParameterList *Params,
TemplateArgumentLoc &Arg);
@@ -10224,7 +10230,8 @@
/// deferred errors.
///
/// - Otherwise, returns true without emitting any diagnostics.
- bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
+ bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee,
+ FunctionDecl *Caller = nullptr);
/// Set __device__ or __host__ __device__ attributes on the given lambda
/// operator() method.
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
