From 77887820ad076a489ed35e4be0bd1c0f840d8feb Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 9 Jun 2020 15:00:20 +0300 Subject: [PATCH 1/4] [SYCL] Implement OpenCL kernel function generation Summary: 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. Reviewers: bader, Naghasan, ABataev, keryell Subscribers: agozillon, mgorny, yaxunl, jfb, ebevhan, Anastasia, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D71016 --- 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/SemaSYCL.cpp | 443 ++++++++++++++++++ .../lib/Sema/SemaTemplateInstantiateDecl.cpp | 17 +- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 86 ++++ .../CodeGenSYCL/address-space-conversions.cpp | 14 +- .../CodeGenSYCL/address-space-deduction.cpp | 22 +- .../CodeGenSYCL/address-space-mangling.cpp | 32 +- .../test/CodeGenSYCL/basic-kernel-wrapper.cpp | 57 +++ clang/test/CodeGenSYCL/device-functions.cpp | 41 ++ .../CodeGenSYCL/field-annotate-addr-space.cpp | 14 +- clang/test/SemaSYCL/Inputs/sycl.hpp | 87 ++++ clang/test/SemaSYCL/accessors-targets.cpp | 33 ++ .../SemaSYCL/built-in-type-kernel-arg.cpp | 70 +++ clang/test/SemaSYCL/fake-accessors.cpp | 56 +++ clang/test/SemaSYCL/mangle-kernel.cpp | 29 ++ 18 files changed, 1010 insertions(+), 18 deletions(-) create mode 100644 clang/test/CodeGenSYCL/Inputs/sycl.hpp create mode 100644 clang/test/CodeGenSYCL/basic-kernel-wrapper.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/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 43ce5d983217f..41e32c2433ac3 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -13116,6 +13116,19 @@ class Sema final { /// Adds Callee to DeviceCallGraph if we don't know if its caller will be /// codegen'ed yet. bool checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee); + +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 294cc20f76c53..4c7a9e6df02ba 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -11072,6 +11072,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 59f3e02705713..2cdd98571817f 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -2960,6 +2960,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 01510e8caf3b7..eee790a109f60 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/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 815463307ecc7..9876182499705 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -8,7 +8,11 @@ // 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" #include "clang/Sema/SemaDiagnostic.h" @@ -48,3 +52,442 @@ bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) { return DiagKind != SemaDiagnosticBuilder::K_Immediate && DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack; } + +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(), FPOptionsOverride()); + 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 std::string(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 27ac2cd08f2a8..8bcec92993a68 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -18,6 +18,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/Basic/SourceManager.h" @@ -6270,6 +6271,8 @@ NamedDecl *Sema::FindInstantiatedDecl(SourceLocation Loc, NamedDecl *D, /// instantiations we have seen until this point. void Sema::PerformPendingInstantiations(bool LocalOnly) { std::deque delayedPCHInstantiations; + std::unique_ptr MangleCtx( + getASTContext().createMangleContext()); while (!PendingLocalImplicitInstantiations.empty() || (!LocalOnly && !PendingInstantiations.empty())) { PendingImplicitInstantiation Inst; @@ -6288,17 +6291,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); + } } // Definition of a PCH-ed template declaration may be available only in the TU. if (!LocalOnly && LangOpts.PCHInstantiateTemplates && diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp new file mode 100644 index 0000000000000..56908fe5f9a3a --- /dev/null +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -0,0 +1,86 @@ +#pragma once + +inline 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__((opencl_global)) dataT *Ptr, + range AccessRange, + range MemRange, id Offset) {} +}; + +} // namespace sycl +} // namespace cl diff --git a/clang/test/CodeGenSYCL/address-space-conversions.cpp b/clang/test/CodeGenSYCL/address-space-conversions.cpp index 3732c4a1b889b..b49ee9d8c1505 100644 --- a/clang/test/CodeGenSYCL/address-space-conversions.cpp +++ b/clang/test/CodeGenSYCL/address-space-conversions.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -I%S/Inputs -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s void bar(int &Data) {} // CHECK-DAG: define{{.*}} spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32 addrspace(4)* align 4 dereferenceable(4) % void bar2(int &Data) {} @@ -136,3 +136,15 @@ void usages() { // CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS3iEvT_(i32 addrspace(3)* % // CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPU3AS0iEvT_(i32* % // CHECK-DAG: define linkonce_odr spir_func void @_Z4tmplIPiEvT_(i32 addrspace(4)* % + +#include "sycl.hpp" + +int main() { + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + cgh.single_task([=]() { + usages(); + }); + }); + return 0; +} diff --git a/clang/test/CodeGenSYCL/address-space-deduction.cpp b/clang/test/CodeGenSYCL/address-space-deduction.cpp index 3453d18787c26..03206c5d4a24b 100644 --- a/clang/test/CodeGenSYCL/address-space-deduction.cpp +++ b/clang/test/CodeGenSYCL/address-space-deduction.cpp @@ -1,7 +1,10 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -I%S/Inputs -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s -// CHECK-LABEL: @_Z4testv( +// Validates SYCL deduction rules compliance. +// See clang/docs/SYCLSupport.rst#address-space-handling for the details. + +// CHECK-LABEL: define {{.*}} @_Z4testv( // CHECK-NEXT: entry: // CHECK-NEXT: [[I:%.*]] = alloca i32, align 4 // CHECK-NEXT: [[PPTR:%.*]] = alloca i32 addrspace(4)*, align 8 @@ -87,7 +90,8 @@ // CHECK-NEXT: store i8 addrspace(4)* getelementptr inbounds ([21 x i8], [21 x i8] addrspace(4)* addrspacecast ([21 x i8] addrspace(1)* @.str.1 to [21 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* addrspace(4)* [[SELECT_STR_TRIVIAL2_ASCAST]], align 8 // CHECK-NEXT: ret void // -void test() { + void test() { + static const int foo = 0x42; @@ -127,3 +131,15 @@ void test() { const char *select_str_trivial2 = false ? str : "Another hello world!"; (void)select_str_trivial2; } + +#include "sycl.hpp" + +int main() { + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + cgh.single_task([=]() { + test(); + }); + }); + return 0; +} diff --git a/clang/test/CodeGenSYCL/address-space-mangling.cpp b/clang/test/CodeGenSYCL/address-space-mangling.cpp index 76feec552fa2b..ceab9414a75c3 100644 --- a/clang/test/CodeGenSYCL/address-space-mangling.cpp +++ b/clang/test/CodeGenSYCL/address-space-mangling.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=SPIR -// RUN: %clang_cc1 -triple x86_64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=X86 +// RUN: %clang_cc1 -I%S/Inputs -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=SPIR +// RUN: %clang_cc1 -I%S/Inputs -triple x86_64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s --check-prefix=X86 // REQUIRES: x86-registered-target @@ -8,15 +8,15 @@ void foo(__attribute__((opencl_local)) int *); void foo(__attribute__((opencl_private)) int *); void foo(int *); -// SPIR: declare spir_func void @_Z3fooPU3AS1i(i32 addrspace(1)*) #1 -// SPIR: declare spir_func void @_Z3fooPU3AS3i(i32 addrspace(3)*) #1 -// SPIR: declare spir_func void @_Z3fooPU3AS0i(i32*) #1 -// SPIR: declare spir_func void @_Z3fooPi(i32 addrspace(4)*) #1 +// SPIR: declare spir_func void @_Z3fooPU3AS1i(i32 addrspace(1)*) +// SPIR: declare spir_func void @_Z3fooPU3AS3i(i32 addrspace(3)*) +// SPIR: declare spir_func void @_Z3fooPU3AS0i(i32*) +// SPIR: declare spir_func void @_Z3fooPi(i32 addrspace(4)*) -// X86: declare void @_Z3fooPU8SYglobali(i32*) #1 -// X86: declare void @_Z3fooPU7SYlocali(i32*) #1 -// X86: declare void @_Z3fooPU9SYprivatei(i32*) #1 -// X86: declare void @_Z3fooPi(i32*) #1 +// X86: declare void @_Z3fooPU8SYglobali(i32*) +// X86: declare void @_Z3fooPU7SYlocali(i32*) +// X86: declare void @_Z3fooPU9SYprivatei(i32*) +// X86: declare void @_Z3fooPi(i32*) void test() { __attribute__((opencl_global)) int *glob; @@ -28,3 +28,15 @@ void test() { foo(priv); foo(def); } + +#include "sycl.hpp" + +int main() { + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + cgh.single_task([=]() { + test(); + }); + }); + return 0; +} diff --git a/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp new file mode 100644 index 0000000000000..f118c931e171a --- /dev/null +++ b/clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp @@ -0,0 +1,57 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s + +// This test checks that compiler generates correct kernel wrapper for basic +// case. + +#include "Inputs/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: [[ANONALLOCA:%[0-9]+]] = alloca %class.anon +// CHECK: [[ARANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range" +// CHECK: [[MRANGEA:%agg.tmp.*]] = alloca %"struct.cl::sycl::range" +// CHECK: [[OIDA:%agg.tmp.*]] = alloca %"struct.cl::sycl::id" +// CHECK: [[ANON:%[0-9]+]] = addrspacecast %class.anon* [[ANONALLOCA]] to %class.anon addrspace(4)* +// CHECK: [[ARANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[ARANGEA]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[MRANGET:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range"* [[MRANGEA]] to %"struct.cl::sycl::range" addrspace(4)* +// CHECK: [[OIDT:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id"* [[OIDA]] to %"struct.cl::sycl::id" addrspace(4)* +// +// Check store of kernel pointer argument to alloca +// CHECK: store i32 addrspace(1)* [[MEM_ARG]], i32 addrspace(1)* addrspace(4)* [[MEM_ARG]].addr.ascast, 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 addrspace(4)* [[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)* addrspace(4)* [[MEM_ARG]].addr.ascast + +// Check accessor __init method call +// CHECK: [[ARANGE:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[ARANGET]] to %"struct.cl::sycl::range"* +// CHECK: [[MRANGE:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::range" addrspace(4)* [[MRANGET]] to %"struct.cl::sycl::range"* +// CHECK: [[OID:%agg.tmp.*]] = addrspacecast %"struct.cl::sycl::id" addrspace(4)* [[OIDT]] to %"struct.cl::sycl::id"* +// CHECK: call spir_func void @{{.*}}__init{{.*}}(%"class.cl::sycl::accessor" addrspace(4)* {{[^,]*}} [[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: call spir_func void @{{.*}}(%class.anon addrspace(4)* {{[^,]*}}) diff --git a/clang/test/CodeGenSYCL/device-functions.cpp b/clang/test/CodeGenSYCL/device-functions.cpp new file mode 100644 index 0000000000000..c8fa8729e29a6 --- /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 dso_local spir_kernel void @{{.*}}test_kernel +// CHECK-LABEL: define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon addrspace(4)* {{.*}}%this) +// CHECK-LABEL: define dso_local spir_func void @{{.*}}foo +// CHECK-LABEL: define linkonce_odr spir_func i32 @{{.*}}bar diff --git a/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp b/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp index 86d6f9a8a9e32..0c0c306ee361d 100644 --- a/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp +++ b/clang/test/CodeGenSYCL/field-annotate-addr-space.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -I%S/Inputs -triple spir64 -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s // CHECK: [[ANNOT:.+]] = private unnamed_addr constant {{.*}}c"my_annotation\00" @@ -17,3 +17,15 @@ void foo(int *b) { // CHECK: bitcast i8 addrspace(4)* %[[CALL]] to i32 addrspace(4)* addrspace(4)* f.a = b; } + +#include "sycl.hpp" + +int main() { + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + cgh.single_task([=]() { + foo(nullptr); + }); + }); + return 0; +} diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp new file mode 100644 index 0000000000000..9663a895243f8 --- /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__((opencl_global)) dataT; +}; + +template +struct DeviceValueType { + using type = __attribute__((opencl_global)) const dataT; +}; + +template +struct DeviceValueType { + using type = __attribute__((opencl_local)) 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..ad6a6106c8f1a --- /dev/null +++ b/clang/test/SemaSYCL/accessors-targets.cpp @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 -fsycl-is-device -ast-dump %s | FileCheck %s + +// This test checks that compiler generates correct OpenCL kernel arguments for +// different accessors targets. + +#include "Inputs/sycl.hpp" + +using namespace cl::sycl; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + + accessor + local_acc; + accessor + global_acc; + kernel( + [=]() { + local_acc.use(); + }); + kernel( + [=]() { + global_acc.use(); + }); +} +// CHECK: {{.*}}use_local 'void (__local int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' +// CHECK: {{.*}}use_global 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>)' 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..acce120e49f68 --- /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 (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_typedef 'void (__global int *, cl::sycl::range<1>, cl::sycl::range<1>, cl::sycl::id<1>, foo::cl::sycl::accessor, accessor) +// CHECK: accessor_alias 'void (__global 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 From b20e9f48b57781f4d48af8d3554d418a72eb1461 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Thu, 21 Jan 2021 14:06:04 +0300 Subject: [PATCH 2/4] TODO: revert this change after CodeGen part is updated The CodeGen patch is outdated and it doesn't work with SYCL 2020 syntax used in this test. --- clang/test/CodeGenSYCL/convergent.cpp | 4 ++-- clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/convergent.cpp b/clang/test/CodeGenSYCL/convergent.cpp index 779f1592da0e0..659e257912379 100644 --- a/clang/test/CodeGenSYCL/convergent.cpp +++ b/clang/test/CodeGenSYCL/convergent.cpp @@ -8,8 +8,8 @@ void foo() { int a = 1; } -template -__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) { +template +[[clang::sycl_kernel]] void kernel_single_task(KernelType kernelFunc) { kernelFunc(); } diff --git a/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp b/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp index e6efa92716fbc..46f0533f0b784 100644 --- a/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp +++ b/clang/test/CodeGenSYCL/debug-info-kernel-variables.cpp @@ -14,7 +14,7 @@ #define KERNEL __attribute__((sycl_kernel)) template -KERNEL void parallel_for(const KernelType &KernelFunc) { +KERNEL void parallel_for(KernelType KernelFunc) { KernelFunc(); } From e343510bc1172da4c2d8b671ee1a4305e118ffe2 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 20 Jan 2021 12:41:20 +0300 Subject: [PATCH 3/4] [SYCL] Enable Open CL types required for implementing the SYCL headers. This patch implements a few of the OpenCL types for SYCL, however doesn't bother handling semantic analysis as these are not intended to be used anywhere but in the SYCL implementation. Signed-off-by: Erich Keane Differential Revision: https://reviews.llvm.org/D77220 --- clang/lib/AST/ASTContext.cpp | 9 + clang/lib/CodeGen/CodeGenModule.cpp | 2 +- clang/lib/Sema/Sema.cpp | 11 + clang/lib/Sema/SemaInit.cpp | 5 +- clang/lib/Sema/SemaSYCL.cpp | 17 +- clang/test/CodeGenSYCL/Inputs/sycl.hpp | 350 +++++++++++++++++- clang/test/CodeGenSYCL/image_accessor.cpp | 111 ++++++ clang/test/CodeGenSYCL/sampler.cpp | 63 ++++ clang/test/SemaSYCL/Inputs/sycl.hpp | 115 +++++- .../test/SemaSYCL/accessors-targets-image.cpp | 74 ++++ clang/test/SemaSYCL/sampler.cpp | 32 ++ 11 files changed, 781 insertions(+), 8 deletions(-) create mode 100644 clang/test/CodeGenSYCL/image_accessor.cpp create mode 100644 clang/test/CodeGenSYCL/sampler.cpp create mode 100644 clang/test/SemaSYCL/accessors-targets-image.cpp create mode 100644 clang/test/SemaSYCL/sampler.cpp diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 4c7a9e6df02ba..f07842ca358d0 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1423,6 +1423,15 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass); InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel); + if (LangOpts.SYCLIsDevice) { + InitBuiltinType(OCLSamplerTy, BuiltinType::OCLSampler); + InitBuiltinType(OCLEventTy, BuiltinType::OCLEvent); +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + InitBuiltinType(SingletonId, BuiltinType::Id); +#include "clang/Basic/OpenCLImageTypes.def" +#undef IMAGE_TYPE + } + if (LangOpts.OpenCL) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ InitBuiltinType(SingletonId, BuiltinType::Id); diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 2cdd98571817f..5851976555445 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -139,7 +139,7 @@ CodeGenModule::CodeGenModule(ASTContext &C, const HeaderSearchOptions &HSO, if (LangOpts.ObjC) createObjCRuntime(); - if (LangOpts.OpenCL) + if (LangOpts.OpenCL || LangOpts.SYCLIsDevice) createOpenCLRuntime(); if (LangOpts.OpenMP) createOpenMPRuntime(); diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index a2b8f475aa8c9..4470ea2a4df29 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -317,6 +317,17 @@ void Sema::Initialize() { addImplicitTypedef("size_t", Context.getSizeType()); } + if (getLangOpts().SYCLIsDevice) { + addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy); + addImplicitTypedef("__ocl_event_t", Context.OCLEventTy); +#define SEMA_STRINGIZE(s) #s +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + addImplicitTypedef(SEMA_STRINGIZE(__ocl_##ImgType##_##Suffix##_t), \ + Context.SingletonId); +#include "clang/Basic/OpenCLImageTypes.def" +#undef SEMA_STRINGIZE + } + // Initialize predefined OpenCL types and supported extensions and (optional) // core features. if (getLangOpts().OpenCL) { diff --git a/clang/lib/Sema/SemaInit.cpp b/clang/lib/Sema/SemaInit.cpp index 119a90deb9c26..dbe132eadfcce 100644 --- a/clang/lib/Sema/SemaInit.cpp +++ b/clang/lib/Sema/SemaInit.cpp @@ -5579,9 +5579,10 @@ static bool TryOCLSamplerInitialization(Sema &S, InitializationSequence &Sequence, QualType DestType, Expr *Initializer) { - if (!S.getLangOpts().OpenCL || !DestType->isSamplerT() || + if ((!S.getLangOpts().OpenCL && !S.getLangOpts().SYCLIsDevice) || + !DestType->isSamplerT() || (!Initializer->isIntegerConstantExpr(S.Context) && - !Initializer->getType()->isSamplerT())) + !Initializer->getType()->isSamplerT())) return false; Sequence.AddOCLSamplerInitStep(DestType); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9876182499705..8dcd756b3e4b3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -63,6 +63,9 @@ class Util { /// 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 a full specialization of the SYCL + /// sampler class. + static bool isSyclSamplerType(const QualType &Ty); /// Checks whether given clang type is declared in the given hierarchy of /// declaration contexts. @@ -258,7 +261,8 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S, CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl(); InitializedEntity Entity = InitializedEntity::InitializeMember(Field, &VarEntity); - if (Util::isSyclAccessorType(FieldType)) { + if (Util::isSyclAccessorType(FieldType) || + Util::isSyclSamplerType(FieldType)) { // Initialize kernel object field with the default constructor and // construct a call of __init method. InitializationKind InitKind = @@ -369,7 +373,7 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, // and add parameter decriptor for them properly. for (const auto *Fld : KernelObj->fields()) { QualType ArgTy = Fld->getType(); - if (Util::isSyclAccessorType(ArgTy)) + if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) createSpecialSYCLObjParamDesc(Fld, ArgTy); else if (ArgTy->isStructureOrClassType()) CreateAndAddPrmDsc(Fld, ArgTy); @@ -455,6 +459,15 @@ bool Util::isSyclAccessorType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isSyclSamplerType(const QualType &Ty) { + static const std::array Scopes = { + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{clang::Decl::Kind::CXXRecord, + "sampler"}}; + return matchQualifiedTypeName(Ty, Scopes); +} + bool Util::matchQualifiedTypeName(const QualType &Ty, ArrayRef Scopes) { // The idea: check the declaration context chain starting from the type diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 56908fe5f9a3a..9100ada1d06a1 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -1,7 +1,32 @@ #pragma once +#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) + +// Dummy runtime classes to model SYCL API. inline namespace cl { namespace sycl { +struct sampler_impl { +#ifdef __SYCL_DEVICE_ONLY__ + __ocl_sampler_t m_Sampler; +#endif +}; + +class sampler { + struct sampler_impl impl; +#ifdef __SYCL_DEVICE_ONLY__ + void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; } +#endif + +public: + void use(void) const {} +}; + +template +class group { +public: + group() = default; // fake constructor +}; + namespace access { enum class target { @@ -36,6 +61,39 @@ enum class address_space : int { }; } // namespace access +namespace property { + +enum prop_type { + use_host_ptr = 0, + use_mutex, + context_bound, + enable_profiling, + base_prop +}; + +struct property_base { + virtual prop_type type() const = 0; +}; +} // namespace property + +class property_list { +public: + template + property_list(propertyTN... props) {} + + template + bool has_property() const { return true; } + + template + propertyT get_property() const { + return propertyT{}; + } + + bool operator==(const property_list &rhs) const { return false; } + + bool operator!=(const property_list &rhs) const { return false; } +}; + template struct id { template @@ -56,6 +114,10 @@ struct range { int Data; }; +template +struct nd_range { +}; + template struct _ImplT { range AccessRange; @@ -77,10 +139,294 @@ class accessor { _ImplT impl; private: - void __init(__attribute__((opencl_global)) dataT *Ptr, - range AccessRange, + void __init(__attribute__((opencl_global)) dataT *Ptr, range AccessRange, range MemRange, id Offset) {} }; +template +struct opencl_image_type; + +#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \ + template <> \ + struct opencl_image_type { \ + using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t; \ + }; + +#define IMAGETY_READ_3_DIM_IMAGE \ + IMAGETY_DEFINE(1, read, ro, image, ) \ + IMAGETY_DEFINE(2, read, ro, image, ) \ + IMAGETY_DEFINE(3, read, ro, image, ) + +#define IMAGETY_WRITE_3_DIM_IMAGE \ + IMAGETY_DEFINE(1, write, wo, image, ) \ + IMAGETY_DEFINE(2, write, wo, image, ) \ + IMAGETY_DEFINE(3, write, wo, image, ) + +#define IMAGETY_READ_2_DIM_IARRAY \ + IMAGETY_DEFINE(1, read, ro, image_array, array_) \ + IMAGETY_DEFINE(2, read, ro, image_array, array_) + +#define IMAGETY_WRITE_2_DIM_IARRAY \ + IMAGETY_DEFINE(1, write, wo, image_array, array_) \ + IMAGETY_DEFINE(2, write, wo, image_array, array_) + +IMAGETY_READ_3_DIM_IMAGE +IMAGETY_WRITE_3_DIM_IMAGE + +IMAGETY_READ_2_DIM_IARRAY +IMAGETY_WRITE_2_DIM_IARRAY + +template +struct _ImageImplT { +#ifdef __SYCL_DEVICE_ONLY__ + typename opencl_image_type::type MImageObj; +#else + range AccessRange; + range MemRange; + id Offset; +#endif +}; + +template +class accessor { +public: + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImageImplT impl; +#ifdef __SYCL_DEVICE_ONLY__ + void __init(typename opencl_image_type::type ImageObj) { impl.MImageObj = ImageObj; } +#endif +}; + +template +class accessor { +public: + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImageImplT impl; +}; + +// TODO: Add support for image_array accessor. +// template +//class accessor + +class kernel {}; +class context {}; +class device {}; +class event {}; + +class queue { +public: + template + event submit(T cgf) { return event{}; } + + void wait() {} + void wait_and_throw() {} + void throw_asynchronous() {} +}; + +class auto_name {}; +template +struct get_kernel_name_t { + using name = Name; +}; +template +struct get_kernel_name_t { + using name = Type; +}; +#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) +template +ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) { + kernelFunc(); +} + +template +ATTR_SYCL_KERNEL void +kernel_parallel_for(KernelType KernelFunc) { + KernelFunc(id()); +} + +template +ATTR_SYCL_KERNEL void +kernel_parallel_for_work_group(KernelType KernelFunc) { + KernelFunc(group()); +} + +class handler { +public: + template + void parallel_for(range numWorkItems, KernelType kernelFunc) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_parallel_for(kernelFunc); +#else + kernelFunc(); +#endif + } + + template + void parallel_for_work_group(range numWorkGroups, range WorkGroupSize, KernelType kernelFunc) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_parallel_for_work_group(kernelFunc); +#else + group G; + kernelFunc(G); +#endif + } + + template + void single_task(KernelType kernelFunc) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_single_task(kernelFunc); +#else + kernelFunc(); +#endif + } +}; + +class stream { +public: + stream(unsigned long BufferSize, unsigned long MaxStatementSize, + handler &CGH) {} + + void __init() {} + + void __finalize() {} +}; + +template +const stream& operator<<(const stream &S, T&&) { + return S; +} + +template +class buffer { +public: + using value_type = T; + using reference = value_type &; + using const_reference = const value_type &; + using allocator_type = AllocatorT; + + template + buffer(ParamTypes... args) {} // fake constructor + + buffer(const range &bufferRange, + const property_list &propList = {}) {} + + buffer(T *hostData, const range &bufferRange, + const property_list &propList = {}) {} + + buffer(const T *hostData, const range &bufferRange, + const property_list &propList = {}) {} + + buffer(const buffer &rhs) = default; + + buffer(buffer &&rhs) = default; + + buffer &operator=(const buffer &rhs) = default; + + buffer &operator=(buffer &&rhs) = default; + + ~buffer() = default; + + range get_range() const { return range{}; } + + template + accessor + get_access(handler &commandGroupHandler) { + return accessor{}; + } + + template + accessor + get_access() { + return accessor{}; + } + + template + void set_final_data(Destination finalData = nullptr) {} +}; + +enum class image_channel_order : unsigned int { + a, + r, + rx, + rg, + rgx, + ra, + rgb, + rgbx, + rgba, + argb, + bgra, + intensity, + luminance, + abgr +}; + +enum class image_channel_type : unsigned int { + snorm_int8, + snorm_int16, + unorm_int8, + unorm_int16, + unorm_short_565, + unorm_short_555, + unorm_int_101010, + signed_int8, + signed_int16, + signed_int32, + unsigned_int8, + unsigned_int16, + unsigned_int32, + fp16, + fp32 +}; + +template +class image { +public: + image(image_channel_order Order, image_channel_type Type, + const range &Range, const property_list &PropList = {}) {} + + /* -- common interface members -- */ + + image(const image &rhs) = default; + + image(image &&rhs) = default; + + image &operator=(const image &rhs) = default; + + image &operator=(image &&rhs) = default; + + ~image() = default; + + template + accessor + get_access(handler &commandGroupHandler) { + return accessor{}; + } + + template + accessor + get_access() { + return accessor{}; + } +}; + } // namespace sycl } // namespace cl diff --git a/clang/test/CodeGenSYCL/image_accessor.cpp b/clang/test/CodeGenSYCL/image_accessor.cpp new file mode 100644 index 0000000000000..5204c55db544d --- /dev/null +++ b/clang/test/CodeGenSYCL/image_accessor.cpp @@ -0,0 +1,111 @@ +// RUN: %clang_cc1 -triple spir64 -I %S/Inputs -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o %t.ll +// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-1DRO +// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-2DRO +// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DRO +// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-1DWO +// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-2DWO +// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO +// +// CHECK-1DRO: %opencl.image1d_ro_t = type opaque +// CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_ro_t* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t* %{{[0-9]+}}) +// +// CHECK-2DRO: %opencl.image2d_ro_t = type opaque +// CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_ro_t* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t* %{{[0-9]+}}) +// +// CHECK-3DRO: %opencl.image3d_ro_t = type opaque +// CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_ro_t* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t* %{{[0-9]+}}) +// +// CHECK-1DWO: %opencl.image1d_wo_t = type opaque +// CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_wo_t* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t* %{{[0-9]+}}) +// +// CHECK-2DWO: %opencl.image2d_wo_t = type opaque +// CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_wo_t* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t* %{{[0-9]+}}) +// +// CHECK-3DWO: %opencl.image3d_wo_t = type opaque +// CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_wo_t* [[IMAGE_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t* %{{[0-9]+}}) +// +// TODO: Add tests for the image_array opencl datatype support. +#include "sycl.hpp" + +int main() { + + { + cl::sycl::image<1> MyImage1d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<1>(3)); + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = MyImage1d.get_access(cgh); + + cgh.single_task([=]() { + Acc.use(); + }); + }); + } + + { + cl::sycl::image<2> MyImage2d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<2>(3, 2)); + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = MyImage2d.get_access(cgh); + + cgh.single_task([=]() { + Acc.use(); + }); + }); + } + + { + cl::sycl::image<3> MyImage3d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<3>(3, 2, 4)); + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = MyImage3d.get_access(cgh); + + cgh.single_task([=]() { + Acc.use(); + }); + }); + } + + { + cl::sycl::image<1> MyImage1d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<1>(3)); + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = MyImage1d.get_access(cgh); + + cgh.single_task([=]() { + Acc.use(); + }); + }); + } + + { + cl::sycl::image<2> MyImage2d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<2>(3, 2)); + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = MyImage2d.get_access(cgh); + + cgh.single_task([=]() { + Acc.use(); + }); + }); + } + + { + cl::sycl::image<3> MyImage3d(cl::sycl::image_channel_order::rgbx, cl::sycl::image_channel_type::unorm_short_565, cl::sycl::range<3>(3, 2, 4)); + cl::sycl::queue Q; + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = MyImage3d.get_access(cgh); + + cgh.single_task([=]() { + Acc.use(); + }); + }); + } + + return 0; +} diff --git a/clang/test/CodeGenSYCL/sampler.cpp b/clang/test/CodeGenSYCL/sampler.cpp new file mode 100644 index 0000000000000..bc1c361147a64 --- /dev/null +++ b/clang/test/CodeGenSYCL/sampler.cpp @@ -0,0 +1,63 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck --enable-var-scope %s +// CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]]) +// CHECK-NEXT: entry: +// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 +// CHECK: [[ANON:%[0-9]+]] = alloca %class.anon, align 8 +// CHECK: [[ANONCAST:%[0-9]+]] = addrspacecast %class.anon* [[ANON]] to %class.anon addrspace(4)* +// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8 +// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8* +// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4 +// CHECK-NEXT: [[GEP:%[0-9]+]] = getelementptr inbounds %class.anon, %class.anon addrspace(4)* [[ANONCAST]], i32 0, i32 0 +// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG]].addr.ascast, align 8 +// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler" addrspace(4)* {{[^,]*}} [[GEP]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]]) +// + +// CHECK: define {{.*}}spir_kernel void @{{[a-zA-Z0-9_]+}}(%opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]], i32 [[ARG_A:%[a-zA-Z0-9_]+]]) + +// Check alloca +// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8 +// CHECK: [[ARG_A]].addr = alloca i32, align 4 +// CHECK: [[LAMBDAA:%[0-9]+]] = alloca %class.anon.0, align 8 +// CHECK: [[LAMBDA:%[0-9]+]] = addrspacecast %class.anon.0* [[LAMBDAA]] to %class.anon.0 addrspace(4)* + +// Check argument store +// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8 +// CHECK: store i32 [[ARG_A]], i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4 + +// Initialize 'a' +// CHECK: [[GEP_LAMBDA:%[0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0 +// CHECK: [[GEP_A:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.sampler_wrapper, %struct.sampler_wrapper addrspace(4)* [[GEP_LAMBDA]], i32 0, i32 1 +// CHECK: [[LOAD_A:%[0-9]+]] = load i32, i32 addrspace(4)* [[ARG_A]].addr.ascast, align 4 +// CHECK: store i32 [[LOAD_A]], i32 addrspace(4)* [[GEP_A]], align 8 + +// Initialize wrapped sampler 'smpl' +// CHECK: [[GEP_LAMBDA_0:%[0-9]+]] = getelementptr inbounds %class.anon.0, %class.anon.0 addrspace(4)* [[LAMBDA]], i32 0, i32 0 +// CHECK: [[GEP_SMPL:%[a-zA-Z0-9]+]] = getelementptr inbounds %struct.sampler_wrapper, %struct.sampler_wrapper addrspace(4)* [[GEP_LAMBDA_0]], i32 0, i32 0 +// CHECK: [[LOAD_SMPL:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)* addrspace(4)* [[SAMPLER_ARG_WRAPPED]].addr.ascast, align 8 +// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SMPL]]) +// +#include "Inputs/sycl.hpp" + +struct sampler_wrapper { + cl::sycl::sampler smpl; + int a; +}; + +template +__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) { + kernelFunc(); +} + +int main() { + cl::sycl::sampler smplr; + kernel_single_task([=]() { + smplr.use(); + }); + + sampler_wrapper wrappedSampler = {smplr, 1}; + kernel_single_task([=]() { + wrappedSampler.smpl.use(); + }); + + return 0; +} diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 9663a895243f8..889f4f67bc7f0 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -1,4 +1,7 @@ -#pragma once +#ifndef SYCL_HPP +#define SYCL_HPP + +// Shared code for SYCL tests namespace cl { namespace sycl { @@ -83,5 +86,115 @@ class accessor { range MemRange, id Offset) {} }; +template +struct opencl_image_type; + +#define IMAGETY_DEFINE(dim, accessmode, amsuffix, Target, ifarray_) \ + template <> \ + struct opencl_image_type { \ + using type = __ocl_image##dim##d_##ifarray_##amsuffix##_t; \ + }; + +#define IMAGETY_READ_3_DIM_IMAGE \ + IMAGETY_DEFINE(1, read, ro, image, ) \ + IMAGETY_DEFINE(2, read, ro, image, ) \ + IMAGETY_DEFINE(3, read, ro, image, ) + +#define IMAGETY_WRITE_3_DIM_IMAGE \ + IMAGETY_DEFINE(1, write, wo, image, ) \ + IMAGETY_DEFINE(2, write, wo, image, ) \ + IMAGETY_DEFINE(3, write, wo, image, ) + +#define IMAGETY_READ_2_DIM_IARRAY \ + IMAGETY_DEFINE(1, read, ro, image_array, array_) \ + IMAGETY_DEFINE(2, read, ro, image_array, array_) + +#define IMAGETY_WRITE_2_DIM_IARRAY \ + IMAGETY_DEFINE(1, write, wo, image_array, array_) \ + IMAGETY_DEFINE(2, write, wo, image_array, array_) + +IMAGETY_READ_3_DIM_IMAGE +IMAGETY_WRITE_3_DIM_IMAGE + +IMAGETY_READ_2_DIM_IARRAY +IMAGETY_WRITE_2_DIM_IARRAY + +template +struct _ImageImplT { +#ifdef __SYCL_DEVICE_ONLY__ + typename opencl_image_type::type MImageObj; +#else + range AccessRange; + range MemRange; + id Offset; +#endif +}; + +template +class accessor { +public: + void use(void) const {} + template + void use(T... args) {} + template + void use(T... args) const {} + _ImageImplT impl; +#ifdef __SYCL_DEVICE_ONLY__ + void __init(typename opencl_image_type::type ImageObj) { impl.MImageObj = ImageObj; } +#endif +}; + +struct sampler_impl { +#ifdef __SYCL_DEVICE_ONLY__ + __ocl_sampler_t m_Sampler; +#endif +}; + +class sampler { + struct sampler_impl impl; +#ifdef __SYCL_DEVICE_ONLY__ + void __init(__ocl_sampler_t Sampler) { impl.m_Sampler = Sampler; } +#endif + +public: + void use(void) const {} +}; + +class event {}; +class queue { +public: + template + event submit(T cgf) { return event{}; } +}; +class auto_name {}; +template +struct get_kernel_name_t { + using name = Name; +}; +template +struct get_kernel_name_t { + using name = Type; +}; +#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel)) +template +ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) { + kernelFunc(); +} +class handler { +public: + template + void single_task(KernelType kernelFunc) { + using NameT = typename get_kernel_name_t::name; +#ifdef __SYCL_DEVICE_ONLY__ + kernel_single_task(kernelFunc); +#else + kernelFunc(); +#endif + } +}; + } // namespace sycl } // namespace cl + +#endif diff --git a/clang/test/SemaSYCL/accessors-targets-image.cpp b/clang/test/SemaSYCL/accessors-targets-image.cpp new file mode 100644 index 0000000000000..ead3e8a385c4c --- /dev/null +++ b/clang/test/SemaSYCL/accessors-targets-image.cpp @@ -0,0 +1,74 @@ +// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s + +// This test checks that compiler generates correct kernel wrapper arguments for +// image accessors targets. + +#include + +using namespace cl::sycl; + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + + accessor + image_acc1d_read; + kernel( + [=]() { + image_acc1d_read.use(); + }); + + accessor + image_acc2d_read; + kernel( + [=]() { + image_acc2d_read.use(); + }); + + accessor + image_acc3d_read; + kernel( + [=]() { + image_acc3d_read.use(); + }); + + accessor + image_acc1d_write; + kernel( + [=]() { + image_acc1d_write.use(); + }); + + accessor + image_acc2d_write; + kernel( + [=]() { + image_acc2d_write.use(); + }); + + accessor + image_acc3d_write; + kernel( + [=]() { + image_acc3d_write.use(); + }); +} + +// CHECK: {{.*}}use_image1d_r 'void (__read_only image1d_t)' +// CHECK: {{.*}}use_image2d_r 'void (__read_only image2d_t)' +// CHECK: {{.*}}use_image3d_r 'void (__read_only image3d_t)' +// CHECK: {{.*}}use_image1d_w 'void (__write_only image1d_t)' +// CHECK: {{.*}}use_image2d_w 'void (__write_only image2d_t)' +// CHECK: {{.*}}use_image3d_w 'void (__write_only image3d_t)' + +// TODO: SYCL specific fail - analyze and enable +// XFAIL: windows-msvc diff --git a/clang/test/SemaSYCL/sampler.cpp b/clang/test/SemaSYCL/sampler.cpp new file mode 100644 index 0000000000000..eaa214c9541a5 --- /dev/null +++ b/clang/test/SemaSYCL/sampler.cpp @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 -I %S/Inputs -fsycl-is-device -ast-dump %s | FileCheck %s + +#include + +template +__attribute__((sycl_kernel)) void kernel(Func kernelFunc) { + kernelFunc(); +} + +int main() { + cl::sycl::sampler Sampler; + kernel([=]() { + Sampler.use(); + }); + return 0; +} + +// Check declaration of the test kernel +// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (sampler_t)' +// +// Check parameters of the test kernel +// CHECK: ParmVarDecl {{.*}} used [[_arg_sampler:[0-9a-zA-Z_]+]] 'sampler_t' +// +// Check that sampler field of the test kernel object is initialized using __init method +// CHECK: CXXMemberCallExpr {{.*}} 'void' +// CHECK-NEXT: MemberExpr {{.*}} 'void (__ocl_sampler_t)' lvalue .__init +// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::sampler':'cl::sycl::sampler' lvalue +// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' +// +// Check the parameters of __init method +// CHECK-NEXT: ImplicitCastExpr {{.*}} '__ocl_sampler_t':'sampler_t' +// CHECK-NEXT: DeclRefExpr {{.*}} 'sampler_t' lvalue ParmVar {{.*}} '[[_arg_sampler]]' 'sampler_t' From e90c927d50b615fd180d0951dc2beb16cf2f8bf4 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Mon, 6 Jul 2020 10:59:49 +0100 Subject: [PATCH 4/4] [SPIR-V] Add SPIR-V builtin functions and types * Make the OpenCL Builtin lookup agnostic to programming models Rename the OpenCL Builtin lookup system to make it agnostic. The Builtin emitter is now outputting the builtin information into a class so that different programming models can be produced and used. * Add SPIR-V variants of TypeSampledImage as clang builtin type. This patch adds SPIR-V sampled image types as derivative of the builtin OpenCL Image types. For each OpenCL image type, clang defines a Sampled variant and lowered as a "spirv.SampledImage." llvm opaque type. * Enable SPIR-V builtin lookup Add flag -fdeclare-spirv-builtins to enable lookup of SPIR-V builtins. If -fdeclare-spirv-builtins is passed to clang, the compiler will try to lookup for the builtin described in SPIRVBuiltins.td for any match. If a match is found, overloads are build for the match. This will enforce a stable way to express SPIR-V builtins and make them closer to how the translator mangles them. This will help ensuring builtin for CUDA does not break easily. This will also support any changes suggested by the SPIRV-LLVM people on how to represent builtins. Define __SPIRV_BUILTIN_DECLARATIONS__ when passing -fdeclare-spirv-builtins to clang. Added OpenCL SPIR-V extended set builtins bindings and part of the core SPIR-V (mostly missing Images and Pipes) TODO: Known vendor extensions are not implemented yet. Signed-off-by: Victor Lomuller Differential Revision: https://reviews.llvm.org/D108034 --- clang/include/clang-c/Index.h | 16 +- clang/include/clang/AST/ASTContext.h | 5 + clang/include/clang/AST/Type.h | 28 + clang/include/clang/AST/TypeProperties.td | 6 + clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Driver/Options.td | 2 + .../include/clang/Serialization/ASTBitCodes.h | 5 + clang/lib/AST/ASTContext.cpp | 30 +- clang/lib/AST/ASTImporter.cpp | 6 + clang/lib/AST/ExprConstant.cpp | 5 + clang/lib/AST/ItaniumMangle.cpp | 8 + clang/lib/AST/MicrosoftMangle.cpp | 7 + clang/lib/AST/NSAPI.cpp | 5 + clang/lib/AST/PrintfFormatString.cpp | 5 + clang/lib/AST/Type.cpp | 11 + clang/lib/AST/TypeLoc.cpp | 5 + clang/lib/CodeGen/CGDebugInfo.cpp | 7 + clang/lib/CodeGen/CGDebugInfo.h | 5 + clang/lib/CodeGen/CGOpenCLRuntime.cpp | 9 + clang/lib/CodeGen/CodeGenTypes.cpp | 5 + clang/lib/CodeGen/ItaniumCXXABI.cpp | 5 + clang/lib/Frontend/CompilerInvocation.cpp | 5 + clang/lib/Frontend/InitPreprocessor.cpp | 4 + clang/lib/Index/USRGeneration.cpp | 5 + clang/lib/Sema/CMakeLists.txt | 6 + clang/lib/Sema/OpenCLBuiltins.td | 2 + clang/lib/Sema/SPIRVBuiltins.td | 952 ++++++++++++++++++ clang/lib/Sema/Sema.cpp | 19 +- clang/lib/Sema/SemaDecl.cpp | 2 +- clang/lib/Sema/SemaExpr.cpp | 10 + clang/lib/Sema/SemaLookup.cpp | 150 ++- clang/lib/Sema/SemaType.cpp | 5 +- clang/lib/Serialization/ASTCommon.cpp | 7 + clang/lib/Serialization/ASTReader.cpp | 7 + clang/test/CodeGenOpenCL/sampled_image.cl | 12 + .../CodeGenSPIRV/spirv-builtin-lookup-win.cpp | 17 + .../CodeGenSPIRV/spirv-builtin-lookup.cpp | 17 + clang/test/CodeGenSYCL/unique_stable_name.cpp | 90 +- .../unique_stable_name_windows_diff.cpp | 35 +- clang/test/Preprocessor/spirv-macro.cpp | 6 + .../test/SemaOpenCL/sampled_image_overload.cl | 14 + .../SemaSYCL/spirv-builtin-lookup-invalid.cpp | 12 + clang/test/SemaSYCL/spirv-builtin-lookup.cpp | 29 + clang/tools/libclang/CIndex.cpp | 5 + clang/tools/libclang/CXType.cpp | 11 +- clang/utils/TableGen/CMakeLists.txt | 2 +- ...r.cpp => ClangProgModelBuiltinEmitter.cpp} | 221 ++-- clang/utils/TableGen/TableGen.cpp | 6 + clang/utils/TableGen/TableGenBackends.h | 1 + 49 files changed, 1630 insertions(+), 198 deletions(-) create mode 100644 clang/lib/Sema/SPIRVBuiltins.td create mode 100644 clang/test/CodeGenOpenCL/sampled_image.cl create mode 100644 clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp create mode 100644 clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp create mode 100644 clang/test/Preprocessor/spirv-macro.cpp create mode 100644 clang/test/SemaOpenCL/sampled_image_overload.cl create mode 100644 clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp create mode 100644 clang/test/SemaSYCL/spirv-builtin-lookup.cpp rename clang/utils/TableGen/{ClangOpenCLBuiltinEmitter.cpp => ClangProgModelBuiltinEmitter.cpp} (86%) diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h index b0d7ef509c26f..85fa0fa430fed 100644 --- a/clang/include/clang-c/Index.h +++ b/clang/include/clang-c/Index.h @@ -3401,7 +3401,21 @@ enum CXTypeKind { CXType_OCLIntelSubgroupAVCImeDualRefStreamin = 175, CXType_ExtVector = 176, - CXType_Atomic = 177 + CXType_Atomic = 177, + + /* SPIRV builtin types. */ + CXType_SampledOCLImage1dRO = 178, + CXType_SampledOCLImage1dArrayRO = 179, + CXType_SampledOCLImage1dBufferRO = 180, + CXType_SampledOCLImage2dRO = 181, + CXType_SampledOCLImage2dArrayRO = 182, + CXType_SampledOCLImage2dDepthRO = 183, + CXType_SampledOCLImage2dArrayDepthRO = 184, + CXType_SampledOCLImage2dMSAARO = 185, + CXType_SampledOCLImage2dArrayMSAARO = 186, + CXType_SampledOCLImage2dMSAADepthRO = 187, + CXType_SampledOCLImage2dArrayMSAADepthRO = 188, + CXType_SampledOCLImage3dRO = 189 }; /** diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index d336342e4cda6..6a3e6ed7b684a 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -1107,6 +1107,11 @@ class ASTContext : public RefCountedBase { CanQualType ObjCBuiltinBoolTy; #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ CanQualType SingletonId; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + CanQualType Sampled##SingletonId; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" CanQualType OCLSamplerTy, OCLEventTy, OCLClkEventTy; CanQualType OCLQueueTy, OCLReserveIDTy; diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h index 4c89c297bf340..bf58450860099 100644 --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -2110,9 +2110,15 @@ class alignas(8) Type : public ExtQualsTypeCommonBase { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ bool is##Id##Type() const; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + bool isSampled##Id##Type() const; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" bool isImageType() const; // Any OpenCL image type + bool isSampledImageType() const; // Any SPIR-V Sampled image type bool isSamplerT() const; // OpenCL sampler_t bool isEventT() const; // OpenCL event_t @@ -2496,6 +2502,10 @@ class BuiltinType : public Type { // OpenCL image types #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) Id, #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) Sampled##Id, +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" // OpenCL extension types #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) Id, #include "clang/Basic/OpenCLExtensionTypes.def" @@ -6860,6 +6870,14 @@ inline bool Type::isDecltypeType() const { } #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + inline bool Type::isSampled##Id##Type() const { \ + return isSpecificBuiltinType(BuiltinType::Sampled##Id); \ + } +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" + inline bool Type::isSamplerT() const { return isSpecificBuiltinType(BuiltinType::OCLSampler); } @@ -6882,7 +6900,17 @@ inline bool Type::isReserveIDT() const { inline bool Type::isImageType() const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) is##Id##Type() || + return isSampledImageType() || +#include "clang/Basic/OpenCLImageTypes.def" + false; // end boolean or operation +} + +inline bool Type::isSampledImageType() const { +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + isSampled##Id##Type() || return +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" false; // end boolean or operation } diff --git a/clang/include/clang/AST/TypeProperties.td b/clang/include/clang/AST/TypeProperties.td index 438d5af5a2e26..ee5135633ad81 100644 --- a/clang/include/clang/AST/TypeProperties.td +++ b/clang/include/clang/AST/TypeProperties.td @@ -757,6 +757,12 @@ let Class = BuiltinType in { case BuiltinType::ID: return ctx.SINGLETON_ID; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(IMGTYPE, ID, SINGLETON_ID, ACCESS, SUFFIX) \ + case BuiltinType::Sampled##ID: return ctx.Sampled##SINGLETON_ID; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" + #define EXT_OPAQUE_TYPE(EXTTYPE, ID, EXT) \ case BuiltinType::ID: return ctx.ID##Ty; #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index 4651f4fff6aa0..132e573a8cc19 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -261,6 +261,7 @@ LANGOPT(GPUExcludeWrongSideOverloads, 1, 0, "always exclude wrong side overloads LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device") LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation") ENUM_LANGOPT(SYCLVersion , SYCLMajorVersion, 2, SYCL_None, "Version of the SYCL standard used") +LANGOPT(DeclareSPIRVBuiltins, 1, 0, "Declare SPIR-V builtin functions") LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 7730b7d1915e4..a7c8b15ecae55 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -5771,6 +5771,8 @@ def finclude_default_header : Flag<["-"], "finclude-default-header">, HelpText<"Include default header file for OpenCL">; def fdeclare_opencl_builtins : Flag<["-"], "fdeclare-opencl-builtins">, HelpText<"Add OpenCL builtin function declarations (experimental)">; +def fdeclare_spirv_builtins : Flag<["-"], "fdeclare-spirv-builtins">, + HelpText<"Add SPIR-V builtin function declarations (experimental)">; def fpreserve_vec3_type : Flag<["-"], "fpreserve-vec3-type">, HelpText<"Preserve 3-component vector type">, diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index 341da5bd1d62e..9b22249a8d603 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1074,6 +1074,11 @@ enum PredefinedTypeIDs { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ PREDEF_TYPE_##Id##_ID, #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + PREDEF_TYPE_SAMPLED_##Id##_ID, +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" /// \brief OpenCL extension types with auto numeration #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) PREDEF_TYPE_##Id##_ID, #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index f07842ca358d0..5ecb794cb5b2e 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1423,18 +1423,14 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass); InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel); - if (LangOpts.SYCLIsDevice) { - InitBuiltinType(OCLSamplerTy, BuiltinType::OCLSampler); - InitBuiltinType(OCLEventTy, BuiltinType::OCLEvent); + if (LangOpts.OpenCL || LangOpts.SYCLIsDevice) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ InitBuiltinType(SingletonId, BuiltinType::Id); #include "clang/Basic/OpenCLImageTypes.def" -#undef IMAGE_TYPE - } - - if (LangOpts.OpenCL) { -#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ - InitBuiltinType(SingletonId, BuiltinType::Id); +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + InitBuiltinType(Sampled##SingletonId, BuiltinType::Sampled##Id); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" InitBuiltinType(OCLSamplerTy, BuiltinType::OCLSampler); @@ -2193,6 +2189,11 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" @@ -6860,6 +6861,12 @@ OpenCLTypeKind ASTContext::getOpenCLTypeKind(const Type *T) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: \ return OCLTK_Image; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return OCLTK_Image; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLClkEvent: @@ -7446,6 +7453,11 @@ static char getObjCEncodingForPrimitiveType(const ASTContext *C, #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/ASTImporter.cpp b/clang/lib/AST/ASTImporter.cpp index 710e40bbb4b72..50ccea32127a4 100644 --- a/clang/lib/AST/ASTImporter.cpp +++ b/clang/lib/AST/ASTImporter.cpp @@ -1041,6 +1041,12 @@ ExpectedType ASTNodeImporter::VisitBuiltinType(const BuiltinType *T) { case BuiltinType::Id: \ return Importer.getToContext().SingletonId; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return Importer.getToContext().Sampled##SingletonId; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: \ return Importer.getToContext().Id##Ty; diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp index 99babd58b0276..94827dd24277d 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -11018,6 +11018,11 @@ EvaluateBuiltinClassifyType(QualType T, const LangOptions &LangOpts) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp index 07579d04e2754..ee91f23f0d9ac 100644 --- a/clang/lib/AST/ItaniumMangle.cpp +++ b/clang/lib/AST/ItaniumMangle.cpp @@ -3027,6 +3027,14 @@ void CXXNameMangler::mangleType(const BuiltinType *T) { type_name = "ocl_" #ImgType "_" #Suffix; \ Out << type_name.size() << type_name; \ break; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + type_name = "__spirv_SampledImage__" #ImgType "_" #Suffix; \ + Out << type_name.size() << type_name; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: Out << "11ocl_sampler"; diff --git a/clang/lib/AST/MicrosoftMangle.cpp b/clang/lib/AST/MicrosoftMangle.cpp index 79a448a2435cd..16a88819c3a29 100644 --- a/clang/lib/AST/MicrosoftMangle.cpp +++ b/clang/lib/AST/MicrosoftMangle.cpp @@ -2423,6 +2423,13 @@ void MicrosoftCXXNameMangler::mangleType(const BuiltinType *T, Qualifiers, case BuiltinType::Id: \ Out << "PAUocl_" #ImgType "_" #Suffix "@@"; \ break; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + Out << "PAU__spirv_SampledImage__" #ImgType "_" #Suffix "@@"; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: Out << "PA"; diff --git a/clang/lib/AST/NSAPI.cpp b/clang/lib/AST/NSAPI.cpp index db7878e18c42d..d6906e41ac76a 100644 --- a/clang/lib/AST/NSAPI.cpp +++ b/clang/lib/AST/NSAPI.cpp @@ -464,6 +464,11 @@ NSAPI::getNSNumberFactoryMethodKind(QualType T) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/PrintfFormatString.cpp b/clang/lib/AST/PrintfFormatString.cpp index e2569c9e20df7..014572c5d25de 100644 --- a/clang/lib/AST/PrintfFormatString.cpp +++ b/clang/lib/AST/PrintfFormatString.cpp @@ -786,6 +786,11 @@ bool PrintfSpecifier::fixType(QualType QT, const LangOptions &LangOpt, #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index e0ac3f5b1351d..f46900064dd25 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -3072,6 +3072,12 @@ StringRef BuiltinType::getName(const PrintingPolicy &Policy) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case Id: \ return "__" #Access " " #ImgType "_t"; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case Sampled##Id: \ + return "__ocl_sampled_" #ImgType "_" #Suffix "_t"; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case OCLSampler: return "sampler_t"; @@ -4117,6 +4123,11 @@ bool Type::canHaveNullability(bool ResultIfUnknown) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/TypeLoc.cpp b/clang/lib/AST/TypeLoc.cpp index c3ed08d5a8b3e..b23d460e8c66d 100644 --- a/clang/lib/AST/TypeLoc.cpp +++ b/clang/lib/AST/TypeLoc.cpp @@ -402,6 +402,11 @@ TypeSpecifierType BuiltinTypeLoc::getWrittenTypeSpec() const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index af651e6f44b7c..2c20ff410d96b 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -689,6 +689,13 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) { case BuiltinType::Id: \ return getOrCreateStructPtrType("opencl_" #ImgType "_" #Suffix "_t", \ SingletonId); +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return getOrCreateStructPtrType( \ + "spirv_sampled_" #ImgType "_" #Suffix "_t", Sampled##SingletonId); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: return getOrCreateStructPtrType("opencl_sampler_t", OCLSamplerDITy); diff --git a/clang/lib/CodeGen/CGDebugInfo.h b/clang/lib/CodeGen/CGDebugInfo.h index a7b72fa5f5a65..33eed1503bf2f 100644 --- a/clang/lib/CodeGen/CGDebugInfo.h +++ b/clang/lib/CodeGen/CGDebugInfo.h @@ -71,6 +71,11 @@ class CGDebugInfo { llvm::DIType *SelTy = nullptr; #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ llvm::DIType *SingletonId = nullptr; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + llvm::DIType *Sampled##SingletonId = nullptr; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" llvm::DIType *OCLSamplerDITy = nullptr; llvm::DIType *OCLEventDITy = nullptr; diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.cpp b/clang/lib/CodeGen/CGOpenCLRuntime.cpp index dbe375294d179..78ccb0790e0ec 100644 --- a/clang/lib/CodeGen/CGOpenCLRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenCLRuntime.cpp @@ -46,6 +46,15 @@ llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) { return llvm::PointerType::get( \ llvm::StructType::create(Ctx, "opencl." #ImgType "_" #Suffix "_t"), \ AddrSpc); +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return llvm::PointerType::get( \ + llvm::StructType::create(Ctx, "spirv.SampledImage." #ImgType \ + "_" #Suffix "_t"), \ + AddrSpc); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: return getSamplerType(T); diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index fb05475a4e8ca..2e7ea8d5dbbd8 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -531,6 +531,11 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp index 04163aeaddc52..8aeae7e9b0561 100644 --- a/clang/lib/CodeGen/ItaniumCXXABI.cpp +++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp @@ -3293,6 +3293,11 @@ static bool TypeInfoIsInStandardLibrary(const BuiltinType *Ty) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index c104a6f40e20f..d703904adf368 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -3366,6 +3366,8 @@ void CompilerInvocation::GenerateLangArgs(const LangOptions &Opts, GenerateArg(Args, OPT_finclude_default_header, SA); if (Opts.DeclareOpenCLBuiltins) GenerateArg(Args, OPT_fdeclare_opencl_builtins, SA); + if (Opts.DeclareSPIRVBuiltins) + GenerateArg(Args, OPT_fdeclare_spirv_builtins, SA); const LangOptions *LangOpts = &Opts; @@ -3678,6 +3680,9 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args, Opts.IncludeDefaultHeader = Args.hasArg(OPT_finclude_default_header); Opts.DeclareOpenCLBuiltins = Args.hasArg(OPT_fdeclare_opencl_builtins); + Opts.SYCLIsDevice = Args.hasArg(options::OPT_fsycl_is_device); + Opts.DeclareSPIRVBuiltins = Args.hasArg(OPT_fdeclare_spirv_builtins); + CompilerInvocation::setLangDefaults(Opts, IK, T, Includes, LangStd); // The key paths of codegen options defined in Options.td start with diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 0c153446142ef..81981ca25a168 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -497,6 +497,10 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI, Builder.defineMacro("SYCL_LANGUAGE_VERSION", "202001"); } + if (LangOpts.DeclareSPIRVBuiltins) { + Builder.defineMacro("__SPIRV_BUILTIN_DECLARATIONS__"); + } + // Not "standard" per se, but available even with the -undef flag. if (LangOpts.AsmPreprocessor) Builder.defineMacro("__ASSEMBLER__"); diff --git a/clang/lib/Index/USRGeneration.cpp b/clang/lib/Index/USRGeneration.cpp index 41edd431dd5b8..b931d6b98dfd3 100644 --- a/clang/lib/Index/USRGeneration.cpp +++ b/clang/lib/Index/USRGeneration.cpp @@ -719,6 +719,11 @@ void USRGenerator::VisitType(QualType T) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt index 0e0681a8e2927..a82b4f64c4ef7 100644 --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -10,6 +10,11 @@ clang_tablegen(OpenCLBuiltins.inc -gen-clang-opencl-builtins TARGET ClangOpenCLBuiltinsImpl ) +clang_tablegen(SPIRVBuiltins.inc -gen-clang-spirv-builtins + SOURCE SPIRVBuiltins.td + TARGET ClangSPIRVBuiltinsImpl + ) + add_clang_library(clangSema AnalysisBasedWarnings.cpp CodeCompleteConsumer.cpp @@ -65,6 +70,7 @@ add_clang_library(clangSema DEPENDS ClangOpenCLBuiltinsImpl + ClangSPIRVBuiltinsImpl omp_gen LINK_LIBS diff --git a/clang/lib/Sema/OpenCLBuiltins.td b/clang/lib/Sema/OpenCLBuiltins.td index 8cf7ec58eff56..c870a050983cc 100644 --- a/clang/lib/Sema/OpenCLBuiltins.td +++ b/clang/lib/Sema/OpenCLBuiltins.td @@ -298,6 +298,8 @@ class Builtin _Signature, list _Attributes = Attr. bit IsConst = _Attributes[1]; // Function attribute __attribute__((convergent)) bit IsConv = _Attributes[2]; + // Is function a variadic one + bit IsVariadic = 0; // OpenCL extensions to which the function belongs. FunctionExtension Extension = FuncExtNone; // Version of OpenCL from which the function is available (e.g.: CL10). diff --git a/clang/lib/Sema/SPIRVBuiltins.td b/clang/lib/Sema/SPIRVBuiltins.td new file mode 100644 index 0000000000000..8e4c2175d3aa0 --- /dev/null +++ b/clang/lib/Sema/SPIRVBuiltins.td @@ -0,0 +1,952 @@ +//==--- SPIRVBuiltins.td - SPIRV builtin declarations -------------------===// +// +// The LLVM Compiler Infrastructure +// +// 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 file contains TableGen definitions for SPIR-V builtin function +// declarations. In case of an unresolved function name, Clang will check for +// a function described in this file when -fdeclare-spirv-builtins is specified. +// +//===----------------------------------------------------------------------===// + +//===----------------------------------------------------------------------===// +// Definitions of miscellaneous basic entities. +//===----------------------------------------------------------------------===// +// TODO: basic entities declaration with OpenCLBuiltins.td + +// TODO: Manage version using the JSON grammar. Unused for now. +class Version { + int ID = _Version; +} +def SPIRVAll : Version< 0>; + +// Address spaces +// Pointer types need to be assigned an address space. +class AddressSpace { + string Name = _AS; +} +// Default is important for the frontend as there is not necessarily +// an automatic conversion from this address space to +// the one it will be lowered to. +// This file assumes it will get lowered to generic or private. +def DefaultAS : AddressSpace<"clang::LangAS::Default">; +def PrivateAS : AddressSpace<"clang::LangAS::sycl_private">; +def GlobalAS : AddressSpace<"clang::LangAS::sycl_global">; +def ConstantAS : AddressSpace<"clang::LangAS::opencl_constant">; +def LocalAS : AddressSpace<"clang::LangAS::sycl_local">; +def GenericAS : AddressSpace<"clang::LangAS::opencl_generic">; + +// TODO: Manage capabilities. Unused for now. +class AbstractExtension { + string ExtName = _Ext; +} + +// Extension associated to a builtin function. +class FunctionExtension : AbstractExtension<_Ext>; + +// FunctionExtension definitions. +def FuncExtNone : FunctionExtension<"">; + +// Extension associated to a type. This enables implicit conditionalization of +// builtin function overloads containing a type that depends on an extension. +// During overload resolution, when a builtin function overload contains a type +// with a TypeExtension, those overloads are skipped when the extension is +// disabled. +class TypeExtension : AbstractExtension<_Ext>; + +// TypeExtension definitions. +def NoTypeExt : TypeExtension<"">; + +// Qualified Type. These map to ASTContext::QualType. +// TODO: Create a QualTypeFromASTContext. +// To fully make sense here, this class should represent +// the QualType only. How the QualType is accessed should be separated. +class QualType { + // Name of the field or function in a clang::ASTContext + // E.g. Name="IntTy" for the int type, and "getIntPtrType()" for an intptr_t + string TypeExpr = _TypeExpr; + // Some QualTypes in this file represent an abstract type for which there is + // no corresponding AST QualType, e.g. a GenType or an `image2d_t` type + // without access qualifiers. + bit IsAbstract = _IsAbstract; + bit IsSigned = _IsSigned; +} + +// Qualified Type. These map to a function taking an ASTContext +// and returning a QualType. +// Instead of direclty accessing ASTContext fields, the builtin lookup can +// call a function to extract the correct type for the call. +// The name will be interpreted as the function to call +// rather than the field to access. +class QualTypeFromFunction : + QualType<_Name, _IsAbstract, _IsSigned> { +// TODO: At the moment the user is expected to write the function outside this file. +// Although they could be generated in the .inc file and +// the user would only have to provide the body here +// (like it can be done for attributes for instance). +} + +// List of integers. +class IntList _List> { + string Name = _Name; + list List = _List; +} + +// Basic data types (int, float, image2d_t, ...). +// Its child classes can represent concrete types (e.g. VectorType) or +// abstract types (e.g. GenType). +class Type { + // Name of the Type. + string Name = _Name; + // QualType associated with this type. + QualType QTExpr = _QTExpr; + // Size of the vector (if applicable). + int VecWidth = 1; + // Size of the element in bits. + int ElementSize = 1; + // Is a integer. + bit IsInteger = 0; + // Is a signed integer. + bit IsSigned = 1; + // Is a float. + bit IsFloat = 0; + // Is a pointer. + bit IsPointer = 0; + // "const" qualifier. + bit IsConst = 0; + // "volatile" qualifier. + bit IsVolatile = 0; + // Access qualifier. Must be one of ("RO", "WO", "RW"). + string AccessQualifier = ""; + // Address space. + string AddrSpace = DefaultAS.Name; + // Extension that needs to be enabled to expose a builtin that uses this type. + TypeExtension Extension = NoTypeExt; +} + +class FundamentalType : Type<_Name, _QTName> { + // Inherited fields + let ElementSize = _Size; +} + +// Integer Type. +class IntType : FundamentalType<_Name, _QTName, _Size> { + // Inherited fields + let IsInteger = 1; + let IsSigned = 1; +} + +// Unsigned integer Type. +class UIntType : FundamentalType<_Name, _QTName, _Size> { + // Inherited fields + let IsInteger = 1; + let IsSigned = 0; +} + +// Floating Type. +class FPType : FundamentalType<_Name, _QTName, _Size> { + // Inherited fields + let IsFloat = 1; +} + +class CompoundType : Type<_Ty.Name, _Ty.QTExpr> { + // Inherited fields + let VecWidth = _Ty.VecWidth; + let ElementSize = _Ty.ElementSize; + let IsInteger = _Ty.IsInteger; + let IsSigned = _Ty.IsSigned; + let IsFloat = _Ty.IsFloat; + let IsPointer = _Ty.IsPointer; + let IsConst = _Ty.IsConst; + let IsVolatile = _Ty.IsVolatile; + let AccessQualifier = _Ty.AccessQualifier; + let AddrSpace = _Ty.AddrSpace; + let Extension = _Ty.Extension; + + Type ElementType = _Ty; +} + +// Vector types (e.g. int2, int3, int16, float8, ...). +class VectorType : Type<_Ty.Name, _Ty.QTExpr> { + let VecWidth = _VecWidth; + let AccessQualifier = ""; + // Inherited fields + let ElementSize = _Ty.ElementSize; + let IsInteger = _Ty.IsInteger; + let IsSigned = _Ty.IsSigned; + let IsFloat = _Ty.IsFloat; + let IsPointer = _Ty.IsPointer; + let IsConst = _Ty.IsConst; + let IsVolatile = _Ty.IsVolatile; + let AccessQualifier = _Ty.AccessQualifier; + let AddrSpace = _Ty.AddrSpace; + let Extension = _Ty.Extension; +} + +// Pointer types (e.g. int*, float*, ...). +class PointerType : + CompoundType<_Ty> { + // Inherited fields + let IsPointer = 1; + let AddrSpace = _AS.Name; + let Extension = _Ty.Extension; +} + +// Const types (e.g. const int). +class ConstType : CompoundType<_Ty> { + // Inherited fields + let IsConst = 1; + let Extension = _Ty.Extension; +} + +// Volatile types (e.g. volatile int). +class VolatileType : CompoundType<_Ty> { + // Inherited fields + let IsVolatile = 1; + let Extension = _Ty.Extension; +} + +// Image types (e.g. image2d). +class ImageType : + Type<_Ty.Name, QualType<_Ty.QTExpr.TypeExpr # _AccessQualifier # "Ty", 0>> { + let VecWidth = 0; + let AccessQualifier = _AccessQualifier; + // Inherited fields + let ElementSize = _Ty.ElementSize; + let IsInteger = _Ty.IsInteger; + let IsSigned = _Ty.IsSigned; + let IsFloat = _Ty.IsFloat; + let IsPointer = _Ty.IsPointer; + let IsConst = _Ty.IsConst; + let IsVolatile = _Ty.IsVolatile; + let AddrSpace = _Ty.AddrSpace; + let Extension = _Ty.Extension; +} + +// List of Types. +class TypeList _Type> { + list List = _Type; +} + +// A GenericType is an abstract type that defines a set of types as a +// combination of Types and vector sizes. +// +// For example, if TypeList = and VectorList = <1, 2, 4>, then it +// represents . +// +// Some rules apply when using multiple GenericType arguments in a declaration: +// 1. The number of vector sizes must be equal or 1 for all gentypes in a +// declaration. +// 2. The number of Types must be equal or 1 for all gentypes in a +// declaration. +// 3. Generic types are combined by iterating over all generic types at once. +// For example, for the following GenericTypes +// GenT1 = GenericType and +// GenT2 = GenericType +// A declaration f(GenT1, GenT2) results in the combinations +// f(half, float), f(half2, float2), f(half, int), f(half2, int2) . +// 4. "sgentype" from the OpenCL specification is supported by specifying +// a single vector size. +// For example, for the following GenericTypes +// GenT = GenericType and +// SGenT = GenericType +// A declaration f(GenT, SGenT) results in the combinations +// f(half, half), f(half2, half), f(int, int), f(int2, int) . +class GenericType : + Type<_Ty, QualType<"null", 1>> { + // Possible element types of the generic type. + TypeList TypeList = _TypeList; + // Possible vector sizes of the types in the TypeList. + IntList VectorList = _VectorList; + // The VecWidth field is ignored for GenericTypes. Use VectorList instead. + let VecWidth = 0; +} + +// Builtin function attributes. +def Attr { + list None = [0, 0, 0]; + list Pure = [1, 0, 0]; + list Const = [0, 1, 0]; + list Convergent = [0, 0, 1]; +} + +//===----------------------------------------------------------------------===// +// Class for builtin functions +//===----------------------------------------------------------------------===// +class Builtin _Signature, list _Attributes = Attr.None> { + // Name of the builtin function + string Name = _Name; + // List of types used by the function. The first one is the return type and + // the following are the arguments. The list must have at least one element + // (the return type). + list Signature = _Signature; + // Function attribute __attribute__((pure)) + bit IsPure = _Attributes[0]; + // Function attribute __attribute__((const)) + bit IsConst = _Attributes[1]; + // Function attribute __attribute__((convergent)) + bit IsConv = _Attributes[2]; + // Is function a variadic one + bit IsVariadic = 0; + // OpenCL extensions to which the function belongs. + FunctionExtension Extension = FuncExtNone; + // Version from which the function is available. + // MinVersion is inclusive. + Version MinVersion = SPIRVAll; + // Version from which the function is not supported anymore. + // MaxVersion is exclusive. + // SPIRVAll makes the function available for all versions. + Version MaxVersion = SPIRVAll; +} + +// Helper to declare SPIR-V Core builtins. +class SPVBuiltin _Signature, list _Attributes = Attr.None> : + Builtin<"__spirv_" # _Name, _Signature, _Attributes> {} + +// Helper to declare OpenCL SPIR-V extended set builtins. +class OCLSPVBuiltin _Signature, list _Attributes = Attr.None> : + SPVBuiltin<"ocl_" # _Name, _Signature, _Attributes> {} + +class ConstOCLSPVBuiltin _Signature> : + OCLSPVBuiltin<_Name, _Signature, Attr.Const> {} + +//===----------------------------------------------------------------------===// +// Definitions of types +//===----------------------------------------------------------------------===// + +// OpenCL v1.0/1.2/2.0 s6.1.1: Built-in Scalar Data Types. +def Bool : IntType<"bool", QualType<"Context.BoolTy">, 1>; +def Char : IntType<"char", QualType<"Context.CharTy", 0, 1>, 8>; +def SChar : IntType<"schar", QualType<"Context.SignedCharTy", 0, 1>, 8>; +def UChar : UIntType<"uchar", QualType<"Context.UnsignedCharTy">, 8>; +def Short : IntType<"short", QualType<"Context.ShortTy", 0, 1>, 16>; +def UShort : UIntType<"ushort", QualType<"Context.UnsignedShortTy">, 16>; +def Int : IntType<"int", QualType<"Context.IntTy", 0, 1>, 32>; +def UInt : UIntType<"uint", QualType<"Context.UnsignedIntTy">, 32>; +def Long : IntType<"long", QualType<"Context.getIntTypeForBitwidth(64, true)", 0, 1>, 64>; +def ULong : UIntType<"ulong", QualType<"Context.getIntTypeForBitwidth(64, false)">, 64>; +def Float : FPType<"float", QualType<"Context.FloatTy">, 32>; +def Double : FPType<"double", QualType<"Context.DoubleTy">, 64>; +def Half : FPType<"half", QualTypeFromFunction<"GetFloat16Type">, 16>; +def Void : Type<"void", QualType<"Context.VoidTy">>; +// FIXME: ensure this is portable... +def Size : Type<"size_t", QualType<"Context.getSizeType()">>; + +def Sampler : Type<"sampler_t", QualType<"Context.OCLSamplerTy">>; +def Event : Type<"event_t", QualType<"Context.OCLEventTy">>; + +//===----------------------------------------------------------------------===// +// Definitions of gentype variants +//===----------------------------------------------------------------------===// + +// Vector width lists. +def VecAndScalar: IntList<"VecAndScalar", [1, 2, 3, 4, 8, 16]>; +def VecNoScalar : IntList<"VecNoScalar", [2, 3, 4, 8, 16]>; +def Vec1 : IntList<"Vec1", [1]>; +def Vec2 : IntList<"Vec2", [2]>; +def Vec4 : IntList<"Vec4", [4]>; +def Vec8 : IntList<"Vec8", [8]>; +def Vec16 : IntList<"Vec16", [16]>; +def Vec1234 : IntList<"Vec1234", [1, 2, 3, 4]>; + +// Type lists. +def TLAll : TypeList<[Char, UChar, Short, UShort, Int, UInt, Long, ULong, Float, Double, Half]>; +def TLAllUnsigned : TypeList<[UChar, UChar, UShort, UShort, UInt, UInt, ULong, ULong, UInt, ULong, UShort]>; +def TLFloat : TypeList<[Float, Double, Half]>; +def TLSignedInts : TypeList<[Char, Short, Int, Long]>; +def TLUnsignedInts : TypeList<[UChar, UShort, UInt, ULong]>; + +// Signed to Unsigned conversion +def TLSToUSignedInts : TypeList<[Char, Short, Int, Long]>; +def TLSToUUnsignedInts : TypeList<[UChar, UShort, UInt, ULong]>; + +def TLIntLongFloats : TypeList<[Int, UInt, Long, ULong, Float, Double, Half]>; + +// All unsigned integer types twice, to facilitate unsigned return types for e.g. +// uchar abs(char) and +// uchar abs(uchar). +def TLAllUIntsTwice : TypeList<[UChar, UChar, UChar, UShort, UShort, UInt, UInt, ULong, ULong]>; + +def TLAllInts : TypeList<[Char, SChar, UChar, Short, UShort, Int, UInt, Long, ULong]>; + +// GenType definitions for multiple base types (e.g. all floating point types, +// or all integer types). +// All types +def AGenType1 : GenericType<"AGenType1", TLAll, Vec1>; +def AGenTypeN : GenericType<"AGenTypeN", TLAll, VecAndScalar>; +def AGenTypeNNoScalar : GenericType<"AGenTypeNNoScalar", TLAll, VecNoScalar>; +// All integer +def AIGenType1 : GenericType<"AIGenType1", TLAllInts, Vec1>; +def AIGenTypeN : GenericType<"AIGenTypeN", TLAllInts, VecAndScalar>; +def AUIGenTypeN : GenericType<"AUIGenTypeN", TLUnsignedInts, VecAndScalar>; +def ASIGenTypeN : GenericType<"ASIGenTypeN", TLSignedInts, VecAndScalar>; +def AIGenTypeNNoScalar : GenericType<"AIGenTypeNNoScalar", TLAllInts, VecNoScalar>; +// All integer to unsigned +def AI2UGenTypeN : GenericType<"AI2UGenTypeN", TLAllUIntsTwice, VecAndScalar>; +// Signed integer +def SGenTypeN : GenericType<"SGenTypeN", TLSignedInts, VecAndScalar>; +// Unsigned integer +def UGenTypeN : GenericType<"UGenTypeN", TLUnsignedInts, VecAndScalar>; +// Float +def FGenTypeN : GenericType<"FGenTypeN", TLFloat, VecAndScalar>; +// (u)int, (u)long, and all floats +def IntLongFloatGenType1 : GenericType<"IntLongFloatGenType1", TLIntLongFloats, Vec1>; + +// GenType definitions for every single base type (e.g. fp32 only). +// Names are like: GenTypeFloatVecAndScalar. +foreach Type = [Char, SChar, UChar, Short, UShort, + Int, UInt, Long, ULong, + Float, Double, Half] in { + foreach VecSizes = [VecAndScalar, VecNoScalar] in { + def "GenType" # Type # VecSizes : + GenericType<"GenType" # Type # VecSizes, + TypeList<[Type]>, VecSizes>; + } +} + +// GenType definitions for vec1234. +foreach Type = [Float, Double, Half] in { + def "GenType" # Type # Vec1234 : + GenericType<"GenType" # Type # Vec1234, + TypeList<[Type]>, Vec1234>; +} + +//===----------------------------------------------------------------------===// +// Definitions of builtins +// extinst.opencl.std.100.grammar.json +//===----------------------------------------------------------------------===// + +// 2.1. Math extended instructions + + +foreach name = ["acos", "acosh", "acospi", + "asin", "asinh", "asinpi", + "atan", "atanh", "atanpi", + "cbrt", "ceil", "cos", + "cosh", "cospi", + "erfc", "erf", + "exp", "exp2", "exp10", + "expm1", "fabs", "floor", "lgamma", + "log", "log2", "log10", "log1p", "logb", + "rint", "round", "rsqrt", + "sin", "sinh", "sinpi", + "sqrt", + "tan", "tanh", "tanpi", + "tgamma", "trunc"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["fmax", "fmin", "fmod", + "atan2", "atan2pi", + "copysign", "fdim", "hypot", + "maxmag", "minmag", "nextafter", + "pow", "powr", "remainder"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["fma", "mad"] in { + def : ConstOCLSPVBuiltin; +} + +foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in { + foreach name = ["fract", "modf"] in { + def : OCLSPVBuiltin]>; + } + + foreach name = ["frexp", "lgamma_r"] in { + foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in { + def : OCLSPVBuiltin]>; + } + } +} + +foreach name = ["ilogb"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +foreach name = ["ldexp"] in { + foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + } +} + +foreach name = ["nan"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +foreach name = ["pown"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in { + foreach name = ["remquo"] in { + foreach Type = [GenTypeFloatVecAndScalar, GenTypeDoubleVecAndScalar, GenTypeHalfVecAndScalar] in { + def : OCLSPVBuiltin]>; + } + } +} + +foreach name = ["rootn"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in { + foreach name = ["sincos"] in { + def : OCLSPVBuiltin]>; + } +} + +foreach name = ["half_cos", + "half_exp", "half_exp2", "half_exp10", + "half_log", "half_log2", "half_log10", + "half_recip", "half_rsqrt", + "half_sin", "half_sqrt", "half_tan", + "native_cos", "native_exp", "native_exp2", "native_exp10", + "native_log", "native_log2", "native_log10", + "native_recip", "native_rsqrt", + "native_sin", "native_sqrt", "native_tan"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["half_divide", "half_powr", "native_divide", "native_powr"] in { + def : ConstOCLSPVBuiltin; +} + +// 2.2. Integer instructions + +foreach name = ["clz", "ctz", "popcount"] in { + def : ConstOCLSPVBuiltin; +} + +def : ConstOCLSPVBuiltin<"rotate", [AIGenTypeN, AIGenTypeN, AIGenTypeN]>; + +def : ConstOCLSPVBuiltin<"s_abs", [AUIGenTypeN, ASIGenTypeN]>; + +def : ConstOCLSPVBuiltin<"s_abs_diff", [AUIGenTypeN, ASIGenTypeN, ASIGenTypeN]>; + +foreach name = ["s_add_sat", + "s_hadd", "s_rhadd", + "s_max", "s_min", + "s_mul_hi", "s_sub_sat"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["s_clamp", "s_mad_hi", "s_mad_sat"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["s_upsample"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +def : ConstOCLSPVBuiltin<"s_mad24", [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar]>; + +def : ConstOCLSPVBuiltin<"s_mul24", [GenTypeIntVecAndScalar, GenTypeIntVecAndScalar, GenTypeIntVecAndScalar]>; + +foreach name = ["u_add_sat", "u_hadd", + "u_rhadd", + "u_max", "u_min", "u_sub_sat", + "u_abs_diff", "u_mul_hi"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["u_clamp", "u_mad_sat", "u_mad_hi"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["u_upsample"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +def : ConstOCLSPVBuiltin<"u_mad24", [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar]>; + +def : ConstOCLSPVBuiltin<"u_mul24", [GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar, GenTypeUIntVecAndScalar]>; + +def : ConstOCLSPVBuiltin<"u_abs", [AUIGenTypeN, AUIGenTypeN]>; + +// 2.3. Common instructions + +foreach name = ["degrees", "radians", "sign"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["fmax_common", "fmin_common", "step"] in { + def : ConstOCLSPVBuiltin; +} + +foreach name = ["fclamp", "mix", "smoothstep"] in { + def : ConstOCLSPVBuiltin; +} + +// 2.4. Geometric instructions + +foreach name = ["cross"] in { + foreach VSize = [3, 4] in { + def : ConstOCLSPVBuiltin, VectorType, VectorType]>; + def : ConstOCLSPVBuiltin, VectorType, VectorType]>; + def : ConstOCLSPVBuiltin, VectorType, VectorType]>; + } +} + +foreach name = ["distance"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +foreach name = ["length"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +foreach name = ["normalize"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +def : ConstOCLSPVBuiltin<"fast_distance", [Float, GenTypeFloatVec1234, GenTypeFloatVec1234]>; + +def : ConstOCLSPVBuiltin<"fast_length", [Float, GenTypeFloatVec1234]>; + +def : ConstOCLSPVBuiltin<"fast_normalize", [GenTypeFloatVec1234, GenTypeFloatVec1234]>; + +// 2.5. Relational instructions + +def : ConstOCLSPVBuiltin<"bitselect", [AGenTypeN, AGenTypeN, AGenTypeN, AGenTypeN]>; + +foreach name = ["select"] in { + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; + def : ConstOCLSPVBuiltin; +} + +// 2.6. Vector Data Load and Store instructions + +foreach VSize = [2, 3, 4, 8, 16] in { + foreach AS = [GlobalAS, LocalAS, PrivateAS, ConstantAS, GenericAS, DefaultAS] in { + foreach Ty = TLAll.List in { + foreach name = ["vloadn"] in { + def : OCLSPVBuiltin, Size, PointerType, AS>]>; + } + } + foreach name = ["vloada_halfn", "vload_halfn"] in { + def : OCLSPVBuiltin, Size, PointerType, AS>]>; + } + } + foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in { + foreach Ty = TLAll.List in { + foreach name = ["vstoren"] in { + def : OCLSPVBuiltin, Size, PointerType]>; + } + } + foreach name = ["vstore_halfn", "vstorea_halfn"] in { + def : OCLSPVBuiltin, Size, PointerType]>; + def : OCLSPVBuiltin, Size, PointerType]>; + } + foreach name = ["vstore_halfn_r", "vstorea_halfn_r"] in { + def : OCLSPVBuiltin, Size, PointerType, UInt]>; + def : OCLSPVBuiltin, Size, PointerType, UInt]>; + } + } +} + +foreach AS = [GlobalAS, LocalAS, PrivateAS, ConstantAS, GenericAS, DefaultAS] in { + foreach name = ["vload_half"] in { + def : OCLSPVBuiltin, AS>]>; + } +} + +foreach AS = [GlobalAS, LocalAS, PrivateAS, GenericAS, DefaultAS] in { + foreach name = ["vstore_half"] in { + def : OCLSPVBuiltin]>; + def : OCLSPVBuiltin]>; + } + foreach name = ["vstore_half_r"] in { + def : OCLSPVBuiltin, UInt]>; + def : OCLSPVBuiltin, UInt]>; + } +} + +// 2.7. Miscellaneous Vector instructions + +foreach VSize1 = [Vec2, Vec4, Vec8, Vec16] in { + foreach VSize2 = [Vec2, Vec4, Vec8, Vec16] in { + def : OCLSPVBuiltin<"shuffle", [GenericType<"TLAll" # VSize1.Name, TLAll, VSize1>, + GenericType<"TLAll" # VSize2.Name, TLAll, VSize2>, + GenericType<"TLAllUnsigned" # VSize1.Name, TLAllUnsigned, VSize1>], + Attr.Const>; + } +} +foreach VSize1 = [Vec2, Vec4, Vec8, Vec16] in { + foreach VSize2 = [Vec2, Vec4, Vec8, Vec16] in { + def : OCLSPVBuiltin<"shuffle2", [GenericType<"TLAll" # VSize1.Name, TLAll, VSize1>, + GenericType<"TLAll" # VSize2.Name, TLAll, VSize2>, + GenericType<"TLAll" # VSize2.Name, TLAll, VSize2>, + GenericType<"TLAllUnsigned" # VSize1.Name, TLAllUnsigned, VSize1>], + Attr.Const>; + } +} + +// 2.8. Misc instructions + +let IsVariadic = 1 in { + foreach name = ["printf"] in { + def : OCLSPVBuiltin, ConstantAS>]>; + } +} + +foreach name = ["prefetch"] in { + def : OCLSPVBuiltin, GlobalAS>, Size]>; +} + + +// Core builtins + +// 3.32.8. Memory Instructions + +foreach name = ["GenericPtrMemSemantics"] in { + def : SPVBuiltin, GenericAS>], Attr.Const>; +} + +// 3.32.11. Conversion Instructions +foreach rnd = ["", "_rte", "_rtn", "_rtp", "_rtz"] in { + foreach IType = TLUnsignedInts.List in { + foreach FType = TLFloat.List in { + foreach sat = ["", "_sat"] in { + def : SPVBuiltin<"ConvertFToU_R" # IType.Name # sat # rnd, [IType, FType], Attr.Const>; + } + def : SPVBuiltin<"ConvertUToF_R" # FType.Name # rnd, [FType, IType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + foreach sat = ["", "_sat"] in { + def : SPVBuiltin<"ConvertFToU_R" # IType.Name # v # sat # rnd, + [VectorType, VectorType], + Attr.Const>; + } + def : SPVBuiltin<"ConvertUToF_R" # FType.Name # v # rnd, + [VectorType, VectorType], + Attr.Const>; + } + } + } + + foreach IType = TLSignedInts.List in { + foreach FType = TLFloat.List in { + foreach sat = ["", "_sat"] in { + def : SPVBuiltin<"ConvertFToS_R" # IType.Name # sat # rnd, [IType, FType], Attr.Const>; + } + def : SPVBuiltin<"ConvertSToF_R" # FType.Name # rnd, [FType, IType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + foreach sat = ["", "_sat"] in { + def : SPVBuiltin<"ConvertFToS_R" # IType.Name # v # sat # rnd, + [VectorType, VectorType], + Attr.Const>; + } + def : SPVBuiltin<"ConvertSToF_R" # FType.Name # v # rnd, + [VectorType, VectorType], + Attr.Const>; + } + } + } + + foreach InType = TLFloat.List in { + foreach OutType = TLFloat.List in { + if !ne(OutType.ElementSize, InType.ElementSize) then { + def : SPVBuiltin<"FConvert_R" # OutType.Name # rnd, [OutType, InType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + def : SPVBuiltin<"FConvert_R" # OutType.Name # v # rnd, + [VectorType, VectorType], + Attr.Const>; + } + } + } + } +} + +foreach sat = ["", "_sat"] in { + foreach InType = TLAllInts.List in { + foreach OutType = TLUnsignedInts.List in { + if !ne(OutType.ElementSize, InType.ElementSize) then { + def : SPVBuiltin<"UConvert_R" # OutType.Name # sat, [OutType, InType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + def : SPVBuiltin<"UConvert_R" # OutType.Name # v # sat, + [VectorType, VectorType], + Attr.Const>; + } + } + } + foreach OutType = TLSignedInts.List in { + if !ne(OutType.ElementSize, InType.ElementSize) then { + def : SPVBuiltin<"SConvert_R" # OutType.Name # sat, [OutType, InType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + def : SPVBuiltin<"SConvert_R" # OutType.Name # v # sat, + [VectorType, VectorType], + Attr.Const>; + } + } + } + } +} + +foreach InType = TLSignedInts.List in { + foreach OutType = TLUnsignedInts.List in { + def : SPVBuiltin<"SatConvertSToU_R" # OutType.Name, [OutType, InType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + def : SPVBuiltin<"SatConvertSToU_R" # OutType.Name # v, + [VectorType, VectorType], + Attr.Const>; + } + } +} + +foreach InType = TLUnsignedInts.List in { + foreach OutType = TLSignedInts.List in { + def : SPVBuiltin<"SatConvertUToS_R" # OutType.Name, [OutType, InType], Attr.Const>; + foreach v = [2, 3, 4, 8, 16] in { + def : SPVBuiltin<"SatConvertUToS_R" # OutType.Name # v, + [VectorType, VectorType], + Attr.Const>; + } + } +} + +foreach AS = [GlobalAS, LocalAS, PrivateAS] in { + def : SPVBuiltin<"GenericCastToPtrExplicit", [PointerType, PointerType], Attr.Const>; +} + +foreach Type = TLFloat.List in { + foreach v = [2, 3, 4, 8, 16] in { + def : SPVBuiltin<"VectorTimesScalar", [VectorType, VectorType, Type], Attr.Const>; + } +} + +foreach name = ["Dot"] in { + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; +} + +foreach name = ["Any", "All"] in { + def : SPVBuiltin; +} + +foreach name = ["IsNan", "IsInf", "IsFinite", "IsNormal", "SignBitSet"] in { + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; +} + +foreach name = ["LessOrGreater", + "Ordered", "Unordered", + "FOrdEqual", "FUnordEqual", + "FOrdNotEqual", "FUnordNotEqual", + "FOrdLessThan", "FUnordLessThan", + "FOrdGreaterThan", "FUnordGreaterThan", + "FOrdLessThanEqual", "FUnordLessThanEqual", + "FOrdGreaterThanEqual", "FUnordGreaterThanEqual"] in { + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; + def : SPVBuiltin; +} + +foreach name = ["BitCount"] in { + def : SPVBuiltin; +} + +// 3.32.20. Barrier Instructions + +foreach name = ["ControlBarrier"] in { + // TODO: Allow enum flags instead of UInt ? + // TODO: We should enforce that the UInt must be a literal. + def : SPVBuiltin; +} + +foreach name = ["MemoryBarrier"] in { + // TODO: Allow enum flags instead of UInt ? + // TODO: We should enforce that the UInt must be a literal. + def : SPVBuiltin; +} + +// 3.32.21. Group and Subgroup Instructions + +foreach name = ["GroupAsyncCopy"] in { + // TODO: Allow enum flags instead of UInt ? + // TODO: We should enforce that the UInt must be a literal. + def : SPVBuiltin, PointerType, GlobalAS>, Size, Size, Event], Attr.Convergent>; + def : SPVBuiltin, PointerType, LocalAS>, Size, Size, Event], Attr.Convergent>; +} + +foreach name = ["GroupWaitEvents"] in { + def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin], Attr.Convergent>; +} + +foreach name = ["GroupAll", "GroupAny"] in { + def : SPVBuiltin; +} + +foreach name = ["GroupBroadcast"] in { + foreach IDType = TLAllInts.List in { + def : SPVBuiltin; + def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin; + def : SPVBuiltin], Attr.Convergent>; + def : SPVBuiltin], Attr.Convergent>; + } +} + +foreach name = ["GroupIAdd", "GroupNonUniformIMul", "GroupNonUniformBitwiseOr", + "GroupNonUniformBitwiseXor", "GroupNonUniformBitwiseAnd"] in { + def : SPVBuiltin; +} + +foreach name = ["GroupFAdd", "GroupFMin", "GroupFMax", + "GroupNonUniformFMul"] in { + def : SPVBuiltin; +} + +foreach name = ["GroupUMin", "GroupUMax"] in { + def : SPVBuiltin; +} + +foreach name = ["GroupSMin", "GroupSMax"] in { + def : SPVBuiltin; +} diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 4470ea2a4df29..37be02255a521 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -318,8 +318,11 @@ void Sema::Initialize() { } if (getLangOpts().SYCLIsDevice) { - addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy); addImplicitTypedef("__ocl_event_t", Context.OCLEventTy); + addImplicitTypedef("__ocl_sampler_t", Context.OCLSamplerTy); +#ifdef SEMA_STRINGIZE +#error "Undefine SEMA_STRINGIZE macro." +#endif #define SEMA_STRINGIZE(s) #s #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ addImplicitTypedef(SEMA_STRINGIZE(__ocl_##ImgType##_##Suffix##_t), \ @@ -328,6 +331,20 @@ void Sema::Initialize() { #undef SEMA_STRINGIZE } + if (getLangOpts().SYCLIsDevice || getLangOpts().OpenCL) { +#ifdef SEMA_STRINGIZE +#error "Undefine SEMA_STRINGIZE macro." +#endif +#define SEMA_STRINGIZE(s) #s +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + addImplicitTypedef(SEMA_STRINGIZE(__ocl_sampled_##ImgType##_##Suffix##_t), \ + Context.Sampled##SingletonId); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" +#undef SEMA_STRINGIZE + } + // Initialize predefined OpenCL types and supported extensions and (optional) // core features. if (getLangOpts().OpenCL) { diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index af174ac1ca1a7..32201da0234b6 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -6833,7 +6833,7 @@ static bool diagnoseOpenCLTypes(Sema &Se, VarDecl *NewVD) { // OpenCL v2.0 s6.9.b - Image type can only be used as a function argument. // OpenCL v2.0 s6.13.16.1 - Pipe type can only be used as a function // argument. - if (R->isImageType() || R->isPipeType()) { + if (!R->isSampledImageType() && (R->isImageType() || R->isPipeType())) { Se.Diag(NewVD->getLocation(), diag::err_opencl_type_can_only_be_used_as_function_parameter) << R; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 8592335e20d31..f80ac30a92628 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6102,6 +6102,11 @@ static bool isPlaceholderToRemoveAsArg(QualType type) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" @@ -19816,6 +19821,11 @@ ExprResult Sema::CheckPlaceholderExpr(Expr *E) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp index 0711e6d89383c..cab9f6707693c 100644 --- a/clang/lib/Sema/SemaLookup.cpp +++ b/clang/lib/Sema/SemaLookup.cpp @@ -47,7 +47,10 @@ #include #include +static inline clang::QualType GetFloat16Type(clang::ASTContext &Context); + #include "OpenCLBuiltins.inc" +#include "SPIRVBuiltins.inc" using namespace clang; using namespace sema; @@ -677,6 +680,10 @@ LLVM_DUMP_METHOD void LookupResult::dump() { D->dump(); } +static inline QualType GetFloat16Type(clang::ASTContext &Context) { + return Context.getLangOpts().OpenCL ? Context.HalfTy : Context.Float16Ty; +} + /// Diagnose a missing builtin type. static QualType diagOpenCLBuiltinTypeError(Sema &S, llvm::StringRef TypeClass, llvm::StringRef Name) { @@ -711,10 +718,10 @@ static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name) { return S.Context.getTypedefType(Decl); } -/// Get the QualType instances of the return type and arguments for an OpenCL +/// Get the QualType instances of the return type and arguments for a ProgModel /// builtin function signature. -/// \param S (in) The Sema instance. -/// \param OpenCLBuiltin (in) The signature currently handled. +/// \param Context (in) The Context instance. +/// \param Builtin (in) The signature currently handled. /// \param GenTypeMaxCnt (out) Maximum number of types contained in a generic /// type used as return type or as argument. /// Only meaningful for generic types, otherwise equals 1. @@ -722,27 +729,31 @@ static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name) { /// \param ArgTypes (out) List of the possible argument types. For each /// argument, ArgTypes contains QualTypes for the Cartesian product /// of (vector sizes) x (types) . -static void GetQualTypesForOpenCLBuiltin( - Sema &S, const OpenCLBuiltinStruct &OpenCLBuiltin, unsigned &GenTypeMaxCnt, - SmallVector &RetTypes, +template +static void GetQualTypesForProgModelBuiltin( + Sema &S, const typename ProgModel::BuiltinStruct &Builtin, + unsigned &GenTypeMaxCnt, SmallVector &RetTypes, SmallVector, 5> &ArgTypes) { // Get the QualType instances of the return types. - unsigned Sig = SignatureTable[OpenCLBuiltin.SigTableIndex]; - OCL2Qual(S, TypeTable[Sig], RetTypes); + unsigned Sig = ProgModel::SignatureTable[Builtin.SigTableIndex]; + ProgModel::Bultin2Qual(S, ProgModel::TypeTable[Sig], RetTypes); GenTypeMaxCnt = RetTypes.size(); // Get the QualType instances of the arguments. // First type is the return type, skip it. - for (unsigned Index = 1; Index < OpenCLBuiltin.NumTypes; Index++) { + for (unsigned Index = 1; Index < Builtin.NumTypes; Index++) { SmallVector Ty; - OCL2Qual(S, TypeTable[SignatureTable[OpenCLBuiltin.SigTableIndex + Index]], - Ty); + ProgModel::Bultin2Qual( + S, + ProgModel::TypeTable[ProgModel::SignatureTable[Builtin.SigTableIndex + + Index]], + Ty); GenTypeMaxCnt = (Ty.size() > GenTypeMaxCnt) ? Ty.size() : GenTypeMaxCnt; ArgTypes.push_back(std::move(Ty)); } } -/// Create a list of the candidate function overloads for an OpenCL builtin +/// Create a list of the candidate function overloads for a ProgModel builtin /// function. /// \param Context (in) The ASTContext instance. /// \param GenTypeMaxCnt (in) Maximum number of types contained in a generic @@ -751,13 +762,13 @@ static void GetQualTypesForOpenCLBuiltin( /// \param FunctionList (out) List of FunctionTypes. /// \param RetTypes (in) List of the possible return types. /// \param ArgTypes (in) List of the possible types for the arguments. -static void GetOpenCLBuiltinFctOverloads( +static void GetProgModelBuiltinFctOverloads( ASTContext &Context, unsigned GenTypeMaxCnt, std::vector &FunctionList, SmallVector &RetTypes, - SmallVector, 5> &ArgTypes) { + SmallVector, 5> &ArgTypes, bool IsVariadic) { FunctionProtoType::ExtProtoInfo PI( Context.getDefaultCallingConvention(false, false, true)); - PI.Variadic = false; + PI.Variadic = IsVariadic; // Do not attempt to create any FunctionTypes if there are no return types, // which happens when a type belongs to a disabled extension. @@ -787,8 +798,22 @@ static void GetOpenCLBuiltinFctOverloads( } } -/// When trying to resolve a function name, if isOpenCLBuiltin() returns a -/// non-null pair, then the name is referencing an OpenCL +template +static bool isVersionInMask(const LangOptions &O, unsigned Mask); +template <> +bool isVersionInMask(const LangOptions &LO, unsigned Mask) { + return isOpenCLVersionContainedInMask(LO, Mask); +} + +// SPIRV Builtins are always permitted, since all builtins are 'SPIRV_ALL'. We +// have no corresponding language option to check, so we always include them. +template <> +bool isVersionInMask(const LangOptions &LO, unsigned Mask) { + return true; +} + +/// When trying to resolve a function name, if ProgModel::isBuiltin() returns a +/// non-null pair, then the name is referencing a /// builtin function. Add all candidate signatures to the LookUpResult. /// /// \param S (in) The Sema instance. @@ -796,10 +821,13 @@ static void GetOpenCLBuiltinFctOverloads( /// \param II (in) The identifier being resolved. /// \param FctIndex (in) Starting index in the BuiltinTable. /// \param Len (in) The signature list has Len elements. -static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, - IdentifierInfo *II, - const unsigned FctIndex, - const unsigned Len) { +template +static void InsertBuiltinDeclarationsFromTable( + Sema &S, LookupResult &LR, IdentifierInfo *II, const unsigned FctIndex, + const unsigned Len, + std::function + ProgModelFinalizer) { // The builtin function declaration uses generic types (gentype). bool HasGenType = false; @@ -810,19 +838,18 @@ static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, ASTContext &Context = S.Context; for (unsigned SignatureIndex = 0; SignatureIndex < Len; SignatureIndex++) { - const OpenCLBuiltinStruct &OpenCLBuiltin = - BuiltinTable[FctIndex + SignatureIndex]; + const typename ProgModel::BuiltinStruct &Builtin = + ProgModel::BuiltinTable[FctIndex + SignatureIndex]; // Ignore this builtin function if it is not available in the currently // selected language version. - if (!isOpenCLVersionContainedInMask(Context.getLangOpts(), - OpenCLBuiltin.Versions)) + if (!isVersionInMask(Context.getLangOpts(), Builtin.Versions)) continue; // Ignore this builtin function if it carries an extension macro that is // not defined. This indicates that the extension is not supported by the // target, so the builtin function should not be available. - StringRef Extensions = FunctionExtensionTable[OpenCLBuiltin.Extension]; + StringRef Extensions = ProgModel::FunctionExtensionTable[Builtin.Extension]; if (!Extensions.empty()) { SmallVector ExtVec; Extensions.split(ExtVec, " "); @@ -841,27 +868,27 @@ static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, SmallVector, 5> ArgTypes; // Obtain QualType lists for the function signature. - GetQualTypesForOpenCLBuiltin(S, OpenCLBuiltin, GenTypeMaxCnt, RetTypes, - ArgTypes); + GetQualTypesForProgModelBuiltin(S, Builtin, GenTypeMaxCnt, + RetTypes, ArgTypes); if (GenTypeMaxCnt > 1) { HasGenType = true; } // Create function overload for each type combination. std::vector FunctionList; - GetOpenCLBuiltinFctOverloads(Context, GenTypeMaxCnt, FunctionList, RetTypes, - ArgTypes); + GetProgModelBuiltinFctOverloads(Context, GenTypeMaxCnt, FunctionList, + RetTypes, ArgTypes, Builtin.IsVariadic); SourceLocation Loc = LR.getNameLoc(); DeclContext *Parent = Context.getTranslationUnitDecl(); - FunctionDecl *NewOpenCLBuiltin; + FunctionDecl *NewBuiltin; for (const auto &FTy : FunctionList) { - NewOpenCLBuiltin = FunctionDecl::Create( - Context, Parent, Loc, Loc, II, FTy, /*TInfo=*/nullptr, SC_Extern, - S.getCurFPFeatures().isFPConstrained(), false, - FTy->isFunctionProtoType()); - NewOpenCLBuiltin->setImplicit(); + NewBuiltin = FunctionDecl::Create(Context, Parent, Loc, Loc, II, FTy, + /*TInfo=*/nullptr, SC_Extern, + S.getCurFPFeatures().isFPConstrained(), + false, FTy->isFunctionProtoType()); + NewBuiltin->setImplicit(); // Create Decl objects for each parameter, adding them to the // FunctionDecl. @@ -869,25 +896,25 @@ static void InsertOCLBuiltinDeclarationsFromTable(Sema &S, LookupResult &LR, SmallVector ParmList; for (unsigned IParm = 0, e = FP->getNumParams(); IParm != e; ++IParm) { ParmVarDecl *Parm = ParmVarDecl::Create( - Context, NewOpenCLBuiltin, SourceLocation(), SourceLocation(), - nullptr, FP->getParamType(IParm), nullptr, SC_None, nullptr); + Context, NewBuiltin, SourceLocation(), SourceLocation(), nullptr, + FP->getParamType(IParm), nullptr, SC_None, nullptr); Parm->setScopeInfo(0, IParm); ParmList.push_back(Parm); } - NewOpenCLBuiltin->setParams(ParmList); + NewBuiltin->setParams(ParmList); // Add function attributes. - if (OpenCLBuiltin.IsPure) - NewOpenCLBuiltin->addAttr(PureAttr::CreateImplicit(Context)); - if (OpenCLBuiltin.IsConst) - NewOpenCLBuiltin->addAttr(ConstAttr::CreateImplicit(Context)); - if (OpenCLBuiltin.IsConv) - NewOpenCLBuiltin->addAttr(ConvergentAttr::CreateImplicit(Context)); - + if (Builtin.IsPure) + NewBuiltin->addAttr(PureAttr::CreateImplicit(Context)); + if (Builtin.IsConst) + NewBuiltin->addAttr(ConstAttr::CreateImplicit(Context)); + if (Builtin.IsConv) + NewBuiltin->addAttr(ConvergentAttr::CreateImplicit(Context)); if (!S.getLangOpts().OpenCLCPlusPlus) - NewOpenCLBuiltin->addAttr(OverloadableAttr::CreateImplicit(Context)); + NewBuiltin->addAttr(OverloadableAttr::CreateImplicit(Context)); - LR.addDecl(NewOpenCLBuiltin); + ProgModelFinalizer(Builtin, *NewBuiltin); + LR.addDecl(NewBuiltin); } } @@ -920,10 +947,31 @@ bool Sema::LookupBuiltin(LookupResult &R) { // Check if this is an OpenCL Builtin, and if so, insert its overloads. if (getLangOpts().OpenCL && getLangOpts().DeclareOpenCLBuiltins) { - auto Index = isOpenCLBuiltin(II->getName()); + auto Index = OpenCLBuiltin::isBuiltin(II->getName()); + if (Index.first) { + InsertBuiltinDeclarationsFromTable( + *this, R, II, Index.first - 1, Index.second, + [this](const OpenCLBuiltin::BuiltinStruct &OpenCLBuiltin, + FunctionDecl &NewOpenCLBuiltin) { + if (!this->getLangOpts().OpenCLCPlusPlus) + NewOpenCLBuiltin.addAttr( + OverloadableAttr::CreateImplicit(Context)); + }); + return true; + } + } + + // Check if this is a SPIR-V Builtin, and if so, insert its overloads. + if (getLangOpts().DeclareSPIRVBuiltins) { + auto Index = SPIRVBuiltin::isBuiltin(II->getName()); if (Index.first) { - InsertOCLBuiltinDeclarationsFromTable(*this, R, II, Index.first - 1, - Index.second); + InsertBuiltinDeclarationsFromTable( + *this, R, II, Index.first - 1, Index.second, + [this](const SPIRVBuiltin::BuiltinStruct &, + FunctionDecl &NewBuiltin) { + if (!this->getLangOpts().CPlusPlus) + NewBuiltin.addAttr(OverloadableAttr::CreateImplicit(Context)); + }); return true; } } diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index d2ee669debd0c..2be31fe525101 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -5084,8 +5084,9 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state, if (LangOpts.OpenCL) { // OpenCL v2.0 s6.12.5 - A block cannot be the return value of a // function. - if (T->isBlockPointerType() || T->isImageType() || T->isSamplerT() || - T->isPipeType()) { + if (!T->isSampledImageType() && + (T->isBlockPointerType() || T->isImageType() || T->isSamplerT() || + T->isPipeType())) { S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return) << T << 1 /*hint off*/; D.setInvalidType(true); diff --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp index c60f87a239857..e9c168a0a4d79 100644 --- a/clang/lib/Serialization/ASTCommon.cpp +++ b/clang/lib/Serialization/ASTCommon.cpp @@ -215,6 +215,13 @@ serialization::TypeIdxFromBuiltin(const BuiltinType *BT) { ID = PREDEF_TYPE_##Id##_ID; \ break; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + ID = PREDEF_TYPE_SAMPLED_##Id##_ID; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: \ ID = PREDEF_TYPE_##Id##_ID; \ diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index a033bccbe5061..d7f70e64369aa 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -6996,6 +6996,13 @@ QualType ASTReader::GetType(TypeID ID) { T = Context.SingletonId; \ break; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case PREDEF_TYPE_SAMPLED_##Id##_ID: \ + T = Context.Sampled##SingletonId; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case PREDEF_TYPE_##Id##_ID: \ T = Context.Id##Ty; \ diff --git a/clang/test/CodeGenOpenCL/sampled_image.cl b/clang/test/CodeGenOpenCL/sampled_image.cl new file mode 100644 index 0000000000000..f9ccd754a74aa --- /dev/null +++ b/clang/test/CodeGenOpenCL/sampled_image.cl @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - -cl-std=clc++ | FileCheck %s + +__attribute__((overloadable)) void my_read_image(__ocl_sampled_image1d_ro_t img); +__attribute__((overloadable)) void my_read_image(__ocl_sampled_image2d_ro_t img); + +void test_read_image(__ocl_sampled_image1d_ro_t img_ro, __ocl_sampled_image2d_ro_t img_2d) { + // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image1d_ro(%spirv.SampledImage.image1d_ro_t* %{{[0-9]+}}) + my_read_image(img_ro); + // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image2d_ro(%spirv.SampledImage.image2d_ro_t* %{{[0-9]+}}) + my_read_image(img_2d); +} diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp new file mode 100644 index 0000000000000..1d477a20dc6fc --- /dev/null +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup-win.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple x86_64-windows-msvc -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s + +float acos(float val) { + // CHECK: @"?acos@@YAMM@Z" + // CHECK: call float @"?__spirv_ocl_acos@@YAMM@Z" + return __spirv_ocl_acos(val); +} + +// CHECK: declare dso_local float @"?__spirv_ocl_acos@@YAMM@Z"(float) + +double acos(double val) { + // CHECK: @"?acos@@YANN@Z" + // CHECK: call double @"?__spirv_ocl_acos@@YANN@Z" + return __spirv_ocl_acos(val); +} + +// CHECK: declare dso_local double @"?__spirv_ocl_acos@@YANN@Z"(double) diff --git a/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp b/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp new file mode 100644 index 0000000000000..a6805c12aa55e --- /dev/null +++ b/clang/test/CodeGenSPIRV/spirv-builtin-lookup.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fdeclare-spirv-builtins -fsyntax-only -emit-llvm %s -o - | FileCheck %s + +float acos(float val) { + // CHECK: @_Z4acosf + // CHECK: call float @_Z16__spirv_ocl_acosf + return __spirv_ocl_acos(val); +} + +// CHECK: declare float @_Z16__spirv_ocl_acosf(float) + +double acos(double val) { + // CHECK: @_Z4acosd + // CHECK: call double @_Z16__spirv_ocl_acosd + return __spirv_ocl_acos(val); +} + +// CHECK: declare double @_Z16__spirv_ocl_acosd(double) diff --git a/clang/test/CodeGenSYCL/unique_stable_name.cpp b/clang/test/CodeGenSYCL/unique_stable_name.cpp index 462541a29fc2a..50b3f04fba139 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-linux-pc -fsycl-is-host -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s // CHECK: @[[LAMBDA_KERNEL3:[^\w]+]] = private unnamed_addr constant [[LAMBDA_K3_SIZE:\[[0-9]+ x i8\]]] c"_ZTSZ4mainEUlPZ4mainEUlvE_E_\00" // CHECK: @[[INT1:[^\w]+]] = private unnamed_addr constant [[INT_SIZE:\[[0-9]+ x i8\]]] c"_ZTSi\00" // CHECK: @[[STRING:[^\w]+]] = private unnamed_addr constant [[STRING_SIZE:\[[0-9]+ x i8\]]] c"_ZTSAppL_ZZ4mainE1jE_i\00", @@ -65,95 +65,105 @@ template kernelFunc(); } +template +void unnamed_kernel_single_task(KernelType kernelFunc) { + kernel_single_task(kernelFunc); +} + +template +void not_kernel_single_task(KernelType kernelFunc) { + kernelFunc(); +} + int main() { - kernel_single_task(func); - // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8 addrspace(4)* ()* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) + not_kernel_single_task(func); + // CHECK: call void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_(i8* ()* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv) auto l1 = []() { return 1; }; auto l2 = [](decltype(l1) *l = nullptr) { return 2; }; - kernel_single_task(l2); + kernel_single_task(l2); puts(__builtin_sycl_unique_stable_name(decltype(l2))); - // CHECK: call spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_ - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]]* @[[LAMBDA_KERNEL3]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_ + // CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_K3_SIZE]], [[LAMBDA_K3_SIZE]]* @[[LAMBDA_KERNEL3]], i32 0, i32 0)) constexpr const char str[] = "lalala"; static_assert(__builtin_strcmp(__builtin_sycl_unique_stable_name(decltype(str)), "_ZTSA7_Kc\0") == 0, "unexpected mangling"); int i = 0; puts(__builtin_sycl_unique_stable_name(decltype(i++))); - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT1]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT1]], i32 0, i32 0)) // FIXME: Ensure that j is incremented because VLAs are terrible. int j = 55; puts(__builtin_sycl_unique_stable_name(int[++j])); - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]]* @[[STRING]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[STRING_SIZE]], [[STRING_SIZE]]* @[[STRING]], i32 0, i32 0)) - // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_ - // CHECK: declare spir_func i8 addrspace(4)* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv - // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE7kernel3Z4mainEUlPZ4mainEUlvE_E_EvT0_ - // CHECK: define internal spir_func void @_Z18kernel_single_taskIZ4mainE6kernelZ4mainEUlvE0_EvT0_ + // CHECK: define internal void @_Z22not_kernel_single_taskIZ4mainE7kernel2PFPKcvEEvT0_ + // CHECK: declare i8* @_Z4funcI4DerpEDTu33__builtin_sycl_unique_stable_nameDtsrT_3strEEEv + // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlPZ4mainEUlvE_E_S2_EvT0_ + // CHECK: define internal void @_Z18kernel_single_taskIZ4mainEUlvE0_S0_EvT0_ - kernel_single_task( + unnamed_kernel_single_task( []() { puts(__builtin_sycl_unique_stable_name(int)); - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT2]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT2]], i32 0, i32 0)) auto x = []() {}; puts(__builtin_sycl_unique_stable_name(decltype(x))); - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_X_SIZE]], [[LAMBDA_X_SIZE]]* @[[LAMBDA_X]], i32 0, i32 0)) DEF_IN_MACRO(); - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*)) - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_X]], i32 0, i32 0)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_SIZE]], [[MACRO_SIZE]]* @[[MACRO_Y]], i32 0, i32 0)) MACRO_CALLS_MACRO(); - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]], i32 0, i32 0) to i8 addrspace(4)*)) - // CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]], i32 0, i32 0) to i8 addrspace(4)*)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_X]], i32 0, i32 0)) + // CHECK: call void @puts(i8* getelementptr inbounds ([[MACRO_MACRO_SIZE]], [[MACRO_MACRO_SIZE]]* @[[MACRO_MACRO_Y]], i32 0, i32 0)) template_param(); - // CHECK: call spir_func void @_Z14template_paramIiEvv + // CHECK: call void @_Z14template_paramIiEvv template_param(); - // CHECK: call spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv + // CHECK: call void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv lambda_in_dependent_function(); - // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIiEvv + // CHECK: call void @_Z28lambda_in_dependent_functionIiEvv lambda_in_dependent_function(); - // CHECK: call spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv + // CHECK: call void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv lambda_no_dep(3, 5.5); - // CHECK: call spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 3, double 5.500000e+00) + // CHECK: call void @_Z13lambda_no_depIidEvT_T0_(i32 3, double 5.500000e+00) int a = 5; double b = 10.7; auto y = [](int a) { return a; }; auto z = [](double b) { return b; }; lambda_two_dep(); - // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv + // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv lambda_two_dep(); - // CHECK: call spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv + // CHECK: call void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv }); } -// CHECK: define linkonce_odr spir_func void @_Z14template_paramIiEvv -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT3]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define linkonce_odr void @_Z14template_paramIiEvv +// CHECK: call void @puts(i8* getelementptr inbounds ([[INT_SIZE]], [[INT_SIZE]]* @[[INT3]], i32 0, i32 0)) -// CHECK: define internal spir_func void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[LAMBDA_SIZE]], [[LAMBDA_SIZE]]* @[[LAMBDA]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define internal void @_Z14template_paramIZZ4mainENKUlvE0_clEvEUlvE_Evv +// CHECK: call void @puts(i8* getelementptr inbounds ([[LAMBDA_SIZE]], [[LAMBDA_SIZE]]* @[[LAMBDA]], i32 0, i32 0)) -// CHECK: define linkonce_odr spir_func void @_Z28lambda_in_dependent_functionIiEvv -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define linkonce_odr void @_Z28lambda_in_dependent_functionIiEvv +// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_INT_SIZE]], [[DEP_INT_SIZE]]* @[[LAMBDA_IN_DEP_INT]], i32 0, i32 0)) -// CHECK: define internal spir_func void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define internal void @_Z28lambda_in_dependent_functionIZZ4mainENKUlvE0_clEvEUlvE_Evv +// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA_SIZE]], [[DEP_LAMBDA_SIZE]]* @[[LAMBDA_IN_DEP_X]], i32 0, i32 0)) -// CHECK: define linkonce_odr spir_func void @_Z13lambda_no_depIidEvT_T0_(i32 %a, double %b) -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[NO_DEP_LAMBDA_SIZE]], [[NO_DEP_LAMBDA_SIZE]]* @[[LAMBDA_NO_DEP]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define linkonce_odr void @_Z13lambda_no_depIidEvT_T0_(i32 %a, double %b) +// CHECK: call void @puts(i8* getelementptr inbounds ([[NO_DEP_LAMBDA_SIZE]], [[NO_DEP_LAMBDA_SIZE]]* @[[LAMBDA_NO_DEP]], i32 0, i32 0)) -// CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA1_SIZE]], [[DEP_LAMBDA1_SIZE]]* @[[LAMBDA_TWO_DEP]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUliE_ZZ4mainENKS0_clEvEUldE_Evv +// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA1_SIZE]], [[DEP_LAMBDA1_SIZE]]* @[[LAMBDA_TWO_DEP]], i32 0, i32 0)) -// CHECK: define internal spir_func void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv -// CHECK: call spir_func void @puts(i8 addrspace(4)* addrspacecast (i8* getelementptr inbounds ([[DEP_LAMBDA2_SIZE]], [[DEP_LAMBDA2_SIZE]]* @[[LAMBDA_TWO_DEP2]], i32 0, i32 0) to i8 addrspace(4)*)) +// CHECK: define internal void @_Z14lambda_two_depIZZ4mainENKUlvE0_clEvEUldE_ZZ4mainENKS0_clEvEUliE_Evv +// CHECK: call void @puts(i8* getelementptr inbounds ([[DEP_LAMBDA2_SIZE]], [[DEP_LAMBDA2_SIZE]]* @[[LAMBDA_TWO_DEP2]], i32 0, i32 0)) diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp index 7c7979f712f05..fb02c1b876106 100644 --- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp +++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp @@ -1,22 +1,36 @@ -// RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s -// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s - +// RUN: %clang_cc1 -triple spir64-unknown-unknown -aux-triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsycl-is-device -disable-llvm-passes -fsycl-is-device -emit-llvm %s -o - | FileCheck %s --check-prefixes=WIN,CHECK template __attribute__((sycl_kernel)) void kernel(Func F){ F(); } +template +void kernel_wrapper(Func F) { + kernel(F); +} + template __attribute__((sycl_kernel)) void kernel2(Func F){ F(1); } +template +void kernel2_wrapper(Func F) { + kernel2(F); +} + template __attribute__((sycl_kernel)) void kernel3(Func F){ F(1.1); } +template +void kernel3_wrapper(Func F) { + kernel3(F); +} + int main() { int i; double d; @@ -25,15 +39,17 @@ int main() { auto lambda2 = [](int){}; auto lambda3 = [](double){}; - kernel(lambda1); - kernel2(lambda2); - kernel3(lambda3); + kernel_wrapper(lambda1); + kernel2_wrapper(lambda2); + kernel3_wrapper(lambda3); // Ensure the kernels are named the same between the device and host // invocations. + kernel_wrapper([](){ (void)__builtin_sycl_unique_stable_name(decltype(lambda1)); (void)__builtin_sycl_unique_stable_name(decltype(lambda2)); (void)__builtin_sycl_unique_stable_name(decltype(lambda3)); + }); // Make sure the following 3 are the same between the host and device compile. // Note that these are NOT the same value as eachother, they differ by the @@ -41,4 +57,11 @@ int main() { // CHECK: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUlvE_\00" // CHECK: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUliE_\00" // CHECK: private unnamed_addr constant [17 x i8] c"_ZTSZ4mainEUldE_\00" + + // On Windows, ensure that we haven't broken the 'lambda numbering' for thex + // lambda itself. + // WIN: define internal void @"??R line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t)' +void kernel ker(__ocl_sampled_image1d_ro_t src1, __ocl_sampled_image2d_ro_t src2) { + // CHECK: CallExpr + // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image1d_ro_t)' + foo(src1); + // CHECK: CallExpr + // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image2d_ro_t)' + foo(src2); +} diff --git a/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp new file mode 100644 index 0000000000000..b30dfc9ae3b6f --- /dev/null +++ b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify %s + +// Verify that invalid call to __spirv_ocl_acos (no viable overloads) get diagnosed + +struct InvalidType {}; +void acos(InvalidType Invalid) { + __spirv_ocl_acos(Invalid); // expected-error {{no matching function for call to '__spirv_ocl_acos'}} + // expected-note@-1 + {{candidate function not viable: no known conversion from}} + // too many params + __spirv_ocl_acos(42.f, 42.f); // expected-error {{no matching function for call to '__spirv_ocl_acos'}} + // expected-note@-1 + {{candidate function not viable: requires 1 argument, but 2 were provided}} +} diff --git a/clang/test/SemaSYCL/spirv-builtin-lookup.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup.cpp new file mode 100644 index 0000000000000..df28ec6b4337c --- /dev/null +++ b/clang/test/SemaSYCL/spirv-builtin-lookup.cpp @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -fdeclare-spirv-builtins -fsyntax-only -verify %s +// expected-no-diagnostics + +// Verify that __spirv_ocl_acos is recognized as a builtin + +float acos(float val) { + return __spirv_ocl_acos(val); +} + +double acos(double val) { + return __spirv_ocl_acos(val); +} + +typedef int int4 __attribute__((ext_vector_type(4))); +typedef float float4 __attribute__((ext_vector_type(4))); + +int4 ilogb() { + float4 f4 = {0.f, 0.f, 0.f, 0.f}; + int4 i4 = __spirv_ocl_ilogb(f4); + return i4; +} + +double sincos(double val, double *res) { + return __spirv_ocl_sincos(val, res); +} + +double dot(float4 v1, float4 v2) { + return __spirv_Dot(v1, v2); +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index ab3bf222430a9..be025b0c64044 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -1538,6 +1538,11 @@ bool CursorVisitor::VisitBuiltinTypeLoc(BuiltinTypeLoc TL) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtTYpe, Id, Ext) case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" case BuiltinType::OCLSampler: diff --git a/clang/tools/libclang/CXType.cpp b/clang/tools/libclang/CXType.cpp index 822ab3bb64b8e..35cf1cd79a03a 100644 --- a/clang/tools/libclang/CXType.cpp +++ b/clang/tools/libclang/CXType.cpp @@ -69,7 +69,11 @@ static CXTypeKind GetBuiltinTypeKind(const BuiltinType *BT) { BTCASE(ObjCSel); #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) BTCASE(Id); #include "clang/Basic/OpenCLImageTypes.def" -#undef IMAGE_TYPE +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + BTCASE(Sampled##Id); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) BTCASE(Id); #include "clang/Basic/OpenCLExtensionTypes.def" BTCASE(OCLSampler); @@ -614,6 +618,11 @@ CXString clang_getTypeKindSpelling(enum CXTypeKind K) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) TKIND(Id); #include "clang/Basic/OpenCLImageTypes.def" #undef IMAGE_TYPE +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) TKIND(Sampled##Id); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" +#undef IMAGE_TYPE #define EXT_OPAQUE_TYPE(ExtTYpe, Id, Ext) TKIND(Id); #include "clang/Basic/OpenCLExtensionTypes.def" TKIND(OCLSampler); diff --git a/clang/utils/TableGen/CMakeLists.txt b/clang/utils/TableGen/CMakeLists.txt index 6379cc4e11e83..a1f13a223b35a 100644 --- a/clang/utils/TableGen/CMakeLists.txt +++ b/clang/utils/TableGen/CMakeLists.txt @@ -11,7 +11,7 @@ add_tablegen(clang-tblgen CLANG ClangDataCollectorsEmitter.cpp ClangDiagnosticsEmitter.cpp ClangOpcodesEmitter.cpp - ClangOpenCLBuiltinEmitter.cpp + ClangProgModelBuiltinEmitter.cpp ClangOptionDocEmitter.cpp ClangSACheckersEmitter.cpp ClangSyntaxEmitter.cpp diff --git a/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp similarity index 86% rename from clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp rename to clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp index 4795b008dda3c..4cdfbac8c2a0d 100644 --- a/clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp +++ b/clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp @@ -1,4 +1,4 @@ -//===- ClangOpenCLBuiltinEmitter.cpp - Generate Clang OpenCL Builtin handling +//===- ClangProgModelBuiltinEmitter.cpp - Generate Clang Builtin handling // // The LLVM Compiler Infrastructure // @@ -8,8 +8,8 @@ // //===----------------------------------------------------------------------===// // -// These backends consume the definitions of OpenCL builtin functions in -// clang/lib/Sema/OpenCLBuiltins.td and produce builtin handling code for +// These backends consume the definitions of builtin functions in +// clang/lib/Sema/*Builtins.td and produce builtin handling code for // inclusion in SemaLookup.cpp, or a test file that calls all declared builtins. // //===----------------------------------------------------------------------===// @@ -39,19 +39,19 @@ struct BuiltinTableEntries { std::vector> Signatures; }; -// This tablegen backend emits code for checking whether a function is an -// OpenCL builtin function. If so, all overloads of this function are -// added to the LookupResult. The generated include file is used by -// SemaLookup.cpp +// This tablegen backend emits code for checking whether a function is a +// builtin function of a programming model. If so, all overloads of this +// function are added to the LookupResult. The generated include file is +// used by SemaLookup.cpp // -// For a successful lookup of e.g. the "cos" builtin, isOpenCLBuiltin("cos") +// For a successful lookup of e.g. the "cos" builtin, isBuiltin("cos") // returns a pair . // BuiltinTable[Index] to BuiltinTable[Index + Len] contains the pairs // of the overloads of "cos". // SignatureTable[SigIndex] to SignatureTable[SigIndex + SigLen] contains // one of the signatures of "cos". The SignatureTable entry can be // referenced by other functions, e.g. "sin", to exploit the fact that -// many OpenCL builtins share the same signature. +// many builtins may share the same signature. // // The file generated by this TableGen emitter contains the following: // @@ -62,7 +62,7 @@ struct BuiltinTableEntries { // entry in this table when the builtin requires a particular (set of) // extension(s) to be enabled. // -// * OpenCLTypeStruct TypeTable[] +// * ProgModelTypeStruct TypeTable[] // Type information for return types and arguments. // // * unsigned SignatureTable[] @@ -71,39 +71,47 @@ struct BuiltinTableEntries { // signature, where the first entry is the return type and subsequent // entries are the argument types. // -// * OpenCLBuiltinStruct BuiltinTable[] -// Each entry represents one overload of an OpenCL builtin function and +// * BuiltinStruct BuiltinTable[] +// Each entry represents one overload of a builtin function and // consists of an index into the SignatureTable and the number of arguments. // -// * std::pair isOpenCLBuiltin(llvm::StringRef Name) -// Find out whether a string matches an existing OpenCL builtin function +// * std::pair isBuiltin(llvm::StringRef Name) +// Find out whether a string matches an existing builtin function // name and return an index into BuiltinTable and the number of overloads. // -// * void OCL2Qual(Sema&, OpenCLTypeStruct, std::vector&) -// Convert an OpenCLTypeStruct type to a list of QualType instances. -// One OpenCLTypeStruct can represent multiple types, primarily when using +// * void Bultin2Qual(Sema&, ProgModelTypeStruct, std::vector&) +// Convert an ProgModelTypeStruct type to a list of QualType instances. +// One ProgModelTypeStruct can represent multiple types, primarily when using // GenTypes. // class BuiltinNameEmitter { public: - BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS) - : Records(Records), OS(OS) {} + BuiltinNameEmitter(RecordKeeper &Records, raw_ostream &OS, + llvm::StringRef Family) + : Records(Records), OS(OS), Family(Family), + ClassName((Family + "Builtin").str()) {} // Entrypoint to generate the functions and structures for checking - // whether a function is an OpenCL builtin function. + // whether a function is a builtin function. void Emit(); private: // A list of indices into the builtin function table. using BuiltinIndexListTy = SmallVector; - // Contains OpenCL builtin functions and related information, stored as + // Contains builtin functions and related information, stored as // Record instances. They are coming from the associated TableGen file. RecordKeeper &Records; // The output file. raw_ostream &OS; + // Family for which the builtin are for. + llvm::StringRef Family; + + // Class for which the builtin are for. + std::string ClassName; + // Helper function for BuiltinNameEmitter::EmitDeclarations. Generate enum // definitions in the Output string parameter, and save their Record instances // in the List parameter. @@ -124,7 +132,7 @@ class BuiltinNameEmitter { // FctOverloadMap and TypeMap. void GetOverloads(); - // Compare two lists of signatures and check that e.g. the OpenCL version, + // Compare two lists of signatures and check that e.g. the version, // function attributes, and extension are equal for each signature. // \param Candidate (in) Entry in the SignatureListMap to check. // \param SignatureList (in) List of signatures of the considered function. @@ -137,14 +145,14 @@ class BuiltinNameEmitter { // SignatureListMap. // Some builtin functions have the same list of signatures, for example the // "sin" and "cos" functions. To save space in the BuiltinTable, the - // "isOpenCLBuiltin" function will have the same output for these two + // "isBuiltin" function will have the same output for these two // function names. void GroupBySignature(); // Emit the FunctionExtensionTable that lists all function extensions. void EmitExtensionTable(); - // Emit the TypeTable containing all types used by OpenCL builtins. + // Emit the TypeTable containing all types used by the builtins. void EmitTypeTable(); // Emit the SignatureTable. This table contains all the possible signatures. @@ -158,7 +166,7 @@ class BuiltinNameEmitter { void EmitSignatureTable(); // Emit the BuiltinTable table. This table contains all the overloads of - // each function, and is a struct OpenCLBuiltinDecl. + // each function, and is a struct BuiltinDecl. // E.g.: // // 891 convert_float2_rtn // { 58, 2, 3, 100, 0 }, @@ -166,12 +174,12 @@ class BuiltinNameEmitter { // 1 argument (+1 for the return type), stored at index 58 in // the SignatureTable. This prototype requires extension "3" in the // FunctionExtensionTable. The last two values represent the minimum (1.0) - // and maximum (0, meaning no max version) OpenCL version in which this + // and maximum (0, meaning no max version) version in which this // overload is supported. void EmitBuiltinTable(); - // Emit a StringMatcher function to check whether a function name is an - // OpenCL builtin function name. + // Emit a StringMatcher function to check whether a function name is a + // builtin function name. void EmitStringMatcher(); // Emit a function returning the clang QualType instance associated with @@ -200,14 +208,14 @@ class BuiltinNameEmitter { MapVector>> FctOverloadMap; - // Contains the map of OpenCL types to their index in the TypeTable. + // Contains the map of types to their index in the TypeTable. MapVector TypeMap; - // List of OpenCL function extensions mapping extension strings to + // List of function extensions mapping extension strings to // an index into the FunctionExtensionTable. StringMap FunctionExtensionIndex; - // List of OpenCL type names in the same order as in enum OpenCLTypeID. + // List of type names in the same order as in enum TypeID. // This list does not contain generic types. std::vector TypeList; @@ -317,7 +325,8 @@ class OpenCLBuiltinTestEmitter : public OpenCLBuiltinFileEmitterBase { } // namespace void BuiltinNameEmitter::Emit() { - emitSourceFileHeader("OpenCL Builtin handling", OS); + std::string Banner = (Family + " Builtin handling").str(); + emitSourceFileHeader(Banner, OS); OS << "#include \"llvm/ADT/StringRef.h\"\n"; OS << "using namespace clang;\n\n"; @@ -348,7 +357,7 @@ void BuiltinNameEmitter::ExtractEnumTypes(std::vector &Types, for (const auto *T : Types) { if (TypesSeen.find(T->getValueAsString("Name")) == TypesSeen.end()) { - SS << " OCLT_" + T->getValueAsString("Name") << ",\n"; + SS << " TID_" + T->getValueAsString("Name") << ",\n"; // Save the type names in the same order as their enum value. Note that // the Record can be a VectorType or something else, only the name is // important. @@ -360,8 +369,11 @@ void BuiltinNameEmitter::ExtractEnumTypes(std::vector &Types, } void BuiltinNameEmitter::EmitDeclarations() { + OS << "class " << ClassName << " {\n\n" + << "public:\n\n"; + // Enum of scalar type names (float, int, ...) and generic type sets. - OS << "enum OpenCLTypeID {\n"; + OS << "enum TypeID {\n"; StringMap TypesSeen; std::string GenTypeEnums; @@ -384,17 +396,17 @@ void BuiltinNameEmitter::EmitDeclarations() { // Structure definitions. OS << R"( // Image access qualifier. -enum OpenCLAccessQual : unsigned char { - OCLAQ_None, - OCLAQ_ReadOnly, - OCLAQ_WriteOnly, - OCLAQ_ReadWrite +enum AccessQual : unsigned char { + AQ_None, + AQ_ReadOnly, + AQ_WriteOnly, + AQ_ReadWrite }; // Represents a return type or argument type. -struct OpenCLTypeStruct { +struct ProgModelTypeStruct { // A type (e.g. float, int, ...). - const OpenCLTypeID ID; + const TypeID ID; // Vector size (if applicable; 0 for scalars and generic types). const unsigned VectorWidth; // 0 if the type is not a pointer. @@ -404,14 +416,14 @@ struct OpenCLTypeStruct { // 0 if the type is not volatile. const bool IsVolatile : 1; // Access qualifier. - const OpenCLAccessQual AccessQualifier; + const AccessQual AccessQualifier; // Address space of the pointer (if applicable). const LangAS AS; }; -// One overload of an OpenCL builtin function. -struct OpenCLBuiltinStruct { - // Index of the signature in the OpenCLTypeStruct table. +// One overload of a builtin function. +struct BuiltinStruct { + // Index of the signature in the ProgModelTypeStruct table. const unsigned SigTableIndex; // Entries between index SigTableIndex and (SigTableIndex + NumTypes - 1) in // the SignatureTable represent the complete signature. The first type at @@ -423,13 +435,26 @@ struct OpenCLBuiltinStruct { const bool IsConst : 1; // Function attribute __attribute__((convergent)) const bool IsConv : 1; + // 0 if the function is not variadic. + const bool IsVariadic : 1; // OpenCL extension(s) required for this overload. const unsigned short Extension; // OpenCL versions in which this overload is available. const unsigned short Versions; }; +static const char *FunctionExtensionTable[]; +static const ProgModelTypeStruct TypeTable[]; +static const unsigned short SignatureTable[]; +static const BuiltinStruct BuiltinTable[]; + +static std::pair isBuiltin(llvm::StringRef Name); +static void Bultin2Qual(Sema &Sema, const ProgModelTypeStruct &Ty, + llvm::SmallVectorImpl &QT); + )"; + + OS << "}; // class " << ClassName << "\n"; } // Verify that the combination of GenTypes in a signature is supported. @@ -452,7 +477,8 @@ static void VerifySignature(const std::vector &Signature, if (NVecSizes != GenTypeVecSizes && NVecSizes != 1) { if (GenTypeVecSizes > 1) { // We already saw a gentype with a different number of vector sizes. - PrintFatalError(BuiltinRec->getLoc(), + PrintFatalError( + BuiltinRec->getLoc(), "number of vector sizes should be equal or 1 for all gentypes " "in a declaration"); } @@ -465,7 +491,8 @@ static void VerifySignature(const std::vector &Signature, if (NTypes != GenTypeTypes && NTypes != 1) { if (GenTypeTypes > 1) { // We already saw a gentype with a different number of types. - PrintFatalError(BuiltinRec->getLoc(), + PrintFatalError( + BuiltinRec->getLoc(), "number of types should be equal or 1 for all gentypes " "in a declaration"); } @@ -514,7 +541,7 @@ void BuiltinNameEmitter::GetOverloads() { } void BuiltinNameEmitter::EmitExtensionTable() { - OS << "static const char *FunctionExtensionTable[] = {\n"; + OS << "const char * " << ClassName << "::FunctionExtensionTable[] = {\n"; unsigned Index = 0; std::vector FuncExtensions = Records.getAllDerivedDefinitions("FunctionExtension"); @@ -531,22 +558,22 @@ void BuiltinNameEmitter::EmitExtensionTable() { } void BuiltinNameEmitter::EmitTypeTable() { - OS << "static const OpenCLTypeStruct TypeTable[] = {\n"; + OS << "const " << ClassName << "::ProgModelTypeStruct " << ClassName + << "::TypeTable[] = {\n"; for (const auto &T : TypeMap) { const char *AccessQual = StringSwitch(T.first->getValueAsString("AccessQualifier")) - .Case("RO", "OCLAQ_ReadOnly") - .Case("WO", "OCLAQ_WriteOnly") - .Case("RW", "OCLAQ_ReadWrite") - .Default("OCLAQ_None"); + .Case("RO", "AQ_ReadOnly") + .Case("WO", "AQ_WriteOnly") + .Case("RW", "AQ_ReadWrite") + .Default("AQ_None"); OS << " // " << T.second << "\n" - << " {OCLT_" << T.first->getValueAsString("Name") << ", " + << " {TID_" << T.first->getValueAsString("Name") << ", " << T.first->getValueAsInt("VecWidth") << ", " << T.first->getValueAsBit("IsPointer") << ", " << T.first->getValueAsBit("IsConst") << ", " - << T.first->getValueAsBit("IsVolatile") << ", " - << AccessQual << ", " + << T.first->getValueAsBit("IsVolatile") << ", " << AccessQual << ", " << T.first->getValueAsString("AddrSpace") << "},\n"; } OS << "};\n\n"; @@ -554,9 +581,9 @@ void BuiltinNameEmitter::EmitTypeTable() { void BuiltinNameEmitter::EmitSignatureTable() { // Store a type (e.g. int, float, int2, ...). The type is stored as an index - // of a struct OpenCLType table. Multiple entries following each other form a - // signature. - OS << "static const unsigned short SignatureTable[] = {\n"; + // of a struct ProgModelTypeStruct table. Multiple entries following each + // other form a signature. + OS << "const unsigned short " << ClassName << "::SignatureTable[] = {\n"; for (const auto &P : SignaturesList) { OS << " // " << P.second << "\n "; for (const Record *R : P.first) { @@ -600,7 +627,8 @@ static unsigned short EncodeVersions(unsigned int MinVersion, void BuiltinNameEmitter::EmitBuiltinTable() { unsigned Index = 0; - OS << "static const OpenCLBuiltinStruct BuiltinTable[] = {\n"; + OS << "const " << ClassName << "::BuiltinStruct " << ClassName + << "::BuiltinTable[] = {\n"; for (const auto &SLM : SignatureListMap) { OS << " // " << (Index + 1) << ": "; @@ -621,6 +649,7 @@ void BuiltinNameEmitter::EmitBuiltinTable() { << (Overload.first->getValueAsBit("IsPure")) << ", " << (Overload.first->getValueAsBit("IsConst")) << ", " << (Overload.first->getValueAsBit("IsConv")) << ", " + << (Overload.first->getValueAsBit("IsVariadic")) << ", " << FunctionExtensionIndex[ExtName] << ", " << EncodeVersions(MinVersion, MaxVersion) << " },\n"; Index++; @@ -635,14 +664,14 @@ bool BuiltinNameEmitter::CanReuseSignature( assert(Candidate->size() == SignatureList.size() && "signature lists should have the same size"); - auto &CandidateSigs = - SignatureListMap.find(Candidate)->second.Signatures; + auto &CandidateSigs = SignatureListMap.find(Candidate)->second.Signatures; for (unsigned Index = 0; Index < Candidate->size(); Index++) { const Record *Rec = SignatureList[Index].first; const Record *Rec2 = CandidateSigs[Index].first; if (Rec->getValueAsBit("IsPure") == Rec2->getValueAsBit("IsPure") && Rec->getValueAsBit("IsConst") == Rec2->getValueAsBit("IsConst") && Rec->getValueAsBit("IsConv") == Rec2->getValueAsBit("IsConv") && + Rec->getValueAsBit("IsVariadic") == Rec2->getValueAsBit("IsVariadic") && Rec->getValueAsDef("MinVersion")->getValueAsInt("ID") == Rec2->getValueAsDef("MinVersion")->getValueAsInt("ID") && Rec->getValueAsDef("MaxVersion")->getValueAsInt("ID") == @@ -719,27 +748,29 @@ void BuiltinNameEmitter::EmitStringMatcher() { } OS << R"( -// Find out whether a string matches an existing OpenCL builtin function name. +// Find out whether a string matches an existing builtin function name. // Returns: A pair <0, 0> if no name matches. // A pair indexing the BuiltinTable if the name is -// matching an OpenCL builtin function. -static std::pair isOpenCLBuiltin(llvm::StringRef Name) { - +// matching a builtin function. )"; + OS << "std::pair " << ClassName + << "::isBuiltin(llvm::StringRef Name) {\n\n"; + StringMatcher("Name", ValidBuiltins, OS).Emit(0, true); OS << " return std::make_pair(0, 0);\n"; - OS << "} // isOpenCLBuiltin\n"; + OS << "} // isBuiltin\n"; } void BuiltinNameEmitter::EmitQualTypeFinder() { OS << R"( +// Convert an ProgModelTypeStruct type to a list of QualTypes. static QualType getOpenCLEnumType(Sema &S, llvm::StringRef Name); static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name); -// Convert an OpenCLTypeStruct type to a list of QualTypes. +// Convert an ProgModelTypeStruct type to a list of QualTypes. // Generic types represent multiple types and vector sizes, thus a vector // is returned. The conversion is done in two steps: // Step 1: A switch statement fills a vector with scalar base types for the @@ -747,8 +778,13 @@ static QualType getOpenCLTypedefType(Sema &S, llvm::StringRef Name); // or a single scalar type for non generic types. // Step 2: Qualifiers and other type properties such as vector size are // applied. -static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, - llvm::SmallVectorImpl &QT) { +)"; + + OS << "void " << ClassName + << "::Bultin2Qual(Sema &S, const ProgModelTypeStruct &Ty, " + "llvm::SmallVectorImpl &QT) {\n"; + + OS << R"( ASTContext &Context = S.Context; // Number of scalar types in the GenType. unsigned GenTypeNumTypes; @@ -758,8 +794,8 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, // Generate list of vector sizes for each generic type. for (const auto *VectList : Records.getAllDerivedDefinitions("IntList")) { - OS << " constexpr unsigned List" - << VectList->getValueAsString("Name") << "[] = {"; + OS << " constexpr unsigned List" << VectList->getValueAsString("Name") + << "[] = {"; for (const auto V : VectList->getValueAsListOfInts("List")) { OS << V << ", "; } @@ -793,16 +829,16 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, // tells which one is needed. Emit a switch statement that puts the // corresponding QualType into "QT". for (const auto &ITE : ImageTypesMap) { - OS << " case OCLT_" << ITE.getKey() << ":\n" + OS << " case TID_" << ITE.getKey() << ":\n" << " switch (Ty.AccessQualifier) {\n" - << " case OCLAQ_None:\n" + << " case AQ_None:\n" << " llvm_unreachable(\"Image without access qualifier\");\n"; for (const auto &Image : ITE.getValue()) { OS << StringSwitch( Image->getValueAsString("AccessQualifier")) - .Case("RO", " case OCLAQ_ReadOnly:\n") - .Case("WO", " case OCLAQ_WriteOnly:\n") - .Case("RW", " case OCLAQ_ReadWrite:\n") + .Case("RO", " case AQ_ReadOnly:\n") + .Case("WO", " case AQ_WriteOnly:\n") + .Case("RW", " case AQ_ReadWrite:\n") << " QT.push_back(" << Image->getValueAsDef("QTExpr")->getValueAsString("TypeExpr") << ");\n" @@ -814,7 +850,7 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, // Switch cases for generic types. for (const auto *GenType : Records.getAllDerivedDefinitions("GenericType")) { - OS << " case OCLT_" << GenType->getValueAsString("Name") << ": {\n"; + OS << " case TID_" << GenType->getValueAsString("Name") << ": {\n"; // Build the Cartesian product of (vector sizes) x (types). Only insert // the plain scalar types for now; other type information such as vector @@ -831,8 +867,14 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, OS << " if (S.getPreprocessor().isMacroDefined(\"" << Ext << "\")) {\n "; } - OS << " TypeList.push_back(" - << T->getValueAsDef("QTExpr")->getValueAsString("TypeExpr") << ");\n"; + if (T->getValueAsDef("QTExpr")->isSubClassOf("QualTypeFromFunction")) + OS << " TypeList.push_back(" + << T->getValueAsDef("QTExpr")->getValueAsString("TypeExpr") + << "(Context));\n"; + else + OS << " TypeList.push_back(" + << T->getValueAsDef("QTExpr")->getValueAsString("TypeExpr") + << ");\n"; if (!Ext.empty()) { OS << " }\n"; } @@ -875,7 +917,7 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, if (QT->getValueAsBit("IsAbstract") == 1) continue; // Emit the cases for non generic, non image types. - OS << " case OCLT_" << T->getValueAsString("Name") << ":\n"; + OS << " case TID_" << T->getValueAsString("Name") << ":\n"; StringRef Ext = T->getValueAsDef("Extension")->getValueAsString("ExtName"); // If this type depends on an extension, ensure the extension macro is @@ -884,7 +926,11 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, OS << " if (S.getPreprocessor().isMacroDefined(\"" << Ext << "\")) {\n "; } - OS << " QT.push_back(" << QT->getValueAsString("TypeExpr") << ");\n"; + if (QT->isSubClassOf("QualTypeFromFunction")) + OS << " QT.push_back(" << QT->getValueAsString("TypeExpr") + << "(Context));\n"; + else + OS << " QT.push_back(" << QT->getValueAsString("TypeExpr") << ");\n"; if (!Ext.empty()) { OS << " }\n"; } @@ -943,8 +989,8 @@ static void OCL2Qual(Sema &S, const OpenCLTypeStruct &Ty, } )"; - // End of the "OCL2Qual" function. - OS << "\n} // OCL2Qual\n"; + // End of the "Bultin2Qual" function. + OS << "\n} // Bultin2Qual\n"; } std::string OpenCLBuiltinFileEmitterBase::getTypeString(const Record *Type, @@ -1185,7 +1231,12 @@ void OpenCLBuiltinTestEmitter::emit() { } void clang::EmitClangOpenCLBuiltins(RecordKeeper &Records, raw_ostream &OS) { - BuiltinNameEmitter NameChecker(Records, OS); + BuiltinNameEmitter NameChecker(Records, OS, "OpenCL"); + NameChecker.Emit(); +} + +void clang::EmitClangSPIRVBuiltins(RecordKeeper &Records, raw_ostream &OS) { + BuiltinNameEmitter NameChecker(Records, OS, "SPIRV"); NameChecker.Emit(); } diff --git a/clang/utils/TableGen/TableGen.cpp b/clang/utils/TableGen/TableGen.cpp index bb9366e2b7fc0..fbc48008523ef 100644 --- a/clang/utils/TableGen/TableGen.cpp +++ b/clang/utils/TableGen/TableGen.cpp @@ -65,6 +65,7 @@ enum ActionType { GenClangCommentCommandList, GenClangOpenCLBuiltins, GenClangOpenCLBuiltinTests, + GenClangSPIRVBuiltins, GenArmNeon, GenArmFP16, GenArmBF16, @@ -200,6 +201,8 @@ cl::opt Action( "Generate OpenCL builtin declaration handlers"), clEnumValN(GenClangOpenCLBuiltinTests, "gen-clang-opencl-builtin-tests", "Generate OpenCL builtin declaration tests"), + clEnumValN(GenClangSPIRVBuiltins, "gen-clang-spirv-builtins", + "Generate SPIR-V builtin declaration handlers"), clEnumValN(GenArmNeon, "gen-arm-neon", "Generate arm_neon.h for clang"), clEnumValN(GenArmFP16, "gen-arm-fp16", "Generate arm_fp16.h for clang"), clEnumValN(GenArmBF16, "gen-arm-bf16", "Generate arm_bf16.h for clang"), @@ -383,6 +386,9 @@ bool ClangTableGenMain(raw_ostream &OS, RecordKeeper &Records) { case GenClangOpenCLBuiltinTests: EmitClangOpenCLBuiltinTests(Records, OS); break; + case GenClangSPIRVBuiltins: + EmitClangSPIRVBuiltins(Records, OS); + break; case GenClangSyntaxNodeList: EmitClangSyntaxNodeList(Records, OS); break; diff --git a/clang/utils/TableGen/TableGenBackends.h b/clang/utils/TableGen/TableGenBackends.h index fd8b9fcda20f0..256d033aa1db4 100644 --- a/clang/utils/TableGen/TableGenBackends.h +++ b/clang/utils/TableGen/TableGenBackends.h @@ -125,6 +125,7 @@ void EmitClangOpenCLBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); void EmitClangOpenCLBuiltinTests(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); +void EmitClangSPIRVBuiltins(llvm::RecordKeeper &Records, llvm::raw_ostream &OS); void EmitClangDataCollectors(llvm::RecordKeeper &Records, llvm::raw_ostream &OS);