[clang] [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
Fri Aug 21 09:01:43 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 01/15] 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 02/15] 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 03/15] 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;
>From 5f32870b9615aea37b6863720afe7d11241300c7 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 4 Aug 2026 12:07:24 -0700
Subject: [PATCH 04/15] WIP
---
clang/lib/Sema/SemaSYCL.cpp | 108 +++++++++++++-----
.../sycl-kernel-param-is-device-copyable.cpp | 11 +-
2 files changed, 90 insertions(+), 29 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 54895983d76da..192221de2a19d 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -706,14 +706,14 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
if (IdentResult.isAmbiguous()) {
- // TODO error or let go?
+ // TODO warn or error
llvm::errs() << "debug2\n";
return false;
}
ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
if (nullptr == IDCDecl) {
- // TODO error or let go?
+ // TODO simply let go; it's undefined
llvm::errs() << "debug3\n";
return false;
}
@@ -728,13 +728,13 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
/*Scope=*/nullptr, /*ForNestedNameSpecifier=*/false);
if (IDCTrait.isNull()) {
- // TODO error or let go?
+ // TODO simply let go; it's undefined
llvm::errs() << "debug4\n";
return false;
}
if (SemaRef.RequireCompleteType(Loc, IDCTrait,
diag::err_sycl_incomplete_type_trait)) {
- // TODO error or let go?
+ // TODO do I need this if?
llvm::errs() << "debug5\n";
return false;
}
@@ -742,12 +742,13 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
assert(RD && "specialization of class template is not a class?");
- // Look up the ::promise_type member.
+ // Look up the ::value 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?
+ // TODO should I error or let go?
+ // definitely error on ambiguous, but what about empty?
llvm::errs() << "debug6\n";
return false;
}
@@ -755,7 +756,7 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
if (ValueExpr.isInvalid()) {
- // TODO error or let go?
+ // TODO compiler error: this shouldn't happen?
llvm::errs() << "debug7\n";
return false;
}
@@ -775,7 +776,7 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
ValueExpr = SemaRef.VerifyIntegerConstantExpression(ValueExpr.get(),
&IDCValue, Diagnoser);
if (ValueExpr.isInvalid()) {
- // TODO error or let go?
+ // TODO compiler error: this shouldn't happen?
llvm::errs() << "debug8\n";
return false;
}
@@ -791,6 +792,22 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
const FieldDecl *>;
SmallVector<ObjectAccess, 4> ObjectAccessPath;
+ struct DiagDetails {
+ QualType Type;
+ SourceLocation Loc;
+ };
+
+ // Return diagnostics info for an 'ObjectAccess' stored on ObjectAccessPath
+ DiagDetails getObjectAccessDiagDetails(ObjectAccess o) {
+ if (auto *PVD = dyn_cast<const ParmVarDecl *>(o))
+ return {PVD->getType(), PVD->getLocation()};
+ if (auto *FD = dyn_cast<const FieldDecl *>(o))
+ return {FD->getType(), FD->getLocation()};
+ if (auto *BS = dyn_cast<const CXXBaseSpecifier *>(o))
+ return {BS->getType(), BS->getBaseTypeLoc()};
+ llvm_unreachable("Unexpected type in ObjectAccess");
+ }
+
void emitObjectAccessPathNotes() {
for (auto Parent : llvm::reverse(ObjectAccessPath)) {
if (auto *FD = Parent.dyn_cast<const FieldDecl *>()) {
@@ -869,29 +886,64 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
return false;
}
- auto DirectParent = ObjectAccessPath.back();
+ DiagDetails Detail = getObjectAccessDiagDetails(ObjectAccessPath.back());
+
+ if (!lookupIsDeviceCopyable(SemaSYCLRef.SemaRef, Ty, Detail.Loc)) {
+ SemaSYCLRef.Diag(Detail.Loc,
+ diag::err_sycl_kernel_param_not_device_copyable)
+ << Ty;
+ emitObjectAccessPathNotes();
+
+ IsValid = false;
+ return false;
+ }
+ // TODO the issue with this logic:
+ // - if a class is marked as copyable, I should STOP descending into the class's subfields
+ // - we STOP because the user's declaration that a class is copyable should overwrite
+ // the results of future traversals; we shouldn't descend any further
+ // - although, because we want to catch classes that are obviously not copyable, we'll
+ // need to do a shallow traversal in the future, checking that there aren't data members
+ // in the class that are obviously not copyable: this will need to be its own function
+ // - the stop condition should be if there are stuff that obviously breaks SYCL spec
+ // for "is device copyable"
+ // - if a class is not marked as copyable, _then_ descend into the class's subfields
+ // - this is current behavior
+
+
+ // TODO Issue warning if type is obviously not copyable
+ // ... you can do this by making some sort of dict for memoizing whether or
+ // not a field/known type is not copyable
+
// 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?
- }
- }
+
+ // TODO I am not sure that parmVar guarantees base argument; we might be
+ // better off checking the objectaccess vector instead, althoug there has
+ // to be a better way to do this
+
+ //// ???: I don't think I need this check, it should be up to the logic of
+ //// this function to stop traversing
+ // if (const ParmVarDecl *parmVar =
+ // dyn_cast<const ParmVarDecl *>(DirectParent)) {
+ // // TODO don't need this check, we can check after it failed I guess
+ // 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 786561c18013e..2de2f479c7de3 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -48,6 +48,11 @@ class DeviceCopyable : public NotTriviallyCopyable {};
template<>
struct sycl::is_device_copyable<DeviceCopyable> : std::true_type {};
+struct DefinitelyCopyable {
+ int foo = 0;
+};
+static_assert(std::is_trivially_copyable_v<DefinitelyCopyable>,
+ "DefinitelyCopyable should be trivially copyable");
// Check that sycl::is_device_copyable is respected
namespace iscopyable1 {
@@ -58,8 +63,12 @@ void kernel_single_task(T t) {}
void test() {
DeviceCopyable a;
- NotTriviallyCopyable b;
kernel_single_task<KN<1>>(a);
+
+ NotTriviallyCopyable b;
kernel_single_task<KN<2>>(b);
+
+ DefinitelyCopyable c;
+ kernel_single_task<KN<3>>(c);
}
} // namespace iscopyable1
\ No newline at end of file
>From 343eb71121e84e99e925a7697b85f52c5348b42d Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 14 Aug 2026 09:21:25 -0700
Subject: [PATCH 05/15] rethink how device copyable checks worked
---
clang/lib/Sema/SemaSYCL.cpp | 115 ++++++++++--------
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 17 +++
.../sycl-kernel-param-is-device-copyable.cpp | 13 +-
3 files changed, 91 insertions(+), 54 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 192221de2a19d..a4a9f3e6bec03 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -681,16 +681,20 @@ NamespaceDecl *getSyclNamespace(Sema &SemaRef, SourceLocation Loc) {
return NamespaceResult.getAsSingle<NamespaceDecl>();
}
-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;
- }
+struct DeviceCopyableResult {
+ bool copyable;
+ // PD contains a(n):
+ // - error, if queried type is not copyable.
+ // - warning, if queried type is marked copyable but detected to be non-device
+ // copyable.
+ std::optional<PartialDiagnostic> PD;
+};
+bool isMarkedDeviceCopyable(Sema &SemaRef, const QualType &Ty, SourceLocation Loc) {
+ ASTContext &Ctx = SemaRef.getASTContext();
// TODO too slow; cache this
NamespaceDecl *SyclNamespace = getSyclNamespace(SemaRef, Loc);
- if (nullptr == SyclNamespace) {
+ if (!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
@@ -784,6 +788,35 @@ bool lookupIsDeviceCopyable(Sema &SemaRef, QualType &Ty, SourceLocation Loc) {
return IDCValue.getBoolValue();
}
+// TODO should these results be cached?
+DeviceCopyableResult isDeviceCopyable(Sema &SemaRef, const QualType &Ty, SourceLocation Loc) {
+ // No need to lookup anything if trivially copyable already
+ ASTContext &Ctx = SemaRef.getASTContext();
+ if (Ty.isTriviallyCopyableType(Ctx))
+ return {true, /*PD=*/std::nullopt};
+
+ bool markedCopyable = isMarkedDeviceCopyable(SemaRef, Ty, Loc);
+ if (const CXXRecordDecl *RD = Ty->getAsCXXRecordDecl()) {
+ // Set all lambdas as copyable: Future traversal passes will determine
+ // whether or not the parameters/capture of the lambda are actually
+ // copyable.
+ if (RD->isLambda())
+ return {true, /*PD=*/std::nullopt};
+ // TODO confirm:
+ // - does RD have at least one eligible copy constructor, move constructor, copy assignment operator, or move assignment operator?
+ // - for each of aforementioned, ensure it is public
+ // - confirm each does a bitwise copy (perhaps not possible, up to the user to enforce)
+ // - confirm it has a non deleted destructor
+ // - does the destructor have "no effect"? (perhaps not possible, up to user to enforce)
+ }
+ // TODO subsequent base classes shouldn't be checked for not device copyable
+
+ if (!markedCopyable)
+ return { false, PartialDiagnostic(diag::err_sycl_kernel_param_not_device_copyable, SemaRef.Context.getDiagAllocator()) << Ty };
+ return {true, /*PD=*/std::nullopt};
+}
+
+
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
@@ -841,7 +874,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
void checkParameter(const ParmVarDecl *PVD) {
ObjectAccessPath.push_back(PVD);
// Check the immediate type of the parameter.
- if (checkType(PVD->getType())) {
+ if (checkType(PVD->getType()) && checkDeviceCopyable(PVD->getType())) {
// If type checking wasn't short circuited, visit subobjects to check
// them.
visit(PVD->getType());
@@ -857,7 +890,8 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
bool visitFieldDeclPre(const FieldDecl *FD) {
ObjectAccessPath.push_back(FD);
- return checkType(FD->getType());
+ return checkType(FD->getType()) && checkDeviceCopyable(FD->getType());
+ // TODO: do we need to check if a class's fields are deviceCopyable?
}
// Returns true if subobjects should be visited and false otherwise.
@@ -886,18 +920,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
return false;
}
- DiagDetails Detail = getObjectAccessDiagDetails(ObjectAccessPath.back());
-
- if (!lookupIsDeviceCopyable(SemaSYCLRef.SemaRef, Ty, Detail.Loc)) {
- SemaSYCLRef.Diag(Detail.Loc,
- diag::err_sycl_kernel_param_not_device_copyable)
- << Ty;
- emitObjectAccessPathNotes();
-
- IsValid = false;
- return false;
- }
- // TODO the issue with this logic:
+ // TODO
// - if a class is marked as copyable, I should STOP descending into the class's subfields
// - we STOP because the user's declaration that a class is copyable should overwrite
// the results of future traversals; we shouldn't descend any further
@@ -909,7 +932,6 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// - if a class is not marked as copyable, _then_ descend into the class's subfields
// - this is current behavior
-
// TODO Issue warning if type is obviously not copyable
// ... you can do this by making some sort of dict for memoizing whether or
// not a field/known type is not copyable
@@ -917,33 +939,30 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// 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
+ return true;
+ }
- // TODO I am not sure that parmVar guarantees base argument; we might be
- // better off checking the objectaccess vector instead, althoug there has
- // to be a better way to do this
-
- //// ???: I don't think I need this check, it should be up to the logic of
- //// this function to stop traversing
- // if (const ParmVarDecl *parmVar =
- // dyn_cast<const ParmVarDecl *>(DirectParent)) {
- // // TODO don't need this check, we can check after it failed I guess
- // 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?
- // }
- // }
+ bool checkDeviceCopyable(QualType Ty) {
+ auto DirectParent = ObjectAccessPath.back();
+ QualType Type = Ty;
+
+ DiagDetails Detail = getObjectAccessDiagDetails(DirectParent);
+ // Since references are allowed as direct kernel parameters, we need to
+ // explicitly check the referenced type:
+ if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent)) {
+ Type = Ty->getPointeeType();
+ }
+ const DeviceCopyableResult DCR = isDeviceCopyable(SemaSYCLRef.SemaRef, Type, Detail.Loc);
+ if (DCR.PD) {
+ // Emit diagnostics if any were generated.
+ SemaSYCLRef.Diag(Detail.Loc, DCR.PD.value());
+ emitObjectAccessPathNotes();
+ }
+ if (!DCR.copyable) {
+ assert(DCR.PD && "DeviceCopyableResult must emit an explanatory diagnostic if Ty is not device copyable");
+ IsValid = false;
+ return false;
+ }
return true;
}
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index cc751cc683b59..006e5fade82a9 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -68,6 +68,23 @@ struct copyable {
~copyable();
};
+namespace sycl {
+
+template <bool B>
+struct bool_constant {
+ static constexpr bool value = B;
+};
+
+template <typename T>
+struct is_device_copyable : bool_constant<false> {};
+
+// define copyable as is_device_copyable since copyable isn't actually
+// is_trivially_copyable due to custom destructor:
+template<>
+struct is_device_copyable<copyable> : bool_constant<true> {};
+
+} // namespace sycl
+
namespace std {
template<typename T> constexpr T &&move(T &val) { return static_cast<T&&>(val); }
template<class T>
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 2de2f479c7de3..4ef716ec78a37 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -59,16 +59,17 @@ 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 kernel_single_task(T t) {} // expected-error {{'NotTriviallyCopyable' is not device copyable}} \
+ // expected-note {{within parameter 't' of type 'NotTriviallyCopyable' declared here}}
void test() {
- DeviceCopyable a;
- kernel_single_task<KN<1>>(a);
+ DefinitelyCopyable a;
+ kernel_single_task<KN<3>>(a);
NotTriviallyCopyable b;
- kernel_single_task<KN<2>>(b);
+ kernel_single_task<KN<2>>(b); // expected-note {{in instantiation of function template specialization 'iscopyable1::kernel_single_task<KN<2>, NotTriviallyCopyable>' requested here}}
- DefinitelyCopyable c;
- kernel_single_task<KN<3>>(c);
+ DeviceCopyable c;
+ kernel_single_task<KN<1>>(c);
}
} // namespace iscopyable1
\ No newline at end of file
>From 1c70f5690634c94a668544e85150ce4a6e173804 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 14 Aug 2026 09:26:18 -0700
Subject: [PATCH 06/15] clang-format
---
clang/lib/Sema/SemaSYCL.cpp | 53 ++++++++++++++++++++++++-------------
1 file changed, 34 insertions(+), 19 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index a4a9f3e6bec03..9a28c60a745f5 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -690,7 +690,8 @@ struct DeviceCopyableResult {
std::optional<PartialDiagnostic> PD;
};
-bool isMarkedDeviceCopyable(Sema &SemaRef, const QualType &Ty, SourceLocation Loc) {
+bool isMarkedDeviceCopyable(Sema &SemaRef, const QualType &Ty,
+ SourceLocation Loc) {
ASTContext &Ctx = SemaRef.getASTContext();
// TODO too slow; cache this
NamespaceDecl *SyclNamespace = getSyclNamespace(SemaRef, Loc);
@@ -789,34 +790,40 @@ bool isMarkedDeviceCopyable(Sema &SemaRef, const QualType &Ty, SourceLocation Lo
}
// TODO should these results be cached?
-DeviceCopyableResult isDeviceCopyable(Sema &SemaRef, const QualType &Ty, SourceLocation Loc) {
+DeviceCopyableResult isDeviceCopyable(Sema &SemaRef, const QualType &Ty,
+ SourceLocation Loc) {
// No need to lookup anything if trivially copyable already
ASTContext &Ctx = SemaRef.getASTContext();
if (Ty.isTriviallyCopyableType(Ctx))
return {true, /*PD=*/std::nullopt};
-
+
bool markedCopyable = isMarkedDeviceCopyable(SemaRef, Ty, Loc);
if (const CXXRecordDecl *RD = Ty->getAsCXXRecordDecl()) {
// Set all lambdas as copyable: Future traversal passes will determine
- // whether or not the parameters/capture of the lambda are actually
+ // whether or not the parameters/capture of the lambda are actually
// copyable.
if (RD->isLambda())
return {true, /*PD=*/std::nullopt};
// TODO confirm:
- // - does RD have at least one eligible copy constructor, move constructor, copy assignment operator, or move assignment operator?
+ // - does RD have at least one eligible copy constructor, move constructor,
+ // copy assignment operator, or move assignment operator?
// - for each of aforementioned, ensure it is public
- // - confirm each does a bitwise copy (perhaps not possible, up to the user to enforce)
+ // - confirm each does a bitwise copy (perhaps not possible, up to the
+ // user to enforce)
// - confirm it has a non deleted destructor
- // - does the destructor have "no effect"? (perhaps not possible, up to user to enforce)
+ // - does the destructor have "no effect"? (perhaps not possible, up to
+ // user to enforce)
}
// TODO subsequent base classes shouldn't be checked for not device copyable
if (!markedCopyable)
- return { false, PartialDiagnostic(diag::err_sycl_kernel_param_not_device_copyable, SemaRef.Context.getDiagAllocator()) << Ty };
+ return {false,
+ PartialDiagnostic(diag::err_sycl_kernel_param_not_device_copyable,
+ SemaRef.Context.getDiagAllocator())
+ << Ty};
return {true, /*PD=*/std::nullopt};
}
-
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
@@ -921,15 +928,21 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
}
// TODO
- // - if a class is marked as copyable, I should STOP descending into the class's subfields
- // - we STOP because the user's declaration that a class is copyable should overwrite
- // the results of future traversals; we shouldn't descend any further
- // - although, because we want to catch classes that are obviously not copyable, we'll
- // need to do a shallow traversal in the future, checking that there aren't data members
- // in the class that are obviously not copyable: this will need to be its own function
- // - the stop condition should be if there are stuff that obviously breaks SYCL spec
+ // - if a class is marked as copyable, I should STOP descending into the
+ // class's subfields
+ // - we STOP because the user's declaration that a class is copyable
+ // should overwrite
+ // the results of future traversals; we shouldn't descend any further
+ // - although, because we want to catch classes that are obviously not
+ // copyable, we'll
+ // need to do a shallow traversal in the future, checking that there
+ // aren't data members in the class that are obviously not copyable:
+ // this will need to be its own function
+ // - the stop condition should be if there are stuff that obviously
+ // breaks SYCL spec
// for "is device copyable"
- // - if a class is not marked as copyable, _then_ descend into the class's subfields
+ // - if a class is not marked as copyable, _then_ descend into the class's
+ // subfields
// - this is current behavior
// TODO Issue warning if type is obviously not copyable
@@ -952,14 +965,16 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent)) {
Type = Ty->getPointeeType();
}
- const DeviceCopyableResult DCR = isDeviceCopyable(SemaSYCLRef.SemaRef, Type, Detail.Loc);
+ const DeviceCopyableResult DCR =
+ isDeviceCopyable(SemaSYCLRef.SemaRef, Type, Detail.Loc);
if (DCR.PD) {
// Emit diagnostics if any were generated.
SemaSYCLRef.Diag(Detail.Loc, DCR.PD.value());
emitObjectAccessPathNotes();
}
if (!DCR.copyable) {
- assert(DCR.PD && "DeviceCopyableResult must emit an explanatory diagnostic if Ty is not device copyable");
+ assert(DCR.PD && "DeviceCopyableResult must emit an explanatory "
+ "diagnostic if Ty is not device copyable");
IsValid = false;
return false;
}
>From 60c1691f2d9f3684453f3e024cba99375b4c0849 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 18 Aug 2026 13:33:14 -0700
Subject: [PATCH 07/15] Change behavior to stop checking is_device_copyable
params and structs
---
clang/include/clang/Basic/DiagnosticGroups.td | 1 +
.../clang/Basic/DiagnosticSemaKinds.td | 5 +-
clang/lib/Sema/SemaSYCL.cpp | 440 ++++++++++--------
.../sycl-kernel-param-is-device-copyable.cpp | 59 ++-
.../sycl-kernel-param-restrictions.cpp | 12 +-
5 files changed, 298 insertions(+), 219 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index b7072634cccf3..4a20dcd283eb5 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -856,6 +856,7 @@ def NonPODVarargs : DiagGroup<"non-pod-varargs">;
def ClassVarargs : DiagGroup<"class-varargs", [NonPODVarargs]>;
def : DiagGroup<"nonportable-cfstrings">;
def NonPortableSYCL : DiagGroup<"nonportable-sycl">;
+def SyclDeviceCopyable : DiagGroup<"sycl-device-copyable">;
def NonVirtualDtor : DiagGroup<"non-virtual-dtor">;
def GNUNullPointerArithmetic : DiagGroup<"gnu-null-pointer-arithmetic">;
def NullPointerArithmetic
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index fd1bbee14498d..584b8cabc48fa 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13584,8 +13584,9 @@ 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_sycl_kernel_param_not_device_copyable : Warning<
+ "%0 is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter">,
+ InGroup<SyclDeviceCopyable>, DefaultError;
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 9a28c60a745f5..6fee4989177c6 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -666,164 +666,6 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
-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, Loc,
- Sema::LookupNamespaceName);
- SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
-
- if (NamespaceResult.isAmbiguous())
- return nullptr;
-
- return NamespaceResult.getAsSingle<NamespaceDecl>();
-}
-
-struct DeviceCopyableResult {
- bool copyable;
- // PD contains a(n):
- // - error, if queried type is not copyable.
- // - warning, if queried type is marked copyable but detected to be non-device
- // copyable.
- std::optional<PartialDiagnostic> PD;
-};
-
-bool isMarkedDeviceCopyable(Sema &SemaRef, const QualType &Ty,
- SourceLocation Loc) {
- ASTContext &Ctx = SemaRef.getASTContext();
- // TODO too slow; cache this
- NamespaceDecl *SyclNamespace = getSyclNamespace(SemaRef, Loc);
- if (!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;
- }
-
- // is_device_copyable Identifier
- IdentifierInfo const &IDCIdent = Ctx.Idents.get("is_device_copyable");
-
- LookupResult IdentResult(SemaRef, &IDCIdent, Loc, Sema::LookupOrdinaryName);
- SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
-
- if (IdentResult.isAmbiguous()) {
- // TODO warn or error
- llvm::errs() << "debug2\n";
- return false;
- }
-
- ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
- if (nullptr == IDCDecl) {
- // TODO simply let go; it's undefined
- 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 simply let go; it's undefined
- llvm::errs() << "debug4\n";
- return false;
- }
- if (SemaRef.RequireCompleteType(Loc, IDCTrait,
- diag::err_sycl_incomplete_type_trait)) {
- // TODO do I need this if?
- llvm::errs() << "debug5\n";
- return false;
- }
-
- CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
- assert(RD && "specialization of class template is not a class?");
-
- // Look up the ::value 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 should I error or let go?
- // definitely error on ambiguous, but what about empty?
- llvm::errs() << "debug6\n";
- return false;
- }
-
- ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
- CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
- if (ValueExpr.isInvalid()) {
- // TODO compiler error: this shouldn't happen?
- llvm::errs() << "debug7\n";
- 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 compiler error: this shouldn't happen?
- llvm::errs() << "debug8\n";
- return false;
- }
-
- return IDCValue.getBoolValue();
-}
-
-// TODO should these results be cached?
-DeviceCopyableResult isDeviceCopyable(Sema &SemaRef, const QualType &Ty,
- SourceLocation Loc) {
- // No need to lookup anything if trivially copyable already
- ASTContext &Ctx = SemaRef.getASTContext();
- if (Ty.isTriviallyCopyableType(Ctx))
- return {true, /*PD=*/std::nullopt};
-
- bool markedCopyable = isMarkedDeviceCopyable(SemaRef, Ty, Loc);
- if (const CXXRecordDecl *RD = Ty->getAsCXXRecordDecl()) {
- // Set all lambdas as copyable: Future traversal passes will determine
- // whether or not the parameters/capture of the lambda are actually
- // copyable.
- if (RD->isLambda())
- return {true, /*PD=*/std::nullopt};
- // TODO confirm:
- // - does RD have at least one eligible copy constructor, move constructor,
- // copy assignment operator, or move assignment operator?
- // - for each of aforementioned, ensure it is public
- // - confirm each does a bitwise copy (perhaps not possible, up to the
- // user to enforce)
- // - confirm it has a non deleted destructor
- // - does the destructor have "no effect"? (perhaps not possible, up to
- // user to enforce)
- }
- // TODO subsequent base classes shouldn't be checked for not device copyable
-
- if (!markedCopyable)
- return {false,
- PartialDiagnostic(diag::err_sycl_kernel_param_not_device_copyable,
- SemaRef.Context.getDiagAllocator())
- << Ty};
- return {true, /*PD=*/std::nullopt};
-}
-
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
@@ -873,6 +715,179 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
}
}
+ // Return true if err_sycl_kernel_param_not_device_copyable has not been
+ // downgraded from an error via -Wsycl-device-copyable or
+ // -Wno-sycl-device-copyable. Loc is needed incase pragmas are used to set
+ // -Wsycl-device-copyable or -Wno-sycl-device-copyable.
+ bool enforcingDeviceCopyable(SourceLocation Loc) {
+ DiagnosticsEngine::Level DiagLvl =
+ SemaSYCLRef.getDiagnostics().getDiagnosticLevel(
+ diag::warn_sycl_kernel_param_not_device_copyable, Loc);
+ return DiagLvl >= DiagnosticsEngine::Level::Error;
+ }
+
+ // SyclNamespaceCache has value:
+ // - std::nullopt, if lookup hasn't been performed yet.
+ // - A pointer, if lookup was successful and sycl namespace was defined.
+ // - nullptr, if lookup failed.
+ std::optional<NamespaceDecl*> SyclNamespaceCache = std::nullopt;
+
+ NamespaceDecl *getSyclNamespace(SourceLocation Loc) {
+ if (SyclNamespaceCache.has_value())
+ return SyclNamespaceCache.value();
+
+ ASTContext &Ctx = SemaSYCLRef.getASTContext();
+ IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
+
+ Sema &SemaRef = SemaSYCLRef.SemaRef;
+ LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, Loc,
+ Sema::LookupNamespaceName);
+ SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
+
+ if (NamespaceResult.isAmbiguous())
+ SyclNamespaceCache = (NamespaceDecl*) nullptr;
+ else
+ SyclNamespaceCache = NamespaceResult.getAsSingle<NamespaceDecl>();
+
+ return SyclNamespaceCache.value();
+ }
+
+ llvm::DenseMap<QualType, bool> MarkedDeviceCopyableMap;
+
+ bool isMarkedDeviceCopyable(const QualType &Ty) {
+ // TODO: do I even bother caching lambdas?
+ // Lambda signatures are weird, and odds are lambdas are a one-time lookup anyway
+ QualType CanonicalTy = SemaSYCLRef.getASTContext().getCanonicalType(Ty);
+ auto It = MarkedDeviceCopyableMap.find(CanonicalTy);
+ if (It != MarkedDeviceCopyableMap.end()) {
+ return It->second;
+ }
+
+ // VMT's are not legal template parameters and cannot be marked device
+ // copyable.
+ if (Ty->isVariablyModifiedType()) {
+ llvm::errs() << "VMT shortpath\n";
+ return false;
+ }
+
+ auto DirectParent = ObjectAccessPath.back();
+ DiagDetails Detail = getObjectAccessDiagDetails(DirectParent);
+ SourceLocation Loc = Detail.Loc;
+ // If we are not checking device-copyability, completely ignore
+ // is_device_copyable.
+ if (!enforcingDeviceCopyable(Loc))
+ return false;
+
+ ASTContext &Ctx = SemaSYCLRef.getASTContext();
+ NamespaceDecl *SyclNamespace = getSyclNamespace(Loc);
+ if (!SyclNamespace) {
+ llvm::errs() << "Debug: SYCL namespace undefined / not found.\n";
+ return false;
+ }
+
+ // is_device_copyable Identifier
+ IdentifierInfo const &IDCIdent = Ctx.Idents.get("is_device_copyable");
+
+ Sema &SemaRef = SemaSYCLRef.SemaRef;
+ LookupResult IdentResult(SemaRef, &IDCIdent, Loc, Sema::LookupOrdinaryName);
+ SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
+
+ if (IdentResult.isAmbiguous()) {
+ // TODO warn or error
+ llvm::errs() << "debug2\n";
+ return false;
+ }
+
+ ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
+ if (nullptr == IDCDecl) {
+ // TODO simply let go; it's undefined
+ // TODO perhaps consider a warning: is_device_copyable *should* be defined
+ llvm::errs() << "debug3\n";
+ return false;
+ }
+
+ TemplateArgumentListInfo Args{};
+ TemplateArgument TyArg{Ty};
+ Args.addArgument(
+ SemaRef.getTrivialTemplateArgumentLoc(TyArg, QualType{}, Loc));
+
+ QualType IDCTrait;
+ {
+ // CheckTemplateIdType tries to diagnose illegal template argument types.
+ // An SFINAETrap is used here to catch said errors if Ty is an illegal
+ // template argument type.
+ Sema::SFINAETrap Trap(SemaRef);
+ IDCTrait = SemaRef.CheckTemplateIdType(
+ ElaboratedTypeKeyword::None, TemplateName{IDCDecl}, Loc, Args,
+ /*Scope=*/nullptr, /*ForNestedNameSpecifier=*/false);
+
+ if (IDCTrait.isNull()) {
+ // TODO simply let go; it's undefined
+ llvm::errs() << "debug4\n";
+ return false;
+ }
+
+ // IDCTrait must be checked for null before calling RequireCompleteType:
+ // it isn't valid to require completeness of a null QualType.
+ if (SemaRef.RequireCompleteType(Loc, IDCTrait,
+ diag::err_sycl_incomplete_type_trait)) {
+ llvm::errs() << "debug5\n";
+ return false;
+ }
+ }
+
+ CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
+ assert(RD && "specialization of class template is not a class?");
+
+ if (!RD->hasDefinition()) {
+ llvm::errs() << "RD undefined\n";
+ return false;
+ }
+
+ // Look up the ::value 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 should I error or let go?
+ // definitely error on ambiguous, but what about empty?
+ llvm::errs() << "debug6\n";
+ return false;
+ }
+
+ ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
+ CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
+ if (ValueExpr.isInvalid()) {
+ // TODO compiler error: this shouldn't happen?
+ llvm::errs() << "debug7\n";
+ 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 compiler error: this shouldn't happen?
+ llvm::errs() << "debug8\n";
+ return false;
+ }
+
+ bool Marked = IDCValue.getBoolValue();
+ MarkedDeviceCopyableMap.try_emplace(Ty, Marked);
+ return Marked;
+ }
+
public:
KernelParamsChecker(SemaSYCL &SR, SourceLocation Loc)
: ConstSubobjectVisitor<KernelParamsChecker>(SR.getASTContext()),
@@ -880,11 +895,15 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
void checkParameter(const ParmVarDecl *PVD) {
ObjectAccessPath.push_back(PVD);
- // Check the immediate type of the parameter.
- if (checkType(PVD->getType()) && checkDeviceCopyable(PVD->getType())) {
+ QualType Ty = PVD->getType();
+ // If type is explicitly marked as sycl::is_device_copyable, don't check the
+ // type further: Defining a non device-copyable type as device-copyable is
+ // UB.
+ // Otherwise, check the immediate type of the parameter.
+ if (!isMarkedDeviceCopyable(Ty) && checkType(Ty) && checkDeviceCopyable(Ty)) {
// If type checking wasn't short circuited, visit subobjects to check
// them.
- visit(PVD->getType());
+ visit(Ty);
}
ObjectAccessPath.pop_back();
assert(ObjectAccessPath.empty());
@@ -892,13 +911,27 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
bool visitBaseSpecifierPre(const CXXBaseSpecifier *BS) {
ObjectAccessPath.push_back(BS);
- return checkType(BS->getType());
+ QualType Ty = BS->getType();
+ // If type is explicitly marked as sycl::is_device_copyable, don't check the
+ // type further: Defining a non device-copyable type as device-copyable is
+ // UB.
+ if (isMarkedDeviceCopyable(Ty))
+ return false;
+ return checkType(Ty) && checkDeviceCopyable(Ty);
+ // TODO: if a class inherits an is_device_copyable class, does that make it device copyable?
}
bool visitFieldDeclPre(const FieldDecl *FD) {
ObjectAccessPath.push_back(FD);
- return checkType(FD->getType()) && checkDeviceCopyable(FD->getType());
- // TODO: do we need to check if a class's fields are deviceCopyable?
+ QualType Ty = FD->getType();
+ // If type is explicitly marked as sycl::is_device_copyable, don't check the
+ // type further: Defining a non device-copyable type as device-copyable is
+ // UB.
+ if (isMarkedDeviceCopyable(Ty))
+ return false;
+ return checkType(Ty) && checkDeviceCopyable(Ty);
+ // TODO: if a class has a field of type is_device_copyable, should I still check the field type?
+ // can I guarantee e.g. copy constructors and other special functions are going to be safe on the device?
}
// Returns true if subobjects should be visited and false otherwise.
@@ -927,54 +960,59 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
return false;
}
- // TODO
- // - if a class is marked as copyable, I should STOP descending into the
- // class's subfields
- // - we STOP because the user's declaration that a class is copyable
- // should overwrite
- // the results of future traversals; we shouldn't descend any further
- // - although, because we want to catch classes that are obviously not
- // copyable, we'll
- // need to do a shallow traversal in the future, checking that there
- // aren't data members in the class that are obviously not copyable:
- // this will need to be its own function
- // - the stop condition should be if there are stuff that obviously
- // breaks SYCL spec
- // for "is device copyable"
- // - if a class is not marked as copyable, _then_ descend into the class's
- // subfields
- // - this is current behavior
-
- // TODO Issue warning if type is obviously not copyable
- // ... you can do this by making some sort of dict for memoizing whether or
- // not a field/known type is not copyable
-
- // 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
+ // TODO Make sure I'm descending into Lambdas properly
return true;
}
+ // TODO this function needs more context to produce useful diagnostics
+ // - i.e. if traversal class is already marked is_device_copyable, then issue a
+ // warning instead
+ // TODO should these results be cached?
bool checkDeviceCopyable(QualType Ty) {
auto DirectParent = ObjectAccessPath.back();
- QualType Type = Ty;
-
DiagDetails Detail = getObjectAccessDiagDetails(DirectParent);
// Since references are allowed as direct kernel parameters, we need to
// explicitly check the referenced type:
- if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent)) {
+ QualType Type = Ty;
+ if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent))
Type = Ty->getPointeeType();
+ // TODO exit earlier if it's a reference but not a parmvardecl
+
+ // No need to lookup anything if trivially copyable already
+ if (Type.isTriviallyCopyableType(SemaSYCLRef.getASTContext()))
+ return true;
+
+ bool markedCopyable = isMarkedDeviceCopyable(Type);
+
+ if (const CXXRecordDecl *RD = Type->getAsCXXRecordDecl()) {
+ // Set all lambdas as copyable: Future traversal deeper into the lambda
+ // will determine whether or not the parameters/capture of the lambda are
+ // actually copyable.
+ if (RD->isLambda())
+ return true;
+ // TODO confirm:
+ // - does RD have at least one eligible copy constructor, move constructor,
+ // copy assignment operator, or move assignment operator?
+ // - for each of aforementioned, ensure it is public
+ // - confirm each does a bitwise copy (perhaps not possible, up to the
+ // user to enforce)
+ // - confirm it has a non deleted destructor
+ // - does the destructor have "no effect"? (perhaps not possible, up to
+ // user to enforce)
+
+ // MAKE SURE THE SPECIAL FUNCTION CHECKS PROPAGATE TO THE BASE CLASSES
}
- const DeviceCopyableResult DCR =
- isDeviceCopyable(SemaSYCLRef.SemaRef, Type, Detail.Loc);
- if (DCR.PD) {
- // Emit diagnostics if any were generated.
- SemaSYCLRef.Diag(Detail.Loc, DCR.PD.value());
+ // TODO subsequent base classes shouldn't be checked for not device copyable
+
+ if (!markedCopyable) {
+ // TODO Type is culprit but Ty is original caller: fix diagnostics to include both
+ SemaSYCLRef.Diag(Detail.Loc, diag::warn_sycl_kernel_param_not_device_copyable) << Type;
emitObjectAccessPathNotes();
- }
- if (!DCR.copyable) {
- assert(DCR.PD && "DeviceCopyableResult must emit an explanatory "
- "diagnostic if Ty is not device copyable");
+
+ // Continue traversal if not enforcing device-copyability.
+ if (!enforcingDeviceCopyable(Detail.Loc))
+ return true;
+ // Do not continue traversing if enforcing device-copyability.
IsValid = false;
return false;
}
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 4ef716ec78a37..6327b450e80ac 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -59,17 +59,64 @@ namespace iscopyable1 {
// Kernel entry point template definition.
template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
-void kernel_single_task(T t) {} // expected-error {{'NotTriviallyCopyable' is not device copyable}} \
- // expected-note {{within parameter 't' of type 'NotTriviallyCopyable' declared here}}
+void kernel_single_task(T t) {} // expected-error {{'NotTriviallyCopyable' is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter}}
+ // expected-note-re at -1 2{{within parameter 't' of type '{{.*}}' declared here}}
void test() {
DefinitelyCopyable a;
- kernel_single_task<KN<3>>(a);
+ kernel_single_task<KN<1>>(a);
NotTriviallyCopyable b;
- kernel_single_task<KN<2>>(b); // expected-note {{in instantiation of function template specialization 'iscopyable1::kernel_single_task<KN<2>, NotTriviallyCopyable>' requested here}}
+ kernel_single_task<KN<2>>(b);
+ // expected-note-re at -1 {{in instantiation of function template specialization 'iscopyable1::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
DeviceCopyable c;
- kernel_single_task<KN<1>>(c);
+ kernel_single_task<KN<3>>(c);
+
+
+ kernel_single_task<KN<4>>([=] { (void) a; });
+
+ kernel_single_task<KN<5>>([=] { (void) c; });
+
+ kernel_single_task<KN<6>>([=] { (void) b; });
+ // expected-error at -1 {{'NotTriviallyCopyable' is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'iscopyable1::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note at -3 {{within capture 'b' of lambda expression here}}
+
+ auto notCopyableLambda = [](NotTriviallyCopyable notCopyable) { (void) notCopyable; };
+ kernel_single_task<KN<7>>(notCopyableLambda);
+ kernel_single_task<KN<8>>([=](NotTriviallyCopyable NC) { notCopyableLambda(NC); });
+ // TODO this isn't firing; shouldn't this create an error
+
+}
+
+} // namespace iscopyable1
+
+struct HasNonDeviceCopyableMember {
+ NotTriviallyCopyable member;
+};
+
+struct ExplicitlyDeviceCopyableWithMember {
+ NotTriviallyCopyable member;
+};
+template<>
+struct sycl::is_device_copyable<ExplicitlyDeviceCopyableWithMember>
+ : std::true_type {};
+
+namespace iscopyable2 {
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} // expected-note-re {{within parameter 't' of type '{{.*}}' declared here}}
+
+void test() {
+ HasNonDeviceCopyableMember a;
+ kernel_single_task<KN<9>>([=] { (void) a; });
+ // expected-error at -1 {{'HasNonDeviceCopyableMember' is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'iscopyable2::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note at -3 {{within capture 'a' of lambda expression here}}
+
+ ExplicitlyDeviceCopyableWithMember b;
+ kernel_single_task<KN<10>>([=] { (void) b; });
}
-} // namespace iscopyable1
\ No newline at end of file
+} // namespace iscopyable2
+
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 429aefde2a016..4c72941e5e57e 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -Wno-vla-cxx-extension -fsycl-is-host -verify %s
-// RUN: %clang_cc1 -triple spirv64 -std=c++17 -fsyntax-only -Wno-vla-cxx-extension -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -Wno-vla-cxx-extension -Wno-sycl-device-copyable -fsycl-is-host -verify %s
+// RUN: %clang_cc1 -triple spirv64 -std=c++17 -fsyntax-only -Wno-vla-cxx-extension -Wno-sycl-device-copyable -fsycl-is-device -verify %s
// A unique kernel name type is required for each declared kernel entry point.
template<int, int = 0> struct KN;
@@ -8,14 +8,6 @@ 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.
>From 62d71f6086f192fe51f0b53eb43a35b2bf3fa514 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 18 Aug 2026 15:12:32 -0700
Subject: [PATCH 08/15] Move checkExplicitDeviceCopyable outside of
KernelParamsChecker so that the caches actually do something
---
clang/include/clang/Sema/SemaSYCL.h | 29 +++
clang/lib/Sema/SemaSYCL.cpp | 341 +++++++++++++---------------
2 files changed, 192 insertions(+), 178 deletions(-)
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 4980aa44c3012..c2cc9544c7898 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -18,10 +18,12 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Sema/Ownership.h"
#include "clang/Sema/SemaBase.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
namespace clang {
class Decl;
+class NamespaceDecl;
class ParsedAttr;
class SemaSYCL : public SemaBase {
@@ -96,6 +98,33 @@ class SemaSYCL : public SemaBase {
/// BuildSYCLKernelLaunchIdExpr().
StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
Expr *LaunchIdExpr);
+
+
+ /// Lookup the sycl namespace declaration. Return nullptr if lookup fails.
+ NamespaceDecl *getSyclNamespace(SourceLocation Loc);
+
+ /// Determines whether Ty has been explicitly marked as device-copyable via
+ /// sycl::is_device_copyable<Ty>
+ bool checkExplicitDeviceCopyable(QualType Ty, SourceLocation Loc);
+
+ // TODO: perhaps a good idea to add a normal checkDeviceCopyable as well, and
+ // not just checking if something's explicitly device-copyable.
+
+ /// Returns true if sycl-device-copyable diagnostics has not been downgraded
+ /// from an error via -Wsycl-device-copyable or -Wno-sycl-device-copyable.
+ /// 'Loc' is needed in case pragmas are used to set -Wsycl-device-copyable or
+ /// -Wno-sycl-device-copyable.
+ bool isEnforcingDeviceCopyable(SourceLocation Loc) const;
+
+private:
+ // Cache used by getSyclNamespace as the lookup is expensive. Since
+ // NamespaceDecls always track a "primary" DeclContext chain, saving this
+ // pointer is fine as future changes to the namespace will still be reachable
+ // via the decl chain.
+ NamespaceDecl *SyclNamespacePtr = nullptr;
+
+ /// Cache used by checkExplicitDEviceCopyable().
+ llvm::DenseMap<QualType, bool> MarkedDeviceCopyableCache;
};
} // namespace clang
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 6fee4989177c6..51bdeb1e11410 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -493,6 +493,163 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
return IdExpr;
}
+NamespaceDecl *SemaSYCL::getSyclNamespace(SourceLocation Loc) {
+ if (SyclNamespacePtr)
+ return SyclNamespacePtr;
+
+ ASTContext &Ctx = getASTContext();
+ IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
+
+ LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, Loc,
+ Sema::LookupNamespaceName);
+ SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
+
+ if (NamespaceResult.isAmbiguous())
+ return nullptr;
+
+ SyclNamespacePtr = NamespaceResult.getAsSingle<NamespaceDecl>();
+ return SyclNamespacePtr;
+ // Don't cache the sycl NamespaceDecl pointer if we can't find it, since
+ // namespace declarations could have been declared later on.
+}
+
+bool SemaSYCL::isEnforcingDeviceCopyable(SourceLocation Loc) const {
+ DiagnosticsEngine::Level DiagLvl =
+ getDiagnostics().getDiagnosticLevel(
+ diag::warn_sycl_kernel_param_not_device_copyable, Loc);
+ return DiagLvl >= DiagnosticsEngine::Level::Error;
+}
+
+bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty, SourceLocation Loc) {
+ ASTContext &Ctx = getASTContext();
+
+ // TODO: do I even bother caching lambdas?
+ // Lambda signatures are weird, and odds are lambdas are a one-time lookup anyway
+ QualType CanonicalTy = Ctx.getCanonicalType(Ty);
+ auto It = MarkedDeviceCopyableCache.find(CanonicalTy);
+ if (It != MarkedDeviceCopyableCache.end()) {
+ return It->second;
+ }
+
+ // VMT's are not legal template parameters and cannot be marked device
+ // copyable.
+ if (Ty->isVariablyModifiedType()) {
+ llvm::errs() << "VMT shortpath\n";
+ return false;
+ }
+
+ // If we are not checking device-copyability, completely ignore
+ // is_device_copyable.
+ if (!isEnforcingDeviceCopyable(Loc))
+ return false;
+
+ NamespaceDecl *SyclNamespace = getSyclNamespace(Loc);
+ if (!SyclNamespace) {
+ llvm::errs() << "Debug: SYCL namespace undefined / not found.\n";
+ return false;
+ }
+
+ // is_device_copyable Identifier
+ IdentifierInfo const &IDCIdent = Ctx.Idents.get("is_device_copyable");
+ LookupResult IdentResult(SemaRef, &IDCIdent, Loc, Sema::LookupOrdinaryName);
+ SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
+
+ if (IdentResult.isAmbiguous()) {
+ // TODO warn or error
+ llvm::errs() << "debug2\n";
+ return false;
+ }
+
+ ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
+ if (nullptr == IDCDecl) {
+ // TODO simply let go; it's undefined
+ // TODO perhaps consider a warning: is_device_copyable *should* be defined
+ llvm::errs() << "debug3\n";
+ return false;
+ }
+
+ TemplateArgumentListInfo Args{};
+ TemplateArgument TyArg{Ty};
+ Args.addArgument(
+ SemaRef.getTrivialTemplateArgumentLoc(TyArg, QualType{}, Loc));
+
+ QualType IDCTrait;
+ {
+ // CheckTemplateIdType tries to diagnose illegal template argument types.
+ // An SFINAETrap is used here to catch said errors if Ty is an illegal
+ // template argument type.
+ Sema::SFINAETrap Trap(SemaRef);
+ IDCTrait = SemaRef.CheckTemplateIdType(
+ ElaboratedTypeKeyword::None, TemplateName{IDCDecl}, Loc, Args,
+ /*Scope=*/nullptr, /*ForNestedNameSpecifier=*/false);
+
+ if (IDCTrait.isNull()) {
+ // TODO simply let go; it's undefined
+ llvm::errs() << "debug4\n";
+ return false;
+ }
+
+ // IDCTrait must be checked for null before calling RequireCompleteType:
+ // it isn't valid to require completeness of a null QualType.
+ if (SemaRef.RequireCompleteType(Loc, IDCTrait,
+ diag::err_sycl_incomplete_type_trait)) {
+ llvm::errs() << "debug5\n";
+ return false;
+ }
+ }
+
+ CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
+ assert(RD && "specialization of class template is not a class?");
+
+ if (!RD->hasDefinition()) {
+ llvm::errs() << "RD undefined\n";
+ return false;
+ }
+
+ // Look up the ::value 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 should I error or let go?
+ // definitely error on ambiguous, but what about empty?
+ llvm::errs() << "debug6\n";
+ return false;
+ }
+
+ ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
+ CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
+ if (ValueExpr.isInvalid()) {
+ // TODO compiler error: this shouldn't happen?
+ llvm::errs() << "debug7\n";
+ 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 compiler error: this shouldn't happen?
+ llvm::errs() << "debug8\n";
+ return false;
+ }
+
+ bool Marked = IDCValue.getBoolValue();
+ MarkedDeviceCopyableCache.try_emplace(Ty, Marked);
+ return Marked;
+}
+
namespace {
// Constructs the arguments to be passed for the SYCL kernel launch call.
@@ -715,179 +872,6 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
}
}
- // Return true if err_sycl_kernel_param_not_device_copyable has not been
- // downgraded from an error via -Wsycl-device-copyable or
- // -Wno-sycl-device-copyable. Loc is needed incase pragmas are used to set
- // -Wsycl-device-copyable or -Wno-sycl-device-copyable.
- bool enforcingDeviceCopyable(SourceLocation Loc) {
- DiagnosticsEngine::Level DiagLvl =
- SemaSYCLRef.getDiagnostics().getDiagnosticLevel(
- diag::warn_sycl_kernel_param_not_device_copyable, Loc);
- return DiagLvl >= DiagnosticsEngine::Level::Error;
- }
-
- // SyclNamespaceCache has value:
- // - std::nullopt, if lookup hasn't been performed yet.
- // - A pointer, if lookup was successful and sycl namespace was defined.
- // - nullptr, if lookup failed.
- std::optional<NamespaceDecl*> SyclNamespaceCache = std::nullopt;
-
- NamespaceDecl *getSyclNamespace(SourceLocation Loc) {
- if (SyclNamespaceCache.has_value())
- return SyclNamespaceCache.value();
-
- ASTContext &Ctx = SemaSYCLRef.getASTContext();
- IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
-
- Sema &SemaRef = SemaSYCLRef.SemaRef;
- LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, Loc,
- Sema::LookupNamespaceName);
- SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
-
- if (NamespaceResult.isAmbiguous())
- SyclNamespaceCache = (NamespaceDecl*) nullptr;
- else
- SyclNamespaceCache = NamespaceResult.getAsSingle<NamespaceDecl>();
-
- return SyclNamespaceCache.value();
- }
-
- llvm::DenseMap<QualType, bool> MarkedDeviceCopyableMap;
-
- bool isMarkedDeviceCopyable(const QualType &Ty) {
- // TODO: do I even bother caching lambdas?
- // Lambda signatures are weird, and odds are lambdas are a one-time lookup anyway
- QualType CanonicalTy = SemaSYCLRef.getASTContext().getCanonicalType(Ty);
- auto It = MarkedDeviceCopyableMap.find(CanonicalTy);
- if (It != MarkedDeviceCopyableMap.end()) {
- return It->second;
- }
-
- // VMT's are not legal template parameters and cannot be marked device
- // copyable.
- if (Ty->isVariablyModifiedType()) {
- llvm::errs() << "VMT shortpath\n";
- return false;
- }
-
- auto DirectParent = ObjectAccessPath.back();
- DiagDetails Detail = getObjectAccessDiagDetails(DirectParent);
- SourceLocation Loc = Detail.Loc;
- // If we are not checking device-copyability, completely ignore
- // is_device_copyable.
- if (!enforcingDeviceCopyable(Loc))
- return false;
-
- ASTContext &Ctx = SemaSYCLRef.getASTContext();
- NamespaceDecl *SyclNamespace = getSyclNamespace(Loc);
- if (!SyclNamespace) {
- llvm::errs() << "Debug: SYCL namespace undefined / not found.\n";
- return false;
- }
-
- // is_device_copyable Identifier
- IdentifierInfo const &IDCIdent = Ctx.Idents.get("is_device_copyable");
-
- Sema &SemaRef = SemaSYCLRef.SemaRef;
- LookupResult IdentResult(SemaRef, &IDCIdent, Loc, Sema::LookupOrdinaryName);
- SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
-
- if (IdentResult.isAmbiguous()) {
- // TODO warn or error
- llvm::errs() << "debug2\n";
- return false;
- }
-
- ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
- if (nullptr == IDCDecl) {
- // TODO simply let go; it's undefined
- // TODO perhaps consider a warning: is_device_copyable *should* be defined
- llvm::errs() << "debug3\n";
- return false;
- }
-
- TemplateArgumentListInfo Args{};
- TemplateArgument TyArg{Ty};
- Args.addArgument(
- SemaRef.getTrivialTemplateArgumentLoc(TyArg, QualType{}, Loc));
-
- QualType IDCTrait;
- {
- // CheckTemplateIdType tries to diagnose illegal template argument types.
- // An SFINAETrap is used here to catch said errors if Ty is an illegal
- // template argument type.
- Sema::SFINAETrap Trap(SemaRef);
- IDCTrait = SemaRef.CheckTemplateIdType(
- ElaboratedTypeKeyword::None, TemplateName{IDCDecl}, Loc, Args,
- /*Scope=*/nullptr, /*ForNestedNameSpecifier=*/false);
-
- if (IDCTrait.isNull()) {
- // TODO simply let go; it's undefined
- llvm::errs() << "debug4\n";
- return false;
- }
-
- // IDCTrait must be checked for null before calling RequireCompleteType:
- // it isn't valid to require completeness of a null QualType.
- if (SemaRef.RequireCompleteType(Loc, IDCTrait,
- diag::err_sycl_incomplete_type_trait)) {
- llvm::errs() << "debug5\n";
- return false;
- }
- }
-
- CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
- assert(RD && "specialization of class template is not a class?");
-
- if (!RD->hasDefinition()) {
- llvm::errs() << "RD undefined\n";
- return false;
- }
-
- // Look up the ::value 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 should I error or let go?
- // definitely error on ambiguous, but what about empty?
- llvm::errs() << "debug6\n";
- return false;
- }
-
- ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
- CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
- if (ValueExpr.isInvalid()) {
- // TODO compiler error: this shouldn't happen?
- llvm::errs() << "debug7\n";
- 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 compiler error: this shouldn't happen?
- llvm::errs() << "debug8\n";
- return false;
- }
-
- bool Marked = IDCValue.getBoolValue();
- MarkedDeviceCopyableMap.try_emplace(Ty, Marked);
- return Marked;
- }
-
public:
KernelParamsChecker(SemaSYCL &SR, SourceLocation Loc)
: ConstSubobjectVisitor<KernelParamsChecker>(SR.getASTContext()),
@@ -900,7 +884,8 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// type further: Defining a non device-copyable type as device-copyable is
// UB.
// Otherwise, check the immediate type of the parameter.
- if (!isMarkedDeviceCopyable(Ty) && checkType(Ty) && checkDeviceCopyable(Ty)) {
+ if (!SemaSYCLRef.checkExplicitDeviceCopyable(Ty, PVD->getLocation()) &&
+ checkType(Ty) && checkDeviceCopyable(Ty)) {
// If type checking wasn't short circuited, visit subobjects to check
// them.
visit(Ty);
@@ -915,7 +900,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// If type is explicitly marked as sycl::is_device_copyable, don't check the
// type further: Defining a non device-copyable type as device-copyable is
// UB.
- if (isMarkedDeviceCopyable(Ty))
+ if (SemaSYCLRef.checkExplicitDeviceCopyable(Ty, BS->getBaseTypeLoc()))
return false;
return checkType(Ty) && checkDeviceCopyable(Ty);
// TODO: if a class inherits an is_device_copyable class, does that make it device copyable?
@@ -927,7 +912,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// If type is explicitly marked as sycl::is_device_copyable, don't check the
// type further: Defining a non device-copyable type as device-copyable is
// UB.
- if (isMarkedDeviceCopyable(Ty))
+ if (SemaSYCLRef.checkExplicitDeviceCopyable(Ty, FD->getLocation()))
return false;
return checkType(Ty) && checkDeviceCopyable(Ty);
// TODO: if a class has a field of type is_device_copyable, should I still check the field type?
@@ -982,7 +967,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (Type.isTriviallyCopyableType(SemaSYCLRef.getASTContext()))
return true;
- bool markedCopyable = isMarkedDeviceCopyable(Type);
+ bool markedCopyable = SemaSYCLRef.checkExplicitDeviceCopyable(Type, Detail.Loc);
if (const CXXRecordDecl *RD = Type->getAsCXXRecordDecl()) {
// Set all lambdas as copyable: Future traversal deeper into the lambda
@@ -1010,7 +995,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
emitObjectAccessPathNotes();
// Continue traversal if not enforcing device-copyability.
- if (!enforcingDeviceCopyable(Detail.Loc))
+ if (!SemaSYCLRef.isEnforcingDeviceCopyable(Detail.Loc))
return true;
// Do not continue traversing if enforcing device-copyability.
IsValid = false;
>From 98031ba3db761c5d0e673331a43ed53918761336 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Tue, 18 Aug 2026 15:15:27 -0700
Subject: [PATCH 09/15] 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 51bdeb1e11410..99945171bd5f4 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -501,7 +501,7 @@ NamespaceDecl *SemaSYCL::getSyclNamespace(SourceLocation Loc) {
IdentifierInfo const &SyclNamespaceID = Ctx.Idents.get("sycl");
LookupResult NamespaceResult(SemaRef, &SyclNamespaceID, Loc,
- Sema::LookupNamespaceName);
+ Sema::LookupNamespaceName);
SemaRef.LookupQualifiedName(NamespaceResult, Ctx.getTranslationUnitDecl());
if (NamespaceResult.isAmbiguous())
@@ -514,17 +514,18 @@ NamespaceDecl *SemaSYCL::getSyclNamespace(SourceLocation Loc) {
}
bool SemaSYCL::isEnforcingDeviceCopyable(SourceLocation Loc) const {
- DiagnosticsEngine::Level DiagLvl =
- getDiagnostics().getDiagnosticLevel(
- diag::warn_sycl_kernel_param_not_device_copyable, Loc);
+ DiagnosticsEngine::Level DiagLvl = getDiagnostics().getDiagnosticLevel(
+ diag::warn_sycl_kernel_param_not_device_copyable, Loc);
return DiagLvl >= DiagnosticsEngine::Level::Error;
}
-bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty, SourceLocation Loc) {
+bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty,
+ SourceLocation Loc) {
ASTContext &Ctx = getASTContext();
// TODO: do I even bother caching lambdas?
- // Lambda signatures are weird, and odds are lambdas are a one-time lookup anyway
+ // Lambda signatures are weird, and odds are lambdas are a one-time lookup
+ // anyway
QualType CanonicalTy = Ctx.getCanonicalType(Ty);
auto It = MarkedDeviceCopyableCache.find(CanonicalTy);
if (It != MarkedDeviceCopyableCache.end()) {
@@ -630,9 +631,9 @@ bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty, SourceLocation Loc
Expr *GotExpr;
ICEDiagnoser(QualType &TT, Expr *E) : TraitTy(TT), GotExpr(E) {}
Sema::SemaDiagnosticBuilder diagnoseNotICE(Sema &S,
- SourceLocation Loc) override {
+ SourceLocation Loc) override {
return S.Diag(Loc, diag::err_sycl_unexpected_type_trait_val)
- << TraitTy << "std::true_type or std::false_type" << GotExpr;
+ << TraitTy << "std::true_type or std::false_type" << GotExpr;
}
} Diagnoser(IDCTrait, ValueExpr.get());
@@ -898,25 +899,27 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
ObjectAccessPath.push_back(BS);
QualType Ty = BS->getType();
// If type is explicitly marked as sycl::is_device_copyable, don't check the
- // type further: Defining a non device-copyable type as device-copyable is
+ // type further: Defining a non device-copyable type as device-copyable is
// UB.
if (SemaSYCLRef.checkExplicitDeviceCopyable(Ty, BS->getBaseTypeLoc()))
return false;
return checkType(Ty) && checkDeviceCopyable(Ty);
- // TODO: if a class inherits an is_device_copyable class, does that make it device copyable?
+ // TODO: if a class inherits an is_device_copyable class, does that make it
+ // device copyable?
}
bool visitFieldDeclPre(const FieldDecl *FD) {
ObjectAccessPath.push_back(FD);
QualType Ty = FD->getType();
// If type is explicitly marked as sycl::is_device_copyable, don't check the
- // type further: Defining a non device-copyable type as device-copyable is
+ // type further: Defining a non device-copyable type as device-copyable is
// UB.
if (SemaSYCLRef.checkExplicitDeviceCopyable(Ty, FD->getLocation()))
return false;
return checkType(Ty) && checkDeviceCopyable(Ty);
- // TODO: if a class has a field of type is_device_copyable, should I still check the field type?
- // can I guarantee e.g. copy constructors and other special functions are going to be safe on the device?
+ // TODO: if a class has a field of type is_device_copyable, should I still
+ // check the field type? can I guarantee e.g. copy constructors and other
+ // special functions are going to be safe on the device?
}
// Returns true if subobjects should be visited and false otherwise.
@@ -945,13 +948,13 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
return false;
}
- // TODO Make sure I'm descending into Lambdas properly
+ // TODO Make sure I'm descending into Lambdas properly
return true;
}
// TODO this function needs more context to produce useful diagnostics
- // - i.e. if traversal class is already marked is_device_copyable, then issue a
- // warning instead
+ // - i.e. if traversal class is already marked is_device_copyable, then issue
+ // a warning instead
// TODO should these results be cached?
bool checkDeviceCopyable(QualType Ty) {
auto DirectParent = ObjectAccessPath.back();
@@ -961,13 +964,14 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
QualType Type = Ty;
if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent))
Type = Ty->getPointeeType();
- // TODO exit earlier if it's a reference but not a parmvardecl
+ // TODO exit earlier if it's a reference but not a parmvardecl
// No need to lookup anything if trivially copyable already
if (Type.isTriviallyCopyableType(SemaSYCLRef.getASTContext()))
return true;
- bool markedCopyable = SemaSYCLRef.checkExplicitDeviceCopyable(Type, Detail.Loc);
+ bool markedCopyable =
+ SemaSYCLRef.checkExplicitDeviceCopyable(Type, Detail.Loc);
if (const CXXRecordDecl *RD = Type->getAsCXXRecordDecl()) {
// Set all lambdas as copyable: Future traversal deeper into the lambda
@@ -976,8 +980,8 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (RD->isLambda())
return true;
// TODO confirm:
- // - does RD have at least one eligible copy constructor, move constructor,
- // copy assignment operator, or move assignment operator?
+ // - does RD have at least one eligible copy constructor, move
+ // constructor, copy assignment operator, or move assignment operator?
// - for each of aforementioned, ensure it is public
// - confirm each does a bitwise copy (perhaps not possible, up to the
// user to enforce)
@@ -990,8 +994,11 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// TODO subsequent base classes shouldn't be checked for not device copyable
if (!markedCopyable) {
- // TODO Type is culprit but Ty is original caller: fix diagnostics to include both
- SemaSYCLRef.Diag(Detail.Loc, diag::warn_sycl_kernel_param_not_device_copyable) << Type;
+ // TODO Type is culprit but Ty is original caller: fix diagnostics to
+ // include both
+ SemaSYCLRef.Diag(Detail.Loc,
+ diag::warn_sycl_kernel_param_not_device_copyable)
+ << Type;
emitObjectAccessPathNotes();
// Continue traversal if not enforcing device-copyability.
>From cf10787a34ba7f838f8ebbe74f049eae449fd314 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 20 Aug 2026 17:07:26 -0700
Subject: [PATCH 10/15] Add class warnings
---
clang/include/clang/Basic/DiagnosticGroups.td | 1 +
.../clang/Basic/DiagnosticSemaKinds.td | 20 +++
clang/include/clang/Sema/Sema.h | 3 +
clang/include/clang/Sema/SemaSYCL.h | 3 +-
clang/lib/Sema/SemaLookup.cpp | 16 ++
clang/lib/Sema/SemaSYCL.cpp | 170 +++++++++++-------
.../sycl-kernel-param-is-device-copyable.cpp | 160 ++++++++++++++++-
7 files changed, 299 insertions(+), 74 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index 4a20dcd283eb5..9d442752b8515 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -857,6 +857,7 @@ def ClassVarargs : DiagGroup<"class-varargs", [NonPODVarargs]>;
def : DiagGroup<"nonportable-cfstrings">;
def NonPortableSYCL : DiagGroup<"nonportable-sycl">;
def SyclDeviceCopyable : DiagGroup<"sycl-device-copyable">;
+def PedanticSycl : DiagGroup<"pedantic-sycl">;
def NonVirtualDtor : DiagGroup<"non-virtual-dtor">;
def GNUNullPointerArithmetic : DiagGroup<"gnu-null-pointer-arithmetic">;
def NullPointerArithmetic
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 584b8cabc48fa..b0c683a16752d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13522,6 +13522,8 @@ 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_type_trait_bad_value : Error<
+ "SYCL type trait %0 does not provide a valid 'value' member">;
def err_sycl_unexpected_type_trait_val : Error<
"expected %1 for SYCL type trait %0, but found %2 instead">;
@@ -13587,6 +13589,24 @@ def err_sycl_entry_point_device_use : Error<
def warn_sycl_kernel_param_not_device_copyable : Warning<
"%0 is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter">,
InGroup<SyclDeviceCopyable>, DefaultError;
+def warn_sycl_device_copyable_no_eligible_smf : Warning<
+ "%0 is explicitly marked as device copyable (sycl::is_device_copyable) but"
+ " has no eligible copy constructor, move constructor, copy assignment"
+ " operator, or move assignment operator">,
+ InGroup<PedanticSycl>, DefaultIgnore;
+def warn_sycl_device_copyable_smf_not_public : Warning<
+ "%0 is explicitly marked as device copyable (sycl::is_device_copyable) but"
+ " its eligible %enum_select<EligibleSMFKind>{"
+ "%CopyCtor{copy constructor}|"
+ "%MoveCtor{move constructor}|"
+ "%CopyAssign{copy assignment operator}|"
+ "%MoveAssign{move assignment operator}"
+ "}1 is not public">,
+ InGroup<PedanticSycl>, DefaultIgnore;
+def warn_sycl_device_copyable_bad_destructor : Warning<
+ "%0 is explicitly marked as device copyable (sycl::is_device_copyable) but"
+ " does not have a public, non-deleted destructor">,
+ InGroup<PedanticSycl>;
def warn_cuda_maxclusterrank_sm_90 : Warning<
"maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 8a30f6319bcef..0921198b98316 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -9637,6 +9637,9 @@ class Sema final : public SemaBase {
/// Look up the constructors for the given class.
DeclContextLookupResult LookupConstructors(CXXRecordDecl *Class);
+ /// Look up the copy and move assignment operators for the given class.
+ DeclContextLookupResult LookupAssignmentOperators(CXXRecordDecl *Class);
+
/// Look up the default constructor for the given class.
CXXConstructorDecl *LookupDefaultConstructor(CXXRecordDecl *Class);
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index c2cc9544c7898..e36868d21f9d9 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -99,7 +99,6 @@ class SemaSYCL : public SemaBase {
StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
Expr *LaunchIdExpr);
-
/// Lookup the sycl namespace declaration. Return nullptr if lookup fails.
NamespaceDecl *getSyclNamespace(SourceLocation Loc);
@@ -120,7 +119,7 @@ class SemaSYCL : public SemaBase {
// Cache used by getSyclNamespace as the lookup is expensive. Since
// NamespaceDecls always track a "primary" DeclContext chain, saving this
// pointer is fine as future changes to the namespace will still be reachable
- // via the decl chain.
+ // via the decl chain.
NamespaceDecl *SyclNamespacePtr = nullptr;
/// Cache used by checkExplicitDEviceCopyable().
diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp
index 43129800e9813..5811d99c0c3f0 100644
--- a/clang/lib/Sema/SemaLookup.cpp
+++ b/clang/lib/Sema/SemaLookup.cpp
@@ -3651,6 +3651,22 @@ DeclContext::lookup_result Sema::LookupConstructors(CXXRecordDecl *Class) {
return Class->lookup(Name);
}
+DeclContext::lookup_result Sema::LookupAssignmentOperators(CXXRecordDecl *Class) {
+ // If the implicit copy or move assignment operators have not yet been
+ // declared, do so now.
+ if (CanDeclareSpecialMemberFunction(Class)) {
+ runWithSufficientStackSpace(Class->getLocation(), [&] {
+ if (Class->needsImplicitCopyAssignment())
+ DeclareImplicitCopyAssignment(Class);
+ if (getLangOpts().CPlusPlus11 && Class->needsImplicitMoveAssignment())
+ DeclareImplicitMoveAssignment(Class);
+ });
+ }
+
+ DeclarationName Name = Context.DeclarationNames.getCXXOperatorName(OO_Equal);
+ return Class->lookup(Name);
+}
+
CXXMethodDecl *Sema::LookupCopyingAssignment(CXXRecordDecl *Class,
unsigned Quals, bool RValueThis,
unsigned ThisQuals) {
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 99945171bd5f4..3f8d05bb57291 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -523,9 +523,6 @@ bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty,
SourceLocation Loc) {
ASTContext &Ctx = getASTContext();
- // TODO: do I even bother caching lambdas?
- // Lambda signatures are weird, and odds are lambdas are a one-time lookup
- // anyway
QualType CanonicalTy = Ctx.getCanonicalType(Ty);
auto It = MarkedDeviceCopyableCache.find(CanonicalTy);
if (It != MarkedDeviceCopyableCache.end()) {
@@ -534,21 +531,16 @@ bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty,
// VMT's are not legal template parameters and cannot be marked device
// copyable.
- if (Ty->isVariablyModifiedType()) {
- llvm::errs() << "VMT shortpath\n";
+ if (Ty->isVariablyModifiedType())
return false;
- }
-
// If we are not checking device-copyability, completely ignore
// is_device_copyable.
if (!isEnforcingDeviceCopyable(Loc))
return false;
NamespaceDecl *SyclNamespace = getSyclNamespace(Loc);
- if (!SyclNamespace) {
- llvm::errs() << "Debug: SYCL namespace undefined / not found.\n";
+ if (!SyclNamespace)
return false;
- }
// is_device_copyable Identifier
IdentifierInfo const &IDCIdent = Ctx.Idents.get("is_device_copyable");
@@ -556,18 +548,13 @@ bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty,
SemaRef.LookupQualifiedName(IdentResult, SyclNamespace);
if (IdentResult.isAmbiguous()) {
- // TODO warn or error
- llvm::errs() << "debug2\n";
+ SemaRef.DiagnoseAmbiguousLookup(IdentResult);
return false;
}
ClassTemplateDecl *IDCDecl = IdentResult.getAsSingle<ClassTemplateDecl>();
- if (nullptr == IDCDecl) {
- // TODO simply let go; it's undefined
- // TODO perhaps consider a warning: is_device_copyable *should* be defined
- llvm::errs() << "debug3\n";
+ if (!IDCDecl)
return false;
- }
TemplateArgumentListInfo Args{};
TemplateArgument TyArg{Ty};
@@ -583,46 +570,38 @@ bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty,
IDCTrait = SemaRef.CheckTemplateIdType(
ElaboratedTypeKeyword::None, TemplateName{IDCDecl}, Loc, Args,
/*Scope=*/nullptr, /*ForNestedNameSpecifier=*/false);
-
- if (IDCTrait.isNull()) {
- // TODO simply let go; it's undefined
- llvm::errs() << "debug4\n";
+ if (IDCTrait.isNull())
return false;
- }
// IDCTrait must be checked for null before calling RequireCompleteType:
// it isn't valid to require completeness of a null QualType.
if (SemaRef.RequireCompleteType(Loc, IDCTrait,
- diag::err_sycl_incomplete_type_trait)) {
- llvm::errs() << "debug5\n";
+ diag::err_sycl_incomplete_type_trait))
return false;
- }
}
CXXRecordDecl *RD = IDCTrait->getAsCXXRecordDecl();
assert(RD && "specialization of class template is not a class?");
-
- if (!RD->hasDefinition()) {
- llvm::errs() << "RD undefined\n";
- return false;
- }
+ assert(RD->hasDefinition() &&
+ "RequireCompleteType should have guaranteed a definition exists");
// Look up the ::value 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 should I error or let go?
- // definitely error on ambiguous, but what about empty?
- llvm::errs() << "debug6\n";
+ if (ValueResult.isAmbiguous()) {
+ SemaRef.DiagnoseAmbiguousLookup(ValueResult);
+ return false;
+ }
+ if (ValueResult.empty()) {
+ SemaRef.Diag(Loc, diag::err_sycl_type_trait_bad_value) << IDCTrait;
return false;
}
ExprResult ValueExpr = SemaRef.BuildDeclarationNameExpr(
CXXScopeSpec{}, ValueResult, /*NeedsADL=*/false);
if (ValueExpr.isInvalid()) {
- // TODO compiler error: this shouldn't happen?
- llvm::errs() << "debug7\n";
+ SemaRef.Diag(Loc, diag::err_sycl_type_trait_bad_value) << IDCTrait;
return false;
}
@@ -640,11 +619,8 @@ bool SemaSYCL::checkExplicitDeviceCopyable(const QualType Ty,
llvm::APSInt IDCValue;
ValueExpr = SemaRef.VerifyIntegerConstantExpression(ValueExpr.get(),
&IDCValue, Diagnoser);
- if (ValueExpr.isInvalid()) {
- // TODO compiler error: this shouldn't happen?
- llvm::errs() << "debug8\n";
+ if (ValueExpr.isInvalid())
return false;
- }
bool Marked = IDCValue.getBoolValue();
MarkedDeviceCopyableCache.try_emplace(Ty, Marked);
@@ -952,9 +928,6 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
return true;
}
- // TODO this function needs more context to produce useful diagnostics
- // - i.e. if traversal class is already marked is_device_copyable, then issue
- // a warning instead
// TODO should these results be cached?
bool checkDeviceCopyable(QualType Ty) {
auto DirectParent = ObjectAccessPath.back();
@@ -964,38 +937,99 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
QualType Type = Ty;
if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent))
Type = Ty->getPointeeType();
- // TODO exit earlier if it's a reference but not a parmvardecl
// No need to lookup anything if trivially copyable already
if (Type.isTriviallyCopyableType(SemaSYCLRef.getASTContext()))
return true;
-
- bool markedCopyable =
+ // FIXME: isTriviallyCopyableType allows explicitly deleted destructors,
+ // but the SYCL spec stipulates that deleted destructors on an explicitly-
+ // marked device-copyable class is UB.
+ // This shortcut path ends up resulting in missed -Wpendantic-sycl deleted
+ // destructor warnings. Consider moving -Wpendantic-sycl deleted destructor
+ // test before this shortcut.
+
+ bool MarkedCopyable =
SemaSYCLRef.checkExplicitDeviceCopyable(Type, Detail.Loc);
- if (const CXXRecordDecl *RD = Type->getAsCXXRecordDecl()) {
- // Set all lambdas as copyable: Future traversal deeper into the lambda
- // will determine whether or not the parameters/capture of the lambda are
- // actually copyable.
- if (RD->isLambda())
- return true;
- // TODO confirm:
- // - does RD have at least one eligible copy constructor, move
- // constructor, copy assignment operator, or move assignment operator?
- // - for each of aforementioned, ensure it is public
- // - confirm each does a bitwise copy (perhaps not possible, up to the
- // user to enforce)
- // - confirm it has a non deleted destructor
- // - does the destructor have "no effect"? (perhaps not possible, up to
- // user to enforce)
-
- // MAKE SURE THE SPECIAL FUNCTION CHECKS PROPAGATE TO THE BASE CLASSES
+ CXXRecordDecl *RD = Type->getAsCXXRecordDecl();
+ // Set all lambdas as copyable: Future traversal deeper into the lambda
+ // will determine whether or not the parameters/capture of the lambda are
+ // actually copyable.
+ if (RD && RD->isLambda())
+ return true;
+
+ // Checking SYCL 2020 3.13.1, when explicitly declaring certain class types
+ // as device copyable:
+ if (MarkedCopyable && RD) {
+ // * Type T has a public non-deleted destructor; and
+ CXXDestructorDecl *DD = SemaSYCLRef.SemaRef.LookupDestructor(RD);
+ if (!DD || DD->isDeleted() || DD->getAccess() != AS_public) {
+ SemaSYCLRef.Diag(DD ? DD->getLocation() : RD->getLocation(),
+ diag::warn_sycl_device_copyable_bad_destructor)
+ << Type;
+ emitObjectAccessPathNotes();
+ }
+
+ DiagnosticsEngine &Diags = SemaSYCLRef.getDiagnostics();
+ bool CheckSMFEligible = !Diags.isIgnored(
+ diag::warn_sycl_device_copyable_smf_not_public, Detail.Loc);
+ if (CheckSMFEligible) {
+ // * Each eligible copy constructor, move constructor, copy assignment
+ // operator, and move assignment operator is public;
+ bool HasEligibleSMF = false;
+ for (const NamedDecl *ND : SemaSYCLRef.SemaRef.LookupConstructors(RD)) {
+ auto *CD = dyn_cast<CXXConstructorDecl>(ND);
+ if (CD && CD->isCopyOrMoveConstructor() && !CD->isDeleted() &&
+ !CD->isIneligibleOrNotSelected()) {
+ if (CD->getAccess() != AS_public) {
+ SemaSYCLRef.Diag(CD->getLocation(),
+ diag::warn_sycl_device_copyable_smf_not_public)
+ << Type
+ << (CD->isCopyConstructor()
+ ? diag::EligibleSMFKind::CopyCtor
+ : diag::EligibleSMFKind::MoveCtor);
+ emitObjectAccessPathNotes();
+ }
+ HasEligibleSMF = true;
+ }
+ }
+ for (const NamedDecl *ND :
+ SemaSYCLRef.SemaRef.LookupAssignmentOperators(RD)) {
+ auto *MD = dyn_cast<CXXMethodDecl>(ND);
+ if (MD &&
+ (MD->isCopyAssignmentOperator() ||
+ MD->isMoveAssignmentOperator()) &&
+ !MD->isDeleted() && !MD->isIneligibleOrNotSelected()) {
+ if (MD->getAccess() != AS_public) {
+ SemaSYCLRef.Diag(MD->getLocation(),
+ diag::warn_sycl_device_copyable_smf_not_public)
+ << Type
+ << (MD->isCopyAssignmentOperator()
+ ? diag::EligibleSMFKind::CopyAssign
+ : diag::EligibleSMFKind::MoveAssign);
+ emitObjectAccessPathNotes();
+ }
+ HasEligibleSMF = true;
+ }
+ }
+ // * Type T has at least one eligible copy constructor, move
+ // constructor,
+ // copy assignment operator, or move assignment operator;
+ if (!HasEligibleSMF) {
+ SemaSYCLRef.Diag(RD->getLocation(),
+ diag::warn_sycl_device_copyable_no_eligible_smf)
+ << Type;
+ emitObjectAccessPathNotes();
+ }
+ }
+ // Not possible to check for the following:
+ // * The effect of each eligible copy constructor, move constructor, copy
+ // assignment operator, and move assignment operator is the same as a
+ // bitwise copy of the object;
+ // * The destructor has no effect.
}
- // TODO subsequent base classes shouldn't be checked for not device copyable
- if (!markedCopyable) {
- // TODO Type is culprit but Ty is original caller: fix diagnostics to
- // include both
+ if (!MarkedCopyable) {
SemaSYCLRef.Diag(Detail.Loc,
diag::warn_sycl_kernel_param_not_device_copyable)
<< Type;
@@ -1004,7 +1038,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// Continue traversal if not enforcing device-copyability.
if (!SemaSYCLRef.isEnforcingDeviceCopyable(Detail.Loc))
return true;
- // Do not continue traversing if enforcing device-copyability.
+ // Otherwise, do not continue if enforcing device-copyability.
IsValid = false;
return false;
}
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 6327b450e80ac..213183bafd383 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -1,5 +1,7 @@
// 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
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -Wpedantic-sycl -DCHECK_PEDANTIC_SYCL -verify %s
+// RUN: %clang_cc1 -triple spirv64 -std=c++17 -fsyntax-only -fsycl-is-device -Wpedantic-sycl -DCHECK_PEDANTIC_SYCL -verify %s
namespace std {
@@ -21,10 +23,6 @@ inline constexpr bool is_trivially_copyable_v = is_trivially_copyable<T>::value;
// 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>
@@ -56,6 +54,10 @@ static_assert(std::is_trivially_copyable_v<DefinitelyCopyable>,
// Check that sycl::is_device_copyable is respected
namespace iscopyable1 {
+
+template<typename KNT, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
// Kernel entry point template definition.
template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
@@ -104,6 +106,10 @@ struct sycl::is_device_copyable<ExplicitlyDeviceCopyableWithMember>
: std::true_type {};
namespace iscopyable2 {
+
+template<typename KNT, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
void kernel_single_task(T t) {} // expected-note-re {{within parameter 't' of type '{{.*}}' declared here}}
@@ -120,3 +126,149 @@ void test() {
}
} // namespace iscopyable2
+#ifdef CHECK_PEDANTIC_SYCL
+
+struct BadDestructorDeleted {
+ // Classes with deleted destructors are still is_trivially_copyable: This
+ // copy constructor is added to make the class not trivially copyable.
+ // FIXME: SYCL spec simultaneously stipulates deleted destructors are UB
+ // for device-copyable classes while also stipulates that all
+ // is_trivially_copyable classes are device-copyable. Which is it?
+ BadDestructorDeleted(const BadDestructorDeleted &) {}
+ ~BadDestructorDeleted() = delete;
+ // expected-warning at -1 {{'BadDestructorDeleted' is explicitly marked as device copyable (sycl::is_device_copyable) but does not have a public, non-deleted destructor}}
+};
+template<>
+struct sycl::is_device_copyable<BadDestructorDeleted> : std::true_type {};
+
+struct BadDestructorPrivate {
+ // Classes with deleted destructors are still is_trivially_copyable: This
+ // copy constructor is added to make the class not trivially copyable.
+ BadDestructorPrivate(const BadDestructorPrivate &) {}
+private:
+ ~BadDestructorPrivate() {}
+ // expected-warning at -1 {{'BadDestructorPrivate' is explicitly marked as device copyable (sycl::is_device_copyable) but does not have a public, non-deleted destructor}}
+};
+template<>
+struct sycl::is_device_copyable<BadDestructorPrivate> : std::true_type {};
+
+struct PrivateCopyCtor {
+ PrivateCopyCtor() {}
+private:
+ PrivateCopyCtor(const PrivateCopyCtor &) {}
+ // expected-warning at -1 {{'PrivateCopyCtor' is explicitly marked as device copyable (sycl::is_device_copyable) but its eligible copy constructor is not public}}
+};
+template<>
+struct sycl::is_device_copyable<PrivateCopyCtor> : std::true_type {};
+
+struct PrivateMoveAssign {
+ PrivateMoveAssign() {}
+private:
+ PrivateMoveAssign& operator=(PrivateMoveAssign &&other) { return other; }
+ // expected-warning at -1 {{'PrivateMoveAssign' is explicitly marked as device copyable (sycl::is_device_copyable) but its eligible move assignment operator is not public}}
+};
+template<>
+struct sycl::is_device_copyable<PrivateMoveAssign> : std::true_type {};
+
+struct NoEligibleSMF {
+ // expected-warning at -1 {{'NoEligibleSMF' is explicitly marked as device copyable (sycl::is_device_copyable) but has no eligible copy constructor, move constructor, copy assignment operator, or move assignment operator}}
+ NoEligibleSMF() {}
+ ~NoEligibleSMF() {}
+ NoEligibleSMF(const NoEligibleSMF &) = delete;
+ NoEligibleSMF(NoEligibleSMF &&) = delete;
+ NoEligibleSMF &operator=(const NoEligibleSMF &) = delete;
+ NoEligibleSMF &operator=(NoEligibleSMF &&) = delete;
+};
+template<>
+struct sycl::is_device_copyable<NoEligibleSMF> : std::true_type {};
+
+namespace pedanticsycl {
+
+// Custom sycl_kernel_launch that forwards its arguments directly, preventing
+// passing by value and the resulting additional copy + decay/deletion.
+// Although technically incorrect, this version of sycl_kernel_launch allows us
+// to pass "objects" without triggering its destructor (that we delete for test
+// purposes).
+template<typename KNT, typename... Ts>
+void sycl_kernel_launch(const char *, Ts &&...) {}
+
+// Custom kernel_single_task that takes a ref instead
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T &t) {}
+// expected-note-re at -1 5{{within parameter 't' of type '{{.*}}' declared here}}
+
+// Used to obtain an T& argument without ever constructing a real T object,
+// preventing destructors (that we deleted for test purposes) from triggering.
+template <typename T> T &getRef();
+
+void test() {
+ kernel_single_task<KN<11>>(getRef<BadDestructorDeleted>());
+ // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+
+ kernel_single_task<KN<12>>(getRef<BadDestructorPrivate>());
+ // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+
+ PrivateCopyCtor c1;
+ kernel_single_task<KN<13>>(c1);
+ // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+
+ PrivateMoveAssign m1;
+ kernel_single_task<KN<14>>(m1);
+ // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+
+ kernel_single_task<KN<15>>(getRef<NoEligibleSMF>());
+ // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+}
+
+} // namespace pedanticsycl
+#endif // CHECK_PEDANTIC_SYCL
+
+// Same as previous pendanticsycl testcase, but this time checking warnings are
+// not thrown if -Wpendantic-sycl is not enabled
+#ifndef CHECK_PEDANTIC_SYCL
+
+struct BadDestructorDeleted {
+ // Classes with deleted destructors are still is_trivially_copyable: This
+ // copy constructor is added to make the class not trivially copyable.
+ BadDestructorDeleted(const BadDestructorDeleted &) {}
+ ~BadDestructorDeleted() = delete;
+ // expected-warning at -1 {{'BadDestructorDeleted' is explicitly marked as device copyable (sycl::is_device_copyable) but does not have a public, non-deleted destructor}}
+};
+template<>
+struct sycl::is_device_copyable<BadDestructorDeleted> : std::true_type {};
+
+struct PrivateCopyCtor {
+ PrivateCopyCtor() {}
+private:
+ PrivateCopyCtor(const PrivateCopyCtor &) {}
+};
+template<>
+struct sycl::is_device_copyable<PrivateCopyCtor> : std::true_type {};
+
+namespace pedanticsycl {
+
+template<typename KNT, typename... Ts>
+void sycl_kernel_launch(const char *, Ts &&...) {}
+
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T &t) {}
+// expected-note-re at -1 {{within parameter 't' of type '{{.*}}' declared here}}
+
+template <typename T> T &getRef();
+
+void test() {
+ // Destructor tests are cheap and still enabled without -Wpendantic-sycl:
+ kernel_single_task<KN<16>>(getRef<BadDestructorDeleted>());
+ // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+
+ // Eligible special member function tests are expensive, and thus should not
+ // trigger without -Wpendantic-sycl:
+ PrivateCopyCtor c2;
+ kernel_single_task<KN<17>>(c2);
+}
+
+} // namespace pedanticsycl
+
+#endif // !CHECK_PEDANTIC_SYCL
\ No newline at end of file
>From 8af7f31f88613429372a8bf488448ef46b9ed1da Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 20 Aug 2026 17:10:06 -0700
Subject: [PATCH 11/15] remove todos
---
clang/lib/Sema/SemaSYCL.cpp | 5 -----
1 file changed, 5 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3f8d05bb57291..fd4102ccf2f4d 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -880,8 +880,6 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (SemaSYCLRef.checkExplicitDeviceCopyable(Ty, BS->getBaseTypeLoc()))
return false;
return checkType(Ty) && checkDeviceCopyable(Ty);
- // TODO: if a class inherits an is_device_copyable class, does that make it
- // device copyable?
}
bool visitFieldDeclPre(const FieldDecl *FD) {
@@ -893,9 +891,6 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (SemaSYCLRef.checkExplicitDeviceCopyable(Ty, FD->getLocation()))
return false;
return checkType(Ty) && checkDeviceCopyable(Ty);
- // TODO: if a class has a field of type is_device_copyable, should I still
- // check the field type? can I guarantee e.g. copy constructors and other
- // special functions are going to be safe on the device?
}
// Returns true if subobjects should be visited and false otherwise.
>From 9121c73dab2a7def1832335dd2d3eceb5a07d42e Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 20 Aug 2026 17:11:00 -0700
Subject: [PATCH 12/15] clang-format
---
clang/lib/Sema/SemaLookup.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp
index 5811d99c0c3f0..96d70996f9cf6 100644
--- a/clang/lib/Sema/SemaLookup.cpp
+++ b/clang/lib/Sema/SemaLookup.cpp
@@ -3651,7 +3651,8 @@ DeclContext::lookup_result Sema::LookupConstructors(CXXRecordDecl *Class) {
return Class->lookup(Name);
}
-DeclContext::lookup_result Sema::LookupAssignmentOperators(CXXRecordDecl *Class) {
+DeclContext::lookup_result
+Sema::LookupAssignmentOperators(CXXRecordDecl *Class) {
// If the implicit copy or move assignment operators have not yet been
// declared, do so now.
if (CanDeclareSpecialMemberFunction(Class)) {
>From e886528a36ea9078cb531396379bf3a5cd6943b7 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 20 Aug 2026 22:25:38 -0700
Subject: [PATCH 13/15] Solve fixmes
---
clang/lib/Sema/SemaSYCL.cpp | 121 +++++++++---------
.../sycl-kernel-param-is-device-copyable.cpp | 15 ---
2 files changed, 60 insertions(+), 76 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index fd4102ccf2f4d..195d879eaef97 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -919,11 +919,9 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
return false;
}
- // TODO Make sure I'm descending into Lambdas properly
return true;
}
- // TODO should these results be cached?
bool checkDeviceCopyable(QualType Ty) {
auto DirectParent = ObjectAccessPath.back();
DiagDetails Detail = getObjectAccessDiagDetails(DirectParent);
@@ -933,16 +931,6 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent))
Type = Ty->getPointeeType();
- // No need to lookup anything if trivially copyable already
- if (Type.isTriviallyCopyableType(SemaSYCLRef.getASTContext()))
- return true;
- // FIXME: isTriviallyCopyableType allows explicitly deleted destructors,
- // but the SYCL spec stipulates that deleted destructors on an explicitly-
- // marked device-copyable class is UB.
- // This shortcut path ends up resulting in missed -Wpendantic-sycl deleted
- // destructor warnings. Consider moving -Wpendantic-sycl deleted destructor
- // test before this shortcut.
-
bool MarkedCopyable =
SemaSYCLRef.checkExplicitDeviceCopyable(Type, Detail.Loc);
@@ -953,10 +941,16 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
if (RD && RD->isLambda())
return true;
+ // These SMF checks are not cheap and should only be enabled with
+ // -Wpendantic-sycl.
+ DiagnosticsEngine &Diags = SemaSYCLRef.getDiagnostics();
+ bool CheckSMFs = !Diags.isIgnored(
+ diag::warn_sycl_device_copyable_smf_not_public, Detail.Loc);
+
// Checking SYCL 2020 3.13.1, when explicitly declaring certain class types
// as device copyable:
- if (MarkedCopyable && RD) {
- // * Type T has a public non-deleted destructor; and
+ if (CheckSMFs && MarkedCopyable && RD) {
+ // * Type T has a public non-deleted destructor;
CXXDestructorDecl *DD = SemaSYCLRef.SemaRef.LookupDestructor(RD);
if (!DD || DD->isDeleted() || DD->getAccess() != AS_public) {
SemaSYCLRef.Diag(DD ? DD->getLocation() : RD->getLocation(),
@@ -964,58 +958,63 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
<< Type;
emitObjectAccessPathNotes();
}
+ }
- DiagnosticsEngine &Diags = SemaSYCLRef.getDiagnostics();
- bool CheckSMFEligible = !Diags.isIgnored(
- diag::warn_sycl_device_copyable_smf_not_public, Detail.Loc);
- if (CheckSMFEligible) {
- // * Each eligible copy constructor, move constructor, copy assignment
- // operator, and move assignment operator is public;
- bool HasEligibleSMF = false;
- for (const NamedDecl *ND : SemaSYCLRef.SemaRef.LookupConstructors(RD)) {
- auto *CD = dyn_cast<CXXConstructorDecl>(ND);
- if (CD && CD->isCopyOrMoveConstructor() && !CD->isDeleted() &&
- !CD->isIneligibleOrNotSelected()) {
- if (CD->getAccess() != AS_public) {
- SemaSYCLRef.Diag(CD->getLocation(),
- diag::warn_sycl_device_copyable_smf_not_public)
- << Type
- << (CD->isCopyConstructor()
- ? diag::EligibleSMFKind::CopyCtor
- : diag::EligibleSMFKind::MoveCtor);
- emitObjectAccessPathNotes();
- }
- HasEligibleSMF = true;
+ // SYCL 2020 3.13.1: trivially copyable implies device-copyability.
+ // However, due to Clang not implementing DR 1734, Clang treats classes
+ // with deleted destructors as trivially copyable. Thus, checking for
+ // is_trivially_copyable must happen after deleted destructors are checked.
+ // FIXME: DR1734
+ if (Type.isTriviallyCopyableType(SemaSYCLRef.getASTContext()))
+ return true;
+
+ if (CheckSMFs && MarkedCopyable && RD) {
+ // * Each eligible copy constructor, move constructor, copy assignment
+ // operator, and move assignment operator is public;
+ bool HasEligibleSMF = false;
+ for (const NamedDecl *ND : SemaSYCLRef.SemaRef.LookupConstructors(RD)) {
+ auto *CD = dyn_cast<CXXConstructorDecl>(ND);
+ if (CD && CD->isCopyOrMoveConstructor() && !CD->isDeleted() &&
+ !CD->isIneligibleOrNotSelected()) {
+ if (CD->getAccess() != AS_public) {
+ SemaSYCLRef.Diag(CD->getLocation(),
+ diag::warn_sycl_device_copyable_smf_not_public)
+ << Type
+ << (CD->isCopyConstructor()
+ ? diag::EligibleSMFKind::CopyCtor
+ : diag::EligibleSMFKind::MoveCtor);
+ emitObjectAccessPathNotes();
}
+ HasEligibleSMF = true;
}
- for (const NamedDecl *ND :
- SemaSYCLRef.SemaRef.LookupAssignmentOperators(RD)) {
- auto *MD = dyn_cast<CXXMethodDecl>(ND);
- if (MD &&
- (MD->isCopyAssignmentOperator() ||
- MD->isMoveAssignmentOperator()) &&
- !MD->isDeleted() && !MD->isIneligibleOrNotSelected()) {
- if (MD->getAccess() != AS_public) {
- SemaSYCLRef.Diag(MD->getLocation(),
- diag::warn_sycl_device_copyable_smf_not_public)
- << Type
- << (MD->isCopyAssignmentOperator()
- ? diag::EligibleSMFKind::CopyAssign
- : diag::EligibleSMFKind::MoveAssign);
- emitObjectAccessPathNotes();
- }
- HasEligibleSMF = true;
+ }
+ for (const NamedDecl *ND :
+ SemaSYCLRef.SemaRef.LookupAssignmentOperators(RD)) {
+ auto *MD = dyn_cast<CXXMethodDecl>(ND);
+ if (MD &&
+ (MD->isCopyAssignmentOperator() ||
+ MD->isMoveAssignmentOperator()) &&
+ !MD->isDeleted() && !MD->isIneligibleOrNotSelected()) {
+ if (MD->getAccess() != AS_public) {
+ SemaSYCLRef.Diag(MD->getLocation(),
+ diag::warn_sycl_device_copyable_smf_not_public)
+ << Type
+ << (MD->isCopyAssignmentOperator()
+ ? diag::EligibleSMFKind::CopyAssign
+ : diag::EligibleSMFKind::MoveAssign);
+ emitObjectAccessPathNotes();
}
+ HasEligibleSMF = true;
}
- // * Type T has at least one eligible copy constructor, move
- // constructor,
- // copy assignment operator, or move assignment operator;
- if (!HasEligibleSMF) {
- SemaSYCLRef.Diag(RD->getLocation(),
- diag::warn_sycl_device_copyable_no_eligible_smf)
- << Type;
- emitObjectAccessPathNotes();
- }
+ }
+ // * Type T has at least one eligible copy constructor, move
+ // constructor,
+ // copy assignment operator, or move assignment operator;
+ if (!HasEligibleSMF) {
+ SemaSYCLRef.Diag(RD->getLocation(),
+ diag::warn_sycl_device_copyable_no_eligible_smf)
+ << Type;
+ emitObjectAccessPathNotes();
}
// Not possible to check for the following:
// * The effect of each eligible copy constructor, move constructor, copy
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 213183bafd383..c0ed5424c5c81 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-is-device-copyable.cpp
@@ -84,12 +84,6 @@ void test() {
// expected-error at -1 {{'NotTriviallyCopyable' is not device copyable (sycl::is_device_copyable) and cannot be used as a kernel parameter}}
// expected-note-re at -2 {{in instantiation of function template specialization 'iscopyable1::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
// expected-note at -3 {{within capture 'b' of lambda expression here}}
-
- auto notCopyableLambda = [](NotTriviallyCopyable notCopyable) { (void) notCopyable; };
- kernel_single_task<KN<7>>(notCopyableLambda);
- kernel_single_task<KN<8>>([=](NotTriviallyCopyable NC) { notCopyableLambda(NC); });
- // TODO this isn't firing; shouldn't this create an error
-
}
} // namespace iscopyable1
@@ -131,9 +125,6 @@ void test() {
struct BadDestructorDeleted {
// Classes with deleted destructors are still is_trivially_copyable: This
// copy constructor is added to make the class not trivially copyable.
- // FIXME: SYCL spec simultaneously stipulates deleted destructors are UB
- // for device-copyable classes while also stipulates that all
- // is_trivially_copyable classes are device-copyable. Which is it?
BadDestructorDeleted(const BadDestructorDeleted &) {}
~BadDestructorDeleted() = delete;
// expected-warning at -1 {{'BadDestructorDeleted' is explicitly marked as device copyable (sycl::is_device_copyable) but does not have a public, non-deleted destructor}}
@@ -233,7 +224,6 @@ struct BadDestructorDeleted {
// copy constructor is added to make the class not trivially copyable.
BadDestructorDeleted(const BadDestructorDeleted &) {}
~BadDestructorDeleted() = delete;
- // expected-warning at -1 {{'BadDestructorDeleted' is explicitly marked as device copyable (sycl::is_device_copyable) but does not have a public, non-deleted destructor}}
};
template<>
struct sycl::is_device_copyable<BadDestructorDeleted> : std::true_type {};
@@ -254,17 +244,12 @@ void sycl_kernel_launch(const char *, Ts &&...) {}
template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
void kernel_single_task(T &t) {}
-// expected-note-re at -1 {{within parameter 't' of type '{{.*}}' declared here}}
template <typename T> T &getRef();
void test() {
- // Destructor tests are cheap and still enabled without -Wpendantic-sycl:
kernel_single_task<KN<16>>(getRef<BadDestructorDeleted>());
- // expected-note-re at -1 {{in instantiation of function template specialization 'pedanticsycl::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
- // Eligible special member function tests are expensive, and thus should not
- // trigger without -Wpendantic-sycl:
PrivateCopyCtor c2;
kernel_single_task<KN<17>>(c2);
}
>From 6d8eab3b2a754999c1cdec41a02ef17e99b7bad0 Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Thu, 20 Aug 2026 22:29:18 -0700
Subject: [PATCH 14/15] clang-format
---
clang/lib/Sema/SemaSYCL.cpp | 15 +++++++--------
1 file changed, 7 insertions(+), 8 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 195d879eaef97..d08be1e4a8b5e 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -978,26 +978,25 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
!CD->isIneligibleOrNotSelected()) {
if (CD->getAccess() != AS_public) {
SemaSYCLRef.Diag(CD->getLocation(),
- diag::warn_sycl_device_copyable_smf_not_public)
+ diag::warn_sycl_device_copyable_smf_not_public)
<< Type
- << (CD->isCopyConstructor()
- ? diag::EligibleSMFKind::CopyCtor
- : diag::EligibleSMFKind::MoveCtor);
+ << (CD->isCopyConstructor() ? diag::EligibleSMFKind::CopyCtor
+ : diag::EligibleSMFKind::MoveCtor);
emitObjectAccessPathNotes();
}
HasEligibleSMF = true;
}
}
for (const NamedDecl *ND :
- SemaSYCLRef.SemaRef.LookupAssignmentOperators(RD)) {
+ SemaSYCLRef.SemaRef.LookupAssignmentOperators(RD)) {
auto *MD = dyn_cast<CXXMethodDecl>(ND);
if (MD &&
(MD->isCopyAssignmentOperator() ||
- MD->isMoveAssignmentOperator()) &&
+ MD->isMoveAssignmentOperator()) &&
!MD->isDeleted() && !MD->isIneligibleOrNotSelected()) {
if (MD->getAccess() != AS_public) {
SemaSYCLRef.Diag(MD->getLocation(),
- diag::warn_sycl_device_copyable_smf_not_public)
+ diag::warn_sycl_device_copyable_smf_not_public)
<< Type
<< (MD->isCopyAssignmentOperator()
? diag::EligibleSMFKind::CopyAssign
@@ -1012,7 +1011,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// copy assignment operator, or move assignment operator;
if (!HasEligibleSMF) {
SemaSYCLRef.Diag(RD->getLocation(),
- diag::warn_sycl_device_copyable_no_eligible_smf)
+ diag::warn_sycl_device_copyable_no_eligible_smf)
<< Type;
emitObjectAccessPathNotes();
}
>From 62c40c1bfb76d384034f5fdc1ef0374a2d8c096b Mon Sep 17 00:00:00 2001
From: "Li, Ian" <ian.li at intel.com>
Date: Fri, 21 Aug 2026 09:01:07 -0700
Subject: [PATCH 15/15] Remove unneccessary code
---
clang/include/clang/Sema/SemaSYCL.h | 3 ---
clang/lib/Sema/SemaSYCL.cpp | 32 +++++++++++------------------
2 files changed, 12 insertions(+), 23 deletions(-)
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index e36868d21f9d9..6010246253ded 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -106,9 +106,6 @@ class SemaSYCL : public SemaBase {
/// sycl::is_device_copyable<Ty>
bool checkExplicitDeviceCopyable(QualType Ty, SourceLocation Loc);
- // TODO: perhaps a good idea to add a normal checkDeviceCopyable as well, and
- // not just checking if something's explicitly device-copyable.
-
/// Returns true if sycl-device-copyable diagnostics has not been downgraded
/// from an error via -Wsycl-device-copyable or -Wno-sycl-device-copyable.
/// 'Loc' is needed in case pragmas are used to set -Wsycl-device-copyable or
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index d08be1e4a8b5e..a30f17ea58432 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -808,19 +808,13 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
const FieldDecl *>;
SmallVector<ObjectAccess, 4> ObjectAccessPath;
- struct DiagDetails {
- QualType Type;
- SourceLocation Loc;
- };
-
- // Return diagnostics info for an 'ObjectAccess' stored on ObjectAccessPath
- DiagDetails getObjectAccessDiagDetails(ObjectAccess o) {
- if (auto *PVD = dyn_cast<const ParmVarDecl *>(o))
- return {PVD->getType(), PVD->getLocation()};
- if (auto *FD = dyn_cast<const FieldDecl *>(o))
- return {FD->getType(), FD->getLocation()};
- if (auto *BS = dyn_cast<const CXXBaseSpecifier *>(o))
- return {BS->getType(), BS->getBaseTypeLoc()};
+ SourceLocation getObjectAccessLoc(ObjectAccess O) {
+ if (auto *PVD = dyn_cast<const ParmVarDecl *>(O))
+ return PVD->getLocation();
+ if (auto *FD = dyn_cast<const FieldDecl *>(O))
+ return FD->getLocation();
+ if (auto *BS = dyn_cast<const CXXBaseSpecifier *>(O))
+ return BS->getBaseTypeLoc();
llvm_unreachable("Unexpected type in ObjectAccess");
}
@@ -924,15 +918,14 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
bool checkDeviceCopyable(QualType Ty) {
auto DirectParent = ObjectAccessPath.back();
- DiagDetails Detail = getObjectAccessDiagDetails(DirectParent);
+ SourceLocation Loc = getObjectAccessLoc(DirectParent);
// Since references are allowed as direct kernel parameters, we need to
// explicitly check the referenced type:
QualType Type = Ty;
if (Ty->isReferenceType() && isa<const ParmVarDecl *>(DirectParent))
Type = Ty->getPointeeType();
- bool MarkedCopyable =
- SemaSYCLRef.checkExplicitDeviceCopyable(Type, Detail.Loc);
+ bool MarkedCopyable = SemaSYCLRef.checkExplicitDeviceCopyable(Type, Loc);
CXXRecordDecl *RD = Type->getAsCXXRecordDecl();
// Set all lambdas as copyable: Future traversal deeper into the lambda
@@ -945,7 +938,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
// -Wpendantic-sycl.
DiagnosticsEngine &Diags = SemaSYCLRef.getDiagnostics();
bool CheckSMFs = !Diags.isIgnored(
- diag::warn_sycl_device_copyable_smf_not_public, Detail.Loc);
+ diag::warn_sycl_device_copyable_smf_not_public, Loc);
// Checking SYCL 2020 3.13.1, when explicitly declaring certain class types
// as device copyable:
@@ -1023,13 +1016,12 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
}
if (!MarkedCopyable) {
- SemaSYCLRef.Diag(Detail.Loc,
- diag::warn_sycl_kernel_param_not_device_copyable)
+ SemaSYCLRef.Diag(Loc, diag::warn_sycl_kernel_param_not_device_copyable)
<< Type;
emitObjectAccessPathNotes();
// Continue traversal if not enforcing device-copyability.
- if (!SemaSYCLRef.isEnforcingDeviceCopyable(Detail.Loc))
+ if (!SemaSYCLRef.isEnforcingDeviceCopyable(Loc))
return true;
// Otherwise, do not continue if enforcing device-copyability.
IsValid = false;
More information about the cfe-commits
mailing list