diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h --- a/clang/include/clang-c/Index.h +++ b/clang/include/clang-c/Index.h @@ -3392,7 +3392,21 @@ 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 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -1079,6 +1079,11 @@ 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 --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -2109,9 +2109,15 @@ #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 @@ -2495,6 +2501,10 @@ // 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" @@ -6864,6 +6874,14 @@ } #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); } @@ -6886,7 +6904,17 @@ 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 --- a/clang/include/clang/AST/TypeProperties.td +++ b/clang/include/clang/AST/TypeProperties.td @@ -757,6 +757,12 @@ 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 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -257,6 +257,7 @@ 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 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -5681,6 +5681,8 @@ HelpText<"Include default header file for OpenCL">; def fdeclare_opencl_builtins : Flag<["-"], "fdeclare-opencl-builtins">, HelpText<"Add OpenCL builtin function declarations (experimental)">; +def fdeclare_spirv_builtins : Flag<["-"], "fdeclare-spirv-builtins">, + HelpText<"Add SPIR-V builtin function declarations (experimental)">; def fpreserve_vec3_type : Flag<["-"], "fpreserve-vec3-type">, HelpText<"Preserve 3-component vector type">, diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1068,6 +1068,11 @@ #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 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1413,9 +1413,14 @@ 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); @@ -2172,6 +2177,11 @@ #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" @@ -6756,6 +6766,12 @@ #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: @@ -7341,6 +7357,11 @@ #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/ASTImporter.cpp b/clang/lib/AST/ASTImporter.cpp --- a/clang/lib/AST/ASTImporter.cpp +++ b/clang/lib/AST/ASTImporter.cpp @@ -1046,6 +1046,12 @@ 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 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -11003,6 +11003,11 @@ #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 --- a/clang/lib/AST/ItaniumMangle.cpp +++ b/clang/lib/AST/ItaniumMangle.cpp @@ -3014,6 +3014,14 @@ 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 --- a/clang/lib/AST/MicrosoftMangle.cpp +++ b/clang/lib/AST/MicrosoftMangle.cpp @@ -2394,6 +2394,13 @@ 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 --- a/clang/lib/AST/NSAPI.cpp +++ b/clang/lib/AST/NSAPI.cpp @@ -463,6 +463,11 @@ #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 --- a/clang/lib/AST/PrintfFormatString.cpp +++ b/clang/lib/AST/PrintfFormatString.cpp @@ -785,6 +785,11 @@ #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 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -3064,6 +3064,12 @@ #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"; @@ -4109,6 +4115,11 @@ #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 --- a/clang/lib/AST/TypeLoc.cpp +++ b/clang/lib/AST/TypeLoc.cpp @@ -392,6 +392,11 @@ #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.h b/clang/lib/CodeGen/CGDebugInfo.h --- a/clang/lib/CodeGen/CGDebugInfo.h +++ b/clang/lib/CodeGen/CGDebugInfo.h @@ -71,6 +71,11 @@ 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/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -683,6 +683,13 @@ 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/CGOpenCLRuntime.cpp b/clang/lib/CodeGen/CGOpenCLRuntime.cpp --- a/clang/lib/CodeGen/CGOpenCLRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenCLRuntime.cpp @@ -46,6 +46,15 @@ return llvm::PointerType::get( \ llvm::StructType::create(Ctx, "opencl." #ImgType "_" #Suffix "_t"), \ AddrSpc); +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return llvm::PointerType::get( \ + llvm::StructType::create(Ctx, "spirv.SampledImage." #ImgType \ + "_" #Suffix "_t"), \ + AddrSpc); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: return getSamplerType(T); diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -530,6 +530,11 @@ #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 --- a/clang/lib/CodeGen/ItaniumCXXABI.cpp +++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp @@ -3284,6 +3284,11 @@ #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 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -3626,6 +3626,9 @@ 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 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -487,6 +487,10 @@ 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 --- a/clang/lib/Index/USRGeneration.cpp +++ b/clang/lib/Index/USRGeneration.cpp @@ -718,6 +718,11 @@ #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Sema/CMakeLists.txt b/clang/lib/Sema/CMakeLists.txt --- a/clang/lib/Sema/CMakeLists.txt +++ b/clang/lib/Sema/CMakeLists.txt @@ -10,6 +10,11 @@ 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 @@ DEPENDS ClangOpenCLBuiltinsImpl + ClangSPIRVBuiltinsImpl omp_gen LINK_LIBS diff --git a/clang/lib/Sema/OpenCLBuiltins.td b/clang/lib/Sema/OpenCLBuiltins.td --- a/clang/lib/Sema/OpenCLBuiltins.td +++ b/clang/lib/Sema/OpenCLBuiltins.td @@ -277,6 +277,8 @@ 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 --- /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 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -323,6 +323,34 @@ 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 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -6816,7 +6816,7 @@ // 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 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6152,6 +6152,11 @@ #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" @@ -19837,6 +19842,11 @@ #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Sema/SemaLookup.cpp b/clang/lib/Sema/SemaLookup.cpp --- 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 @@ 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 @@ 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 @@ /// \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 @@ /// \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 @@ } } -/// 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 @@ /// \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 @@ 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 @@ 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 @@ 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 @@ // Check if this is an OpenCL Builtin, and if so, insert its overloads. if (getLangOpts().OpenCL && getLangOpts().DeclareOpenCLBuiltins) { - auto Index = isOpenCLBuiltin(II->getName()); + auto Index = OpenCLBuiltin::isBuiltin(II->getName()); + if (Index.first) { + InsertBuiltinDeclarationsFromTable( + *this, R, II, Index.first - 1, Index.second, + [this](const OpenCLBuiltin::BuiltinStruct &OpenCLBuiltin, + FunctionDecl &NewOpenCLBuiltin) { + if (!this->getLangOpts().OpenCLCPlusPlus) + NewOpenCLBuiltin.addAttr( + OverloadableAttr::CreateImplicit(Context)); + }); + return true; + } + } + + // Check if this is a SPIR-V Builtin, and if so, insert its overloads. + if (getLangOpts().DeclareSPIRVBuiltins) { + auto Index = SPIRVBuiltin::isBuiltin(II->getName()); if (Index.first) { - InsertOCLBuiltinDeclarationsFromTable(*this, R, II, Index.first - 1, - Index.second); + InsertBuiltinDeclarationsFromTable( + *this, R, II, Index.first - 1, Index.second, + [this](const SPIRVBuiltin::BuiltinStruct &, + FunctionDecl &NewBuiltin) { + if (!this->getLangOpts().CPlusPlus) + NewBuiltin.addAttr(OverloadableAttr::CreateImplicit(Context)); + }); return true; } } diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -5079,8 +5079,9 @@ 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 --- a/clang/lib/Serialization/ASTCommon.cpp +++ b/clang/lib/Serialization/ASTCommon.cpp @@ -212,6 +212,13 @@ 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 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -7025,6 +7025,13 @@ 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 --- /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 --- /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 --- /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/Preprocessor/spirv-macro.cpp b/clang/test/Preprocessor/spirv-macro.cpp new file mode 100644 --- /dev/null +++ b/clang/test/Preprocessor/spirv-macro.cpp @@ -0,0 +1,6 @@ +// RUN: %clang_cc1 %s -E -dM | FileCheck %s +// RUN: %clang_cc1 %s -fdeclare-spirv-builtins -E -dM | FileCheck --check-prefix=CHECK-SPIRV %s + +// CHECK-NOT:#define __SPIRV_BUILTIN_DECLARATIONS__ + +// CHECK-SPIRV:#define __SPIRV_BUILTIN_DECLARATIONS__ diff --git a/clang/test/SemaOpenCL/sampled_image_overload.cl b/clang/test/SemaOpenCL/sampled_image_overload.cl new file mode 100644 --- /dev/null +++ b/clang/test/SemaOpenCL/sampled_image_overload.cl @@ -0,0 +1,14 @@ +// RUN: %clang_cc1 -ast-dump %s | FileCheck %s + +void __attribute__((overloadable)) foo(__ocl_sampled_image1d_ro_t); +void __attribute__((overloadable)) foo(__ocl_sampled_image2d_ro_t); + +// CHECK: FunctionDecl {{.*}} <{{.*}}> line:{{.*}} ker 'void (__private __ocl_sampled_image1d_ro_t, __private __ocl_sampled_image2d_ro_t)' +void kernel ker(__ocl_sampled_image1d_ro_t src1, __ocl_sampled_image2d_ro_t src2) { + // CHECK: CallExpr + // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image1d_ro_t)' + foo(src1); + // CHECK: CallExpr + // CHECK-NEXT: ImplicitCastExpr {{.*}} <{{.*}}> 'void (*)(__private __ocl_sampled_image2d_ro_t)' + foo(src2); +} diff --git a/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp b/clang/test/SemaSYCL/spirv-builtin-lookup-invalid.cpp new file mode 100644 --- /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 --- /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 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -1538,6 +1538,11 @@ #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 --- a/clang/tools/libclang/CXType.cpp +++ b/clang/tools/libclang/CXType.cpp @@ -68,7 +68,11 @@ 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); @@ -612,6 +616,11 @@ #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 --- a/clang/utils/TableGen/CMakeLists.txt +++ b/clang/utils/TableGen/CMakeLists.txt @@ -11,7 +11,7 @@ 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 rename from clang/utils/TableGen/ClangOpenCLBuiltinEmitter.cpp rename to clang/utils/TableGen/ClangProgModelBuiltinEmitter.cpp --- 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 @@ 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ 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 @@ // 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 @@ 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 @@ } // 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 @@ 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::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 @@ // 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 @@ // 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 @@ 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 @@ 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 @@ 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::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::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::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 @@ 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 @@ << (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 @@ 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 @@ } 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ 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 @@ 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 @@ 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 @@ } )"; - // 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 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 --- a/clang/utils/TableGen/TableGen.cpp +++ b/clang/utils/TableGen/TableGen.cpp @@ -65,6 +65,7 @@ GenClangCommentCommandList, GenClangOpenCLBuiltins, GenClangOpenCLBuiltinTests, + GenClangSPIRVBuiltins, GenArmNeon, GenArmFP16, GenArmBF16, @@ -200,6 +201,8 @@ "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 @@ 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 --- a/clang/utils/TableGen/TableGenBackends.h +++ b/clang/utils/TableGen/TableGenBackends.h @@ -125,6 +125,7 @@ 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);