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