From b305e941e22449484f7e640153461636cdb20c8e Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 7 Mar 2025 07:40:43 -0800 Subject: [PATCH 01/23] [SYCL] add e2e test for namespace --- .../Experimental/free_functions/namespace.cpp | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 sycl/test-e2e/Experimental/free_functions/namespace.cpp diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp new file mode 100644 index 0000000000000..9a3d01cabaf81 --- /dev/null +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -0,0 +1,38 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void iota(float start, float *ptr) { + // Get the ID of this kernel iteration. + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + + ptr[id] = start + static_cast(id); +} + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + + // Get a kernel bundle that contains the free function kernel "iota". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "iota" function from that bundle. + sycl::kernel k_iota = exe_bndl.ext_oneapi_get_kernel(); + + float *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + // Set the values of the kernel arguments. + cgh.set_args(3.14f, ptr); + + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, k_iota); + }).wait(); +} From 5ff7b2cdee767cb28c52e81377f778a12ec3edc0 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 18 Mar 2025 16:17:09 +0100 Subject: [PATCH 02/23] [SYCL] implement namespace support for free function --- clang/lib/Sema/SemaSYCL.cpp | 302 ++++++++++-------- .../Experimental/free_functions/namespace.cpp | 116 ++++++- 2 files changed, 275 insertions(+), 143 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 218d32dab5602..9b8422b17b990 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -15,10 +15,9 @@ #include "clang/AST/QualTypeNames.h" #include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" -#include "clang/AST/TemplateArgumentVisitor.h" -#include "clang/AST/Mangle.h" #include "clang/AST/SYCLKernelInfo.h" #include "clang/AST/StmtSYCL.h" +#include "clang/AST/TemplateArgumentVisitor.h" #include "clang/AST/TypeOrdering.h" #include "clang/AST/TypeVisitor.h" #include "clang/Analysis/CallGraph.h" @@ -27,7 +26,6 @@ #include "clang/Basic/Diagnostic.h" #include "clang/Basic/TargetInfo.h" #include "clang/Basic/Version.h" -#include "clang/AST/SYCLKernelInfo.h" #include "clang/Sema/Attr.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/ParsedAttr.h" @@ -398,8 +396,7 @@ bool SemaSYCL::isDeclAllowedInSYCLDeviceCode(const Decl *D) { return true; const DeclContext *DC = FD->getDeclContext(); - if (II && II->isStr("__spirv_ocl_printf") && - !FD->isDefined() && + if (II && II->isStr("__spirv_ocl_printf") && !FD->isDefined() && FD->getLanguageLinkage() == CXXLanguageLinkage && DC->getEnclosingNamespaceContext()->isTranslationUnit()) return true; @@ -632,19 +629,18 @@ static void collectSYCLAttributes(SemaSYCL &S, FunctionDecl *FD, if (DirectlyCalled) { llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { // FIXME: Make this list self-adapt as new SYCL attributes are added. - return isa(A); + return isa< + IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr, + SYCLReqdWorkGroupSizeAttr, SYCLWorkGroupSizeHintAttr, + SYCLIntelKernelArgsRestrictAttr, SYCLIntelNumSimdWorkItemsAttr, + SYCLIntelSchedulerTargetFmaxMhzAttr, SYCLIntelMaxWorkGroupSizeAttr, + SYCLIntelMaxGlobalWorkDimAttr, + SYCLIntelMinWorkGroupsPerComputeUnitAttr, + SYCLIntelMaxWorkGroupsPerMultiprocessorAttr, + SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr, SYCLIntelLoopFuseAttr, + SYCLIntelMaxConcurrencyAttr, SYCLIntelDisableLoopPipeliningAttr, + SYCLIntelInitiationIntervalAttr, SYCLIntelUseStallEnableClustersAttr, + SYCLDeviceHasAttr, SYCLAddIRAttributesFunctionAttr>(A); }); } } @@ -733,9 +729,7 @@ class DiagDeviceFunction : public RecursiveASTVisitor { // Make sure we skip the condition of the case, since that is a constant // expression. - bool TraverseCaseStmt(CaseStmt *S) { - return TraverseStmt(S->getSubStmt()); - } + bool TraverseCaseStmt(CaseStmt *S) { return TraverseStmt(S->getSubStmt()); } // Skip checking the size expr, since a constant array type loc's size expr is // a constant expression. @@ -1002,7 +996,8 @@ class SingleDeviceFunctionTracker { !KernelBody->hasAttr() && !KernelBody->hasAttr()) { KernelBody->addAttr(AlwaysInlineAttr::CreateImplicit( - KernelBody->getASTContext(), {}, AlwaysInlineAttr::Keyword_forceinline)); + KernelBody->getASTContext(), {}, + AlwaysInlineAttr::Keyword_forceinline)); } } @@ -1094,8 +1089,7 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor { // not a member of sycl::group - continue search return true; auto Name = Callee->getName(); - if (Name != "wait_for" || - Callee->hasAttr()) + if (Name != "wait_for" || Callee->hasAttr()) return true; // it is a call to sycl::group::wait_for - mark the callee Callee->addAttr( @@ -1314,15 +1308,21 @@ static bool isReadOnlyAccessor(const TemplateArgument &AccessModeArg) { // anonymous namespace so these don't get linkage. namespace { -template struct bind_param { using type = T; }; +template struct bind_param { + using type = T; +}; template <> struct bind_param { using type = const CXXBaseSpecifier &; }; -template <> struct bind_param { using type = FieldDecl *; }; +template <> struct bind_param { + using type = FieldDecl *; +}; -template <> struct bind_param { using type = FieldDecl *; }; +template <> struct bind_param { + using type = FieldDecl *; +}; template using bind_param_t = typename bind_param::type; @@ -1331,7 +1331,7 @@ class KernelObjVisitor { template void VisitUnionImpl(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, HandlerTys &... Handlers) { + const CXXRecordDecl *Wrapper, HandlerTys &...Handlers) { (void)std::initializer_list{ (Handlers.enterUnion(Owner, Parent), 0)...}; VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...); @@ -1341,13 +1341,13 @@ class KernelObjVisitor { // These enable handler execution only when previous Handlers succeed. template - bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) { + bool handleField(FieldDecl *FD, QualType FDTy, Tn &&...tn) { bool result = true; (void)std::initializer_list{(result = result && tn(FD, FDTy), 0)...}; return result; } template - bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&... tn) { + bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&...tn) { bool result = true; std::initializer_list{(result = result && tn(BD, BDTy), 0)...}; return result; @@ -1363,7 +1363,7 @@ class KernelObjVisitor { std::ref(Handlers), _1, _2)...) // The following simpler definition works with gcc 8.x and later. - //#define KF_FOR_EACH(FUNC) \ + // #define KF_FOR_EACH(FUNC) \ // handleField(Field, FieldTy, ([&](FieldDecl *FD, QualType FDTy) { \ // return Handlers.f(FD, FDTy); \ // })...) @@ -1392,7 +1392,7 @@ 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...); @@ -1404,7 +1404,7 @@ class KernelObjVisitor { template void visitSimpleRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &... Handlers) { + HandlerTys &...Handlers) { (void)std::initializer_list{ (Handlers.handleNonDecompStruct(Owner, Parent, RecordTy), 0)...}; } @@ -1412,16 +1412,16 @@ class KernelObjVisitor { template void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &... Handlers); + HandlerTys &...Handlers); template void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, HandlerTys &... Handlers); + const CXXRecordDecl *Wrapper, HandlerTys &...Handlers); template void VisitRecordHelper(const CXXRecordDecl *Owner, clang::CXXRecordDecl::base_class_const_range Range, - HandlerTys &... Handlers) { + HandlerTys &...Handlers) { for (const auto &Base : Range) { QualType BaseTy = Base.getType(); // Handle accessor class as base @@ -1438,14 +1438,14 @@ class KernelObjVisitor { template void VisitRecordHelper(const CXXRecordDecl *Owner, RecordDecl::field_range Range, - HandlerTys &... Handlers) { + HandlerTys &...Handlers) { VisitRecordFields(Owner, Handlers...); } template void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - HandlerTys &... Handlers) { + HandlerTys &...Handlers) { (void)std::initializer_list{ (Handlers.nextElement(ElementTy, Index), 0)...}; visitField(Owner, ArrayField, ElementTy, Handlers...); @@ -1453,24 +1453,24 @@ class KernelObjVisitor { template void visitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, HandlerTys &... Handlers) { + QualType ElementTy, HandlerTys &...Handlers) { visitArrayElementImpl(Owner, ArrayField, ElementTy, 0, Handlers...); } template void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - HandlerTys &... Handlers); + HandlerTys &...Handlers); template void visitSimpleArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &... Handlers) { + QualType ArrayTy, HandlerTys &...Handlers) { (void)std::initializer_list{ (Handlers.handleSimpleArrayType(Field, ArrayTy), 0)...}; } template void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &... Handlers) { + QualType ArrayTy, HandlerTys &...Handlers) { // Array workflow is: // handleArrayType // enterArray @@ -1502,7 +1502,7 @@ class KernelObjVisitor { template void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType FieldTy, HandlerTys &... Handlers) { + QualType FieldTy, HandlerTys &...Handlers) { if (isSyclSpecialType(FieldTy, SemaSYCLRef)) KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy); else if (FieldTy->isStructureOrClassType()) { @@ -1577,14 +1577,14 @@ class KernelObjVisitor { template void VisitRecordBases(const CXXRecordDecl *KernelFunctor, - HandlerTys &... Handlers) { + HandlerTys &...Handlers) { VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), Handlers...); } // A visitor function that dispatches to functions as defined in // SyclKernelFieldHandler for the purposes of kernel generation. template - void VisitRecordFields(const CXXRecordDecl *Owner, HandlerTys &... Handlers) { + void VisitRecordFields(const CXXRecordDecl *Owner, HandlerTys &...Handlers) { for (const auto Field : Owner->fields()) visitField(Owner, Field, Field->getType(), Handlers...); } @@ -1770,7 +1770,9 @@ template struct HandlerFilter { template struct AnyTrue; -template struct AnyTrue { static constexpr bool Value = B; }; +template struct AnyTrue { + static constexpr bool Value = B; +}; template struct AnyTrue { static constexpr bool Value = B || AnyTrue::Value; @@ -1778,7 +1780,9 @@ template struct AnyTrue { template struct AllTrue; -template struct AllTrue { static constexpr bool Value = B; }; +template struct AllTrue { + static constexpr bool Value = B; +}; template struct AllTrue { static constexpr bool Value = B && AllTrue::Value; @@ -1787,7 +1791,7 @@ template struct AllTrue { template void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, - Handlers &... handlers) { + Handlers &...handlers) { // Don't continue descending if none of the handlers 'care'. This could be 'if // constexpr' starting in C++17. Until then, we have to count on the // optimizer to realize "if (false)" is a dead branch. @@ -1801,7 +1805,7 @@ template void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - Handlers &... handlers) { + Handlers &...handlers) { // Don't continue descending if none of the handlers 'care'. This could be 'if // constexpr' starting in C++17. Until then, we have to count on the // optimizer to realize "if (false)" is a dead branch. @@ -1815,8 +1819,7 @@ void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner, template void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, - QualType RecordTy, - HandlerTys &... Handlers) { + QualType RecordTy, HandlerTys &...Handlers) { RecordDecl *RD = RecordTy->getAsRecordDecl(); assert(RD && "should not be null."); if (RD->hasAttr()) { @@ -1861,7 +1864,7 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, template void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &... Handlers) { + QualType ArrayTy, HandlerTys &...Handlers) { if (Field->hasAttr()) { visitComplexArray(Owner, Field, ArrayTy, Handlers...); @@ -2061,8 +2064,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { NotForwardDeclarableReason NFDR = isForwardDeclarable(RD, SemaSYCLRef, /*DiagForFreeFunction=*/true); if (NFDR != NotForwardDeclarableReason::None) { - Diag.Report(PD->getLocation(), - diag::err_bad_kernel_param_type) + Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type) << ParamTy; Diag.Report(PD->getLocation(), diag::note_free_function_kernel_param_type_not_fwd_declarable) @@ -2166,8 +2168,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { // experience. CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); if (RD->hasAttr()) { - Diag.Report(PD->getLocation(), - diag::err_bad_kernel_param_type) + Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type) << ParamTy; Diag.Report(PD->getLocation(), diag::note_free_function_kernel_param_type_not_supported) @@ -5313,7 +5314,6 @@ void SemaSYCL::SetSYCLKernelNames() { getSyclIntegrationHeader().updateKernelNames(Pair.first, KernelName, StableName); - // Set name of generated kernel. Pair.second->setDeclName(&getASTContext().Idents.get(KernelName)); // Update the AsmLabel for this generated kernel. @@ -6399,6 +6399,107 @@ static void EmitPragmaDiagnosticPop(raw_ostream &O) { O << "\n"; } +template +static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS, + const DeclContext *DC) { + if (DC->isTranslationUnit()) + return; + + const auto *CurDecl = cast(DC); + // Ensure we are in the canonical version, so that we know we have the 'full' + // name of the thing. + CurDecl = CurDecl->getCanonicalDecl(); + + // We are intentionally skipping linkage decls and record decls. Namespaces + // can appear in a linkage decl, but not a record decl, so we don't have to + // worry about the names getting messed up from that. We handle record decls + // later when printing the name of the thing. + const auto *NS = dyn_cast(CurDecl); + if (NS) + Before(OS, NS); + + if (const DeclContext *NewDC = CurDecl->getDeclContext()) + PrintNSHelper(Before, After, OS, NewDC); + + if (NS) + After(OS, NS); +} + +static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC, bool isPrintNamesOnly = false) { + PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {}, + [isPrintNamesOnly](raw_ostream &OS, const NamespaceDecl *NS) { + if (!isPrintNamesOnly) + { + if (NS->isInline()) + OS << "inline "; + OS << "namespace "; + } + if (!NS->isAnonymousNamespace()) + { + OS << NS->getName(); + if (isPrintNamesOnly) + OS << "::"; + else + OS <<" "; + } + if (!isPrintNamesOnly) + { + OS << "{\n"; + } + }, + OS, DC); +} + +static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) { + PrintNSHelper( + [](raw_ostream &OS, const NamespaceDecl *NS) { + OS << "} // "; + if (NS->isInline()) + OS << "inline "; + + OS << "namespace "; + if (!NS->isAnonymousNamespace()) + OS << NS->getName(); + + OS << '\n'; + }, + [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC); +} + +static bool insertFreeFunctionDeclaration(const PrintingPolicy& Policy, const FunctionDecl * FD, raw_ostream& O) +{ + const auto* DC = FD->getDeclContext(); + bool NSInserted{false}; + if (DC) + { + if (isa(DC)) + { + PrintNamespaces(O, FD); + NSInserted = true; + } + std::string Args; + for(unsigned i = 0; i < FD->getNumParams(); i++) + { + if(i > 0) + { + Args += ","; + } + Args += FD->getParamDecl(i)->getType().getAsString(); + } + O << FD->getReturnType().getAsString(); + O << " "; + O << FD->getNameAsString(); + O << "("; + O << Args; + O << ");\n"; + if (NSInserted) + { + PrintNSClosingBraces(O, FD); + } + } + return NSInserted; +} + void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "// This is auto-generated SYCL integration header.\n"; O << "\n"; @@ -6426,7 +6527,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "#define " << Macro.first << " " << Macro.second << '\n'; O << "#endif //" << Macro.first << "\n\n"; } - switch (S.getLangOpts().getSYCLRangeRounding()) { case LangOptions::SYCLRangeRoundingPreference::Disable: O << "#ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ \n"; @@ -6480,7 +6580,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "};\n"; } } - O << "// Forward declarations of templated kernel function types:\n"; for (const KernelDesc &K : KernelDescs) if (!K.IsUnnamedKernel) @@ -6512,7 +6611,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "\n"; EmitPragmaDiagnosticPop(O); } - // Generate declaration of variable of type __sycl_host_pipe_registration // whose sole purpose is to run its constructor before the application's // main() function. @@ -6533,8 +6631,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "\n"; EmitPragmaDiagnosticPop(O); } - - O << "// names of all kernels defined in the corresponding source\n"; O << "static constexpr\n"; O << "const char* const kernel_names[] = {\n"; @@ -6563,7 +6659,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { } O << "\n"; } - // Sentinel in place for 2 reasons: // 1- to make sure we don't get a warning because this collection is empty. // 2- to provide an obvious value that we can use when debugging to see that @@ -6594,7 +6689,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { Printer.Visit(K.NameType); O << "> {\n"; } - O << " __SYCL_DLL_LOCAL\n"; O << " static constexpr const char* getName() { return \"" << K.Name << "\"; }\n"; @@ -6730,17 +6824,22 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // template arguments that match default template arguments while printing // template-ids, even if the source code doesn't reference them. Policy.EnforceDefaultTemplateArgs = true; + bool NSInserted{false}; if (FTD) { FTD->print(O, Policy); } else { - K.SyclKernel->print(O, Policy); + NSInserted = insertFreeFunctionDeclaration(Policy, K.SyclKernel, O); } - O << ";\n"; // Generate a shim function that returns the address of the free function. O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n"; - O << " return (void (*)(" << ParmList << "))" - << K.SyclKernel->getIdentifier()->getName().data(); + O << " return (void (*)(" << ParmList << "))"; + if (NSInserted) + { + PrintNamespaces(O, K.SyclKernel, true); + } + + O << K.SyclKernel->getIdentifier()->getName().data(); if (FTD) { const TemplateArgumentList *TAL = K.SyclKernel->getTemplateSpecializationArgs(); @@ -6818,7 +6917,14 @@ bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) { return false; } llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); + emit(Out); + + int IntHeaderFD1 = 0; + std::string S{"/tmp/my-files/header.h"}; + llvm::sys::fs::openFileForWrite(S, IntHeaderFD1); + llvm::raw_fd_ostream Out1(IntHeaderFD1, true /*close in destructor*/); + emit(Out1); return true; } @@ -6909,61 +7015,6 @@ bool SYCLIntegrationFooter::emit(StringRef IntHeaderName) { return emit(Out); } -template -static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS, - const DeclContext *DC) { - if (DC->isTranslationUnit()) - return; - - const auto *CurDecl = cast(DC); - // Ensure we are in the canonical version, so that we know we have the 'full' - // name of the thing. - CurDecl = CurDecl->getCanonicalDecl(); - - // We are intentionally skipping linkage decls and record decls. Namespaces - // can appear in a linkage decl, but not a record decl, so we don't have to - // worry about the names getting messed up from that. We handle record decls - // later when printing the name of the thing. - const auto *NS = dyn_cast(CurDecl); - if (NS) - Before(OS, NS); - - if (const DeclContext *NewDC = CurDecl->getDeclContext()) - PrintNSHelper(Before, After, OS, NewDC); - - if (NS) - After(OS, NS); -} - -static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC) { - PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {}, - [](raw_ostream &OS, const NamespaceDecl *NS) { - if (NS->isInline()) - OS << "inline "; - OS << "namespace "; - if (!NS->isAnonymousNamespace()) - OS << NS->getName() << " "; - OS << "{\n"; - }, - OS, DC); -} - -static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) { - PrintNSHelper( - [](raw_ostream &OS, const NamespaceDecl *NS) { - OS << "} // "; - if (NS->isInline()) - OS << "inline "; - - OS << "namespace "; - if (!NS->isAnonymousNamespace()) - OS << NS->getName(); - - OS << '\n'; - }, - [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC); -} - static std::string EmitShim(raw_ostream &OS, unsigned &ShimCounter, const std::string &LastShim, const NamespaceDecl *AnonNS) { @@ -7063,7 +7114,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { for (const VarDecl *VD : GlobalVars) { VD = VD->getCanonicalDecl(); - // Skip if this isn't a SpecIdType, DeviceGlobal, or HostPipe. This + // Skip if this isn't a SpecIdType, DeviceGlobal, or HostPipe. This // can happen if it was a deduced type. if (!SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) && !SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && @@ -7108,8 +7159,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { VD->getNameForDiagnostic(HostPipesOS, Policy, true); } HostPipesOS << ", \""; - HostPipesOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), - VD); + HostPipesOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD); HostPipesOS << "\");\n"; } else { EmittedFirstSpecConstant = true; diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 9a3d01cabaf81..4e57759a957b5 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -1,7 +1,12 @@ +// REQUIRES: aspect-usm_shared_allocations // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include +#include +#include +#include +#include + namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; @@ -9,30 +14,107 @@ static constexpr size_t NUM = 1024; static constexpr size_t WGSIZE = 16; SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void iota(float start, float *ptr) { - // Get the ID of this kernel iteration. +void func_without_ns(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id); } -int main() { - sycl::queue q; - sycl::context ctxt = q.get_context(); +namespace free_functions::tests { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void function_in_ns(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 1); + } +} // namespace free_functions::tests + +namespace free_functions::tests { +inline namespace V1 { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void function_in_inline_ns(float start, float *ptr) + { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 2); + } +} // V1 +} // namespace free_functions::tests - // Get a kernel bundle that contains the free function kernel "iota". - auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); - // Get a kernel object for the "iota" function from that bundle. - sycl::kernel k_iota = exe_bndl.ext_oneapi_get_kernel(); +namespace { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void function_in_anonymous_ns(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 3); + } +} +static void call_kernel_code(sycl::queue& q, sycl::kernel& kernel) { float *ptr = sycl::malloc_shared(NUM, q); q.submit([&](sycl::handler &cgh) { - // Set the values of the kernel arguments. - cgh.set_args(3.14f, ptr); + cgh.set_args(3.14f, ptr); + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, kernel); + }).wait(); +} + +void test_function_without_ns(sycl::queue& q, sycl::context& ctxt) +{ +#ifndef __SYCL_DEVICE_ONLY__ + // Get a kernel bundle that contains the free function kernel "func_without_ns". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + // Get a kernel object for the "func_without_ns" function from that bundle. + sycl::kernel k_func_without_ns = exe_bndl.ext_oneapi_get_kernel(); + call_kernel_code(q, k_func_without_ns); +#endif +} + +void test_function_in_ns(sycl::queue& q, sycl::context& ctxt) +{ +#ifndef __SYCL_DEVICE_ONLY__ + // Get a kernel bundle that contains the free function kernel "function_in_ns". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "function_in_ns" function from that bundle. + sycl::kernel k_function_in_ns = exe_bndl.ext_oneapi_get_kernel(); + call_kernel_code(q, k_function_in_ns); +#endif +} + +void test_function_in_inline_ns(sycl::queue& q, sycl::context& ctxt) +{ +#ifndef __SYCL_DEVICE_ONLY__ + // Get a kernel bundle that contains the free function kernel "function_in_inline_ns". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "function_in_inline_ns" function from that bundle. + sycl::kernel k_function_in_inline_ns = exe_bndl.ext_oneapi_get_kernel(); + call_kernel_code(q, k_function_in_inline_ns); +#endif +} + +void test_function_in_anonimous_ns(sycl::queue& q, sycl::context& ctxt) { + #ifndef __SYCL_DEVICE_ONLY__ + // Get a kernel bundle that contains the free function kernel "function_in_anonymous_ns". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "function_in_anonymous_ns" function from that bundle. + sycl::kernel k_function_in_anonymous_ns = exe_bndl.ext_oneapi_get_kernel(); + call_kernel_code(q, k_function_in_anonymous_ns); +#endif +} + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); - sycl::nd_range ndr{{NUM}, {WGSIZE}}; - cgh.parallel_for(ndr, k_iota); - }).wait(); + test_function_without_ns(q, ctxt); + test_function_in_ns(q, ctxt); + test_function_in_inline_ns(q, ctxt); + test_function_in_anonimous_ns(q, ctxt); } From 12df05eaece35c6d6a8147a2eae02314d40b6b71 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Wed, 19 Mar 2025 16:10:58 +0100 Subject: [PATCH 03/23] [SYCL] fix test issues after adding namespace support [SYCL] fix formatting [SYCL] remove unnecessary space --- clang/lib/Sema/SemaSYCL.cpp | 189 +++++++++--------- ...ee_function_default_template_arguments.cpp | 11 +- .../CodeGenSYCL/free_function_int_header.cpp | 6 +- .../free_function_int_header_rtc_mode.cpp | 4 +- 4 files changed, 105 insertions(+), 105 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9b8422b17b990..f2e03bc741cd5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -15,9 +15,10 @@ #include "clang/AST/QualTypeNames.h" #include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" +#include "clang/AST/TemplateArgumentVisitor.h" +#include "clang/AST/Mangle.h" #include "clang/AST/SYCLKernelInfo.h" #include "clang/AST/StmtSYCL.h" -#include "clang/AST/TemplateArgumentVisitor.h" #include "clang/AST/TypeOrdering.h" #include "clang/AST/TypeVisitor.h" #include "clang/Analysis/CallGraph.h" @@ -396,7 +397,8 @@ bool SemaSYCL::isDeclAllowedInSYCLDeviceCode(const Decl *D) { return true; const DeclContext *DC = FD->getDeclContext(); - if (II && II->isStr("__spirv_ocl_printf") && !FD->isDefined() && + if (II && II->isStr("__spirv_ocl_printf") && + !FD->isDefined() && FD->getLanguageLinkage() == CXXLanguageLinkage && DC->getEnclosingNamespaceContext()->isTranslationUnit()) return true; @@ -629,18 +631,19 @@ static void collectSYCLAttributes(SemaSYCL &S, FunctionDecl *FD, if (DirectlyCalled) { llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { // FIXME: Make this list self-adapt as new SYCL attributes are added. - return isa< - IntelReqdSubGroupSizeAttr, IntelNamedSubGroupSizeAttr, - SYCLReqdWorkGroupSizeAttr, SYCLWorkGroupSizeHintAttr, - SYCLIntelKernelArgsRestrictAttr, SYCLIntelNumSimdWorkItemsAttr, - SYCLIntelSchedulerTargetFmaxMhzAttr, SYCLIntelMaxWorkGroupSizeAttr, - SYCLIntelMaxGlobalWorkDimAttr, - SYCLIntelMinWorkGroupsPerComputeUnitAttr, - SYCLIntelMaxWorkGroupsPerMultiprocessorAttr, - SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr, SYCLIntelLoopFuseAttr, - SYCLIntelMaxConcurrencyAttr, SYCLIntelDisableLoopPipeliningAttr, - SYCLIntelInitiationIntervalAttr, SYCLIntelUseStallEnableClustersAttr, - SYCLDeviceHasAttr, SYCLAddIRAttributesFunctionAttr>(A); + return isa(A); }); } } @@ -729,7 +732,9 @@ class DiagDeviceFunction : public RecursiveASTVisitor { // Make sure we skip the condition of the case, since that is a constant // expression. - bool TraverseCaseStmt(CaseStmt *S) { return TraverseStmt(S->getSubStmt()); } + bool TraverseCaseStmt(CaseStmt *S) { + return TraverseStmt(S->getSubStmt()); + } // Skip checking the size expr, since a constant array type loc's size expr is // a constant expression. @@ -996,8 +1001,7 @@ class SingleDeviceFunctionTracker { !KernelBody->hasAttr() && !KernelBody->hasAttr()) { KernelBody->addAttr(AlwaysInlineAttr::CreateImplicit( - KernelBody->getASTContext(), {}, - AlwaysInlineAttr::Keyword_forceinline)); + KernelBody->getASTContext(), {}, AlwaysInlineAttr::Keyword_forceinline)); } } @@ -1089,7 +1093,8 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor { // not a member of sycl::group - continue search return true; auto Name = Callee->getName(); - if (Name != "wait_for" || Callee->hasAttr()) + if (Name != "wait_for" || + Callee->hasAttr()) return true; // it is a call to sycl::group::wait_for - mark the callee Callee->addAttr( @@ -1308,21 +1313,15 @@ static bool isReadOnlyAccessor(const TemplateArgument &AccessModeArg) { // anonymous namespace so these don't get linkage. namespace { -template struct bind_param { - using type = T; -}; +template struct bind_param { using type = T; }; template <> struct bind_param { using type = const CXXBaseSpecifier &; }; -template <> struct bind_param { - using type = FieldDecl *; -}; +template <> struct bind_param { using type = FieldDecl *; }; -template <> struct bind_param { - using type = FieldDecl *; -}; +template <> struct bind_param { using type = FieldDecl *; }; template using bind_param_t = typename bind_param::type; @@ -1331,7 +1330,7 @@ class KernelObjVisitor { template void VisitUnionImpl(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, HandlerTys &...Handlers) { + const CXXRecordDecl *Wrapper, HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.enterUnion(Owner, Parent), 0)...}; VisitRecordHelper(Wrapper, Wrapper->fields(), Handlers...); @@ -1341,13 +1340,13 @@ class KernelObjVisitor { // These enable handler execution only when previous Handlers succeed. template - bool handleField(FieldDecl *FD, QualType FDTy, Tn &&...tn) { + bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) { bool result = true; (void)std::initializer_list{(result = result && tn(FD, FDTy), 0)...}; return result; } template - bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&...tn) { + bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&... tn) { bool result = true; std::initializer_list{(result = result && tn(BD, BDTy), 0)...}; return result; @@ -1363,7 +1362,7 @@ class KernelObjVisitor { std::ref(Handlers), _1, _2)...) // The following simpler definition works with gcc 8.x and later. - // #define KF_FOR_EACH(FUNC) \ + //#define KF_FOR_EACH(FUNC) \ // handleField(Field, FieldTy, ([&](FieldDecl *FD, QualType FDTy) { \ // return Handlers.f(FD, FDTy); \ // })...) @@ -1392,7 +1391,7 @@ 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...); @@ -1404,7 +1403,7 @@ class KernelObjVisitor { template void visitSimpleRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &...Handlers) { + HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.handleNonDecompStruct(Owner, Parent, RecordTy), 0)...}; } @@ -1412,16 +1411,16 @@ class KernelObjVisitor { template void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, QualType RecordTy, - HandlerTys &...Handlers); + HandlerTys &... Handlers); template void VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, - const CXXRecordDecl *Wrapper, HandlerTys &...Handlers); + const CXXRecordDecl *Wrapper, HandlerTys &... Handlers); template void VisitRecordHelper(const CXXRecordDecl *Owner, clang::CXXRecordDecl::base_class_const_range Range, - HandlerTys &...Handlers) { + HandlerTys &... Handlers) { for (const auto &Base : Range) { QualType BaseTy = Base.getType(); // Handle accessor class as base @@ -1438,14 +1437,14 @@ class KernelObjVisitor { template void VisitRecordHelper(const CXXRecordDecl *Owner, RecordDecl::field_range Range, - HandlerTys &...Handlers) { + HandlerTys &... Handlers) { VisitRecordFields(Owner, Handlers...); } template void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - HandlerTys &...Handlers) { + HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.nextElement(ElementTy, Index), 0)...}; visitField(Owner, ArrayField, ElementTy, Handlers...); @@ -1453,24 +1452,24 @@ class KernelObjVisitor { template void visitFirstArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, - QualType ElementTy, HandlerTys &...Handlers) { + QualType ElementTy, HandlerTys &... Handlers) { visitArrayElementImpl(Owner, ArrayField, ElementTy, 0, Handlers...); } template void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - HandlerTys &...Handlers); + HandlerTys &... Handlers); template void visitSimpleArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &...Handlers) { + QualType ArrayTy, HandlerTys &... Handlers) { (void)std::initializer_list{ (Handlers.handleSimpleArrayType(Field, ArrayTy), 0)...}; } template void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &...Handlers) { + QualType ArrayTy, HandlerTys &... Handlers) { // Array workflow is: // handleArrayType // enterArray @@ -1502,7 +1501,7 @@ class KernelObjVisitor { template void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType FieldTy, HandlerTys &...Handlers) { + QualType FieldTy, HandlerTys &... Handlers) { if (isSyclSpecialType(FieldTy, SemaSYCLRef)) KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy); else if (FieldTy->isStructureOrClassType()) { @@ -1577,14 +1576,14 @@ class KernelObjVisitor { template void VisitRecordBases(const CXXRecordDecl *KernelFunctor, - HandlerTys &...Handlers) { + HandlerTys &... Handlers) { VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), Handlers...); } // A visitor function that dispatches to functions as defined in // SyclKernelFieldHandler for the purposes of kernel generation. template - void VisitRecordFields(const CXXRecordDecl *Owner, HandlerTys &...Handlers) { + void VisitRecordFields(const CXXRecordDecl *Owner, HandlerTys &... Handlers) { for (const auto Field : Owner->fields()) visitField(Owner, Field, Field->getType(), Handlers...); } @@ -1770,9 +1769,7 @@ template struct HandlerFilter { template struct AnyTrue; -template struct AnyTrue { - static constexpr bool Value = B; -}; +template struct AnyTrue { static constexpr bool Value = B; }; template struct AnyTrue { static constexpr bool Value = B || AnyTrue::Value; @@ -1780,9 +1777,7 @@ template struct AnyTrue { template struct AllTrue; -template struct AllTrue { - static constexpr bool Value = B; -}; +template struct AllTrue { static constexpr bool Value = B; }; template struct AllTrue { static constexpr bool Value = B && AllTrue::Value; @@ -1791,7 +1786,7 @@ template struct AllTrue { template void KernelObjVisitor::VisitUnion(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, - Handlers &...handlers) { + Handlers &... handlers) { // Don't continue descending if none of the handlers 'care'. This could be 'if // constexpr' starting in C++17. Until then, we have to count on the // optimizer to realize "if (false)" is a dead branch. @@ -1805,7 +1800,7 @@ template void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField, QualType ElementTy, uint64_t Index, - Handlers &...handlers) { + Handlers &... handlers) { // Don't continue descending if none of the handlers 'care'. This could be 'if // constexpr' starting in C++17. Until then, we have to count on the // optimizer to realize "if (false)" is a dead branch. @@ -1819,7 +1814,8 @@ void KernelObjVisitor::visitNthArrayElement(const CXXRecordDecl *Owner, template void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, const CXXRecordDecl *Wrapper, - QualType RecordTy, HandlerTys &...Handlers) { + QualType RecordTy, + HandlerTys &... Handlers) { RecordDecl *RD = RecordTy->getAsRecordDecl(); assert(RD && "should not be null."); if (RD->hasAttr()) { @@ -1864,7 +1860,7 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent, template void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field, - QualType ArrayTy, HandlerTys &...Handlers) { + QualType ArrayTy, HandlerTys &... Handlers) { if (Field->hasAttr()) { visitComplexArray(Owner, Field, ArrayTy, Handlers...); @@ -2064,7 +2060,8 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { NotForwardDeclarableReason NFDR = isForwardDeclarable(RD, SemaSYCLRef, /*DiagForFreeFunction=*/true); if (NFDR != NotForwardDeclarableReason::None) { - Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type) + Diag.Report(PD->getLocation(), + diag::err_bad_kernel_param_type) << ParamTy; Diag.Report(PD->getLocation(), diag::note_free_function_kernel_param_type_not_fwd_declarable) @@ -2168,7 +2165,8 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { // experience. CXXRecordDecl *RD = ParamTy->getAsCXXRecordDecl(); if (RD->hasAttr()) { - Diag.Report(PD->getLocation(), diag::err_bad_kernel_param_type) + Diag.Report(PD->getLocation(), + diag::err_bad_kernel_param_type) << ParamTy; Diag.Report(PD->getLocation(), diag::note_free_function_kernel_param_type_not_supported) @@ -5314,6 +5312,7 @@ void SemaSYCL::SetSYCLKernelNames() { getSyclIntegrationHeader().updateKernelNames(Pair.first, KernelName, StableName); + // Set name of generated kernel. Pair.second->setDeclName(&getASTContext().Idents.get(KernelName)); // Update the AsmLabel for this generated kernel. @@ -6425,25 +6424,23 @@ static void PrintNSHelper(BeforeFn Before, AfterFn After, raw_ostream &OS, After(OS, NS); } -static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC, bool isPrintNamesOnly = false) { +static void PrintNamespaces(raw_ostream &OS, const DeclContext *DC, + bool isPrintNamesOnly = false) { PrintNSHelper([](raw_ostream &OS, const NamespaceDecl *NS) {}, [isPrintNamesOnly](raw_ostream &OS, const NamespaceDecl *NS) { - if (!isPrintNamesOnly) - { + if (!isPrintNamesOnly) { if (NS->isInline()) OS << "inline "; OS << "namespace "; } - if (!NS->isAnonymousNamespace()) - { + if (!NS->isAnonymousNamespace()) { OS << NS->getName(); if (isPrintNamesOnly) OS << "::"; else - OS <<" "; + OS << " "; } - if (!isPrintNamesOnly) - { + if (!isPrintNamesOnly) { OS << "{\n"; } }, @@ -6466,36 +6463,36 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) { [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC); } -static bool insertFreeFunctionDeclaration(const PrintingPolicy& Policy, const FunctionDecl * FD, raw_ostream& O) -{ - const auto* DC = FD->getDeclContext(); +static bool insertFreeFunctionDeclaration(const PrintingPolicy &Policy, + const FunctionDecl *FD, + raw_ostream &O) { + const auto *DC = FD->getDeclContext(); bool NSInserted{false}; - if (DC) - { - if (isa(DC)) - { + if (DC) { + if (isa(DC)) { PrintNamespaces(O, FD); NSInserted = true; } - std::string Args; - for(unsigned i = 0; i < FD->getNumParams(); i++) - { - if(i > 0) - { - Args += ","; + std::string args; + llvm::raw_string_ostream Args{args}; + for (unsigned i = 0; i < FD->getNumParams(); i++) { + FD->getParamDecl(i)->getType().print(Args, Policy); + Args << " " << FD->getParamDecl(i)->getNameAsString(); + + if (i < FD->getNumParams() - 1) { + Args << ", "; } - Args += FD->getParamDecl(i)->getType().getAsString(); } O << FD->getReturnType().getAsString(); O << " "; O << FD->getNameAsString(); O << "("; - O << Args; - O << ");\n"; - if (NSInserted) - { + O << args; + O << ");"; + if (NSInserted) { PrintNSClosingBraces(O, FD); } + Args.flush(); } return NSInserted; } @@ -6527,6 +6524,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "#define " << Macro.first << " " << Macro.second << '\n'; O << "#endif //" << Macro.first << "\n\n"; } + switch (S.getLangOpts().getSYCLRangeRounding()) { case LangOptions::SYCLRangeRoundingPreference::Disable: O << "#ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ \n"; @@ -6580,6 +6578,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "};\n"; } } + O << "// Forward declarations of templated kernel function types:\n"; for (const KernelDesc &K : KernelDescs) if (!K.IsUnnamedKernel) @@ -6611,6 +6610,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "\n"; EmitPragmaDiagnosticPop(O); } + // Generate declaration of variable of type __sycl_host_pipe_registration // whose sole purpose is to run its constructor before the application's // main() function. @@ -6631,6 +6631,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "\n"; EmitPragmaDiagnosticPop(O); } + + O << "// names of all kernels defined in the corresponding source\n"; O << "static constexpr\n"; O << "const char* const kernel_names[] = {\n"; @@ -6659,6 +6661,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { } O << "\n"; } + // Sentinel in place for 2 reasons: // 1- to make sure we don't get a warning because this collection is empty. // 2- to provide an obvious value that we can use when debugging to see that @@ -6689,6 +6692,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { Printer.Visit(K.NameType); O << "> {\n"; } + O << " __SYCL_DLL_LOCAL\n"; O << " static constexpr const char* getName() { return \"" << K.Name << "\"; }\n"; @@ -6827,19 +6831,20 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { bool NSInserted{false}; if (FTD) { FTD->print(O, Policy); + O << ";\n"; } else { NSInserted = insertFreeFunctionDeclaration(Policy, K.SyclKernel, O); + O << "\n"; } // Generate a shim function that returns the address of the free function. O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n"; O << " return (void (*)(" << ParmList << "))"; - if (NSInserted) - { + if (NSInserted) { PrintNamespaces(O, K.SyclKernel, true); } - O << K.SyclKernel->getIdentifier()->getName().data(); + O << K.SyclKernel->getIdentifier()->getName().data(); if (FTD) { const TemplateArgumentList *TAL = K.SyclKernel->getTemplateSpecializationArgs(); @@ -6917,14 +6922,7 @@ bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) { return false; } llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); - emit(Out); - - int IntHeaderFD1 = 0; - std::string S{"/tmp/my-files/header.h"}; - llvm::sys::fs::openFileForWrite(S, IntHeaderFD1); - llvm::raw_fd_ostream Out1(IntHeaderFD1, true /*close in destructor*/); - emit(Out1); return true; } @@ -7114,7 +7112,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { for (const VarDecl *VD : GlobalVars) { VD = VD->getCanonicalDecl(); - // Skip if this isn't a SpecIdType, DeviceGlobal, or HostPipe. This + // Skip if this isn't a SpecIdType, DeviceGlobal, or HostPipe. This // can happen if it was a deduced type. if (!SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) && !SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && @@ -7159,7 +7157,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { VD->getNameForDiagnostic(HostPipesOS, Policy, true); } HostPipesOS << ", \""; - HostPipesOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD); + HostPipesOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), + VD); HostPipesOS << "\");\n"; } else { EmittedFirstSpecConstant = true; diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp index 62a121d218b8b..0d19924b1ad1a 100644 --- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp +++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp @@ -86,18 +86,19 @@ foo(Arg1 arg) { // CHECK-NEXT: template struct Arg; // CHECK-NEXT: } -// CHECK: void ns::simple(ns::Arg); -// CHECK-NEXT: static constexpr auto __sycl_shim1() { -// CHECK-NEXT: return (void (*)(struct ns::Arg))simple; +// CHECK: namespace ns { +// CHECK-NEXT: void simple(ns::Arg );} // namespace ns +// CHECK: static constexpr auto __sycl_shim1() { +// CHECK-NEXT: return (void (*)(struct ns::Arg))ns::simple; // CHECK-NEXT: } // CHECK: Forward declarations of kernel and its argument types: // CHECK: namespace ns { // CHECK: namespace ns1 { // CHECK-NEXT: template class hasDefaultArg; -// CHECK-NEXT: } +// CHECK-NEXT: }} -// CHECK: void simple1(ns::Arg, int, 12, ns::notatuple>); +// CHECK: void simple1(ns::Arg, int, 12, ns::notatuple> ); // CHECK-NEXT: static constexpr auto __sycl_shim2() { // CHECK-NEXT: return (void (*)(struct ns::Arg, int, 12, struct ns::notatuple>))simple1; // CHECK-NEXT: } diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp index 6a196dedc2fc2..ccd2d73445cf8 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp @@ -163,7 +163,7 @@ void ff_8(sycl::work_group_memory) { // CHECK: Definition of _Z18__sycl_kernel_ff_2Piii as a free function kernel // CHECK: Forward declarations of kernel and its argument types: -// CHECK: void ff_2(int *ptr, int start, int end); +// CHECK: void ff_2(int * ptr, int start, int end); // CHECK-NEXT: static constexpr auto __sycl_shim1() { // CHECK-NEXT: return (void (*)(int *, int, int))ff_2; // CHECK-NEXT: } @@ -180,7 +180,7 @@ void ff_8(sycl::work_group_memory) { // CHECK: Definition of _Z18__sycl_kernel_ff_2Piiii as a free function kernel // CHECK: Forward declarations of kernel and its argument types: -// CHECK: void ff_2(int *ptr, int start, int end, int value); +// CHECK: void ff_2(int * ptr, int start, int end, int value); // CHECK-NEXT: static constexpr auto __sycl_shim2() { // CHECK-NEXT: return (void (*)(int *, int, int, int))ff_2; // CHECK-NEXT: } @@ -309,7 +309,7 @@ void ff_8(sycl::work_group_memory) { // CHECK: Forward declarations of kernel and its argument types: // CHECK: template class work_group_memory; -// CHECK: void ff_8(sycl::work_group_memory); +// CHECK: void ff_8(sycl::work_group_memory ); // CHECK-NEXT: static constexpr auto __sycl_shim9() { // CHECK-NEXT: return (void (*)(class sycl::work_group_memory))ff_8; // CHECK-NEXT: } 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 214318b563fa8..29b697691f445 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header_rtc_mode.cpp @@ -52,7 +52,7 @@ int main(){ // CHECK-RTC-NOT: free_function_single_kernel // CHECK-RTC-NOT: free_function_nd_range -// CHECK-NORTC: void free_function_single(int *ptr, int start, int end); +// CHECK-NORTC: void free_function_single(int * ptr, int start, int end); // CHECK-NORTC: static constexpr auto __sycl_shim[[#FIRST:]]() // CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_single; @@ -63,7 +63,7 @@ int main(){ // CHECK-NORTC-NEXT: static constexpr bool value = true; -// CHECK-NORTC: void free_function_nd_range(int *ptr, int start, int end); +// CHECK-NORTC: void free_function_nd_range(int * ptr, int start, int end); // CHECK-NORTC: static constexpr auto __sycl_shim[[#SECOND:]]() { // CHECK-NORTC-NEXT: return (void (*)(int *, int, int))free_function_nd_range; From 65ef84f1ba84474b211de4aa2d7cb8bcb57da6e3 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Wed, 19 Mar 2025 17:50:14 +0100 Subject: [PATCH 04/23] [SYCL] do not use macros in tests [SYCL] include additional header in tests [SYCL] do not iterate through arguments twice --- clang/lib/Sema/SemaSYCL.cpp | 35 +++++++++---------- .../Experimental/free_functions/namespace.cpp | 11 ++---- 2 files changed, 19 insertions(+), 27 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c0a8b70333648..461a3a2b84e7a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6454,6 +6454,7 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) { static bool insertFreeFunctionDeclaration(const PrintingPolicy &Policy, const FunctionDecl *FD, + const std::string& Args, raw_ostream &O) { const auto *DC = FD->getDeclContext(); bool NSInserted{false}; @@ -6462,26 +6463,11 @@ static bool insertFreeFunctionDeclaration(const PrintingPolicy &Policy, PrintNamespaces(O, FD); NSInserted = true; } - std::string args; - llvm::raw_string_ostream Args{args}; - for (unsigned i = 0; i < FD->getNumParams(); i++) { - FD->getParamDecl(i)->getType().print(Args, Policy); - Args << " " << FD->getParamDecl(i)->getNameAsString(); - - if (i < FD->getNumParams() - 1) { - Args << ", "; - } - } - O << FD->getReturnType().getAsString(); - O << " "; - O << FD->getNameAsString(); - O << "("; - O << args; - O << ");"; + O << FD->getReturnType().getAsString() << " "; + O << FD->getNameAsString() << "(" << Args << ");"; if (NSInserted) { PrintNSClosingBraces(O, FD); } - Args.flush(); } return NSInserted; } @@ -6774,16 +6760,24 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { if (K.SyclKernel->getLanguageLinkage() == CLanguageLinkage) O << "extern \"C\" "; std::string ParmList; + std::string ParmListWithNames; bool FirstParam = true; Policy.SuppressDefaultTemplateArgs = false; Policy.PrintCanonicalTypes = true; + llvm::raw_string_ostream ParmListWithNamesOstream{ParmListWithNames}; for (ParmVarDecl *Param : K.SyclKernel->parameters()) { if (FirstParam) FirstParam = false; else + { ParmList += ", "; + ParmListWithNamesOstream << ", "; + } + Param->getType().print(ParmListWithNamesOstream, Policy); + ParmListWithNamesOstream << " " << Param->getNameAsString(); ParmList += Param->getType().getCanonicalType().getAsString(Policy); } + ParmListWithNamesOstream.flush(); FunctionTemplateDecl *FTD = K.SyclKernel->getPrimaryTemplate(); Policy.PrintCanonicalTypes = false; Policy.SuppressDefinition = true; @@ -6822,7 +6816,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { FTD->print(O, Policy); O << ";\n"; } else { - NSInserted = insertFreeFunctionDeclaration(Policy, K.SyclKernel, O); + NSInserted = insertFreeFunctionDeclaration(Policy, K.SyclKernel, ParmListWithNames, O); O << "\n"; } @@ -6912,6 +6906,11 @@ bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) { } llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); emit(Out); + int IntHeaderFD1 = 0; + std::string S{"/tmp/my-files/header.h"}; + llvm::sys::fs::openFileForWrite(S, IntHeaderFD1); + llvm::raw_fd_ostream Out1(IntHeaderFD1, true /*close in destructor*/); + emit(Out1); return true; } diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 4e57759a957b5..86908b16f2aa8 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -2,10 +2,10 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include #include #include #include +#include namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; @@ -58,19 +58,16 @@ static void call_kernel_code(sycl::queue& q, sycl::kernel& kernel) { void test_function_without_ns(sycl::queue& q, sycl::context& ctxt) { -#ifndef __SYCL_DEVICE_ONLY__ // Get a kernel bundle that contains the free function kernel "func_without_ns". auto exe_bndl = syclexp::get_kernel_bundle(ctxt); // Get a kernel object for the "func_without_ns" function from that bundle. sycl::kernel k_func_without_ns = exe_bndl.ext_oneapi_get_kernel(); call_kernel_code(q, k_func_without_ns); -#endif } void test_function_in_ns(sycl::queue& q, sycl::context& ctxt) { -#ifndef __SYCL_DEVICE_ONLY__ // Get a kernel bundle that contains the free function kernel "function_in_ns". auto exe_bndl = syclexp::get_kernel_bundle(); call_kernel_code(q, k_function_in_ns); -#endif } void test_function_in_inline_ns(sycl::queue& q, sycl::context& ctxt) { -#ifndef __SYCL_DEVICE_ONLY__ // Get a kernel bundle that contains the free function kernel "function_in_inline_ns". auto exe_bndl = syclexp::get_kernel_bundle(); call_kernel_code(q, k_function_in_inline_ns); -#endif } void test_function_in_anonimous_ns(sycl::queue& q, sycl::context& ctxt) { - #ifndef __SYCL_DEVICE_ONLY__ // Get a kernel bundle that contains the free function kernel "function_in_anonymous_ns". auto exe_bndl = syclexp::get_kernel_bundle(); call_kernel_code(q, k_function_in_anonymous_ns); -#endif } int main() { @@ -117,4 +109,5 @@ int main() { test_function_in_ns(q, ctxt); test_function_in_inline_ns(q, ctxt); test_function_in_anonimous_ns(q, ctxt); + return 0; } From 117e97a1c438e1870fef149131d9ac65ab1e75e3 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 21 Mar 2025 17:56:19 +0100 Subject: [PATCH 05/23] [SYCL] fix tests --- clang/lib/Sema/SemaSYCL.cpp | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 461a3a2b84e7a..13dd36afbde0e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6773,7 +6773,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { ParmList += ", "; ParmListWithNamesOstream << ", "; } + Policy.SuppressTagKeyword = true; Param->getType().print(ParmListWithNamesOstream, Policy); + Policy.SuppressTagKeyword = false; ParmListWithNamesOstream << " " << Param->getNameAsString(); ParmList += Param->getType().getCanonicalType().getAsString(Policy); } @@ -6906,11 +6908,6 @@ bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) { } llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); emit(Out); - int IntHeaderFD1 = 0; - std::string S{"/tmp/my-files/header.h"}; - llvm::sys::fs::openFileForWrite(S, IntHeaderFD1); - llvm::raw_fd_ostream Out1(IntHeaderFD1, true /*close in destructor*/); - emit(Out1); return true; } From c53312a427bbd481f07b8b7e8e2ee32db17f4e9b Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 25 Mar 2025 11:46:14 +0100 Subject: [PATCH 06/23] [SYCL] exclude cuda from free function tests --- sycl/test-e2e/Experimental/free_functions/namespace.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 86908b16f2aa8..4287f29d1340d 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -2,6 +2,9 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out +// The name mangling for free function kernels currently does not work with PTX. +// UNSUPPORTED: cuda + #include #include #include From 190ac32df37e9109c433847e24fb82a280ec4ea0 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 25 Mar 2025 12:43:20 +0100 Subject: [PATCH 07/23] [SYCL] update unsupported list of tests --- .../test/e2e_test_requirements/no-unsupported-without-info.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp index 5c86386a52c8e..8a7224f494c14 100644 --- a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp +++ b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp @@ -54,7 +54,7 @@ // tests to match the required format and in that case you should just update // (i.e. reduce) the number and the list below. // -// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 282 +// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 283 // // List of improperly UNSUPPORTED tests. // Remove the CHECK once the test has been properly UNSUPPORTED. @@ -156,6 +156,7 @@ // CHECK-NEXT: ESIMD/slm_alloc_many_kernels_many_funcs.cpp // CHECK-NEXT: ESIMD/slm_alloc_many_kernels_one_func.cpp // CHECK-NEXT: ESIMD/slm_init_no_inline.cpp +// CHECK-NEXT: Experimental/free_functions/namespace.cpp // CHECK-NEXT: Graph/Explicit/buffer_copy_host2target.cpp // CHECK-NEXT: Graph/Explicit/buffer_copy_host2target_2d.cpp // CHECK-NEXT: Graph/Explicit/buffer_copy_host2target_offset.cpp From 0858e70ead56bf8585893c62e5bee1642a60399c Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Wed, 26 Mar 2025 10:52:39 +0100 Subject: [PATCH 08/23] [SYCL] fix typo --- sycl/test-e2e/Experimental/free_functions/namespace.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 4287f29d1340d..7e27144c5366c 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -93,7 +93,7 @@ void test_function_in_inline_ns(sycl::queue& q, sycl::context& ctxt) call_kernel_code(q, k_function_in_inline_ns); } -void test_function_in_anonimous_ns(sycl::queue& q, sycl::context& ctxt) { +void test_function_in_anonymous_ns(sycl::queue& q, sycl::context& ctxt) { // Get a kernel bundle that contains the free function kernel "function_in_anonymous_ns". auto exe_bndl = syclexp::get_kernel_bundle Date: Thu, 27 Mar 2025 18:23:59 +0100 Subject: [PATCH 09/23] [SYCL] rework free functions to use a separate entity --- clang/lib/Sema/SemaSYCL.cpp | 82 +++++++++++++------ ...ee_function_default_template_arguments.cpp | 3 +- 2 files changed, 58 insertions(+), 27 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 13dd36afbde0e..00e9bf5754d5d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6452,25 +6452,55 @@ static void PrintNSClosingBraces(raw_ostream &OS, const DeclContext *DC) { [](raw_ostream &OS, const NamespaceDecl *NS) {}, OS, DC); } -static bool insertFreeFunctionDeclaration(const PrintingPolicy &Policy, - const FunctionDecl *FD, - const std::string& Args, - raw_ostream &O) { - const auto *DC = FD->getDeclContext(); - bool NSInserted{false}; - if (DC) { - if (isa(DC)) { - PrintNamespaces(O, FD); - NSInserted = true; +class FreeFunctionPrinter { + raw_ostream &O; + const PrintingPolicy &Policy; + bool NSInserted = false; + +public: + FreeFunctionPrinter(raw_ostream &O, const PrintingPolicy &Policy) + : O(O), Policy(Policy) {} + + /// Emits the function declaration of a free function. + /// \param FD The function declaration to print. + /// \param Args The arguments of the function. + void printFreeFunctionDeclaration(const FunctionDecl *FD, + const std::string &Args) { + const DeclContext *DC = FD->getDeclContext(); + if (DC) { + // if function in namespace, print namespace + if (isa(DC)) { + PrintNamespaces(O, FD); + // Set flag to print closing braces for namespaces and namespace in shim + // function + NSInserted = true; + } + O << FD->getReturnType().getAsString() << " "; + O << FD->getNameAsString() << "(" << Args << ");"; + if (NSInserted) { + O << "\n"; + PrintNSClosingBraces(O, FD); + } + O << "\n"; } - O << FD->getReturnType().getAsString() << " "; - O << FD->getNameAsString() << "(" << Args << ");"; + } + + /// Emits free function shim function. + /// \param FD The function declaration to print. + /// \param ShimCounter The counter for the shim function. + /// \param ParmList The parameter list of the function. + void printFreeFunctionShim(const FunctionDecl *FD, const unsigned ShimCounter, + const std::string &ParmList) { + // Generate a shim function that returns the address of the free function. + O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n"; + O << " return (void (*)(" << ParmList << "))"; + if (NSInserted) { - PrintNSClosingBraces(O, FD); + PrintNamespaces(O, FD, true); } + O << FD->getIdentifier()->getName().data(); } - return NSInserted; -} +}; void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "// This is auto-generated SYCL integration header.\n"; @@ -6813,23 +6843,16 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // template arguments that match default template arguments while printing // template-ids, even if the source code doesn't reference them. Policy.EnforceDefaultTemplateArgs = true; - bool NSInserted{false}; + FreeFunctionPrinter FFPrinter(O, Policy); + // bool NSInserted{false}; if (FTD) { FTD->print(O, Policy); O << ";\n"; } else { - NSInserted = insertFreeFunctionDeclaration(Policy, K.SyclKernel, ParmListWithNames, O); - O << "\n"; + FFPrinter.printFreeFunctionDeclaration(K.SyclKernel, ParmListWithNames); } - // Generate a shim function that returns the address of the free function. - O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n"; - O << " return (void (*)(" << ParmList << "))"; - if (NSInserted) { - PrintNamespaces(O, K.SyclKernel, true); - } - - O << K.SyclKernel->getIdentifier()->getName().data(); + FFPrinter.printFreeFunctionShim(K.SyclKernel, ShimCounter, ParmList); if (FTD) { const TemplateArgumentList *TAL = K.SyclKernel->getTemplateSpecializationArgs(); @@ -6908,6 +6931,13 @@ bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) { } llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); emit(Out); + + int IntHeaderFD1 = 0; + std::string S{"/tmp/my-files/header.h"}; + llvm::sys::fs::openFileForWrite(S, IntHeaderFD1); + llvm::raw_fd_ostream Out1(IntHeaderFD1, true /*close in destructor*/); + emit(Out1); + return true; } diff --git a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp index 0d19924b1ad1a..f7da6e8f1f772 100644 --- a/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp +++ b/clang/test/CodeGenSYCL/free_function_default_template_arguments.cpp @@ -87,7 +87,8 @@ foo(Arg1 arg) { // CHECK-NEXT: } // CHECK: namespace ns { -// CHECK-NEXT: void simple(ns::Arg );} // namespace ns +// CHECK-NEXT: void simple(ns::Arg ); +// CHECK-NEXT: } // namespace ns // CHECK: static constexpr auto __sycl_shim1() { // CHECK-NEXT: return (void (*)(struct ns::Arg))ns::simple; // CHECK-NEXT: } From 2183658dd851273a81d06ffd66b66666b99b643c Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 10:34:42 +0100 Subject: [PATCH 10/23] [SYCL] Update code style --- clang/lib/Sema/SemaSYCL.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 00e9bf5754d5d..d2aa8a74392f2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6495,9 +6495,8 @@ class FreeFunctionPrinter { O << "static constexpr auto __sycl_shim" << ShimCounter << "() {\n"; O << " return (void (*)(" << ParmList << "))"; - if (NSInserted) { - PrintNamespaces(O, FD, true); - } + if (NSInserted) + PrintNamespaces(O, FD, /*isPrintNamesOnly=*/true); O << FD->getIdentifier()->getName().data(); } }; From 5dd5894a78b78eefe2a8db879b187122d2b117e7 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 11:06:11 +0100 Subject: [PATCH 11/23] [SYCL] update formating --- clang/lib/Sema/SemaSYCL.cpp | 5 +- .../Experimental/free_functions/namespace.cpp | 88 ++++++++++--------- 2 files changed, 49 insertions(+), 44 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d2aa8a74392f2..800007ee55246 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6797,10 +6797,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { for (ParmVarDecl *Param : K.SyclKernel->parameters()) { if (FirstParam) FirstParam = false; - else - { + else { ParmList += ", "; - ParmListWithNamesOstream << ", "; + ParmListWithNamesOstream << ", "; } Policy.SuppressTagKeyword = true; Param->getType().print(ParmListWithNamesOstream, Policy); diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 7e27144c5366c..44c89c77fba89 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -7,8 +7,8 @@ #include #include -#include #include +#include namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; @@ -23,34 +23,32 @@ void func_without_ns(float start, float *ptr) { } namespace free_functions::tests { - SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) - void function_in_ns(float start, float *ptr) { - size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 1); - } +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void function_in_ns(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 1); +} } // namespace free_functions::tests namespace free_functions::tests { inline namespace V1 { - SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) - void function_in_inline_ns(float start, float *ptr) - { - size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 2); - } -} // V1 +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void function_in_inline_ns(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 2); +} +} // namespace V1 } // namespace free_functions::tests - namespace { - SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) - void function_in_anonymous_ns(float start, float *ptr) { - size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 3); - } +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void function_in_anonymous_ns(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 3); } +} // namespace -static void call_kernel_code(sycl::queue& q, sycl::kernel& kernel) { +static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { float *ptr = sycl::malloc_shared(NUM, q); q.submit([&](sycl::handler &cgh) { cgh.set_args(3.14f, ptr); @@ -59,48 +57,56 @@ static void call_kernel_code(sycl::queue& q, sycl::kernel& kernel) { }).wait(); } -void test_function_without_ns(sycl::queue& q, sycl::context& ctxt) -{ - // Get a kernel bundle that contains the free function kernel "func_without_ns". +void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "func_without_ns". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); + syclexp::get_kernel_bundle(ctxt); // Get a kernel object for the "func_without_ns" function from that bundle. - sycl::kernel k_func_without_ns = exe_bndl.ext_oneapi_get_kernel(); + sycl::kernel k_func_without_ns = + exe_bndl.ext_oneapi_get_kernel(); call_kernel_code(q, k_func_without_ns); } -void test_function_in_ns(sycl::queue& q, sycl::context& ctxt) -{ - // Get a kernel bundle that contains the free function kernel "function_in_ns". +void test_function_in_ns(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "function_in_ns". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); // Get a kernel object for the "function_in_ns" function from that bundle. - sycl::kernel k_function_in_ns = exe_bndl.ext_oneapi_get_kernel(); + sycl::kernel k_function_in_ns = + exe_bndl.ext_oneapi_get_kernel(); call_kernel_code(q, k_function_in_ns); } -void test_function_in_inline_ns(sycl::queue& q, sycl::context& ctxt) -{ - // Get a kernel bundle that contains the free function kernel "function_in_inline_ns". +void test_function_in_inline_ns(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "function_in_inline_ns". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); - // Get a kernel object for the "function_in_inline_ns" function from that bundle. - sycl::kernel k_function_in_inline_ns = exe_bndl.ext_oneapi_get_kernel(); + // Get a kernel object for the "function_in_inline_ns" function from that + // bundle. + sycl::kernel k_function_in_inline_ns = exe_bndl.ext_oneapi_get_kernel< + free_functions::tests::function_in_inline_ns>(); call_kernel_code(q, k_function_in_inline_ns); } -void test_function_in_anonymous_ns(sycl::queue& q, sycl::context& ctxt) { - // Get a kernel bundle that contains the free function kernel "function_in_anonymous_ns". +void test_function_in_anonymous_ns(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "function_in_anonymous_ns". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); - // Get a kernel object for the "function_in_anonymous_ns" function from that bundle. - sycl::kernel k_function_in_anonymous_ns = exe_bndl.ext_oneapi_get_kernel(); + // Get a kernel object for the "function_in_anonymous_ns" function from that + // bundle. + sycl::kernel k_function_in_anonymous_ns = + exe_bndl.ext_oneapi_get_kernel(); call_kernel_code(q, k_function_in_anonymous_ns); } From 24999637dd293e4ec3c78cdee7e1d4235281cf45 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 11:22:41 +0100 Subject: [PATCH 12/23] [SYCL] fix includes in SemaSYCL --- clang/lib/Sema/SemaSYCL.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 800007ee55246..ebc175351fd17 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -15,10 +15,9 @@ #include "clang/AST/QualTypeNames.h" #include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" -#include "clang/AST/TemplateArgumentVisitor.h" -#include "clang/AST/Mangle.h" #include "clang/AST/SYCLKernelInfo.h" #include "clang/AST/StmtSYCL.h" +#include "clang/AST/TemplateArgumentVisitor.h" #include "clang/AST/TypeOrdering.h" #include "clang/AST/TypeVisitor.h" #include "clang/Analysis/CallGraph.h" From a5edf00eaf838a8e9b2a64c22fa7e65d6ef2120d Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 12:11:50 +0100 Subject: [PATCH 13/23] [SYCL] include missing header --- clang/lib/Sema/SemaSYCL.cpp | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ebc175351fd17..a4f3c714022d5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -26,6 +26,7 @@ #include "clang/Basic/Diagnostic.h" #include "clang/Basic/TargetInfo.h" #include "clang/Basic/Version.h" +#include "clang/AST/SYCLKernelInfo.h" #include "clang/Sema/Attr.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/ParsedAttr.h" @@ -6928,13 +6929,6 @@ bool SYCLIntegrationHeader::emit(StringRef IntHeaderName) { } llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); emit(Out); - - int IntHeaderFD1 = 0; - std::string S{"/tmp/my-files/header.h"}; - llvm::sys::fs::openFileForWrite(S, IntHeaderFD1); - llvm::raw_fd_ostream Out1(IntHeaderFD1, true /*close in destructor*/); - emit(Out1); - return true; } From 2bb7c2134180d94f97f88b9ffd039a2789755a1c Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 13:54:33 +0100 Subject: [PATCH 14/23] [SYCL] add unit tests of free function namespace support --- .../CodeGenSYCL/free_function_int_header.cpp | 248 ++++++++++++++++++ 1 file changed, 248 insertions(+) diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp index ccd2d73445cf8..f5c702ea2e35a 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp @@ -101,6 +101,47 @@ __attribute__((sycl_device)) void ff_8(sycl::work_group_memory) { } +// function in namespace +namespace free_functions { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_9(int start, int *ptr) { +} +} + +// function in nested namespace +namespace free_functions::tests { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_10(int start, int *ptr) { +} +} + +// function in inline namespace +namespace free_functions::tests { +inline namespace V1 { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_11(int start, int *ptr) { +} +} +} + +//function in anonymous namespace +namespace { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_12(int start, int *ptr) { +} +} + +// functions with the same name but in different namespaces +namespace free_functions { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_13(int start, int *ptr) { +} +} +namespace free_functions::tests { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_13(int start, int *ptr) { +} +} // CHECK: const char* const kernel_names[] = { // CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii @@ -112,6 +153,12 @@ void ff_8(sycl::work_group_memory) { // CHECK-NEXT: {{.*}}__sycl_kernel_ff_6I3Agg7DerivedEvT_T0_i // CHECK-NEXT: {{.*}}__sycl_kernel_ff_7ILi3EEv16KArgWithPtrArrayIXT_EE // CHECK-NEXT: {{.*}}__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions4ff_9EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests5ff_10EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests2V15ff_11EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel__GLOBAL__N_15ff_12EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5ff_13EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests5ff_13EiPi // CHECK-NEXT: "" // CHECK-NEXT: }; @@ -158,6 +205,30 @@ void ff_8(sycl::work_group_memory) { // CHECK: //--- _Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE // CHECK-NEXT: { kernel_param_kind_t::kind_work_group_memory, 8, 0 }, +// CHECK: //--- _ZN28__sycl_kernel_free_functions4ff_9EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5ff_13EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + // CHECK: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 }, // CHECK-NEXT: }; @@ -324,6 +395,135 @@ void ff_8(sycl::work_group_memory) { // CHECK-NEXT: }; // CHECK-NEXT: } +// CHECK: Definition of _ZN28__sycl_kernel_free_functions4ff_9EiPi as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: +// CHECK: namespace free_functions { +// CHECK-NEXT: void ff_9(int start, int * ptr); +// CHECK-NEXT: } // namespace free_functions + +// CHECK: static constexpr auto __sycl_shim10() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::ff_9; +// CHECK-NEXT: } + +// CHECK: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim10()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: + +// CHECK: namespace free_functions { +// CHECK-NEXT: namespace tests { +// CHECK-NEXT: void ff_10(int start, int * ptr); +// CHECK-NEXT: } // namespace tests +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim11() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_10; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim11()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim11()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: + +// CHECK: namespace free_functions { +// CHECK-NEXT: namespace tests { +// CHECK-NEXT: inline namespace V1 { +// CHECK-NEXT: void ff_11(int start, int * ptr); +// CHECK-NEXT: } // inline namespace V1 +// CHECK-NEXT: } // namespace tests +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim12() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::V1::ff_11; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim12()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim12()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: Definition of _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: + +// CHECK: namespace { +// CHECK-NEXT: void ff_12(int start, int * ptr); +// CHECK-NEXT: } // namespace +// CHECK: static constexpr auto __sycl_shim13() { +// CHECK-NEXT: return (void (*)(int, int *))ff_12; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim13()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim13()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: Definition of _ZN28__sycl_kernel_free_functions5ff_13EiPi as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: + +// CHECK: namespace free_functions { +// CHECK-NEXT: void ff_13(int start, int * ptr); +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim14() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::ff_13; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim14()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim14()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: + +// CHECK: namespace free_functions { +// CHECK-NEXT: namespace tests { +// CHECK-NEXT: void ff_13(int start, int * ptr); +// CHECK-NEXT: } // namespace tests +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim15() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_13; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim15()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim15()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + // CHECK: #include // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii @@ -397,3 +597,51 @@ void ff_8(sycl::work_group_memory) { // CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_8N4sycl3_V117work_group_memoryIiEE"}); // CHECK-NEXT: } // CHECK-NEXT: } + +// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions4ff_9EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions4ff_9EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim11()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_10EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim12()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests2V15ff_11EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: Definition of kernel_id of _ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim13()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN26__sycl_kernel__GLOBAL__N_15ff_12EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_13EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim14()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_13EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_13EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim15()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_13EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } From b3bd8ae6642e2fd0956863d91e5968e44e1090ba Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 14:16:34 +0100 Subject: [PATCH 15/23] [SYCL] fix includes --- clang/lib/Sema/SemaSYCL.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a4f3c714022d5..2f94bdfad80aa 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -26,7 +26,6 @@ #include "clang/Basic/Diagnostic.h" #include "clang/Basic/TargetInfo.h" #include "clang/Basic/Version.h" -#include "clang/AST/SYCLKernelInfo.h" #include "clang/Sema/Attr.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/ParsedAttr.h" From 64906f88ca6385d93b6cea1772ee1514aed2fd39 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 28 Mar 2025 16:58:40 +0100 Subject: [PATCH 16/23] [SYCL] fix post merge issue --- clang/test/CodeGenSYCL/free_function_int_header.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp index 8a4c94e356ca0..80fde849163ef 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp @@ -102,7 +102,6 @@ __attribute__((sycl_device)) void ff_8(sycl::work_group_memory) { } -<<<<<<< HEAD // function in namespace namespace free_functions { [[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] From 4cebf72fe61cbc5beb86366db638858413e2be6f Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 1 Apr 2025 11:35:51 +0200 Subject: [PATCH 17/23] [SYCL] check two functions with the same namespace --- .../Experimental/free_functions/namespace.cpp | 33 +++++++++++++++---- 1 file changed, 26 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 44c89c77fba89..8644de76af7c9 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -17,7 +17,7 @@ static constexpr size_t NUM = 1024; static constexpr size_t WGSIZE = 16; SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void func_without_ns(float start, float *ptr) { +void func(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); ptr[id] = start + static_cast(id); } @@ -28,6 +28,12 @@ void function_in_ns(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); ptr[id] = start + static_cast(id + 1); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void func(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id + 2); +} } // namespace free_functions::tests namespace free_functions::tests { @@ -59,14 +65,13 @@ static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { // Get a kernel bundle that contains the free function kernel - // "func_without_ns". + // "func". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); - // Get a kernel object for the "func_without_ns" function from that bundle. - sycl::kernel k_func_without_ns = - exe_bndl.ext_oneapi_get_kernel(); - call_kernel_code(q, k_func_without_ns); + // Get a kernel object for the "func" function from that bundle. + sycl::kernel k_func = exe_bndl.ext_oneapi_get_kernel(); + call_kernel_code(q, k_func); } void test_function_in_ns(sycl::queue &q, sycl::context &ctxt) { @@ -82,6 +87,19 @@ void test_function_in_ns(sycl::queue &q, sycl::context &ctxt) { call_kernel_code(q, k_function_in_ns); } +void test_func_in_ns_with_same_name(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "func". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "func" function from that bundle. + sycl::kernel k_func_in_ns = + exe_bndl.ext_oneapi_get_kernel(); + call_kernel_code(q, k_func_in_ns); +} + void test_function_in_inline_ns(sycl::queue &q, sycl::context &ctxt) { // Get a kernel bundle that contains the free function kernel // "function_in_inline_ns". @@ -118,5 +136,6 @@ int main() { test_function_in_ns(q, ctxt); test_function_in_inline_ns(q, ctxt); test_function_in_anonymous_ns(q, ctxt); + test_func_in_ns_with_same_name(q, ctxt); return 0; } From 157b39a303082f12bb501894870a7f2a744d613e Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 1 Apr 2025 11:47:36 +0200 Subject: [PATCH 18/23] [SYCL][E2E] fix formatting --- sycl/test-e2e/Experimental/free_functions/namespace.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 8644de76af7c9..0cafaa15096bf 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -67,8 +67,7 @@ void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { // Get a kernel bundle that contains the free function kernel // "func". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); + syclexp::get_kernel_bundle(ctxt); // Get a kernel object for the "func" function from that bundle. sycl::kernel k_func = exe_bndl.ext_oneapi_get_kernel(); call_kernel_code(q, k_func); From e3ff53abe485fcc0cdaf3cff697821370a487e51 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 1 Apr 2025 12:52:29 +0200 Subject: [PATCH 19/23] [SYCL][E2E] update free function integration header test --- .../CodeGenSYCL/free_function_int_header.cpp | 47 +++++++++---------- .../Experimental/free_functions/namespace.cpp | 2 +- 2 files changed, 23 insertions(+), 26 deletions(-) diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp index 80fde849163ef..87138427e43e6 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp @@ -423,30 +423,6 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK-NEXT: return (void (*)(int, int *))free_functions::ff_9; // CHECK-NEXT: } -// CHECK: namespace sycl { - -// CHECK: // Definition of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE as a free function kernel -// -// CHECK: Forward declarations of kernel and its argument types: -// CHECK: template class dynamic_work_group_memory; - -// CHECK: void ff_9(sycl::dynamic_work_group_memory); -// CHECK-NEXT: static constexpr auto __sycl_shim10() { -// CHECK-NEXT: return (void (*)(class sycl::dynamic_work_group_memory))ff_9; -// CHECK-NEXT: } -// CHECK-NEXT: namespace sycl { - -// CHECK-NEXT: template <> -// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim10()> { -// CHECK-NEXT: static constexpr bool value = true; -// CHECK-NEXT: }; -// CHECK-NEXT: template <> -// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim10()> { -// CHECK-NEXT: static constexpr bool value = true; -// CHECK-NEXT: }; -// CHECK-NEXT: } - - // CHECK: Definition of _ZN28__sycl_kernel_free_functions5tests5ff_10EiPi as a free function kernel // CHECK: Forward declarations of kernel and its argument types: @@ -556,6 +532,27 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK-NEXT: } +// CHECK: // Definition of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE as a free function kernel +// CHECK: Forward declarations of kernel and its argument types: +// CHECK-NEXT: namespace sycl { inline namespace _V1 { +// CHECK-NEXT: template class dynamic_work_group_memory; +// CHECK-NEXT: }} + +// CHECK: void ff_9(sycl::dynamic_work_group_memory ); +// CHECK-NEXT: static constexpr auto __sycl_shim16() { +// CHECK-NEXT: return (void (*)(class sycl::dynamic_work_group_memory))ff_9; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { + +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim16()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim16()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } // CHECK: #include @@ -684,7 +681,7 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK: // Definition of kernel_id of _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE // CHECK-NEXT: namespace sycl { // CHECK-NEXT: template <> -// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim10()>() { +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim16()>() { // CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE"}); // CHECK-NEXT: } // CHECK-NEXT: } diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 0cafaa15096bf..5b440e7df6fc1 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -67,7 +67,7 @@ void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { // Get a kernel bundle that contains the free function kernel // "func". auto exe_bndl = - syclexp::get_kernel_bundle(ctxt); + syclexp::get_kernel_bundle(ctxt); // Get a kernel object for the "func" function from that bundle. sycl::kernel k_func = exe_bndl.ext_oneapi_get_kernel(); call_kernel_code(q, k_func); From bdb3967bfbebbc0a7731817ffdb6857e7530b9ba Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 1 Apr 2025 17:29:19 +0200 Subject: [PATCH 20/23] [SYCL][E2E] do not update no-unsupported test --- .../Experimental/free_functions/namespace.cpp | 15 +++++++++++---- .../no-unsupported-without-info.cpp | 3 +-- 2 files changed, 12 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 5b440e7df6fc1..6134cac0c2e35 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -4,6 +4,7 @@ // The name mangling for free function kernels currently does not work with PTX. // UNSUPPORTED: cuda +// UNSUPPORTED-INTENDED: Not implemented yet for Nvidia/AMD backends. #include #include @@ -26,13 +27,13 @@ namespace free_functions::tests { SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void function_in_ns(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 1); + ptr[id] = start + static_cast(id); } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void func(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 2); + ptr[id] = start + static_cast(id); } } // namespace free_functions::tests @@ -41,7 +42,7 @@ inline namespace V1 { SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void function_in_inline_ns(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 2); + ptr[id] = start + static_cast(id); } } // namespace V1 } // namespace free_functions::tests @@ -50,7 +51,7 @@ namespace { SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void function_in_anonymous_ns(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); - ptr[id] = start + static_cast(id + 3); + ptr[id] = start + static_cast(id); } } // namespace @@ -61,6 +62,12 @@ static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { sycl::nd_range ndr{{NUM}, {WGSIZE}}; cgh.parallel_for(ndr, kernel); }).wait(); + // Check the result + for (size_t i = 0; i < NUM; ++i) { + const float expected = 3.14f + static_cast(i); + assert(ptr[i] == expected && + "Kernel execution did not produce the expected result"); + } } void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp index 8a7224f494c14..5c86386a52c8e 100644 --- a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp +++ b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp @@ -54,7 +54,7 @@ // tests to match the required format and in that case you should just update // (i.e. reduce) the number and the list below. // -// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 283 +// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 282 // // List of improperly UNSUPPORTED tests. // Remove the CHECK once the test has been properly UNSUPPORTED. @@ -156,7 +156,6 @@ // CHECK-NEXT: ESIMD/slm_alloc_many_kernels_many_funcs.cpp // CHECK-NEXT: ESIMD/slm_alloc_many_kernels_one_func.cpp // CHECK-NEXT: ESIMD/slm_init_no_inline.cpp -// CHECK-NEXT: Experimental/free_functions/namespace.cpp // CHECK-NEXT: Graph/Explicit/buffer_copy_host2target.cpp // CHECK-NEXT: Graph/Explicit/buffer_copy_host2target_2d.cpp // CHECK-NEXT: Graph/Explicit/buffer_copy_host2target_offset.cpp From cae876a493cd8a2545a44edc44f7dcba000eedd2 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 4 Apr 2025 12:17:06 +0200 Subject: [PATCH 21/23] [SYCL] add new tests to cover free fuction namespace support --- .../CodeGenSYCL/free_function_int_header.cpp | 257 ++++++++++++++++++ .../Experimental/free_functions/namespace.cpp | 63 ++++- 2 files changed, 308 insertions(+), 12 deletions(-) diff --git a/clang/test/CodeGenSYCL/free_function_int_header.cpp b/clang/test/CodeGenSYCL/free_function_int_header.cpp index 87138427e43e6..4fe57f8acf34b 100644 --- a/clang/test/CodeGenSYCL/free_function_int_header.cpp +++ b/clang/test/CodeGenSYCL/free_function_int_header.cpp @@ -78,6 +78,7 @@ __attribute__((sycl_device)) template void ff_6(Agg S1, Derived S2, int); constexpr int TestArrSize = 3; +constexpr int TestArrSizeAlias = 50; template struct KArgWithPtrArray { @@ -87,6 +88,18 @@ struct KArgWithPtrArray { constexpr int getArrSize() { return ArrSize; } }; +namespace free_functions { + template + struct KArgWithPtrArray { + float *data[ArrSize]; + float start[ArrSize]; + float end[ArrSize]; + constexpr int getArrSize() { return ArrSize; } + }; + + using AliasStruct = KArgWithPtrArray; +} + template [[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]] void ff_7(KArgWithPtrArray KArg) { @@ -149,7 +162,59 @@ __attribute__((sycl_device)) void ff_9(sycl::dynamic_work_group_memory) { } +typedef int TypedefType; +using AliasType = Derived; + +namespace free_functions::tests { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_14(TypedefType start, TypedefType *ptr) { +} +} + +namespace free_functions::tests { +typedef int NamespaceTypedefType; +using AliasType = Agg; +} + +namespace free_functions { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_15(free_functions::tests::NamespaceTypedefType start, free_functions::tests::NamespaceTypedefType *ptr) { +} +} + +namespace free_functions { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_16(free_functions::tests::AliasType start, free_functions::tests::AliasType *ptr) { +} +} + +namespace free_functions { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_17(AliasType start, AliasType *ptr) { +} +} + +namespace free_functions { + struct Agg { + int a; + float b; + }; +} +namespace free_functions::tests { +[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 0)]] +void ff_18(free_functions::Agg start, free_functions::Agg *ptr) { + ptr->a = start.a + 1; + ptr->b = start.b + 1.1f; +} +} + +[[__sycl_detail__::add_ir_attributes_function("sycl-single-task-kernel", 0)]] +void ff_19(free_functions::AliasStruct KArg) { + for (int j = 0; j < TestArrSizeAlias; j++) + for (int i = KArg.start[j]; i <= KArg.end[j]; i++) + KArg.data[j][i] = KArg.start[j] + KArg.end[j]; +} // CHECK: const char* const kernel_names[] = { // CHECK-NEXT: {{.*}}__sycl_kernel_ff_2Piii @@ -171,6 +236,13 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK-NEXT: {{.*}}__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests5ff_14EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5ff_15EiPi +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5ff_16E3AggPS0_ +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5ff_17E7DerivedPS0_ +// CHECK-NEXT: {{.*}}__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_ +// CHECK-NEXT: {{.*}}__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE + // CHECK-NEXT: "" // CHECK-NEXT: }; @@ -245,6 +317,28 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK: //--- _Z18__sycl_kernel_ff_9N4sycl3_V125dynamic_work_group_memoryIiEE // CHECK-NEXT: { kernel_param_kind_t::kind_dynamic_work_group_memory, 8, 0 }, +// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests5ff_14EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5ff_15EiPi +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 4 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_ +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 32, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 32 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_ +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 40, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 40 }, + +// CHECK: //--- _ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_ +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 8, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_pointer, 8, 8 }, + +// CHECK: //--- _Z19__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 800, 0 }, // CHECK: { kernel_param_kind_t::kind_invalid, -987654321, -987654321 }, // CHECK-NEXT: }; @@ -554,6 +648,130 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK-NEXT: }; // CHECK-NEXT: } +// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5tests5ff_14EiPi as a free function kernel +// CHECK: // Forward declarations of kernel and its argument types: +// CHECK: namespace free_functions { +// CHECK-NEXT: namespace tests { +// CHECK-NEXT: void ff_14(int start, int * ptr); +// CHECK-NEXT: } // namespace tests +// CHECK-NEXT: } // namespace free_functions + +// CHECK: static constexpr auto __sycl_shim17() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::tests::ff_14; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim17()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim17()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5ff_15EiPi as a free function kernel +// CHECK: // Forward declarations of kernel and its argument types: +// CHECK: namespace free_functions { +// CHECK-NEXT: void ff_15(int start, int * ptr); +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim18() { +// CHECK-NEXT: return (void (*)(int, int *))free_functions::ff_15; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim18()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim18()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_ as a free function kernel +// CHECK: // Forward declarations of kernel and its argument types: +// CHECK: namespace free_functions { +// CHECK-NEXT: void ff_16(Agg start, Agg * ptr); +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim19() { +// CHECK-NEXT: return (void (*)(struct Agg, struct Agg *))free_functions::ff_16; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim19()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim19()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_ as a free function kernel +// CHECK: // Forward declarations of kernel and its argument types: +// CHECK: namespace free_functions { +// CHECK-NEXT: void ff_17(Derived start, Derived * ptr); +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim20() { +// CHECK-NEXT: return (void (*)(struct Derived, struct Derived *))free_functions::ff_17; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim20()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim20()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: // Definition of _ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_ as a free function kernel +// CHECK: // Forward declarations of kernel and its argument types: +// CHECK-NEXT: namespace free_functions { +// CHECK-NEXT: struct Agg; +// CHECK-NEXT: } +// CHECK: namespace free_functions { +// CHECK-NEXT: namespace tests { +// CHECK-NEXT: void ff_18(free_functions::Agg start, free_functions::Agg * ptr); +// CHECK-NEXT: } // namespace tests +// CHECK-NEXT: } // namespace free_functions +// CHECK: static constexpr auto __sycl_shim21() { +// CHECK-NEXT: return (void (*)(struct free_functions::Agg, struct free_functions::Agg *))free_functions::tests::ff_18; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim21()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim21()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + +// CHECK: // Definition of _Z19__sycl_kernel_ff_19N14free_functions16KArgWithPtrArrayILi50EEE as a free function kernel +// CHECK: // Forward declarations of kernel and its argument types: +// CHECK-NEXT: namespace free_functions { +// CHECK-NEXT: template struct KArgWithPtrArray; +// CHECK-NEXT: } + +// CHECK: void ff_19(free_functions::KArgWithPtrArray<50> KArg); +// CHECK-NEXT: static constexpr auto __sycl_shim22() { +// CHECK-NEXT: return (void (*)(struct free_functions::KArgWithPtrArray<50>))ff_19; +// CHECK-NEXT: } +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_kernel<__sycl_shim22()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: template <> +// CHECK-NEXT: struct ext::oneapi::experimental::is_single_task_kernel<__sycl_shim22()> { +// CHECK-NEXT: static constexpr bool value = true; +// CHECK-NEXT: }; +// CHECK-NEXT: } + // CHECK: #include // CHECK: Definition of kernel_id of _Z18__sycl_kernel_ff_2Piii @@ -686,3 +904,42 @@ void ff_9(sycl::dynamic_work_group_memory) { // CHECK-NEXT: } // CHECK-NEXT: } +// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_14EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim17()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_14EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_15EiPi +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim18()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_15EiPi"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_ +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim19()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_16E3AggPS0_"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_ +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim20()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5ff_17E7DerivedPS0_"}); +// CHECK-NEXT: } +// CHECK-NEXT: } + +// CHECK: // Definition of kernel_id of _ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_ +// CHECK-NEXT: namespace sycl { +// CHECK-NEXT: template <> +// CHECK-NEXT: kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim21()>() { +// CHECK-NEXT: return sycl::detail::get_kernel_id_impl(std::string_view{"_ZN28__sycl_kernel_free_functions5tests5ff_18ENS_3AggEPS1_"}); +// CHECK-NEXT: } +// CHECK-NEXT: } diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index 6134cac0c2e35..d5a35a2af5eb8 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -55,19 +55,42 @@ void function_in_anonymous_ns(float start, float *ptr) { } } // namespace +struct TestClass { + float data; + TestClass(float d) : data(d) {} +}; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void function_with_test_class(float start, TestClass *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id].data = start + static_cast(id); +} + +template void check_result(T *ptr) { + for (size_t i = 0; i < NUM; ++i) { + const float expected = 3.14f + static_cast(i); + assert(ptr[i] == expected && + "Kernel execution did not produce the expected result"); + } +} + +template <> void check_result(TestClass *ptr) { + for (size_t i = 0; i < NUM; ++i) { + const float expected = 3.14f + static_cast(i); + assert(ptr[i].data == expected && + "Kernel execution did not produce the expected result"); + } +} + +template static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { - float *ptr = sycl::malloc_shared(NUM, q); + T *ptr = sycl::malloc_shared(NUM, q); q.submit([&](sycl::handler &cgh) { cgh.set_args(3.14f, ptr); sycl::nd_range ndr{{NUM}, {WGSIZE}}; cgh.parallel_for(ndr, kernel); }).wait(); - // Check the result - for (size_t i = 0; i < NUM; ++i) { - const float expected = 3.14f + static_cast(i); - assert(ptr[i] == expected && - "Kernel execution did not produce the expected result"); - } + check_result(ptr); } void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { @@ -77,7 +100,7 @@ void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { syclexp::get_kernel_bundle(ctxt); // Get a kernel object for the "func" function from that bundle. sycl::kernel k_func = exe_bndl.ext_oneapi_get_kernel(); - call_kernel_code(q, k_func); + call_kernel_code(q, k_func); } void test_function_in_ns(sycl::queue &q, sycl::context &ctxt) { @@ -90,7 +113,7 @@ void test_function_in_ns(sycl::queue &q, sycl::context &ctxt) { // Get a kernel object for the "function_in_ns" function from that bundle. sycl::kernel k_function_in_ns = exe_bndl.ext_oneapi_get_kernel(); - call_kernel_code(q, k_function_in_ns); + call_kernel_code(q, k_function_in_ns); } void test_func_in_ns_with_same_name(sycl::queue &q, sycl::context &ctxt) { @@ -103,7 +126,7 @@ void test_func_in_ns_with_same_name(sycl::queue &q, sycl::context &ctxt) { // Get a kernel object for the "func" function from that bundle. sycl::kernel k_func_in_ns = exe_bndl.ext_oneapi_get_kernel(); - call_kernel_code(q, k_func_in_ns); + call_kernel_code(q, k_func_in_ns); } void test_function_in_inline_ns(sycl::queue &q, sycl::context &ctxt) { @@ -117,7 +140,7 @@ void test_function_in_inline_ns(sycl::queue &q, sycl::context &ctxt) { // bundle. sycl::kernel k_function_in_inline_ns = exe_bndl.ext_oneapi_get_kernel< free_functions::tests::function_in_inline_ns>(); - call_kernel_code(q, k_function_in_inline_ns); + call_kernel_code(q, k_function_in_inline_ns); } void test_function_in_anonymous_ns(sycl::queue &q, sycl::context &ctxt) { @@ -131,7 +154,22 @@ void test_function_in_anonymous_ns(sycl::queue &q, sycl::context &ctxt) { // bundle. sycl::kernel k_function_in_anonymous_ns = exe_bndl.ext_oneapi_get_kernel(); - call_kernel_code(q, k_function_in_anonymous_ns); + call_kernel_code(q, k_function_in_anonymous_ns); +} + +void test_function_with_class(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "function_with_test_class". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "function_with_test_class" function from that + // bundle. + sycl::kernel k_function_with_test_class = + exe_bndl.template ext_oneapi_get_kernel(); + // call_kernel_code_with_class(q, k_function_with_test_class); + call_kernel_code(q, k_function_with_test_class); } int main() { @@ -143,5 +181,6 @@ int main() { test_function_in_inline_ns(q, ctxt); test_function_in_anonymous_ns(q, ctxt); test_func_in_ns_with_same_name(q, ctxt); + test_function_with_class(q, ctxt); return 0; } From afb8b5957598da88ff673d387ed4b62efc1240c2 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 4 Apr 2025 12:18:40 +0200 Subject: [PATCH 22/23] [SYCL] remove unused variable --- clang/lib/Sema/SemaSYCL.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4eb0c01e9363c..079ef4c3915ea 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6878,7 +6878,6 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { // template-ids, even if the source code doesn't reference them. Policy.EnforceDefaultTemplateArgs = true; FreeFunctionPrinter FFPrinter(O, Policy); - // bool NSInserted{false}; if (FTD) { FTD->print(O, Policy); O << ";\n"; From 09c7786d3522c46514de41691bcb6992b02fd3a4 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 4 Apr 2025 14:26:26 +0200 Subject: [PATCH 23/23] [SYCL][E2E] add new tests for free functions namespace support --- .../Experimental/free_functions/namespace.cpp | 61 +++++++++++++++++-- 1 file changed, 55 insertions(+), 6 deletions(-) diff --git a/sycl/test-e2e/Experimental/free_functions/namespace.cpp b/sycl/test-e2e/Experimental/free_functions/namespace.cpp index d5a35a2af5eb8..472f7bdb15f78 100644 --- a/sycl/test-e2e/Experimental/free_functions/namespace.cpp +++ b/sycl/test-e2e/Experimental/free_functions/namespace.cpp @@ -1,5 +1,5 @@ // REQUIRES: aspect-usm_shared_allocations -// RUN: %{build} -o %t.out +// RUN: %{build} %cxx_std_optionc++20 -o %t.out // RUN: %{run} %t.out // The name mangling for free function kernels currently does not work with PTX. @@ -10,6 +10,7 @@ #include #include #include +#include namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; @@ -60,13 +61,38 @@ struct TestClass { TestClass(float d) : data(d) {} }; +template struct TemplatedTestClass { + T data; + TemplatedTestClass(T d) : data(d) {} +}; + +using IntClassAlias = TemplatedTestClass; +using FloatClassAlias = TemplatedTestClass; + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void function_with_test_class(float start, TestClass *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); ptr[id].data = start + static_cast(id); } -template void check_result(T *ptr) { +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void function_with_int_alias_test_class(float start, IntClassAlias *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id].data = start + static_cast(id); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void function_with_float_alias_test_class(float start, FloatClassAlias *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id].data = start + static_cast(id); +} + +template +concept NumericType = std::is_arithmetic_v>; + +template + requires NumericType +void check_result(T *ptr) { for (size_t i = 0; i < NUM; ++i) { const float expected = 3.14f + static_cast(i); assert(ptr[i] == expected && @@ -74,9 +100,17 @@ template void check_result(T *ptr) { } } -template <> void check_result(TestClass *ptr) { +template +concept HasDataMemeber = requires(T t) { + { t.data } -> NumericType; +}; + +template + requires HasDataMemeber +void check_result(T *ptr) { + using DataType = decltype(ptr->data); for (size_t i = 0; i < NUM; ++i) { - const float expected = 3.14f + static_cast(i); + const DataType expected = 3.14f + static_cast(i); assert(ptr[i].data == expected && "Kernel execution did not produce the expected result"); } @@ -90,7 +124,7 @@ static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { sycl::nd_range ndr{{NUM}, {WGSIZE}}; cgh.parallel_for(ndr, kernel); }).wait(); - check_result(ptr); + check_result(ptr); } void test_function_without_ns(sycl::queue &q, sycl::context &ctxt) { @@ -168,10 +202,24 @@ void test_function_with_class(sycl::queue &q, sycl::context &ctxt) { // bundle. sycl::kernel k_function_with_test_class = exe_bndl.template ext_oneapi_get_kernel(); - // call_kernel_code_with_class(q, k_function_with_test_class); call_kernel_code(q, k_function_with_test_class); } +void test_fucntions_with_int_class_alias(sycl::queue &q, sycl::context &ctxt) { + // Get a kernel bundle that contains the free function kernel + // "function_with_int_alias_test_class". + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + + // Get a kernel object for the "function_with_int_alias_test_class" function + // from that bundle. + sycl::kernel k_function_with_int_alias_test_class = + exe_bndl + .template ext_oneapi_get_kernel(); + call_kernel_code(q, k_function_with_int_alias_test_class); +} + int main() { sycl::queue q; sycl::context ctxt = q.get_context(); @@ -182,5 +230,6 @@ int main() { test_function_in_anonymous_ns(q, ctxt); test_func_in_ns_with_same_name(q, ctxt); test_function_with_class(q, ctxt); + test_fucntions_with_int_class_alias(q, ctxt); return 0; }