From 02e89443764976d4ecd8f5f9d50892b7f1627e09 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 3 Dec 2019 14:23:15 +0300 Subject: [PATCH] [SYCL] Implement OpenCL kernel function generation All SYCL memory objects shared between host and device (buffers/images, these objects map to OpenCL buffers and images) must be accessed through special accessor classes. The "device" side implementation of these classes contain pointers to the device memory. As there is no way in OpenCL to pass structures with pointers inside as kernel arguments, all memory objects shared between host and device must be passed to the kernel as raw pointers. SYCL also has a special mechanism for passing kernel arguments from host to the device. In OpenCL kernel arguments are set by calling `clSetKernelArg` function for each kernel argument, meanwhile in SYCL all the kernel arguments are fields of "SYCL kernel function" which can be defined as a lambda function or a named function object and passed as an argument to SYCL function for invoking kernels (such as `parallel_for` or `single_task`). To facilitate the mapping of SYCL kernel data members to OpenCL kernel arguments and overcome OpenCL limitations we added the generation of an OpenCL kernel function inside the compiler. An OpenCL kernel function contains the body of the SYCL kernel function, receives OpenCL-like parameters and additionally does some manipulation to initialize SYCL kernel data members with these parameters. In some pseudo code the OpenCL kernel function can look like this: ``` // SYCL kernel is defined in SYCL headers: template __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) { // ... KernelFuncObj(); } // Generated OpenCL kernel function __kernel KernelName(global int* a) { KernelType KernelFuncObj; // Actually kernel function object declaration // doesn't have a name in AST. // Let the kernel function object have one captured field - accessor A. // We need to init it with global pointer from arguments: KernelFuncObj.A.__init(a); // Body of the SYCL kernel from SYCL headers: { KernelFuncObj(); } } ``` OpenCL kernel function is generated by the compiler inside the Sema using AST nodes. --- clang/include/clang/Sema/Sema.h | 13 + clang/lib/AST/ASTContext.cpp | 4 + clang/lib/CodeGen/CodeGenModule.cpp | 6 + clang/lib/Parse/ParseAST.cpp | 4 + clang/lib/Sema/CMakeLists.txt | 1 + clang/lib/Sema/SemaSYCL.cpp | 457 ++++++++++++++++++ .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 17 +- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 86 ++++ .../test/CodeGenSYCL/basic-opencl-kernel.cpp | 52 ++ clang/test/CodeGenSYCL/device-functions.cpp | 41 ++ clang/test/SemaSYCL/Inputs/sycl.hpp | 87 ++++ clang/test/SemaSYCL/accessors-targets.cpp | 41 ++ clang/test/SemaSYCL/basic-opencl-kernel.cpp | 74 +++ .../SemaSYCL/built-in-type-kernel-arg.cpp | 70 +++ clang/test/SemaSYCL/fake-accessors.cpp | 56 +++ clang/test/SemaSYCL/mangle-kernel.cpp | 29 ++ 16 files changed, 1035 insertions(+), 3 deletions(-) create mode 100644 clang/lib/Sema/SemaSYCL.cpp create mode 100644 clang/test/CodeGenSYCL/Inputs/sycl.hpp create mode 100644 clang/test/CodeGenSYCL/basic-opencl-kernel.cpp create mode 100644 clang/test/CodeGenSYCL/device-functions.cpp create mode 100644 clang/test/SemaSYCL/Inputs/sycl.hpp create mode 100644 clang/test/SemaSYCL/accessors-targets.cpp create mode 100644 clang/test/SemaSYCL/basic-opencl-kernel.cpp create mode 100644 clang/test/SemaSYCL/built-in-type-kernel-arg.cpp create mode 100644 clang/test/SemaSYCL/fake-accessors.cpp create mode 100644 clang/test/SemaSYCL/mangle-kernel.cpp diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index bb8fd4c9c8570..db734bac0351c 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -11632,6 +11632,19 @@ class Sema final { ConstructorDestructor, BuiltinFunction }; + +private: + /// Contains generated OpenCL kernel functions for SYCL. + SmallVector SYCLKernels; + +public: + void addSYCLKernel(Decl *D) { SYCLKernels.push_back(D); } + /// Access to SYCL kernels. + SmallVectorImpl &getSYCLKernels() { return SYCLKernels; } + + /// Constructs an OpenCL kernel using the KernelCaller function and adds it to + /// the SYCL device code. + void constructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); }; /// RAII object that enters a new expression evaluation context. diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 006eb1e0defb0..5c2d1c504ca51 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -10023,6 +10023,10 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) { if (D->hasAttr() || D->hasAttr()) return true; + // If SYCL, only kernels are required. + if (LangOpts.SYCLIsDevice && !(D->hasAttr())) + return false; + if (const auto *FD = dyn_cast(D)) { // Forward declarations aren't required. if (!FD->doesThisDeclarationHaveABody()) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 4959b80faec7a..0862644806c0c 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2474,6 +2474,12 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) { } } + if (LangOpts.SYCLIsDevice && Global->hasAttr() && + MustBeEmitted(Global)) { + addDeferredDeclToEmit(GD); + return; + } + // Ignore declarations, they will be emitted on their first use. if (const auto *FD = dyn_cast(Global)) { // Forward declarations are emitted lazily on first use. diff --git a/clang/lib/Parse/ParseAST.cpp b/clang/lib/Parse/ParseAST.cpp index 3efd893e499cd..cb59c7e58a738 100644 --- a/clang/lib/Parse/ParseAST.cpp +++ b/clang/lib/Parse/ParseAST.cpp @@ -168,6 +168,10 @@ void clang::ParseAST(Sema &S, bool PrintStats, bool SkipFunctionBodies) { for (Decl *D : S.WeakTopLevelDecls()) Consumer->HandleTopLevelDecl(DeclGroupRef(D)); + if (S.getLangOpts().SYCLIsDevice) + for (Decl *D : S.getSYCLKernels()) + Consumer->HandleTopLevelDecl(DeclGroupRef(D)); + Consumer->HandleTranslationUnit(S.getASTContext()); // Finalize the template instantiation observer chain. diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt index 89c3f6c47b497..7b5b601a2e86e 100644 --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -57,6 +57,7 @@ add_clang_library(clangSema SemaStmt.cpp SemaStmtAsm.cpp SemaStmtAttr.cpp + SemaSYCL.cpp SemaTemplate.cpp SemaTemplateDeduction.cpp SemaTemplateInstantiate.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp new file mode 100644 index 0000000000000..5662cb1bf717f --- /dev/null +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -0,0 +1,457 @@ +//===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// This implements Semantic Analysis for SYCL constructs. +//===----------------------------------------------------------------------===// + +#include "TreeTransform.h" +#include "clang/AST/AST.h" +#include "clang/AST/Mangle.h" +#include "clang/AST/QualTypeNames.h" +#include "clang/Sema/Initialization.h" +#include "clang/Sema/Sema.h" + +using namespace clang; + +using ParamDesc = std::tuple; + +/// Various utilities. +class Util { +public: + using DeclContextDesc = std::pair; + + /// Checks whether given clang type is a full specialization of the SYCL + /// accessor class. + static bool isSyclAccessorType(const QualType &Ty); + + /// Checks whether given clang type is declared in the given hierarchy of + /// declaration contexts. + /// \param Ty the clang type being checked + /// \param Scopes the declaration scopes leading from the type to the + /// translation unit (excluding the latter) + static bool matchQualifiedTypeName(const QualType &Ty, + ArrayRef Scopes); +}; + +static CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { + return (*Caller->param_begin())->getType()->getAsCXXRecordDecl(); +} + +class KernelBodyTransform : public TreeTransform { +public: + KernelBodyTransform(std::pair &MPair, + Sema &S) + : TreeTransform(S), MappingPair(MPair), SemaRef(S) {} + bool AlwaysRebuild() { return true; } + + ExprResult TransformDeclRefExpr(DeclRefExpr *DRE) { + auto Ref = dyn_cast(DRE->getDecl()); + if (Ref && Ref == MappingPair.first) { + auto NewDecl = MappingPair.second; + return DeclRefExpr::Create( + SemaRef.getASTContext(), DRE->getQualifierLoc(), + DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(), + NewDecl->getType(), DRE->getValueKind()); + } + return DRE; + } + +private: + std::pair MappingPair; + Sema &SemaRef; +}; + +static FunctionDecl * +CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, + ArrayRef ParamDescs) { + + DeclContext *DC = Context.getTranslationUnitDecl(); + QualType RetTy = Context.VoidTy; + SmallVector ArgTys; + + // Extract argument types from the descriptor array: + std::transform( + ParamDescs.begin(), ParamDescs.end(), std::back_inserter(ArgTys), + [](const ParamDesc &PD) -> QualType { return std::get<0>(PD); }); + FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + QualType FuncTy = Context.getFunctionType(RetTy, ArgTys, Info); + DeclarationName DN = DeclarationName(&Context.Idents.get(Name)); + + FunctionDecl *OpenCLKernel = FunctionDecl::Create( + Context, DC, SourceLocation(), SourceLocation(), DN, FuncTy, + Context.getTrivialTypeSourceInfo(RetTy), SC_None); + + llvm::SmallVector Params; + int i = 0; + for (const auto &PD : ParamDescs) { + auto P = ParmVarDecl::Create(Context, OpenCLKernel, SourceLocation(), + SourceLocation(), std::get<1>(PD), + std::get<0>(PD), std::get<2>(PD), SC_None, 0); + P->setScopeInfo(0, i++); + P->setIsUsed(); + Params.push_back(P); + } + OpenCLKernel->setParams(Params); + + OpenCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); + OpenCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); + OpenCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); + + // Add kernel to translation unit to see it in AST-dump + DC->addDecl(OpenCLKernel); + return OpenCLKernel; +} + +/// Return __init method +static CXXMethodDecl *getInitMethod(const CXXRecordDecl *CRD) { + CXXMethodDecl *InitMethod; + auto It = std::find_if(CRD->methods().begin(), CRD->methods().end(), + [](const CXXMethodDecl *Method) { + return Method->getNameAsString() == "__init"; + }); + InitMethod = (It != CRD->methods().end()) ? *It : nullptr; + return InitMethod; +} + +// Creates body for new OpenCL kernel. This body contains initialization of SYCL +// kernel object fields with kernel parameters and a little bit transformed body +// of the kernel caller function. +static CompoundStmt *CreateOpenCLKernelBody(Sema &S, + FunctionDecl *KernelCallerFunc, + DeclContext *KernelDecl) { + llvm::SmallVector BodyStmts; + CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc); + assert(LC && "Kernel object must be available"); + TypeSourceInfo *TSInfo = LC->isLambda() ? LC->getLambdaTypeInfo() : nullptr; + + // Create a local kernel object (lambda or functor) assembled from the + // incoming formal parameters. + auto KernelObjClone = VarDecl::Create( + S.Context, KernelDecl, SourceLocation(), SourceLocation(), + LC->getIdentifier(), QualType(LC->getTypeForDecl(), 0), TSInfo, SC_None); + Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), + SourceLocation(), SourceLocation()); + BodyStmts.push_back(DS); + auto KernelObjCloneRef = + DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), SourceLocation(), + KernelObjClone, false, DeclarationNameInfo(), + QualType(LC->getTypeForDecl(), 0), VK_LValue); + + auto KernelFuncDecl = cast(KernelDecl); + auto KernelFuncParam = + KernelFuncDecl->param_begin(); // Iterator to ParamVarDecl (VarDecl) + if (KernelFuncParam) { + llvm::SmallVector InitExprs; + InitializedEntity VarEntity = + InitializedEntity::InitializeVariable(KernelObjClone); + for (auto Field : LC->fields()) { + // Creates Expression for special SYCL object accessor. + // All special SYCL objects must have __init method, here we use it to + // initialize them. We create call of __init method and pass built kernel + // arguments as parameters to the __init method. + auto getExprForSpecialSYCLObj = [&](const QualType ¶mTy, + FieldDecl *Field, + const CXXRecordDecl *CRD, + Expr *Base) { + // All special SYCL objects must have __init method. + CXXMethodDecl *InitMethod = getInitMethod(CRD); + assert(InitMethod && + "__init method is expected."); + unsigned NumParams = InitMethod->getNumParams(); + llvm::SmallVector ParamDREs(NumParams); + auto KFP = KernelFuncParam; + for (size_t I = 0; I < NumParams; ++KFP, ++I) { + QualType ParamType = (*KFP)->getOriginalType(); + ParamDREs[I] = DeclRefExpr::Create( + S.Context, NestedNameSpecifierLoc(), SourceLocation(), *KFP, + false, DeclarationNameInfo(), ParamType, VK_LValue); + } + + if (NumParams) + std::advance(KernelFuncParam, NumParams - 1); + + DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); + // [kernel_obj].special_obj + auto SpecialObjME = MemberExpr::Create( + S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(), + SourceLocation(), Field, FieldDAP, + DeclarationNameInfo(Field->getDeclName(), SourceLocation()), + nullptr, Field->getType(), VK_LValue, OK_Ordinary, NOUR_None); + + // [kernel_obj].special_obj.__init + DeclAccessPair MethodDAP = DeclAccessPair::make(InitMethod, AS_none); + auto ME = MemberExpr::Create( + S.Context, SpecialObjME, false, SourceLocation(), + NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP, + DeclarationNameInfo(InitMethod->getDeclName(), SourceLocation()), + nullptr, InitMethod->getType(), VK_LValue, OK_Ordinary, NOUR_None); + + // Not referenced -> not emitted + S.MarkFunctionReferenced(SourceLocation(), InitMethod, true); + + QualType ResultTy = InitMethod->getReturnType(); + ExprValueKind VK = Expr::getValueKindForType(ResultTy); + ResultTy = ResultTy.getNonLValueExprType(S.Context); + + llvm::SmallVector ParamStmts; + const auto *Proto = cast(InitMethod->getType()); + S.GatherArgumentsForCall(SourceLocation(), InitMethod, Proto, 0, + ParamDREs, ParamStmts); + // [kernel_obj].special_obj.__init(_ValueType*, + // range, range, id) + CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( + S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); + BodyStmts.push_back(Call); + }; + + // Run through kernel object fields and add initialization for them using + // built kernel parameters. There are a several possible cases: + // - Kernel object field is a SYCL special object (SYCL accessor). + // These objects has a special initialization scheme - using + // __init method. + // - Kernel object field has a scalar type. In this case we should add + // simple initialization. + // - Kernel object field has a structure or class type. Same handling as + // a scalar. + QualType FieldType = Field->getType(); + CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl(); + InitializedEntity Entity = + InitializedEntity::InitializeMember(Field, &VarEntity); + if (Util::isSyclAccessorType(FieldType)) { + // Initialize kernel object field with the default constructor and + // construct a call of __init method. + InitializationKind InitKind = + InitializationKind::CreateDefault(SourceLocation()); + InitializationSequence InitSeq(S, Entity, InitKind, None); + ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, None); + InitExprs.push_back(MemberInit.get()); + getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef); + } else if (CRD || FieldType->isScalarType()) { + // If field has built-in or a structure/class type just initialize + // this field with corresponding kernel argument using copy + // initialization. + QualType ParamType = (*KernelFuncParam)->getOriginalType(); + Expr *DRE = + DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), + SourceLocation(), *KernelFuncParam, false, + DeclarationNameInfo(), ParamType, VK_LValue); + + InitializationKind InitKind = + InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); + InitializationSequence InitSeq(S, Entity, InitKind, DRE); + + ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, DRE); + InitExprs.push_back(MemberInit.get()); + + } else + llvm_unreachable("Unsupported field type"); + KernelFuncParam++; + } + Expr *ILE = new (S.Context) + InitListExpr(S.Context, SourceLocation(), InitExprs, SourceLocation()); + ILE->setType(QualType(LC->getTypeForDecl(), 0)); + KernelObjClone->setInit(ILE); + } + + // In the kernel caller function kernel object is a function parameter, so we + // need to replace all refs to this kernel oject with refs to our clone + // declared inside the kernel body. + Stmt *FunctionBody = KernelCallerFunc->getBody(); + ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); + + // DeclRefExpr with a valid source location but with decl which is not marked + // as used becomes invalid. + KernelObjClone->setIsUsed(); + std::pair MappingPair; + MappingPair.first = KernelObjParam; + MappingPair.second = KernelObjClone; + + // Function scope might be empty, so we do push + S.PushFunctionScope(); + KernelBodyTransform KBT(MappingPair, S); + Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); + BodyStmts.push_back(NewBody); + return CompoundStmt::Create(S.Context, BodyStmts, SourceLocation(), + SourceLocation()); +} + +/// Creates a kernel parameter descriptor +/// \param Src field declaration to construct name from +/// \param Ty the desired parameter type +/// \return the constructed descriptor +static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) { + ASTContext &Ctx = Src->getASTContext(); + std::string Name = (Twine("_arg_") + Src->getName()).str(); + return std::make_tuple(Ty, &Ctx.Idents.get(Name), + Ctx.getTrivialTypeSourceInfo(Ty)); +} + +// Creates list of kernel parameters descriptors using KernelObj (kernel +// object). Fields of kernel object must be initialized with SYCL kernel +// arguments so in the following function we extract types of kernel object +// fields and add it to the array with kernel parameters descriptors. +static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, + SmallVectorImpl &ParamDescs) { + auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) { + // Create a parameter descriptor and append it to the result + ParamDescs.push_back(makeParamDesc(Fld, ArgType)); + }; + + // Creates a parameter descriptor for SYCL special object - SYCL accessor. + // All special SYCL objects must have __init method. We extract types for + // kernel parameters from __init method parameters. We will use __init method + // and kernel parameters which we build here to initialize special objects in + // the kernel body. + auto createSpecialSYCLObjParamDesc = [&](const FieldDecl *Fld, + const QualType &ArgTy) { + const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); + assert(RecordDecl && "Special SYCL object must be of a record type"); + + CXXMethodDecl *InitMethod = getInitMethod(RecordDecl); + assert(InitMethod && "__init method is expected."); + unsigned NumParams = InitMethod->getNumParams(); + for (size_t I = 0; I < NumParams; ++I) { + ParmVarDecl *PD = InitMethod->getParamDecl(I); + CreateAndAddPrmDsc(Fld, PD->getType().getCanonicalType()); + } + }; + + // Run through kernel object fields and create corresponding kernel + // parameters descriptors. There are a several possible cases: + // - Kernel object field is a SYCL special object (SYCL accessor). + // These objects has a special initialization scheme - using + // __init method. + // - Kernel object field has a scalar type. In this case we should add + // kernel parameter with the same type. + // - Kernel object field has a structure or class type. Same handling as a + // scalar but we should check if this structure/class contains accessors + // and add parameter decriptor for them properly. + for (const auto *Fld : KernelObj->fields()) { + QualType ArgTy = Fld->getType(); + if (Util::isSyclAccessorType(ArgTy)) + createSpecialSYCLObjParamDesc(Fld, ArgTy); + else if (ArgTy->isStructureOrClassType()) + CreateAndAddPrmDsc(Fld, ArgTy); + else if (ArgTy->isScalarType()) + CreateAndAddPrmDsc(Fld, ArgTy); + else + llvm_unreachable("Unsupported kernel parameter type"); + } +} + +// Creates a mangled kernel name for given kernel name type +static std::string constructKernelName(QualType KernelNameType, + MangleContext &MC) { + SmallString<256> Result; + llvm::raw_svector_ostream Out(Result); + + MC.mangleTypeName(KernelNameType, Out); + return Out.str(); +} + +// Generates the OpenCL kernel using KernelCallerFunc (kernel caller +// function) defined is SYCL headers. +// Generated OpenCL kernel contains the body of the kernel caller function, +// receives OpenCL like parameters and additionally does some manipulation to +// initialize captured lambda/functor fields with these parameters. +// SYCL runtime marks kernel caller function with sycl_kernel attribute. +// To be able to generate OpenCL kernel from KernelCallerFunc we put +// the following requirements to the function which SYCL runtime can mark with +// sycl_kernel attribute: +// - Must be template function with at least two template parameters. +// First parameter must represent "unique kernel name" +// Second parameter must be the function object type +// - Must have only one function parameter - function object. +// +// Example of kernel caller function: +// template +// __attribute__((sycl_kernel)) void kernel_caller_function(KernelType +// KernelFuncObj) { +// KernelFuncObj(); +// } +// +// +void Sema::constructOpenCLKernel(FunctionDecl *KernelCallerFunc, + MangleContext &MC) { + CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc); + assert(LE && "invalid kernel caller"); + + // Build list of kernel arguments. + llvm::SmallVector ParamDescs; + buildArgTys(getASTContext(), LE, ParamDescs); + + // Extract name from kernel caller parameters and mangle it. + const TemplateArgumentList *TemplateArgs = + KernelCallerFunc->getTemplateSpecializationArgs(); + assert(TemplateArgs && "No template argument info"); + QualType KernelNameType = TypeName::getFullyQualifiedType( + TemplateArgs->get(0).getAsType(), getASTContext(), true); + std::string Name = constructKernelName(KernelNameType, MC); + + FunctionDecl *OpenCLKernel = + CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); + + // Let's copy source location of a functor/lambda to emit nicer diagnostics. + OpenCLKernel->setLocation(LE->getLocation()); + + CompoundStmt *OpenCLKernelBody = + CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); + OpenCLKernel->setBody(OpenCLKernelBody); + + addSYCLKernel(OpenCLKernel); +} + +// ----------------------------------------------------------------------------- +// Utility class methods +// ----------------------------------------------------------------------------- + +bool Util::isSyclAccessorType(const QualType &Ty) { + static std::array Scopes = { + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{clang::Decl::Kind::ClassTemplateSpecialization, + "accessor"}}; + return matchQualifiedTypeName(Ty, Scopes); +} + +bool Util::matchQualifiedTypeName(const QualType &Ty, + ArrayRef Scopes) { + // The idea: check the declaration context chain starting from the type + // itself. At each step check the context is of expected kind + // (namespace) and name. + const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); + + if (!RecTy) + return false; // only classes/structs supported + const auto *Ctx = dyn_cast(RecTy); + StringRef Name = ""; + + for (const auto &Scope : llvm::reverse(Scopes)) { + clang::Decl::Kind DK = Ctx->getDeclKind(); + + if (DK != Scope.first) + return false; + + switch (DK) { + case clang::Decl::Kind::ClassTemplateSpecialization: + // ClassTemplateSpecializationDecl inherits from CXXRecordDecl + case clang::Decl::Kind::CXXRecord: + Name = cast(Ctx)->getName(); + break; + case clang::Decl::Kind::Namespace: + Name = cast(Ctx)->getName(); + break; + default: + llvm_unreachable("matchQualifiedTypeName: decl kind not supported"); + } + if (Name != Scope.second) + return false; + Ctx = Ctx->getParent(); + } + return Ctx->isTranslationUnit(); +} + diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index e9cb9f89e0a26..d585eec094318 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -17,6 +17,7 @@ #include "clang/AST/DependentDiagnostic.h" #include "clang/AST/Expr.h" #include "clang/AST/ExprCXX.h" +#include "clang/AST/Mangle.h" #include "clang/AST/PrettyDeclStackTrace.h" #include "clang/AST/TypeLoc.h" #include "clang/Sema/Initialization.h" @@ -5610,6 +5611,8 @@ NamedDecl *Sema::FindInstantiatedDecl(SourceLocation Loc, NamedDecl *D, /// Performs template instantiation for all implicit template /// instantiations we have seen until this point. void Sema::PerformPendingInstantiations(bool LocalOnly) { + std::unique_ptr MangleCtx( + getASTContext().createMangleContext()); while (!PendingLocalImplicitInstantiations.empty() || (!LocalOnly && !PendingInstantiations.empty())) { PendingImplicitInstantiation Inst; @@ -5628,17 +5631,25 @@ void Sema::PerformPendingInstantiations(bool LocalOnly) { TSK_ExplicitInstantiationDefinition; if (Function->isMultiVersion()) { getASTContext().forEachMultiversionedFunctionVersion( - Function, [this, Inst, DefinitionRequired](FunctionDecl *CurFD) { + Function, [this, Inst, DefinitionRequired, + MangleCtx = move(MangleCtx)](FunctionDecl *CurFD) { InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, CurFD, true, DefinitionRequired, true); - if (CurFD->isDefined()) + if (CurFD->isDefined()) { CurFD->setInstantiationIsPending(false); + if (getLangOpts().SYCLIsDevice && + CurFD->hasAttr()) + constructOpenCLKernel(CurFD, *MangleCtx); + } }); } else { InstantiateFunctionDefinition(/*FIXME:*/ Inst.second, Function, true, DefinitionRequired, true); - if (Function->isDefined()) + if (Function->isDefined()) { + if (getLangOpts().SYCLIsDevice && Function->hasAttr()) + constructOpenCLKernel(Function, *MangleCtx); Function->setInstantiationIsPending(false); + } } continue; } diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp new file mode 100644 index 0000000000000..8c01c99c76b44 --- /dev/null +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -0,0 +1,86 @@ +#pragma once + +namespace cl { +namespace sycl { +namespace access { + +enum class target { + global_buffer = 2014, + constant_buffer, + local, + image, + host_buffer, + host_image, + image_array +}; + +enum class mode { + read = 1024, + write, + read_write, + discard_write, + discard_read_write, + atomic +}; + +enum class placeholder { + false_t, + true_t +}; + +enum class address_space : int { + private_space = 0, + global_space, + constant_space, + local_space +}; +} // namespace access + +template +struct id { + template + id(T... args) {} // fake constructor +private: + // Some fake field added to see using of id arguments in the + // kernel wrapper + int Data; +}; + +template +struct range { + template + range(T... args) {} // fake constructor +private: + // Some fake field added to see using of range arguments in the + // kernel wrapper + int Data; +}; + +template +struct _ImplT { + range AccessRange; + range MemRange; + id Offset; +}; + +template +class accessor { + +public: + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImplT impl; + +private: + void __init(__attribute__((address_space(1))) dataT *Ptr, + range AccessRange, + range MemRange, id Offset) {} +}; + +} // namespace sycl +} // namespace cl diff --git a/clang/test/CodeGenSYCL/basic-opencl-kernel.cpp b/clang/test/CodeGenSYCL/basic-opencl-kernel.cpp new file mode 100644 index 0000000000000..842b113fc2252 --- /dev/null +++ b/clang/test/CodeGenSYCL/basic-opencl-kernel.cpp @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-unknown -std=c++11 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s + +// This test checks that compiler generates correct opencl kernel for basic +// case. + +#include "sycl.hpp" + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + cl::sycl::accessor accessorA; + kernel( + [=]() { + accessorA.use(); + }); + return 0; +} + +// CHECK: define spir_kernel void @{{.*}}kernel_function +// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]], +// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]], +// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET:%[a-zA-Z0-9_]+]]) +// Check alloca for pointer argument +// CHECK: [[MEM_ARG]].addr = alloca i32 addrspace(1)* +// Check lambda object alloca +// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon +// Check allocas for ranges +// CHECK: [[ARANGE:%agg.tmp.*]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MRANGE:%agg.tmp.*]] = alloca %"struct.cl::sycl::range" +// CHECK: [[OID:%agg.tmp.*]] = alloca %"struct.cl::sycl::id" +// +// Check store of kernel pointer argument to alloca +// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)** [[MEM_ARG]].addr, align 8 + +// Check for default constructor of accessor +// CHECK: call spir_func {{.*}}accessor + +// Check accessor GEP +// CHECK: [[ACCESSOR:%[a-zA-Z0-9_]+]] = getelementptr inbounds %class.anon, %class.anon* [[ANON]], i32 0, i32 0 + +// Check load from kernel pointer argument alloca +// CHECK: [[MEM_LOAD:%[a-zA-Z0-9_]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[MEM_ARG]].addr + +// Check accessor __init method call +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor"* [[ACCESSOR]], i32 addrspace(1)* [[MEM_LOAD]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[ARANGE]], %"struct.cl::sycl::range"* byval({{.*}}) align 4 [[MRANGE]], %"struct.cl::sycl::id"* byval({{.*}}) align 4 [[OID]]) + +// Check lambda "()" operator call +// CHECK-OLD: call spir_func void @{{.*}}(%class.anon* [[ANON]]) diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp new file mode 100644 index 0000000000000..8cd6a78288c1b --- /dev/null +++ b/clang/test/CodeGenSYCL/device-functions.cpp @@ -0,0 +1,41 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s + +template +T bar(T arg); + +void foo() { + int a = 1 + 1 + bar(1); +} + +template +T bar(T arg) { + return arg; +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} + +// Make sure that definitions for the types not used in SYCL kernels are not +// emitted +// CHECK-NOT: %struct.A +// CHECK-NOT: @a = {{.*}} %struct.A +struct A { + int x = 10; +} a; + +int main() { + a.x = 8; + kernel_single_task([]() { foo(); }); + return 0; +} + +// baz is not called from the SYCL kernel, so it must not be emitted +// CHECK-NOT: define {{.*}} @{{.*}}baz +void baz() {} + +// CHECK-LABEL: define spir_kernel void @{{.*}}test_kernel +// CHECK-LABEL: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%class.anon* %this) +// CHECK-LABEL: define spir_func void @{{.*}}foo +// CHECK-LABEL: define linkonce_odr spir_func i32 @{{.*}}bar diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp new file mode 100644 index 0000000000000..a9ba08e87ce18 --- /dev/null +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -0,0 +1,87 @@ +#pragma once + +namespace cl { +namespace sycl { +namespace access { + +enum class target { + global_buffer = 2014, + constant_buffer, + local, + image, + host_buffer, + host_image, + image_array +}; + +enum class mode { + read = 1024, + write, + read_write, + discard_write, + discard_read_write, + atomic +}; + +enum class placeholder { false_t, + true_t }; + +enum class address_space : int { + private_space = 0, + global_space, + constant_space, + local_space +}; +} // namespace access + +template +struct range { +}; + +template +struct id { +}; + +template +struct _ImplT { + range AccessRange; + range MemRange; + id Offset; +}; + +template +struct DeviceValueType; + +template +struct DeviceValueType { + using type = __attribute__((address_space(1))) dataT; +}; + +template +struct DeviceValueType { + using type = __attribute__((address_space(2))) dataT; +}; + +template +struct DeviceValueType { + using type = __attribute__((address_space(3))) dataT; +}; + +template +class accessor { + +public: + void use(void) const {} + void use(void *) const {} + _ImplT impl; + +private: + using PtrType = typename DeviceValueType::type *; + void __init(PtrType Ptr, range AccessRange, + range MemRange, id Offset) {} +}; + +} // namespace sycl +} // namespace cl diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp new file mode 100644 index 0000000000000..d4e662c54476a --- /dev/null +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -0,0 +1,41 @@ +// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s + +// This test checks that compiler generates correct OpenCL kernel arguments for +// different accessors targets. + +#include + +using namespace cl::sycl; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + + accessor + local_acc; + accessor + global_acc; + accessor + constant_acc; + kernel( + [=]() { + local_acc.use(); + }); + kernel( + [=]() { + global_acc.use(); + }); + kernel( + [=]() { + constant_acc.use(); + }); +} +// CHECK: {{.*}}use_local 'void (__attribute__((address_space(3))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_global 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_constant 'void (__attribute__((address_space(2))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' diff --git a/clang/test/SemaSYCL/basic-opencl-kernel.cpp b/clang/test/SemaSYCL/basic-opencl-kernel.cpp new file mode 100644 index 0000000000000..f509b30a3629a --- /dev/null +++ b/clang/test/SemaSYCL/basic-opencl-kernel.cpp @@ -0,0 +1,74 @@ +// RUN: %clang_cc1 -std=c++11 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s + +// This test checks that compiler generates correct OpenCL kernel for basic +// case. + +#include + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + cl::sycl::accessor acc; + kernel( + [=]() { + acc.use(); + }); +} + +// Check declaration of the kernel + +// CHECK: FunctionDecl {{.*}}kernel{{.*}} 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' + +// Check parameters of the kernel + +// CHECK: ParmVarDecl {{.*}} used [[_arg_Mem:[0-9a-zA-Z_]+]] '__attribute__((address_space(1))) int *' +// CHECK: ParmVarDecl {{.*}} used [[_arg_AccessRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_MemRange:[0-9a-zA-Z_]+]] 'cl::sycl::range<1>' +// CHECK: ParmVarDecl {{.*}} used [[_arg_Offset:[0-9a-zA-Z_]+]] 'cl::sycl::id<1>' + +// Check body of the kernel + +// Check lambda declaration inside the kernel + +// CHECK: DeclStmt +// CHECK-NEXT: VarDecl {{.*}} used '(lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' + +// Check accessor initialization + +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor':'cl::sycl::accessor' lvalue . +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' lvalue Var + +// CHECK-NEXT: ImplicitCastExpr {{.*}} +// CHECK-NEXT: DeclRefExpr {{.*}} '__attribute__((address_space(1))) int *' lvalue ParmVar {{.*}} '[[_arg_Mem]]' '__attribute__((address_space(1))) int *' + +// CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_AccessRange]]' 'cl::sycl::range<1>' + +// CHECK-NEXT: CXXConstructExpr {{.*}} 'range<1>':'cl::sycl::range<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::range<1>' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::range<1>' lvalue ParmVar {{.*}} '[[_arg_MemRange]]' 'cl::sycl::range<1>' + +// CHECK-NEXT: CXXConstructExpr {{.*}} 'id<1>':'cl::sycl::id<1>' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const cl::sycl::id<1>' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'cl::sycl::id<1>' lvalue ParmVar {{.*}} '[[_arg_Offset]]' 'cl::sycl::id<1>' + +// Check that body of the kernel caller function is included into kernel + +// CHECK: CompoundStmt {{.*}} +// CHECK-NEXT: CXXOperatorCallExpr {{.*}} 'void' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'void (*)() const' +// CHECK-NEXT: DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const' +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'const (lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-opencl-kernel.cpp{{.*}})' lvalue Var + +// Check kernel's attributes + +// CHECK: OpenCLKernelAttr {{.*}} Implicit +// CHECK: AsmLabelAttr {{.*}} Implicit "{{.*}}kernel{{.*}}" +// CHECK: ArtificialAttr {{.*}} Implicit diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp new file mode 100644 index 0000000000000..4e78277837f05 --- /dev/null +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -0,0 +1,70 @@ +// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s + +// This test checks that compiler generates correct initialization for arguments +// that have struct or built-in type inside the OpenCL kernel + +#include + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +struct test_struct { + int data; +}; + +void test(const int some_const) { + kernel( + [=]() { + int a = some_const; + }); +} + +int main() { + int data = 5; + test_struct s; + s.data = data; + kernel( + [=]() { + int kernel_data = data; + }); + kernel( + [=]() { + test_struct k_s; + k_s = s; + }); + const int some_const = 10; + test(some_const); + return 0; +} +// Check kernel parameters +// CHECK: FunctionDecl {{.*}}kernel_const{{.*}} 'void (const int)' +// CHECK: ParmVarDecl {{.*}} used _arg_ 'const int' + +// Check that lambda field of const built-in type is initialized +// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: DeclRefExpr {{.*}} 'const int' lvalue ParmVar {{.*}} '_arg_' 'const int' + +// Check kernel parameters +// CHECK: {{.*}}kernel_int{{.*}} 'void (int)' +// CHECK: ParmVarDecl {{.*}} used _arg_ 'int' + +// Check that lambda field of built-in type is initialized +// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' +// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_' 'int' + +// Check kernel parameters +// CHECK: {{.*}}kernel_struct{{.*}} 'void (test_struct)' +// CHECK: ParmVarDecl {{.*}} used _arg_ 'test_struct' + +// Check that lambda field of struct type is initialized +// CHECK: VarDecl {{.*}}'(lambda at {{.*}}built-in-type-kernel-arg.cpp{{.*}})' +// CHECK-NEXT: InitListExpr +// CHECK-NEXT: CXXConstructExpr {{.*}}'test_struct'{{.*}}void (const test_struct &) +// CHECK-NEXT: ImplicitCastExpr {{.*}}'const test_struct' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} 'test_struct' lvalue ParmVar {{.*}} '_arg_' 'test_struct' diff --git a/clang/test/SemaSYCL/fake-accessors.cpp b/clang/test/SemaSYCL/fake-accessors.cpp new file mode 100644 index 0000000000000..0095903c8a6c8 --- /dev/null +++ b/clang/test/SemaSYCL/fake-accessors.cpp @@ -0,0 +1,56 @@ +// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s + +#include + +namespace foo { +namespace cl { +namespace sycl { +class accessor { +public: + int field; +}; +} // namespace sycl +} // namespace cl +} // namespace foo + +class accessor { +public: + int field; +}; + +typedef cl::sycl::accessor + MyAccessorTD; + +using MyAccessorA = cl::sycl::accessor; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + foo::cl::sycl::accessor acc = {1}; + accessor acc1 = {1}; + + cl::sycl::accessor accessorA; + cl::sycl::accessor accessorB; + cl::sycl::accessor accessorC; + kernel( + [=]() { + accessorA.use((void*)(acc.field + acc1.field)); + }); + kernel( + [=]() { + accessorB.use((void*)(acc.field + acc1.field)); + }); + kernel( + [=]() { + accessorC.use((void*)(acc.field + acc1.field)); + }); + return 0; +} +// CHECK: fake_accessors 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias 'void (__attribute__((address_space(1))) int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) diff --git a/clang/test/SemaSYCL/mangle-kernel.cpp b/clang/test/SemaSYCL/mangle-kernel.cpp new file mode 100644 index 0000000000000..4cbdfd56bc5d9 --- /dev/null +++ b/clang/test/SemaSYCL/mangle-kernel.cpp @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -triple spir64-unknown-unknown-unknown -I %S/Inputs -I %S/../Headers/Inputs/include/ -fsycl-is-device -ast-dump %s | FileCheck %s --check-prefix=CHECK-64 +// RUN: %clang_cc1 -triple spir-unknown-unknown-unknown -I %S/Inputs -I %S/../Headers/Inputs/include/ -fsycl-is-device -ast-dump %s | FileCheck %s --check-prefix=CHECK-32 +#include +#include + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +template +class SimpleVadd; + +int main() { + kernel>( + [=](){}); + + kernel>( + [=](){}); + + kernel>( + [=](){}); + return 0; +} + +// CHECK: _ZTS10SimpleVaddIiE +// CHECK: _ZTS10SimpleVaddIdE +// CHECK-64: _ZTS10SimpleVaddImE +// CHECK-32: _ZTS10SimpleVaddIjE