From 34a11ca5577b49d74daa67ed40b866cefed8b45c Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 27 May 2025 19:10:48 -0700 Subject: [PATCH 01/12] Provide supports for decomposable structs --- clang/lib/Sema/SemaSYCL.cpp | 206 +++++++++++------- .../experimental/free_function_traits.hpp | 10 + sycl/include/sycl/handler.hpp | 4 +- 3 files changed, 146 insertions(+), 74 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9a30b3e693ec2..7c2f2762e2133 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -8,6 +8,7 @@ // This implements Semantic Analysis for SYCL constructs. //===----------------------------------------------------------------------===// +#include #include "clang/Sema/SemaSYCL.h" #include "TreeTransform.h" #include "clang/AST/AST.h" @@ -1387,13 +1388,13 @@ class KernelObjVisitor { template void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &... Handlers) { + HandlerTys &...Handlers) { (void)std::initializer_list{ (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; VisitRecordHelper(Wrapper, Wrapper->bases(), Handlers...); - VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...); - (void)std::initializer_list{ - (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; + VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...), + (void)std::initializer_list{ + (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; } template @@ -1499,7 +1500,9 @@ class KernelObjVisitor { void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, QualType FieldTy, HandlerTys &... Handlers) { if (isSyclSpecialType(FieldTy, SemaSYCLRef)) + { FieldTy->dump(); KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy); +} else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); @@ -1526,9 +1529,12 @@ class KernelObjVisitor { void visitParam(ParmVarDecl *Param, QualType ParamTy, HandlerTys &...Handlers) { if (isSyclSpecialType(ParamTy, SemaSYCLRef)) + {ParamTy->dump(); KP_FOR_EACH(handleSyclSpecialType, Param, ParamTy); +} else if (ParamTy->isStructureOrClassType()) { if (KP_FOR_EACH(handleStructType, Param, ParamTy)) { + ParamTy->dump(); CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); visitRecord(nullptr, Param, RD, ParamTy, Handlers...); } @@ -1607,8 +1613,12 @@ class KernelObjVisitor { template void VisitFunctionParameters(FunctionDecl *FreeFunc, HandlerTys &...Handlers) { - for (ParmVarDecl *Param : FreeFunc->parameters()) + for (ParmVarDecl *Param : FreeFunc->parameters()) { +std::cout << "starting!" << std::endl; +Param->getType()->dump(); visitParam(Param, Param->getType(), Handlers...); +std::cout << "ending!" << std::endl; +} } #undef KF_FOR_EACH @@ -1731,10 +1741,6 @@ class SyclKernelFieldHandlerBase { virtual ~SyclKernelFieldHandlerBase() = default; }; - -// A class to act as the direct base for all the SYCL OpenCL Kernel construction -// tasks that contains a reference to Sema (and potentially any other -// universally required data). class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase { protected: SemaSYCL &SemaSYCLRef; @@ -1818,7 +1824,11 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, // If this container requires decomposition, we have to visit it as // 'complex', so all handlers are called in this case with the 'complex' // case. + //RecordTy->dump(); visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); + // 'complex', so all handlers are called in this case with the 'complex' + // case. + //RecordTy->dump(); } else if (AnyTrue:: Value) { // We are currently in PointerHandler visitor. @@ -2141,31 +2151,14 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - // TODO manipulate struct depth once special types are supported for free - // function kernels. - // ++StructFieldDepth; + ++StructFieldDepth; return true; } bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType ParamTy) final { - // TODO manipulate struct depth once special types are supported for free - // function kernels. - // --StructFieldDepth; - // TODO We don't yet support special types and therefore structs that - // require decomposition and leaving/entering. Diagnose for better user - // experience. - CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); - if (RD->hasAttr()) { - Diag.Report(PD->getLocation(), - diag::err_bad_kernel_param_type) - << ParamTy; - Diag.Report(PD->getLocation(), - diag::note_free_function_kernel_param_type_not_supported) - << ParamTy; - IsInvalid = true; - } - return isValid(); + --StructFieldDepth; + return true; } bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, @@ -2269,8 +2262,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool handleSyclSpecialType(ParmVarDecl *, QualType) final { - // TODO We don't support special types in free function kernel parameters, - // but track them to diagnose the case properly. CollectionStack.back() = true; return true; } @@ -2542,7 +2533,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType ParamTy) final { // TODO - unsupportedFreeFunctionParamType(); return true; } @@ -2563,7 +2553,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType ParamTy) final { // TODO - unsupportedFreeFunctionParamType(); return true; } @@ -2660,7 +2649,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool handleNonDecompStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType ParamTy) final { // TODO - unsupportedFreeFunctionParamType(); return true; } @@ -2694,6 +2682,9 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { class SyclKernelDeclCreator : public SyclKernelFieldHandler { FunctionDecl *KernelDecl = nullptr; llvm::SmallVector Params; + // Holds the last handled kernel struct parameter that contains a special type. + // Set in the enterStruct functions. + ParmVarDecl * CurrentStruct; Sema::ContextRAII FuncContext; // Holds the last handled field's first parameter. This doesn't store an // iterator as push_back invalidates iterators. @@ -2711,6 +2702,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { addParam(newParamDesc, ParamTy); } + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { // TODO: There is no name for the base available, but duplicate names are // seemingly already possible, so we'll give them all the same name for now. @@ -2798,7 +2790,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { SourceLocation Loc) { handleAccessorPropertyList(Params.back(), RecordDecl, Loc); - // If "accessor" type check if read only + // If "accessor" type check if read only if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::accessor)) { // Get access mode of accessor. const auto *AccessorSpecializationDecl = @@ -2824,6 +2816,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // lambda kernel by taking the value ParmVarDecl or FieldDecl respectively. template bool handleSpecialType(ParentDecl *decl, QualType Ty) { +std::cout << "Important one!" << std::endl; +Ty->dump(); const auto *RD = Ty->getAsCXXRecordDecl(); assert(RD && "The type must be a RecordDecl"); llvm::StringLiteral MethodName = @@ -2837,7 +2831,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // (if any). size_t ParamIndex = Params.size(); for (const ParmVarDecl *Param : InitMethod->parameters()) { - QualType ParamTy = Param->getType(); + QualType ParamTy = Param->getType(); // For lambda kernels the arguments to the OpenCL kernel are named // based on the position they have as fields in the definition of the // special type structure i.e __arg_field1, __arg_field2 and so on. @@ -2863,6 +2857,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { handleAccessorType(Ty, RD, decl->getBeginLoc()); } LastParamIndex = ParamIndex; + std::cout << LastParamIndex << std::endl; return true; } @@ -2956,6 +2951,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { SYCLKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext())); SemaSYCLRef.addSyclDeviceDecl(KernelDecl); + //KernelDecl->dump(); } bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { @@ -2963,9 +2959,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } - bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - // TODO - // ++StructDepth; + bool enterStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType Ty) final { + ++StructDepth; + //StringRef Name = "_arg_struct"; + //addParam(Name, Ty); + //CurrentStruct = Params.back(); return true; } @@ -2975,8 +2973,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - // TODO - // --StructDepth; + --StructDepth; return true; } @@ -2992,6 +2989,15 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return true; } + bool handleStructType(ParmVarDecl *PD, QualType Ty) final { + StringRef Name = "_arg_struct"; + addParam(Name, Ty); + CurrentStruct = Params.back(); + return true; + } + + bool handleStructType(FieldDecl *, QualType) final { return true; } + bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, QualType FieldTy) final { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); @@ -3166,6 +3172,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return ArrayRef(std::begin(Params) + LastParamIndex, std::end(Params)); } + ParmVarDecl *getParentStructForCurrentField() { return CurrentStruct; } }; // This Visitor traverses the AST of the function with @@ -3619,8 +3626,11 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { SourceLocation LL = NewBody ? NewBody->getBeginLoc() : SourceLocation(); SourceLocation LR = NewBody ? NewBody->getEndLoc() : SourceLocation(); - return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, + CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, + FPOptionsOverride(), LL, LR)->dumpPretty(SemaSYCLRef.getASTContext()); +return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, FPOptionsOverride(), LL, LR); + } void annotateHierarchicalParallelismAPICalls() { @@ -4342,16 +4352,14 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; + llvm::SmallVector CurrentStructs; FunctionDecl *FreeFunc = nullptr; SourceLocation FreeFunctionSrcLoc; // Free function source location. llvm::SmallVector ArgExprs; - // Creates a DeclRefExpr to the ParmVar that represents the current free - // function parameter. - Expr *createParamReferenceExpr() { - ParmVarDecl *FreeFunctionParameter = - DeclCreator.getParamVarDeclsForCurrentField()[0]; - + // Creates a DeclRefExpr to the ParmVar that represents an arbitrary + // free function parameter + Expr *createParamReferenceExpr(ParmVarDecl *FreeFunctionParameter) { QualType FreeFunctionParamType = FreeFunctionParameter->getOriginalType(); Expr *DRE = SemaSYCLRef.SemaRef.BuildDeclRefExpr( FreeFunctionParameter, FreeFunctionParamType, VK_LValue, @@ -4360,6 +4368,14 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { return DRE; } + // Creates a DeclRefExpr to the ParmVar that represents the current free + // function parameter. + Expr *createParamReferenceExpr() { + ParmVarDecl *FreeFunctionParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + return createParamReferenceExpr(FreeFunctionParameter); + } + // Creates a DeclRefExpr to the ParmVar that represents the current pointer // parameter. Expr *createPointerParamReferenceExpr(QualType PointerTy) { @@ -4416,6 +4432,7 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { DRE = createReinterpretCastExpr( createGetAddressOf(DRE), SemaSYCLRef.getASTContext().getPointerType( OrigFunctionParameter->getType())); + DRE = createDerefOp(DRE); } @@ -4450,8 +4467,12 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { auto CallExpr = CallExpr::Create(Context, Fn, ArgExprs, ResultTy, VK, FreeFunctionSrcLoc, FPOptionsOverride()); BodyStmts.push_back(CallExpr); +CompoundStmt::Create(Context, BodyStmts, FPOptionsOverride(), {}, + {})->dumpPretty(Context); + return CompoundStmt::Create(Context, BodyStmts, FPOptionsOverride(), {}, {}); + } MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { @@ -4468,15 +4489,17 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { void createSpecialMethodCall(const CXXRecordDecl *RD, StringRef MethodName, Expr *MemberBaseExpr, SmallVectorImpl &AddTo) { - CXXMethodDecl *Method = getMethodByName(RD, MethodName); +CXXMethodDecl *Method = getMethodByName(RD, MethodName); if (!Method) return; unsigned NumParams = Method->getNumParams(); llvm::SmallVector ParamDREs(NumParams); llvm::ArrayRef KernelParameters = DeclCreator.getParamVarDeclsForCurrentField(); + //std::cout << KernelParameters.size() << std::endl; for (size_t I = 0; I < NumParams; ++I) { QualType ParamType = KernelParameters[I]->getOriginalType(); + //ParamType->dump(); ParamDREs[I] = SemaSYCLRef.SemaRef.BuildDeclRefExpr( KernelParameters[I], ParamType, VK_LValue, FreeFunctionSrcLoc); } @@ -4495,7 +4518,7 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { public: static constexpr const bool VisitInsideSimpleContainers = false; - + FreeFunctionKernelBodyCreator(SemaSYCL &S, SyclKernelDeclCreator &DC, FunctionDecl *FF) : SyclKernelFieldHandler(S), DeclCreator(DC), FreeFunc(FF), @@ -4506,9 +4529,20 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { DeclCreator.setBody(KernelBody); } - bool handleSyclSpecialType(FieldDecl *FD, QualType Ty) final { - // TODO - unsupportedFreeFunctionParamType(); + bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final { + // Being inside this function means there is a struct parameter to the free + // function kernel that contains a special type. +std::cout << "Body!" << std::endl; +FieldTy->dump(); + ParmVarDecl *ParentStruct = DeclCreator.getParentStructForCurrentField(); + // special_type_wrapper_map[ParentStruct->getType()] = true; + Expr *Base = createParamReferenceExpr(ParentStruct); + for (const auto &child : CurrentStructs) { + Base = buildMemberExpr(Base, child); + } + MemberExpr *MemberAccess = buildMemberExpr(Base, FD); + createSpecialMethodCall(FieldTy->getAsCXXRecordDecl(), InitMethodName, + MemberAccess, BodyStmts); return true; } @@ -4527,6 +4561,8 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { // wgm.__init(arg); // user_kernel(some arguments..., wgm, some arguments...); // } + std::cout << "Body!" << std::endl; + ParamTy->dump(); const auto *RecordDecl = ParamTy->getAsCXXRecordDecl(); AccessSpecifier DefaultConstructorAccess; auto DefaultConstructor = @@ -4559,8 +4595,8 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { BodyStmts.push_back(DS); Expr *MemberBaseExpr = SemaSYCLRef.SemaRef.BuildDeclRefExpr( SpecialObjectClone, ParamTy, VK_PRValue, FreeFunctionSrcLoc); - createSpecialMethodCall(RecordDecl, InitMethodName, MemberBaseExpr, - BodyStmts); + createSpecialMethodCall(RecordDecl, InitMethodName, MemberBaseExpr, + BodyStmts); ArgExprs.push_back(MemberBaseExpr); return true; } @@ -4636,26 +4672,24 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { } bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - // TODO - unsupportedFreeFunctionParamType(); + CurrentStructs.push_back(FD); return true; } - bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - // TODO - unsupportedFreeFunctionParamType(); + bool enterStruct(const CXXRecordDecl *RD, ParmVarDecl *PD, + QualType Ty) final { return true; } bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - // TODO - unsupportedFreeFunctionParamType(); + CurrentStructs.pop_back(); return true; } bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - // TODO - unsupportedFreeFunctionParamType(); + ParmVarDecl *ParentStruct = DeclCreator.getParentStructForCurrentField(); + ArgExprs.push_back(SemaSYCLRef.SemaRef.BuildDeclRefExpr( + ParentStruct, ParentStruct->getType(), VK_PRValue, FreeFunctionSrcLoc)); return true; } @@ -4700,6 +4734,11 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { unsupportedFreeFunctionParamType(); return true; } + FieldDecl *getCurrentStruct() { + assert(CurrentStructs.size() && + "Current free function parameter is not inside a structure!"); + return CurrentStructs.back(); + } }; // Kernels are only the unnamed-lambda feature if the feature is enabled, AND @@ -4979,7 +5018,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { // TODO - unsupportedFreeFunctionParamType(); return true; } @@ -4991,7 +5029,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { // TODO - unsupportedFreeFunctionParamType(); return true; } @@ -5488,22 +5525,25 @@ void SemaSYCL::constructFreeFunctionKernel(FunctionDecl *FD, StringRef NameStr) { if (!checkAndAddRegisteredKernelName(*this, FD, NameStr)) return; - SyclKernelArgsSizeChecker argsSizeChecker(*this, FD->getLocation(), false /*IsSIMDKernel*/); SyclKernelDeclCreator kernel_decl(*this, FD->getLocation(), FD->isInlined(), false /*IsSIMDKernel */, FD); - FreeFunctionKernelBodyCreator kernel_body(*this, kernel_decl, FD); - SyclKernelIntHeaderCreator int_header(*this, getSyclIntegrationHeader(), FD->getType(), FD); - SyclKernelIntFooterCreator int_footer(*this, getSyclIntegrationFooter()); KernelObjVisitor Visitor{*this}; - Visitor.VisitFunctionParameters(FD, argsSizeChecker, kernel_decl, kernel_body, - int_header, int_footer); + Visitor.VisitFunctionParameters(FD, argsSizeChecker); + +Visitor.VisitFunctionParameters(FD, kernel_decl); + +Visitor.VisitFunctionParameters(FD, kernel_body); + +Visitor.VisitFunctionParameters(FD, int_header); + +Visitor.VisitFunctionParameters(FD, int_footer); assert(getKernelFDPairs().back().first == FD && "OpenCL Kernel not found for free function entry"); @@ -6984,6 +7024,26 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { } } ParmListWithNamesOstream.flush(); + for (ParmVarDecl *Param : K.SyclKernel->parameters()) { + // if (FreeFunctionKernelBodyCreator::isSpecialTypeWrapper( + // Param->getType())) { + // this is a struct that contains a special type so its neither a + // special type nor a trivially copyable type. We therefore need to + // explicitly communicate to the runtime that this argument should be + // allowed as a free function kernel argument. We do this by defining + // a certain trait recognized by the runtime to be true. + O << "template <>\n"; + O << "struct " + "sycl::ext::oneapi::experimental::detail::is_explicitly_allowed_" + "arg<"; + Policy.SuppressTagKeyword = true; + + Param->getType().print(O, Policy); + Policy.SuppressTagKeyword = false; + O << "> {\n"; + O << " static constexpr bool value = true;\n};\n"; + //} + } FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate(); Policy.PrintCanonicalTypes = false; Policy.SuppressDefinition = true; @@ -7720,7 +7780,7 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD, OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap); Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get(); - OFD->setBody(OFDBody); +OFD->setBody(OFDBody); OFD->setNothrow(); Stmt *NewBody = new (getASTContext()) SYCLKernelCallStmt(Body, OFD); diff --git a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp index 2b5d1f4190d21..0ca1c234c9070 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp @@ -44,6 +44,16 @@ template struct is_kernel { template inline constexpr bool is_kernel_v = is_kernel::value; +namespace detail { +template struct is_explicitly_allowed_arg { + static constexpr bool value = false; +}; + +template +inline constexpr bool is_explicitly_allowed_arg_v = + is_explicitly_allowed_arg::value; + +} // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index d7b304a130c83..54b093f05000a 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -35,6 +35,7 @@ #include #include #include +#include #include #include #include @@ -1766,7 +1767,8 @@ class __SYCL_EXPORT handler { || (!is_same_type::value && std::is_pointer_v>) // USM || is_same_type::value // Interop - || is_same_type::value; // Stream + || is_same_type::value // Stream + || ext::oneapi::experimental::detail::is_explicitly_allowed_arg>::value; }; /// Sets argument for OpenCL interoperability kernels. From c9680a244bdd1f4e8fb979d2032131eba90c4aa3 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 3 Jun 2025 11:46:24 -0700 Subject: [PATCH 02/12] Revert "Provide supports for decomposable structs" This reverts commit 34a11ca5577b49d74daa67ed40b866cefed8b45c. --- clang/lib/Sema/SemaSYCL.cpp | 206 +++++++----------- .../experimental/free_function_traits.hpp | 10 - sycl/include/sycl/handler.hpp | 4 +- 3 files changed, 74 insertions(+), 146 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b1a7b3cedf70f..cf64331198c91 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -8,7 +8,6 @@ // This implements Semantic Analysis for SYCL constructs. //===----------------------------------------------------------------------===// -#include #include "clang/Sema/SemaSYCL.h" #include "TreeTransform.h" #include "clang/AST/AST.h" @@ -1409,13 +1408,13 @@ class KernelObjVisitor { template void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &...Handlers) { + HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...}; VisitRecordHelper(Wrapper, Wrapper->bases(), Handlers...); - VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...), - (void)std::initializer_list{ - (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; + VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...); + (void)std::initializer_list{ + (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...}; } template @@ -1521,9 +1520,7 @@ class KernelObjVisitor { void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, QualType FieldTy, HandlerTys &... Handlers) { if (isSyclSpecialType(FieldTy, SemaSYCLRef)) - { FieldTy->dump(); KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy); -} else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); @@ -1550,12 +1547,9 @@ class KernelObjVisitor { void visitParam(ParmVarDecl *Param, QualType ParamTy, HandlerTys &...Handlers) { if (isSyclSpecialType(ParamTy, SemaSYCLRef)) - {ParamTy->dump(); KP_FOR_EACH(handleSyclSpecialType, Param, ParamTy); -} else if (ParamTy->isStructureOrClassType()) { if (KP_FOR_EACH(handleStructType, Param, ParamTy)) { - ParamTy->dump(); CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); visitRecord(nullptr, Param, RD, ParamTy, Handlers...); } @@ -1634,12 +1628,8 @@ class KernelObjVisitor { template void VisitFunctionParameters(FunctionDecl *FreeFunc, HandlerTys &...Handlers) { - for (ParmVarDecl *Param : FreeFunc->parameters()) { -std::cout << "starting!" << std::endl; -Param->getType()->dump(); + for (ParmVarDecl *Param : FreeFunc->parameters()) visitParam(Param, Param->getType(), Handlers...); -std::cout << "ending!" << std::endl; -} } #undef KF_FOR_EACH @@ -1762,6 +1752,10 @@ class SyclKernelFieldHandlerBase { virtual ~SyclKernelFieldHandlerBase() = default; }; + +// A class to act as the direct base for all the SYCL OpenCL Kernel construction +// tasks that contains a reference to Sema (and potentially any other +// universally required data). class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase { protected: SemaSYCL &SemaSYCLRef; @@ -1845,11 +1839,7 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, // If this container requires decomposition, we have to visit it as // 'complex', so all handlers are called in this case with the 'complex' // case. - //RecordTy->dump(); visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...); - // 'complex', so all handlers are called in this case with the 'complex' - // case. - //RecordTy->dump(); } else if (AnyTrue:: Value) { // We are currently in PointerHandler visitor. @@ -2172,14 +2162,31 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - ++StructFieldDepth; + // TODO manipulate struct depth once special types are supported for free + // function kernels. + // ++StructFieldDepth; return true; } bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType ParamTy) final { - --StructFieldDepth; - return true; + // TODO manipulate struct depth once special types are supported for free + // function kernels. + // --StructFieldDepth; + // TODO We don't yet support special types and therefore structs that + // require decomposition and leaving/entering. Diagnose for better user + // experience. + CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); + if (RD->hasAttr()) { + Diag.Report(PD->getLocation(), + diag::err_bad_kernel_param_type) + << ParamTy; + Diag.Report(PD->getLocation(), + diag::note_free_function_kernel_param_type_not_supported) + << ParamTy; + IsInvalid = true; + } + return isValid(); } bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS, @@ -2283,6 +2290,8 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { } bool handleSyclSpecialType(ParmVarDecl *, QualType) final { + // TODO We don't support special types in free function kernel parameters, + // but track them to diagnose the case properly. CollectionStack.back() = true; return true; } @@ -2554,6 +2563,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType ParamTy) final { // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -2574,6 +2584,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType ParamTy) final { // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -2670,6 +2681,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { bool handleNonDecompStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType ParamTy) final { // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -2703,9 +2715,6 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { class SyclKernelDeclCreator : public SyclKernelFieldHandler { FunctionDecl *KernelDecl = nullptr; llvm::SmallVector Params; - // Holds the last handled kernel struct parameter that contains a special type. - // Set in the enterStruct functions. - ParmVarDecl * CurrentStruct; Sema::ContextRAII FuncContext; // Holds the last handled field's first parameter. This doesn't store an // iterator as push_back invalidates iterators. @@ -2723,7 +2732,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { addParam(newParamDesc, ParamTy); } - void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { // TODO: There is no name for the base available, but duplicate names are // seemingly already possible, so we'll give them all the same name for now. @@ -2811,7 +2819,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { SourceLocation Loc) { handleAccessorPropertyList(Params.back(), RecordDecl, Loc); - // If "accessor" type check if read only + // If "accessor" type check if read only if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::accessor)) { // Get access mode of accessor. const auto *AccessorSpecializationDecl = @@ -2837,8 +2845,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // lambda kernel by taking the value ParmVarDecl or FieldDecl respectively. template bool handleSpecialType(ParentDecl *decl, QualType Ty) { -std::cout << "Important one!" << std::endl; -Ty->dump(); const auto *RD = Ty->getAsCXXRecordDecl(); assert(RD && "The type must be a RecordDecl"); llvm::StringLiteral MethodName = @@ -2852,7 +2858,7 @@ Ty->dump(); // (if any). size_t ParamIndex = Params.size(); for (const ParmVarDecl *Param : InitMethod->parameters()) { - QualType ParamTy = Param->getType(); + QualType ParamTy = Param->getType(); // For lambda kernels the arguments to the OpenCL kernel are named // based on the position they have as fields in the definition of the // special type structure i.e __arg_field1, __arg_field2 and so on. @@ -2878,7 +2884,6 @@ Ty->dump(); handleAccessorType(Ty, RD, decl->getBeginLoc()); } LastParamIndex = ParamIndex; - std::cout << LastParamIndex << std::endl; return true; } @@ -2972,7 +2977,6 @@ Ty->dump(); SYCLKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext())); SemaSYCLRef.addSyclDeviceDecl(KernelDecl); - //KernelDecl->dump(); } bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { @@ -2980,11 +2984,9 @@ Ty->dump(); return true; } - bool enterStruct(const CXXRecordDecl *, ParmVarDecl *PD, QualType Ty) final { - ++StructDepth; - //StringRef Name = "_arg_struct"; - //addParam(Name, Ty); - //CurrentStruct = Params.back(); + bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { + // TODO + // ++StructDepth; return true; } @@ -2994,7 +2996,8 @@ Ty->dump(); } bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - --StructDepth; + // TODO + // --StructDepth; return true; } @@ -3010,15 +3013,6 @@ Ty->dump(); return true; } - bool handleStructType(ParmVarDecl *PD, QualType Ty) final { - StringRef Name = "_arg_struct"; - addParam(Name, Ty); - CurrentStruct = Params.back(); - return true; - } - - bool handleStructType(FieldDecl *, QualType) final { return true; } - bool handleSyclSpecialType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, QualType FieldTy) final { const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); @@ -3193,7 +3187,6 @@ Ty->dump(); return ArrayRef(std::begin(Params) + LastParamIndex, std::end(Params)); } - ParmVarDecl *getParentStructForCurrentField() { return CurrentStruct; } }; // This Visitor traverses the AST of the function with @@ -3647,11 +3640,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { SourceLocation LL = NewBody ? NewBody->getBeginLoc() : SourceLocation(); SourceLocation LR = NewBody ? NewBody->getEndLoc() : SourceLocation(); - CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, - FPOptionsOverride(), LL, LR)->dumpPretty(SemaSYCLRef.getASTContext()); -return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, + return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, FPOptionsOverride(), LL, LR); - } void annotateHierarchicalParallelismAPICalls() { @@ -4373,14 +4363,16 @@ return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; llvm::SmallVector BodyStmts; - llvm::SmallVector CurrentStructs; FunctionDecl *FreeFunc = nullptr; SourceLocation FreeFunctionSrcLoc; // Free function source location. llvm::SmallVector ArgExprs; - // Creates a DeclRefExpr to the ParmVar that represents an arbitrary - // free function parameter - Expr *createParamReferenceExpr(ParmVarDecl *FreeFunctionParameter) { + // Creates a DeclRefExpr to the ParmVar that represents the current free + // function parameter. + Expr *createParamReferenceExpr() { + ParmVarDecl *FreeFunctionParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + QualType FreeFunctionParamType = FreeFunctionParameter->getOriginalType(); Expr *DRE = SemaSYCLRef.SemaRef.BuildDeclRefExpr( FreeFunctionParameter, FreeFunctionParamType, VK_LValue, @@ -4389,14 +4381,6 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { return DRE; } - // Creates a DeclRefExpr to the ParmVar that represents the current free - // function parameter. - Expr *createParamReferenceExpr() { - ParmVarDecl *FreeFunctionParameter = - DeclCreator.getParamVarDeclsForCurrentField()[0]; - return createParamReferenceExpr(FreeFunctionParameter); - } - // Creates a DeclRefExpr to the ParmVar that represents the current pointer // parameter. Expr *createPointerParamReferenceExpr(QualType PointerTy) { @@ -4453,7 +4437,6 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { DRE = createReinterpretCastExpr( createGetAddressOf(DRE), SemaSYCLRef.getASTContext().getPointerType( OrigFunctionParameter->getType())); - DRE = createDerefOp(DRE); } @@ -4488,12 +4471,8 @@ class FreeFunctionKernelBodyCreator : public SyclKernelFieldHandler { auto CallExpr = CallExpr::Create(Context, Fn, ArgExprs, ResultTy, VK, FreeFunctionSrcLoc, FPOptionsOverride()); BodyStmts.push_back(CallExpr); -CompoundStmt::Create(Context, BodyStmts, FPOptionsOverride(), {}, - {})->dumpPretty(Context); - return CompoundStmt::Create(Context, BodyStmts, FPOptionsOverride(), {}, {}); - } MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { @@ -4510,17 +4489,15 @@ CompoundStmt::Create(Context, BodyStmts, FPOptionsOverride(), {}, void createSpecialMethodCall(const CXXRecordDecl *RD, StringRef MethodName, Expr *MemberBaseExpr, SmallVectorImpl &AddTo) { -CXXMethodDecl *Method = getMethodByName(RD, MethodName); + CXXMethodDecl *Method = getMethodByName(RD, MethodName); if (!Method) return; unsigned NumParams = Method->getNumParams(); llvm::SmallVector ParamDREs(NumParams); llvm::ArrayRef KernelParameters = DeclCreator.getParamVarDeclsForCurrentField(); - //std::cout << KernelParameters.size() << std::endl; for (size_t I = 0; I < NumParams; ++I) { QualType ParamType = KernelParameters[I]->getOriginalType(); - //ParamType->dump(); ParamDREs[I] = SemaSYCLRef.SemaRef.BuildDeclRefExpr( KernelParameters[I], ParamType, VK_LValue, FreeFunctionSrcLoc); } @@ -4539,7 +4516,7 @@ CXXMethodDecl *Method = getMethodByName(RD, MethodName); public: static constexpr const bool VisitInsideSimpleContainers = false; - + FreeFunctionKernelBodyCreator(SemaSYCL &S, SyclKernelDeclCreator &DC, FunctionDecl *FF) : SyclKernelFieldHandler(S), DeclCreator(DC), FreeFunc(FF), @@ -4550,20 +4527,9 @@ CXXMethodDecl *Method = getMethodByName(RD, MethodName); DeclCreator.setBody(KernelBody); } - bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final { - // Being inside this function means there is a struct parameter to the free - // function kernel that contains a special type. -std::cout << "Body!" << std::endl; -FieldTy->dump(); - ParmVarDecl *ParentStruct = DeclCreator.getParentStructForCurrentField(); - // special_type_wrapper_map[ParentStruct->getType()] = true; - Expr *Base = createParamReferenceExpr(ParentStruct); - for (const auto &child : CurrentStructs) { - Base = buildMemberExpr(Base, child); - } - MemberExpr *MemberAccess = buildMemberExpr(Base, FD); - createSpecialMethodCall(FieldTy->getAsCXXRecordDecl(), InitMethodName, - MemberAccess, BodyStmts); + bool handleSyclSpecialType(FieldDecl *FD, QualType Ty) final { + // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -4582,8 +4548,6 @@ FieldTy->dump(); // wgm.__init(arg); // user_kernel(some arguments..., wgm, some arguments...); // } - std::cout << "Body!" << std::endl; - ParamTy->dump(); const auto *RecordDecl = ParamTy->getAsCXXRecordDecl(); AccessSpecifier DefaultConstructorAccess; auto DefaultConstructor = @@ -4616,8 +4580,8 @@ FieldTy->dump(); BodyStmts.push_back(DS); Expr *MemberBaseExpr = SemaSYCLRef.SemaRef.BuildDeclRefExpr( SpecialObjectClone, ParamTy, VK_PRValue, FreeFunctionSrcLoc); - createSpecialMethodCall(RecordDecl, InitMethodName, MemberBaseExpr, - BodyStmts); + createSpecialMethodCall(RecordDecl, InitMethodName, MemberBaseExpr, + BodyStmts); ArgExprs.push_back(MemberBaseExpr); return true; } @@ -4693,24 +4657,26 @@ FieldTy->dump(); } bool enterStruct(const CXXRecordDecl *RD, FieldDecl *FD, QualType Ty) final { - CurrentStructs.push_back(FD); + // TODO + unsupportedFreeFunctionParamType(); return true; } - bool enterStruct(const CXXRecordDecl *RD, ParmVarDecl *PD, - QualType Ty) final { + bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { + // TODO + unsupportedFreeFunctionParamType(); return true; } bool leaveStruct(const CXXRecordDecl *, FieldDecl *FD, QualType Ty) final { - CurrentStructs.pop_back(); + // TODO + unsupportedFreeFunctionParamType(); return true; } bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { - ParmVarDecl *ParentStruct = DeclCreator.getParentStructForCurrentField(); - ArgExprs.push_back(SemaSYCLRef.SemaRef.BuildDeclRefExpr( - ParentStruct, ParentStruct->getType(), VK_PRValue, FreeFunctionSrcLoc)); + // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -4755,11 +4721,6 @@ FieldTy->dump(); unsupportedFreeFunctionParamType(); return true; } - FieldDecl *getCurrentStruct() { - assert(CurrentStructs.size() && - "Current free function parameter is not inside a structure!"); - return CurrentStructs.back(); - } }; // Kernels are only the unnamed-lambda feature if the feature is enabled, AND @@ -5049,6 +5010,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool enterStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -5060,6 +5022,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { bool leaveStruct(const CXXRecordDecl *, ParmVarDecl *, QualType) final { // TODO + unsupportedFreeFunctionParamType(); return true; } @@ -5556,25 +5519,22 @@ void SemaSYCL::constructFreeFunctionKernel(FunctionDecl *FD, StringRef NameStr) { if (!checkAndAddRegisteredKernelName(*this, FD, NameStr)) return; + SyclKernelArgsSizeChecker argsSizeChecker(*this, FD->getLocation(), false /*IsSIMDKernel*/); SyclKernelDeclCreator kernel_decl(*this, FD->getLocation(), FD->isInlined(), false /*IsSIMDKernel */, FD); + FreeFunctionKernelBodyCreator kernel_body(*this, kernel_decl, FD); + SyclKernelIntHeaderCreator int_header(*this, getSyclIntegrationHeader(), FD->getType(), FD); + SyclKernelIntFooterCreator int_footer(*this, getSyclIntegrationFooter()); KernelObjVisitor Visitor{*this}; - Visitor.VisitFunctionParameters(FD, argsSizeChecker); - -Visitor.VisitFunctionParameters(FD, kernel_decl); - -Visitor.VisitFunctionParameters(FD, kernel_body); - -Visitor.VisitFunctionParameters(FD, int_header); - -Visitor.VisitFunctionParameters(FD, int_footer); + Visitor.VisitFunctionParameters(FD, argsSizeChecker, kernel_decl, kernel_body, + int_header, int_footer); assert(getKernelFDPairs().back().first == FD && "OpenCL Kernel not found for free function entry"); @@ -7056,26 +7016,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { } } ParmListWithNamesOstream.flush(); - for (ParmVarDecl *Param : K.SyclKernel->parameters()) { - // if (FreeFunctionKernelBodyCreator::isSpecialTypeWrapper( - // Param->getType())) { - // this is a struct that contains a special type so its neither a - // special type nor a trivially copyable type. We therefore need to - // explicitly communicate to the runtime that this argument should be - // allowed as a free function kernel argument. We do this by defining - // a certain trait recognized by the runtime to be true. - O << "template <>\n"; - O << "struct " - "sycl::ext::oneapi::experimental::detail::is_explicitly_allowed_" - "arg<"; - Policy.SuppressTagKeyword = true; - - Param->getType().print(O, Policy); - Policy.SuppressTagKeyword = false; - O << "> {\n"; - O << " static constexpr bool value = true;\n};\n"; - //} - } FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate(); Policy.PrintAsCanonical = false; Policy.SuppressDefinition = true; @@ -7812,7 +7752,7 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD, OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap); Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get(); -OFD->setBody(OFDBody); + OFD->setBody(OFDBody); OFD->setNothrow(); Stmt *NewBody = new (getASTContext()) SYCLKernelCallStmt(Body, OFD); diff --git a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp index 0ca1c234c9070..2b5d1f4190d21 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp @@ -44,16 +44,6 @@ template struct is_kernel { template inline constexpr bool is_kernel_v = is_kernel::value; -namespace detail { -template struct is_explicitly_allowed_arg { - static constexpr bool value = false; -}; - -template -inline constexpr bool is_explicitly_allowed_arg_v = - is_explicitly_allowed_arg::value; - -} // namespace detail } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 3506d1f34f493..fdd27ffe3f5cb 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -36,7 +36,6 @@ #include #include #include -#include #include #include #include @@ -1816,8 +1815,7 @@ class __SYCL_EXPORT handler { || (!is_same_type::value && std::is_pointer_v>) // USM || is_same_type::value // Interop - || is_same_type::value // Stream - || ext::oneapi::experimental::detail::is_explicitly_allowed_arg>::value; + || is_same_type::value; // Stream }; /// Sets argument for OpenCL interoperability kernels. From e7adeb34bc84bbfa96de8e6dbb487e5569ae7261 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Mon, 21 Jul 2025 20:22:01 -0700 Subject: [PATCH 03/12] Add sanity test --- .../free_function_host_compiler.cpp | 26 +++++++++++++++++++ 1 file changed, 26 insertions(+) create mode 100644 sycl/test-e2e/FreeFunctionKernels/free_function_host_compiler.cpp diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_host_compiler.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_host_compiler.cpp new file mode 100644 index 0000000000000..9676c94cab99b --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/free_function_host_compiler.cpp @@ -0,0 +1,26 @@ +// RUN: %{build} -fsycl-host-compiler=g++ -o %t.out +// REQUIRES: linux + +#include +#include +#include + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void kernel() {} + +int main() { + sycl::queue q; + + sycl::kernel_bundle bundle = + sycl::get_kernel_bundle(q.get_context()); + sycl::kernel_id kID = + sycl::ext::oneapi::experimental::get_kernel_id(); + sycl::kernel krn = bundle.get_kernel(kID); + + q.submit([&](sycl::handler &cgh) { + sycl::nd_range<1> ndr; + cgh.parallel_for(ndr, krn); + }); + return 0; +} From b453f78566dc77c908c50bd5b45c96307b06dc8a Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Mon, 21 Jul 2025 20:23:05 -0700 Subject: [PATCH 04/12] Bugfix --- clang/lib/Sema/SemaSYCL.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 07a6c25dced91..b5cc95ab3327f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6946,6 +6946,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { unsigned CurStart = 0; for (const KernelDesc &K : KernelDescs) { + if (S.isFreeFunction(K.SyclKernel)) + continue; const size_t N = K.Params.size(); PresumedLoc PLoc = S.getASTContext().getSourceManager().getPresumedLoc( S.getASTContext() From 68679af9409c89df169c1e615e0fdd34874aefe1 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Mon, 21 Jul 2025 23:59:31 -0400 Subject: [PATCH 05/12] Update free_function_int_header_rtc_mode.cpp --- clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp index 29b697691f445..0f5b2eefcd2e6 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp @@ -44,11 +44,6 @@ int main(){ // CHECK-NEXT: "{{.*}}__sycl_kernel_free_function_nd_rangePiii", // CHECK-NEXT: "{{.*}}Kernel_Function", - -// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_singlePiii"; } -// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_nd_rangePiii"; } -// CHECK: static constexpr const char* getName() { return "{{.*}}Kernel_Function"; } - // CHECK-RTC-NOT: free_function_single_kernel // CHECK-RTC-NOT: free_function_nd_range From 73c76518156d963a0b4fdc81aea60d167b22d0c4 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 22 Jul 2025 06:33:18 -0700 Subject: [PATCH 06/12] Move test to sycl/test and fix precommit failures --- .../test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp | 7 +++---- .../free_function_kernels}/free_function_host_compiler.cpp | 0 2 files changed, 3 insertions(+), 4 deletions(-) rename sycl/{test-e2e/FreeFunctionKernels => test/extensions/free_function_kernels}/free_function_host_compiler.cpp (100%) diff --git a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp index 29b697691f445..2ef457bfa4f9b 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp @@ -44,10 +44,9 @@ int main(){ // CHECK-NEXT: "{{.*}}__sycl_kernel_free_function_nd_rangePiii", // CHECK-NEXT: "{{.*}}Kernel_Function", - -// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_singlePiii"; } -// CHECK: static constexpr const char* getName() { return "{{.*}}__sycl_kernel_free_function_nd_rangePiii"; } -// CHECK: static constexpr const char* getName() { return "{{.*}}Kernel_Function"; } +// CHECK: _Z34__sycl_kernel_free_function_singlePiii +// CHECK: _Z36__sycl_kernel_free_function_nd_rangePiii +// CHECK: _ZTSZ4mainE15Kernel_Function // CHECK-RTC-NOT: free_function_single_kernel // CHECK-RTC-NOT: free_function_nd_range diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp similarity index 100% rename from sycl/test-e2e/FreeFunctionKernels/free_function_host_compiler.cpp rename to sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp From 2c2ff655ccd911fa79ee414eb24ee092e6421867 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 22 Jul 2025 10:38:27 -0400 Subject: [PATCH 07/12] Update free_function_host_compiler.cpp --- .../free_function_kernels/free_function_host_compiler.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp index 9676c94cab99b..acc898e7697e3 100644 --- a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp +++ b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp @@ -1,6 +1,10 @@ -// RUN: %{build} -fsycl-host-compiler=g++ -o %t.out +// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -o %t.out // REQUIRES: linux +// This test serves as a sanity check that compilation succeeds +// for a simple free function kernel when using the host compiler +// flag. + #include #include #include From 80a475b8a25a225f6b2b1a0c47936dc75d24f110 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 22 Jul 2025 11:11:33 -0400 Subject: [PATCH 08/12] Update free_function_host_compiler.cpp --- .../free_function_kernels/free_function_host_compiler.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp index acc898e7697e3..0517c80d091b6 100644 --- a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp +++ b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ -o %t.out +// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s // REQUIRES: linux // This test serves as a sanity check that compilation succeeds From 41daa24fd34efb88b7fb8582191c7f5ab1651d2b Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Tue, 22 Jul 2025 13:05:17 -0700 Subject: [PATCH 09/12] Fix pre-commit failures --- clang/lib/Sema/SemaSYCL.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b5cc95ab3327f..9c47fc88e2c06 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6946,9 +6946,11 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { unsigned CurStart = 0; for (const KernelDesc &K : KernelDescs) { - if (S.isFreeFunction(K.SyclKernel)) - continue; const size_t N = K.Params.size(); + if (S.isFreeFunction(K.SyclKernel)) { + CurStart += N; + continue; + } PresumedLoc PLoc = S.getASTContext().getSourceManager().getPresumedLoc( S.getASTContext() .getSourceManager() From 90c9541a2145c15f57aee1b6bb5699636821eb02 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Mon, 28 Jul 2025 11:41:04 -0700 Subject: [PATCH 10/12] Improve tests by adding cases with unnamed lambda extension disabled --- .../free_function_kernels/free_function_host_compiler.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp index 0517c80d091b6..fed32dc6bc486 100644 --- a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp +++ b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp @@ -1,9 +1,10 @@ // RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s +// RUN: %clangxx -fsycl -fno-sycl-unnamed-lambda -fsycl-host-compiler=g++ %s // REQUIRES: linux // This test serves as a sanity check that compilation succeeds // for a simple free function kernel when using the host compiler -// flag. +// flag with unnamed lambda extension enabled and disabled. #include #include From 9f904f6b5378d9722209a234675e633ad9901730 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Mon, 28 Jul 2025 11:42:49 -0700 Subject: [PATCH 11/12] Improve tests by adding cases with unnamed lambda extension disabled --- .../free_function_kernels/free_function_host_compiler.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp index fed32dc6bc486..e12b5b26a901b 100644 --- a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp +++ b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp @@ -1,10 +1,10 @@ // RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s -// RUN: %clangxx -fsycl -fno-sycl-unnamed-lambda -fsycl-host-compiler=g++ %s +// RUN: %clangxx -fsycl -fno-sycl-unnamed-lambda %s // REQUIRES: linux // This test serves as a sanity check that compilation succeeds // for a simple free function kernel when using the host compiler -// flag with unnamed lambda extension enabled and disabled. +// flag with unnamed lambda extension enabled(default) and disabled. #include #include From 3adcd2abd0203a85614b843aecc179763bf677e2 Mon Sep 17 00:00:00 2001 From: Lorenc Bushi Date: Wed, 30 Jul 2025 13:07:06 -0400 Subject: [PATCH 12/12] Update sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp Co-authored-by: Tom Honermann --- .../free_function_kernels/free_function_host_compiler.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp index e12b5b26a901b..296ee82184ad5 100644 --- a/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp +++ b/sycl/test/extensions/free_function_kernels/free_function_host_compiler.cpp @@ -1,5 +1,6 @@ -// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s +// RUN: %clangxx -fsycl %s // RUN: %clangxx -fsycl -fno-sycl-unnamed-lambda %s +// RUN: %clangxx -fsycl -fsycl-host-compiler=g++ %s // REQUIRES: linux // This test serves as a sanity check that compilation succeeds