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/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/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 294cc20f76c53..5ecb794cb5b2e 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1423,9 +1423,14 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass); InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel); - if (LangOpts.OpenCL) { + if (LangOpts.OpenCL || LangOpts.SYCLIsDevice) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ InitBuiltinType(SingletonId, BuiltinType::Id); +#include "clang/Basic/OpenCLImageTypes.def" +#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); @@ -2184,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" @@ -6851,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: @@ -7437,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" @@ -11072,6 +11093,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/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/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 59f3e02705713..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(); @@ -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/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/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/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 a2b8f475aa8c9..37be02255a521 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -317,6 +317,34 @@ void Sema::Initialize() { addImplicitTypedef("size_t", Context.getSizeType()); } + if (getLangOpts().SYCLIsDevice) { + 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), \ + Context.SingletonId); +#include "clang/Basic/OpenCLImageTypes.def" +#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/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/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/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 815463307ecc7..8dcd756b3e4b3 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,455 @@ 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 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. + /// \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) || + Util::isSyclSamplerType(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) || Util::isSyclSamplerType(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::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 + // 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/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/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp new file mode 100644 index 0000000000000..9100ada1d06a1 --- /dev/null +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -0,0 +1,432 @@ +#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 { + 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 + +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 + 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 nd_range { +}; + +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) {} +}; + +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/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/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(); } 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/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/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/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp new file mode 100644 index 0000000000000..889f4f67bc7f0 --- /dev/null +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -0,0 +1,200 @@ +#ifndef SYCL_HPP +#define SYCL_HPP + +// Shared code for SYCL tests + +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) {} +}; + +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/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 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' 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);