[clang] [WIP][clang][SYCL] Check that SYCL kernel parameters are sycl::is_device_copyable (PR #209658)
Ian Li via cfe-commits
cfe-commits at lists.llvm.org
Tue Jul 21 08:11:46 PDT 2026
https://github.com/ianayl updated https://github.com/llvm/llvm-project/pull/209658
>From a2fe09adf0fcc5b6273002ab56b2649a3bc6a11f Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Mon, 13 Jul 2026 16:13:36 -0700
Subject: [PATCH 1/3] Initial work
---
clang/lib/Sema/SemaSYCL.cpp | 72 +++++++++++++++++++
.../sycl-kernel-param-is-device-copyable.cpp | 47 ++++++++++++
2 files changed, 119 insertions(+)
create mode 100644 clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index b942f19761f40..4eb80bd0713e7 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -666,6 +666,78 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
+NamespaceDecl* getSyclNamespace(Sema &SemaRef) {
+ ASTContext &Ctx = SemaRef.getASTContext();
+ IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
+
+ LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, SourceLocation(),
+ Sema::LookupNamespaceName);
+ SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
+
+ if (NamespaceResult.isAmbiguous())
+ return nullptr;
+
+ return NamespaceResult.getAsSingle<NamespaceDecl>();
+}
+
+bool lookupIsDeviceCopyable(Sema &SemaRef, QualType Ty) {
+ NamespaceDecl* SyclNamespace = getSyclNamespace(SemaRef);
+ if (nullptr == SyclNamespace) {
+ // TODO Decide if I throw error or just let it pass
+ // - Throw error: Assumes the SYCL namespace must exist
+ // - Let it pass: Assumes that the SYCL namespace might not necessarily be declared
+ return false;
+ }
+
+ ASTContext &Ctx = SemaRef.getASTContext();
+ IdentifierInfo const &IsDeviceCopyableIdent =
+ Ctx.Idents.get("is_device_copyable");
+
+ LookupResult Result(SemaRef, &IsDeviceCopyableIdent, SourceLocation(),
+ Sema::LookupOrdinaryName);
+ SemaRef.LookupQualifiedName(Result, SyclNamespace);
+
+ if (Result.isAmbiguous())
+ // TODO error or let go? perhaps letgo here?
+ return false;
+
+ ClassTemplateDecl* IsDeviceCopyable =
+ Result.getAsSingle<ClassTemplateDecl>();
+
+ // TODO: reference this from SemaCoroutine:
+
+ // // Form template argument list for coroutine_traits<R, P1, P2, ...> according
+ // // to [dcl.fct.def.coroutine]3
+ // TemplateArgumentListInfo Args(KwLoc, KwLoc);
+ // auto AddArg = [&](QualType T) {
+ // Args.addArgument(TemplateArgumentLoc(
+ // TemplateArgument(T), S.Context.getTrivialTypeSourceInfo(T, KwLoc)));
+ // };
+ // AddArg(FnType->getReturnType());
+ // // If the function is a non-static member function, add the type
+ // // of the implicit object parameter before the formal parameters.
+ // if (auto *MD = dyn_cast<CXXMethodDecl>(FD)) {
+ // if (MD->isImplicitObjectMemberFunction()) {
+ // // [over.match.funcs]4
+ // // For non-static member functions, the type of the implicit object
+ // // parameter is
+ // // -- "lvalue reference to cv X" for functions declared without a
+ // // ref-qualifier or with the & ref-qualifier
+ // // -- "rvalue reference to cv X" for functions declared with the &&
+ // // ref-qualifier
+ // QualType T = MD->getFunctionObjectParameterType();
+ // T = FnType->getRefQualifier() == RQ_RValue
+ // ? S.Context.getRValueReferenceType(T)
+ // : S.Context.getLValueReferenceType(T, /*SpelledAsLValue*/ true);
+ // AddArg(T);
+ // }
+ // }
+ // for (QualType T : FnType->getParamTypes())
+ // AddArg(T);
+
+ return false;
+}
+
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
new file mode 100644
index 0000000000000..f8bc702da67a5
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -0,0 +1,47 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -verify %s
+// RUN: %clang_cc1 -triple spirv64 -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
+#include <type_traits>
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int = 0> struct KN;
+
+// A generic kernel launch function.
+template<typename KNT, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
+namespace sycl {
+
+template <typename T>
+struct is_device_copyable : std::is_trivially_copyable<T> {};
+
+template <typename T>
+inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;
+
+} // namespace sycl
+
+
+class NotTriviallyCopyable {
+public:
+ NotTriviallyCopyable() {};
+ NotTriviallyCopyable(const NotTriviallyCopyable& x);
+};
+static_assert(!std::is_trivially_copyable_v<NotTriviallyCopyable>,
+ "NotTriviallyCopyable should be not std::is_trivially_copyable");
+
+class DeviceCopyable : public NotTriviallyCopyable {};
+template<>
+struct sycl::is_device_copyable<DeviceCopyable> : std::true_type {};
+
+
+// Check that sycl::is_device_copyable is respected
+namespace iscopyable1 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {}
+
+void test() {
+ DeviceCopyable a;
+ kernel_single_task<KN<1>>([=] { (void)a; });
+}
+} // namespace iscopyable1
\ No newline at end of file
>From 6b68f86588087bf4c2080c824c9211a097ce66ee Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 14 Jul 2026 18:18:26 -0700
Subject: [PATCH 2/3] Initial implementation
---
.../clang/Basic/DiagnosticSemaKinds.td | 6 +
clang/lib/Sema/SemaSYCL.cpp | 150 +++++++++++++-----
.../sycl-kernel-param-is-device-copyable.cpp | 22 ++-
.../sycl-kernel-param-restrictions.cpp | 9 +-
4 files changed, 140 insertions(+), 47 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 046bab580013a..df4c5f0453589 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13521,6 +13521,10 @@ def err_sycl_special_type_num_init_method : Error<
"method defined">;
def err_sycl_device_invalid_target : Error<
"%0 is not a supported SYCL device target">;
+def err_sycl_incomplete_type_trait : Error<
+ "SYCL type trait %0 is declared but incomplete">;
+def err_sycl_unexpected_type_trait_val : Error<
+ "expected %1 for SYCL type trait %0, but found %2 instead">;
// SYCL external attribute diagnostics
def err_sycl_external_invalid_linkage : Error<
@@ -13581,6 +13585,8 @@ def note_sycl_kernel_launch_overload_resolution_here : Note<
def err_sycl_entry_point_device_use : Error<
"function %0 cannot be used in device code because it is declared with the"
" %1 attribute">;
+def err_sycl_kernel_param_not_device_copyable : Error<
+ "%0 is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter">;
def warn_cuda_maxclusterrank_sm_90 : Warning<
"maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 4eb80bd0713e7..9e19263d858a5 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -666,11 +666,12 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
-NamespaceDecl* getSyclNamespace(Sema &SemaRef) {
+NamespaceDecl* getSyclNamespace(Sema &SemaRef, SourceLocation Loc) {
+ // TODO too slow; cache this
ASTContext &Ctx = SemaRef.getASTContext();
IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
- LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, SourceLocation(),
+ LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, Loc,
Sema::LookupNamespaceName);
SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
@@ -680,62 +681,105 @@ NamespaceDecl* getSyclNamespace(Sema &SemaRef) {
return NamespaceResult.getAsSingle<NamespaceDecl>();
}
-bool lookupIsDeviceCopyable(Sema &SemaRef, QualType Ty) {
- NamespaceDecl* SyclNamespace = getSyclNamespace(SemaRef);
+bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
+ // No need to lookup anything if trivially copyable already
+ ASTContext &Ctx = SemaRef.getASTContext();
+ if (Ty.isTriviallyCopyableType(Ctx)) {
+ return true;
+ }
+
+ // TODO too slow; cache this
+ NamespaceDecl* SyclNamespace = getSyclNamespace(SemaRef, Loc);
if (nullptr == SyclNamespace) {
// TODO Decide if I throw error or just let it pass
// - Throw error: Assumes the SYCL namespace must exist
// - Let it pass: Assumes that the SYCL namespace might not necessarily be declared
+ llvm::errs() << "debug1\n";
return false;
}
- ASTContext &Ctx = SemaRef.getASTContext();
- IdentifierInfo const &IsDeviceCopyableIdent =
+ // is_device_copyable Identifier
+ IdentifierInfo const &IDCIdent =
Ctx.Idents.get("is_device_copyable");
- LookupResult Result(SemaRef, &IsDeviceCopyableIdent, SourceLocation(),
+ LookupResult IdentResult(SemaRef, &IDCIdent, Loc,
Sema::LookupOrdinaryName);
- SemaRef.LookupQualifiedName(Result, SyclNamespace);
+ SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
- if (Result.isAmbiguous())
- // TODO error or let go? perhaps letgo here?
+ if (IdentResult.isAmbiguous()) {
+ // TODO error or let go?
+ llvm::errs() << "debug2\n";
+ return false;
+ }
+
+ ClassTemplateDecl* IDCDecl =
+ IdentResult.getAsSingle<ClassTemplateDecl>();
+ if (nullptr == IDCDecl) {
+ // TODO error or let go?
+ llvm::errs() << "debug3\n";
return false;
+ }
+
+ TemplateArgumentListInfo Args{};
+ TemplateArgument TyArg{Ty};
+ Args.addArgument(SemaRef.getTrivialTemplateArgumentLoc(TyArg, QualType{}, Loc));
+
+ QualType IDCTrait = SemaRef.CheckTemplateIdType(
+ ElaboratedTypeKeyword::None, TemplateName{IDCDecl}, Loc, Args,
+ /*Scope=*/nullptr, /*ForNestedNameSpecifier=*/false);
+
+ if (IDCTrait.isNull()) {
+ // TODO error or let go?
+ llvm::errs() << "debug4\n";
+ return false;
+ }
+ if (SemaRef.RequireCompleteType(Loc, IDCTrait, diag::err_sycl_incomplete_type_trait)) {
+ // TODO error or let go?
+ llvm::errs() << "debug5\n";
+ return false;
+ }
- ClassTemplateDecl* IsDeviceCopyable =
- Result.getAsSingle<ClassTemplateDecl>();
-
- // TODO: reference this from SemaCoroutine:
-
- // // Form template argument list for coroutine_traits<R, P1, P2, ...> according
- // // to [dcl.fct.def.coroutine]3
- // TemplateArgumentListInfo Args(KwLoc, KwLoc);
- // auto AddArg = [&](QualType T) {
- // Args.addArgument(TemplateArgumentLoc(
- // TemplateArgument(T), S.Context.getTrivialTypeSourceInfo(T, KwLoc)));
- // };
- // AddArg(FnType->getReturnType());
- // // If the function is a non-static member function, add the type
- // // of the implicit object parameter before the formal parameters.
- // if (auto *MD = dyn_cast<CXXMethodDecl>(FD)) {
- // if (MD->isImplicitObjectMemberFunction()) {
- // // [over.match.funcs]4
- // // For non-static member functions, the type of the implicit object
- // // parameter is
- // // -- "lvalue reference to cv X" for functions declared without a
- // // ref-qualifier or with the & ref-qualifier
- // // -- "rvalue reference to cv X" for functions declared with the &&
- // // ref-qualifier
- // QualType T = MD->getFunctionObjectParameterType();
- // T = FnType->getRefQualifier() == RQ_RValue
- // ? S.Context.getRValueReferenceType(T)
- // : S.Context.getLValueReferenceType(T, /*SpelledAsLValue*/ true);
- // AddArg(T);
- // }
- // }
- // for (QualType T : FnType->getParamTypes())
- // AddArg(T);
+ CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
+ assert(RD && "specialization of class template is not a class?");
+
+ // Look up the ::promise_type member.
+ IdentifierInfo const &ValueIdent = Ctx.Idents.get("value");
+ LookupResult ValueResult(SemaRef, &ValueIdent, Loc, Sema::LookupOrdinaryName);
+ SemaRef.LookupQualifiedName(ValueResult, RD);
+ if (ValueResult.empty() || ValueResult.isAmbiguous()) {
+ // TODO error or let go?
+ llvm::errs() << "debug6\n";
+ return false;
+ }
+
+ ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
+ if (ValueExpr.isInvalid()) {
+ // TODO error or let go?
+ llvm::errs() << "debug7\n";
+ return false;
+ }
- return false;
+ struct ICEDiagnoser : Sema::VerifyICEDiagnoser {
+ QualType &TraitTy;
+ Expr *GotExpr;
+ ICEDiagnoser(QualType &TT, Expr *E)
+ : TraitTy(TT), GotExpr(E) {}
+ Sema::SemaDiagnosticBuilder diagnoseNotICE(Sema &S,
+ SourceLocation Loc) override {
+ return S.Diag(Loc, diag::err_sycl_unexpected_type_trait_val)
+ << TraitTy << "std::true_type or std::false_type" << GotExpr;
+ }
+ } Diagnoser(IDCTrait, ValueExpr.get());
+
+ llvm::APSInt IDCValue;
+ ValueExpr = SemaRef.VerifyIntegerConstantExpression(ValueExpr.get(), &IDCValue, Diagnoser);
+ if (ValueExpr.isInvalid()){
+ // TODO error or let go?
+ llvm::errs() << "debug8\n";
+ return false;
+ }
+
+ return IDCValue.getBoolValue();
}
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
@@ -823,6 +867,24 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
IsValid = false;
return false;
}
+
+ auto DirectParent = ObjectAccessPath.back();
+ // TODO Do I care about deep traversal + checking if every subfield within a class is conformant?
+ // TODO Do I at least need to dive into the lambdas
+ if (const ParmVarDecl *parmVar = dyn_cast<const ParmVarDecl *>(DirectParent)) {
+ const CXXRecordDecl *RD = Ty.getNonReferenceType()->getAsCXXRecordDecl();
+ if (RD && !RD->isLambda() && (RD->isClass() || RD->isStruct())) {
+ if (!lookupIsDeviceCopyable(SemaSYCLRef.SemaRef, Ty, parmVar->getLocation())) {
+ SemaSYCLRef.Diag(parmVar->getLocation(), diag::err_sycl_kernel_param_not_device_copyable) << Ty;
+ emitObjectAccessPathNotes();
+
+ IsValid = false;
+ return false;
+ }
+ // TODO Issue warning if type is obviously not copyable
+ // ... but what does that mean? And how thorough do I want to check, even if the user has already marked it copyable?
+ }
+ }
return true;
}
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
index f8bc702da67a5..786561c18013e 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -1,6 +1,22 @@
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -verify %s
// RUN: %clang_cc1 -triple spirv64 -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
-#include <type_traits>
+
+namespace std {
+
+template <bool B>
+struct bool_constant {
+ static constexpr bool value = B;
+};
+
+using true_type = bool_constant<true>;
+
+template <typename T>
+struct is_trivially_copyable : bool_constant<__is_trivially_copyable(T)> {};
+
+template <typename T>
+inline constexpr bool is_trivially_copyable_v = is_trivially_copyable<T>::value;
+
+} // namespace std
// A unique kernel name type is required for each declared kernel entry point.
template<int, int = 0> struct KN;
@@ -42,6 +58,8 @@ void kernel_single_task(T t) {}
void test() {
DeviceCopyable a;
- kernel_single_task<KN<1>>([=] { (void)a; });
+ NotTriviallyCopyable b;
+ kernel_single_task<KN<1>>(a);
+ kernel_single_task<KN<2>>(b);
}
} // namespace iscopyable1
\ No newline at end of file
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 66aa00da18a04..429aefde2a016 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -8,6 +8,14 @@ template<int, int = 0> struct KN;
template<typename KNT, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+namespace sycl {
+
+// Make everything is_device_copyable for sake of testing
+template <typename T>
+struct is_device_copyable { static constexpr bool value = true; };
+
+} // namespace sycl
+
// Check that reference captures of kernel that defined as lambda are diagnosed.
namespace badref1 {
// Kernel entry point template definition.
@@ -166,7 +174,6 @@ class Base { // expected-note {{within field of type 'Base' declared here}}}
class Derived : virtual Base { // expected-note {{within base class of type 'Base' declared here}}
public:
Derived(int &a) : Base(a) {}
-
};
void test() {
>From 304b00b20fea91542802feb9f5830fff0c4ac11a Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 14 Jul 2026 18:26:43 -0700
Subject: [PATCH 3/3] clang-format
---
clang/lib/Sema/SemaSYCL.cpp | 51 +++++++++++++++++++++----------------
1 file changed, 29 insertions(+), 22 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 9e19263d858a5..54895983d76da 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -666,7 +666,7 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
-NamespaceDecl* getSyclNamespace(Sema &SemaRef, SourceLocation Loc) {
+NamespaceDecl *getSyclNamespace(Sema &SemaRef, SourceLocation Loc) {
// TODO too slow; cache this
ASTContext &Ctx = SemaRef.getASTContext();
IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
@@ -689,21 +689,20 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
}
// TODO too slow; cache this
- NamespaceDecl* SyclNamespace = getSyclNamespace(SemaRef, Loc);
+ NamespaceDecl *SyclNamespace = getSyclNamespace(SemaRef, Loc);
if (nullptr == SyclNamespace) {
// TODO Decide if I throw error or just let it pass
// - Throw error: Assumes the SYCL namespace must exist
- // - Let it pass: Assumes that the SYCL namespace might not necessarily be declared
+ // - Let it pass: Assumes that the SYCL namespace might not necessarily be
+ // declared
llvm::errs() << "debug1\n";
return false;
}
// is_device_copyable Identifier
- IdentifierInfo const &IDCIdent =
- Ctx.Idents.get("is_device_copyable");
+ IdentifierInfo const &IDCIdent = Ctx.Idents.get("is_device_copyable");
- LookupResult IdentResult(SemaRef, &IDCIdent, Loc,
- Sema::LookupOrdinaryName);
+ LookupResult IdentResult(SemaRef, &IDCIdent, Loc, Sema::LookupOrdinaryName);
SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
if (IdentResult.isAmbiguous()) {
@@ -712,8 +711,7 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
return false;
}
- ClassTemplateDecl* IDCDecl =
- IdentResult.getAsSingle<ClassTemplateDecl>();
+ ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
if (nullptr == IDCDecl) {
// TODO error or let go?
llvm::errs() << "debug3\n";
@@ -722,7 +720,8 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
TemplateArgumentListInfo Args{};
TemplateArgument TyArg{Ty};
- Args.addArgument(SemaRef.getTrivialTemplateArgumentLoc(TyArg, QualType{}, Loc));
+ Args.addArgument(
+ SemaRef.getTrivialTemplateArgumentLoc(TyArg, QualType{}, Loc));
QualType IDCTrait = SemaRef.CheckTemplateIdType(
ElaboratedTypeKeyword::None, TemplateName{IDCDecl}, Loc, Args,
@@ -733,7 +732,8 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
llvm::errs() << "debug4\n";
return false;
}
- if (SemaRef.RequireCompleteType(Loc, IDCTrait, diag::err_sycl_incomplete_type_trait)) {
+ if (SemaRef.RequireCompleteType(Loc, IDCTrait,
+ diag::err_sycl_incomplete_type_trait)) {
// TODO error or let go?
llvm::errs() << "debug5\n";
return false;
@@ -752,18 +752,18 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
return false;
}
- ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
+ ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
+ CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
if (ValueExpr.isInvalid()) {
// TODO error or let go?
llvm::errs() << "debug7\n";
return false;
}
-
+
struct ICEDiagnoser : Sema::VerifyICEDiagnoser {
QualType &TraitTy;
Expr *GotExpr;
- ICEDiagnoser(QualType &TT, Expr *E)
- : TraitTy(TT), GotExpr(E) {}
+ ICEDiagnoser(QualType &TT, Expr *E) : TraitTy(TT), GotExpr(E) {}
Sema::SemaDiagnosticBuilder diagnoseNotICE(Sema &S,
SourceLocation Loc) override {
return S.Diag(Loc, diag::err_sycl_unexpected_type_trait_val)
@@ -772,8 +772,9 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
} Diagnoser(IDCTrait, ValueExpr.get());
llvm::APSInt IDCValue;
- ValueExpr = SemaRef.VerifyIntegerConstantExpression(ValueExpr.get(), &IDCValue, Diagnoser);
- if (ValueExpr.isInvalid()){
+ ValueExpr = SemaRef.VerifyIntegerConstantExpression(ValueExpr.get(),
+ &IDCValue, Diagnoser);
+ if (ValueExpr.isInvalid()) {
// TODO error or let go?
llvm::errs() << "debug8\n";
return false;
@@ -869,20 +870,26 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
}
auto DirectParent = ObjectAccessPath.back();
- // TODO Do I care about deep traversal + checking if every subfield within a class is conformant?
+ // TODO Do I care about deep traversal + checking if every subfield within a
+ // class is conformant?
// TODO Do I at least need to dive into the lambdas
- if (const ParmVarDecl *parmVar = dyn_cast<const ParmVarDecl *>(DirectParent)) {
+ if (const ParmVarDecl *parmVar =
+ dyn_cast<const ParmVarDecl *>(DirectParent)) {
const CXXRecordDecl *RD = Ty.getNonReferenceType()->getAsCXXRecordDecl();
if (RD && !RD->isLambda() && (RD->isClass() || RD->isStruct())) {
- if (!lookupIsDeviceCopyable(SemaSYCLRef.SemaRef, Ty, parmVar->getLocation())) {
- SemaSYCLRef.Diag(parmVar->getLocation(), diag::err_sycl_kernel_param_not_device_copyable) << Ty;
+ if (!lookupIsDeviceCopyable(SemaSYCLRef.SemaRef, Ty,
+ parmVar->getLocation())) {
+ SemaSYCLRef.Diag(parmVar->getLocation(),
+ diag::err_sycl_kernel_param_not_device_copyable)
+ << Ty;
emitObjectAccessPathNotes();
IsValid = false;
return false;
}
// TODO Issue warning if type is obviously not copyable
- // ... but what does that mean? And how thorough do I want to check, even if the user has already marked it copyable?
+ // ... but what does that mean? And how thorough do I want to check,
+ // even if the user has already marked it copyable?
}
}
return true;
More information about the cfe-commits
mailing list