From 4b1e7bbc528ec5df02836b41b8a4e9e38617e546 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Mon, 15 Jun 2020 11:28:08 +0100 Subject: [PATCH 1/2] [SYCL-SPIRV] Add SPIR-V variants of TypeSampledImage as builtin type in clang. This patch adds SPIR-V sampled image types as derivative of the builtin OpenCL Image types. For each OpenCL image type, clang defines a Sampled variant and lowered as a "spirv.SampledImage." llvm opaque type. Signed-off-by: Victor Lomuller --- clang/include/clang-c/Index.h | 16 ++++++++- clang/include/clang/AST/ASTContext.h | 5 +++ clang/include/clang/AST/Type.h | 28 +++++++++++++++ clang/include/clang/AST/TypeProperties.td | 6 ++++ .../include/clang/Serialization/ASTBitCodes.h | 5 +++ clang/lib/AST/ASTContext.cpp | 21 +++++++++++ clang/lib/AST/ASTImporter.cpp | 6 ++++ clang/lib/AST/ExprConstant.cpp | 5 +++ clang/lib/AST/ItaniumMangle.cpp | 8 +++++ clang/lib/AST/MicrosoftMangle.cpp | 7 ++++ clang/lib/AST/NSAPI.cpp | 5 +++ clang/lib/AST/PrintfFormatString.cpp | 5 +++ clang/lib/AST/Type.cpp | 11 ++++++ clang/lib/AST/TypeLoc.cpp | 5 +++ clang/lib/CodeGen/CGDebugInfo.cpp | 7 ++++ clang/lib/CodeGen/CGDebugInfo.h | 5 +++ clang/lib/CodeGen/CGOpenCLRuntime.cpp | 9 +++++ clang/lib/CodeGen/CodeGenTypes.cpp | 5 +++ clang/lib/CodeGen/ItaniumCXXABI.cpp | 5 +++ clang/lib/Index/USRGeneration.cpp | 5 +++ clang/lib/Sema/Sema.cpp | 14 ++++++++ clang/lib/Sema/SemaDecl.cpp | 2 +- clang/lib/Sema/SemaExpr.cpp | 10 ++++++ clang/lib/Sema/SemaType.cpp | 8 +++-- clang/lib/Serialization/ASTCommon.cpp | 7 ++++ clang/lib/Serialization/ASTReader.cpp | 7 ++++ clang/test/CodeGenOpenCL/sampled_image.cl | 12 +++++++ .../test/SemaOpenCL/sampled_image_overload.cl | 14 ++++++++ clang/tools/libclang/CIndex.cpp | 5 +++ clang/tools/libclang/CXType.cpp | 11 +++++- sycl/include/CL/__spirv/spirv_types.hpp | 28 +++++++++------ .../CL/sycl/detail/image_ocl_types.hpp | 35 +++++++++++++++---- 32 files changed, 300 insertions(+), 22 deletions(-) create mode 100644 clang/test/CodeGenOpenCL/sampled_image.cl create mode 100644 clang/test/SemaOpenCL/sampled_image_overload.cl diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h index 19471af4a675c..f6b27ad215611 100644 --- a/clang/include/clang-c/Index.h +++ b/clang/include/clang-c/Index.h @@ -3348,7 +3348,21 @@ enum CXTypeKind { CXType_OCLIntelSubgroupAVCImeDualRefStreamin = 175, CXType_ExtVector = 176, - CXType_Atomic = 177 + CXType_Atomic = 177, + + /* SPIRV builtin types. */ + CXType_SampledOCLImage1dRO = 178, + CXType_SampledOCLImage1dArrayRO = 179, + CXType_SampledOCLImage1dBufferRO = 180, + CXType_SampledOCLImage2dRO = 181, + CXType_SampledOCLImage2dArrayRO = 182, + CXType_SampledOCLImage2dDepthRO = 183, + CXType_SampledOCLImage2dArrayDepthRO = 184, + CXType_SampledOCLImage2dMSAARO = 185, + CXType_SampledOCLImage2dArrayMSAARO = 186, + CXType_SampledOCLImage2dMSAADepthRO = 187, + CXType_SampledOCLImage2dArrayMSAADepthRO = 188, + CXType_SampledOCLImage3dRO = 189 }; /** diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 2b988be60da9f..e8c514c204a55 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -976,6 +976,11 @@ class ASTContext : public RefCountedBase { CanQualType ObjCBuiltinBoolTy; #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ CanQualType SingletonId; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + CanQualType Sampled##SingletonId; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" CanQualType OCLSamplerTy, OCLEventTy, OCLClkEventTy; CanQualType OCLQueueTy, OCLReserveIDTy; diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h index d88ad16158671..9e6dd15c21a77 100644 --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -2134,9 +2134,15 @@ class alignas(8) Type : public ExtQualsTypeCommonBase { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ bool is##Id##Type() const; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + bool isSampled##Id##Type() const; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" bool isImageType() const; // Any OpenCL image type + bool isSampledImageType() const; // Any SPIR-V Sampled image type bool isSamplerT() const; // OpenCL sampler_t bool isEventT() const; // OpenCL event_t @@ -2520,6 +2526,10 @@ class BuiltinType : public Type { // OpenCL image types #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) Id, #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) Sampled##Id, +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" // OpenCL extension types #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) Id, #include "clang/Basic/OpenCLExtensionTypes.def" @@ -6847,6 +6857,14 @@ inline bool Type::isDecltypeType() const { } #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + inline bool Type::isSampled##Id##Type() const { \ + return isSpecificBuiltinType(BuiltinType::Sampled##Id); \ + } +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" + inline bool Type::isSamplerT() const { return isSpecificBuiltinType(BuiltinType::OCLSampler); } @@ -6869,7 +6887,17 @@ inline bool Type::isReserveIDT() const { inline bool Type::isImageType() const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) is##Id##Type() || + return isSampledImageType() || +#include "clang/Basic/OpenCLImageTypes.def" + false; // end boolean or operation +} + +inline bool Type::isSampledImageType() const { +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + isSampled##Id##Type() || return +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" false; // end boolean or operation } diff --git a/clang/include/clang/AST/TypeProperties.td b/clang/include/clang/AST/TypeProperties.td index 4540ea0e1952a..4b04bfa0194ca 100644 --- a/clang/include/clang/AST/TypeProperties.td +++ b/clang/include/clang/AST/TypeProperties.td @@ -752,6 +752,12 @@ let Class = BuiltinType in { case BuiltinType::ID: return ctx.SINGLETON_ID; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(IMGTYPE, ID, SINGLETON_ID, ACCESS, SUFFIX) \ + case BuiltinType::Sampled##ID: return ctx.Sampled##SINGLETON_ID; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" + #define EXT_OPAQUE_TYPE(EXTTYPE, ID, EXT) \ case BuiltinType::ID: return ctx.ID##Ty; #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index c6f9f1d1a08f4..726c383ade807 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1072,6 +1072,11 @@ class TypeIdx { /// OpenCL image types with auto numeration #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) \ diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index 0d968688fdb0a..3b036a608c736 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1407,6 +1407,11 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, 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); @@ -2119,6 +2124,11 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" @@ -6474,6 +6484,12 @@ OpenCLTypeKind ASTContext::getOpenCLTypeKind(const Type *T) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: \ return OCLTK_Image; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return OCLTK_Image; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLClkEvent: @@ -7057,6 +7073,11 @@ static char getObjCEncodingForPrimitiveType(const ASTContext *C, #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/ASTImporter.cpp b/clang/lib/AST/ASTImporter.cpp index 8b96e20e374d3..13652d8c27047 100644 --- a/clang/lib/AST/ASTImporter.cpp +++ b/clang/lib/AST/ASTImporter.cpp @@ -1028,6 +1028,12 @@ ExpectedType ASTNodeImporter::VisitBuiltinType(const BuiltinType *T) { case BuiltinType::Id: \ return Importer.getToContext().SingletonId; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return Importer.getToContext().Sampled##SingletonId; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: \ return Importer.getToContext().Id##Ty; diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp index 640bdd0d45e53..c8984f348d533 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -10299,6 +10299,11 @@ EvaluateBuiltinClassifyType(QualType T, const LangOptions &LangOpts) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp index f6dbd944384cd..006dc043b067a 100644 --- a/clang/lib/AST/ItaniumMangle.cpp +++ b/clang/lib/AST/ItaniumMangle.cpp @@ -2819,6 +2819,14 @@ void CXXNameMangler::mangleType(const BuiltinType *T) { type_name = "ocl_" #ImgType "_" #Suffix; \ Out << type_name.size() << type_name; \ break; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + type_name = "__spirv_SampledImage__" #ImgType "_" #Suffix; \ + Out << type_name.size() << type_name; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: Out << "11ocl_sampler"; diff --git a/clang/lib/AST/MicrosoftMangle.cpp b/clang/lib/AST/MicrosoftMangle.cpp index 930d07f6b158c..1e9121c14c47d 100644 --- a/clang/lib/AST/MicrosoftMangle.cpp +++ b/clang/lib/AST/MicrosoftMangle.cpp @@ -2054,6 +2054,13 @@ void MicrosoftCXXNameMangler::mangleType(const BuiltinType *T, Qualifiers, case BuiltinType::Id: \ Out << "PAUocl_" #ImgType "_" #Suffix "@@"; \ break; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + Out << "PAU__spirv_SampledImage__" #ImgType "_" #Suffix "@@"; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: Out << "PA"; diff --git a/clang/lib/AST/NSAPI.cpp b/clang/lib/AST/NSAPI.cpp index ace7f1ceebe74..8991687c162db 100644 --- a/clang/lib/AST/NSAPI.cpp +++ b/clang/lib/AST/NSAPI.cpp @@ -463,6 +463,11 @@ NSAPI::getNSNumberFactoryMethodKind(QualType T) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/PrintfFormatString.cpp b/clang/lib/AST/PrintfFormatString.cpp index f3ac181214ac6..deaa24636aa6c 100644 --- a/clang/lib/AST/PrintfFormatString.cpp +++ b/clang/lib/AST/PrintfFormatString.cpp @@ -785,6 +785,11 @@ bool PrintfSpecifier::fixType(QualType QT, const LangOptions &LangOpt, #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index ba0d86befe1b9..207f812872620 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -3017,6 +3017,12 @@ StringRef BuiltinType::getName(const PrintingPolicy &Policy) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case Id: \ return "__" #Access " " #ImgType "_t"; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case Sampled##Id: \ + return "__ocl_sampled_" #ImgType "_" #Suffix "_t"; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case OCLSampler: return "sampler_t"; @@ -4044,6 +4050,11 @@ bool Type::canHaveNullability(bool ResultIfUnknown) const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/AST/TypeLoc.cpp b/clang/lib/AST/TypeLoc.cpp index 57c11ca5571db..31e35f9968d89 100644 --- a/clang/lib/AST/TypeLoc.cpp +++ b/clang/lib/AST/TypeLoc.cpp @@ -392,6 +392,11 @@ TypeSpecifierType BuiltinTypeLoc::getWrittenTypeSpec() const { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 6965c4a1209c2..a62829cee5323 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -704,6 +704,13 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) { case BuiltinType::Id: \ return getOrCreateStructPtrType("opencl_" #ImgType "_" #Suffix "_t", \ SingletonId); +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return getOrCreateStructPtrType( \ + "spirv_sampled_" #ImgType "_" #Suffix "_t", Sampled##SingletonId); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: return getOrCreateStructPtrType("opencl_sampler_t", OCLSamplerDITy); diff --git a/clang/lib/CodeGen/CGDebugInfo.h b/clang/lib/CodeGen/CGDebugInfo.h index 96ef6c7c1d27d..24d63fad84f52 100644 --- a/clang/lib/CodeGen/CGDebugInfo.h +++ b/clang/lib/CodeGen/CGDebugInfo.h @@ -71,6 +71,11 @@ class CGDebugInfo { llvm::DIType *SelTy = nullptr; #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ llvm::DIType *SingletonId = nullptr; +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + llvm::DIType *Sampled##SingletonId = nullptr; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" llvm::DIType *OCLSamplerDITy = nullptr; llvm::DIType *OCLEventDITy = nullptr; diff --git a/clang/lib/CodeGen/CGOpenCLRuntime.cpp b/clang/lib/CodeGen/CGOpenCLRuntime.cpp index dbe375294d179..78ccb0790e0ec 100644 --- a/clang/lib/CodeGen/CGOpenCLRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenCLRuntime.cpp @@ -46,6 +46,15 @@ llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) { return llvm::PointerType::get( \ llvm::StructType::create(Ctx, "opencl." #ImgType "_" #Suffix "_t"), \ AddrSpc); +#include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + return llvm::PointerType::get( \ + llvm::StructType::create(Ctx, "spirv.SampledImage." #ImgType \ + "_" #Suffix "_t"), \ + AddrSpc); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) #include "clang/Basic/OpenCLImageTypes.def" case BuiltinType::OCLSampler: return getSamplerType(T); diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 906fdc6e86bdd..fab14255ebf7f 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -537,6 +537,11 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp index 2829877cfe5d9..79dbded0edcc1 100644 --- a/clang/lib/CodeGen/ItaniumCXXABI.cpp +++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp @@ -3067,6 +3067,11 @@ static bool TypeInfoIsInStandardLibrary(const BuiltinType *Ty) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Index/USRGeneration.cpp b/clang/lib/Index/USRGeneration.cpp index 0d1e812198234..026f67dc1121f 100644 --- a/clang/lib/Index/USRGeneration.cpp +++ b/clang/lib/Index/USRGeneration.cpp @@ -718,6 +718,11 @@ void USRGenerator::VisitType(QualType T) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 3ce15ff6c931d..eba93b3e6a3f0 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -303,6 +303,20 @@ void Sema::Initialize() { #undef SEMA_STRINGIZE } + if (getLangOpts().SYCLIsDevice || getLangOpts().OpenCL) { +#ifdef SEMA_STRINGIZE +#error "Undefine SEMA_STRINGIZE macro." +#endif +#define SEMA_STRINGIZE(s) #s +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + addImplicitTypedef(SEMA_STRINGIZE(__ocl_sampled_##ImgType##_##Suffix##_t), \ + Context.Sampled##SingletonId); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" +#undef SEMA_STRINGIZE + } + // Initialize predefined OpenCL types and supported extensions and (optional) // core features. if (getLangOpts().OpenCL) { diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index 8064690c2e98d..e7e0614292c5b 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -6674,7 +6674,7 @@ static bool diagnoseOpenCLTypes(Scope *S, Sema &Se, Declarator &D, // 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(D.getIdentifierLoc(), diag::err_opencl_type_can_only_be_used_as_function_parameter) << R; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index d613db8a6b100..ebb23569c425c 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6055,6 +6055,11 @@ static bool isPlaceholderToRemoveAsArg(QualType type) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" @@ -19152,6 +19157,11 @@ ExprResult Sema::CheckPlaceholderExpr(Expr *E) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 44456e4038a8b..6f2679a446e48 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -4953,9 +4953,11 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state, // OpenCL v2.0 s6.12.5 - A block cannot be the return value of a // function. if (T->isBlockPointerType() || T->isImageType() || T->isSamplerT()) { - S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return) - << T << 1 /*hint off*/; - D.setInvalidType(true); + if (!T->isSampledImageType()) { + S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return) + << T << 1 /*hint off*/; + D.setInvalidType(true); + } } // OpenCL doesn't support variadic functions and blocks // (s6.9.e and s6.12.5 OpenCL v2.0) except for printf. diff --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp index bf583b02f96b8..75c510d7aa570 100644 --- a/clang/lib/Serialization/ASTCommon.cpp +++ b/clang/lib/Serialization/ASTCommon.cpp @@ -212,6 +212,13 @@ serialization::TypeIdxFromBuiltin(const BuiltinType *BT) { ID = PREDEF_TYPE_##Id##_ID; \ break; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: \ + ID = PREDEF_TYPE_SAMPLED_##Id##_ID; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case BuiltinType::Id: \ ID = PREDEF_TYPE_##Id##_ID; \ diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 4dd2054d4b079..3da70c0900322 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -6994,6 +6994,13 @@ QualType ASTReader::GetType(TypeID ID) { T = Context.SingletonId; \ break; #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case PREDEF_TYPE_SAMPLED_##Id##_ID: \ + T = Context.Sampled##SingletonId; \ + break; +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) \ case PREDEF_TYPE_##Id##_ID: \ T = Context.Id##Ty; \ diff --git a/clang/test/CodeGenOpenCL/sampled_image.cl b/clang/test/CodeGenOpenCL/sampled_image.cl new file mode 100644 index 0000000000000..f9ccd754a74aa --- /dev/null +++ b/clang/test/CodeGenOpenCL/sampled_image.cl @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - | FileCheck %s +// RUN: %clang_cc1 %s -triple x86_64-unknown-linux-gnu -O0 -emit-llvm -o - -cl-std=clc++ | FileCheck %s + +__attribute__((overloadable)) void my_read_image(__ocl_sampled_image1d_ro_t img); +__attribute__((overloadable)) void my_read_image(__ocl_sampled_image2d_ro_t img); + +void test_read_image(__ocl_sampled_image1d_ro_t img_ro, __ocl_sampled_image2d_ro_t img_2d) { + // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image1d_ro(%spirv.SampledImage.image1d_ro_t* %{{[0-9]+}}) + my_read_image(img_ro); + // CHECK: call void @_Z13my_read_image32__spirv_SampledImage__image2d_ro(%spirv.SampledImage.image2d_ro_t* %{{[0-9]+}}) + my_read_image(img_2d); +} diff --git a/clang/test/SemaOpenCL/sampled_image_overload.cl b/clang/test/SemaOpenCL/sampled_image_overload.cl new file mode 100644 index 0000000000000..4d4a722018a29 --- /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/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 4a65624268d84..5b9e02af3b598 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -1531,6 +1531,11 @@ bool CursorVisitor::VisitBuiltinTypeLoc(BuiltinTypeLoc TL) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ case BuiltinType::Id: #include "clang/Basic/OpenCLImageTypes.def" +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + case BuiltinType::Sampled##Id: +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtTYpe, Id, Ext) case BuiltinType::Id: #include "clang/Basic/OpenCLExtensionTypes.def" case BuiltinType::OCLSampler: diff --git a/clang/tools/libclang/CXType.cpp b/clang/tools/libclang/CXType.cpp index 2d19e13f161a8..99f012f0735b4 100644 --- a/clang/tools/libclang/CXType.cpp +++ b/clang/tools/libclang/CXType.cpp @@ -68,7 +68,11 @@ static CXTypeKind GetBuiltinTypeKind(const BuiltinType *BT) { BTCASE(ObjCSel); #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) BTCASE(Id); #include "clang/Basic/OpenCLImageTypes.def" -#undef IMAGE_TYPE +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \ + BTCASE(Sampled##Id); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" #define EXT_OPAQUE_TYPE(ExtType, Id, Ext) BTCASE(Id); #include "clang/Basic/OpenCLExtensionTypes.def" BTCASE(OCLSampler); @@ -612,6 +616,11 @@ CXString clang_getTypeKindSpelling(enum CXTypeKind K) { #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) TKIND(Id); #include "clang/Basic/OpenCLImageTypes.def" #undef IMAGE_TYPE +#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) TKIND(Sampled##Id); +#define IMAGE_WRITE_TYPE(Type, Id, Ext) +#define IMAGE_READ_WRITE_TYPE(Type, Id, Ext) +#include "clang/Basic/OpenCLImageTypes.def" +#undef IMAGE_TYPE #define EXT_OPAQUE_TYPE(ExtTYpe, Id, Ext) TKIND(Id); #include "clang/Basic/OpenCLExtensionTypes.def" TKIND(OCLSampler); diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp index dd5f91f63834a..386ec2245d64f 100644 --- a/sycl/include/CL/__spirv/spirv_types.hpp +++ b/sycl/include/CL/__spirv/spirv_types.hpp @@ -95,14 +95,22 @@ typedef void* __ocl_event_t; typedef void* __ocl_sampler_t; // Adding only the datatypes that can be currently used in SYCL, // as per SYCL spec 1.2.1 -typedef void *__ocl_image1d_ro_t; -typedef void *__ocl_image2d_ro_t; -typedef void *__ocl_image3d_ro_t; -typedef void *__ocl_image1d_wo_t; -typedef void *__ocl_image2d_wo_t; -typedef void *__ocl_image3d_wo_t; -typedef void *__ocl_image1d_array_ro_t; -typedef void *__ocl_image2d_array_ro_t; -typedef void *__ocl_image1d_array_wo_t; -typedef void *__ocl_image2d_array_wo_t; +#define __SYCL_SPV_IMAGE_TYPE(NAME) typedef void *__ocl_##NAME##_t + +#define __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(NAME) \ + __SYCL_SPV_IMAGE_TYPE(NAME); \ + typedef void *__ocl_sampled_##NAME##_t + +__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image1d_ro); +__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image2d_ro); +__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image3d_ro); +__SYCL_SPV_IMAGE_TYPE(image1d_wo); +__SYCL_SPV_IMAGE_TYPE(image2d_wo); +__SYCL_SPV_IMAGE_TYPE(image3d_wo); +__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image1d_array_ro); +__SYCL_SPV_SAMPLED_AND_IMAGE_TYPE(image2d_array_ro); +__SYCL_SPV_IMAGE_TYPE(image1d_array_wo); +__SYCL_SPV_IMAGE_TYPE(image2d_array_wo); + +#undef __SYCL_SPV_IMAGE_TYPE #endif diff --git a/sycl/include/CL/sycl/detail/image_ocl_types.hpp b/sycl/include/CL/sycl/detail/image_ocl_types.hpp index 2dd20bf60e412..20bfd9655846f 100644 --- a/sycl/include/CL/sycl/detail/image_ocl_types.hpp +++ b/sycl/include/CL/sycl/detail/image_ocl_types.hpp @@ -33,6 +33,17 @@ #include +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { + +// Type trait to get the associated sampled image type for a given image type. +template struct sampled_opencl_image_type; + +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) + #define INVOKE_SPIRV_CALL_ARG1(call) \ template inline R __invoke_##call(T1 ParT1) { \ using Ret = cl::sycl::detail::ConvertToOpenCLType_t; \ @@ -79,7 +90,8 @@ static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords, // Convert from sycl types to builtin types to get correct function mangling. using TempRetT = cl::sycl::detail::ConvertToOpenCLType_t; using TempArgT = cl::sycl::detail::ConvertToOpenCLType_t; - using SampledT = void*; + using SampledT = + typename cl::sycl::detail::sampled_opencl_image_type::type; TempArgT TmpCoords = cl::sycl::detail::convertDataToType(Coords); @@ -191,6 +203,9 @@ struct opencl_image_type { using type = opencl_image_type *; }; +template struct sampled_opencl_image_type { + using type = void *; +}; #define IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, Ifarray_) \ template <> \ @@ -198,11 +213,19 @@ struct opencl_image_type { access::target::Target> { \ using type = __ocl_image##Dim##d_##Ifarray_##AMSuffix##_t; \ }; +#define SAMPLED_AND_IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, \ + Ifarray_) \ + IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, Ifarray_) \ + template <> \ + struct sampled_opencl_image_type::type> { \ + using type = __ocl_sampled_image##Dim##d_##Ifarray_##AMSuffix##_t; \ + }; #define IMAGETY_READ_3_DIM_IMAGE \ - IMAGETY_DEFINE(1, read, ro, image, ) \ - IMAGETY_DEFINE(2, read, ro, image, ) \ - IMAGETY_DEFINE(3, read, ro, image, ) + SAMPLED_AND_IMAGETY_DEFINE(1, read, ro, image, ) \ + SAMPLED_AND_IMAGETY_DEFINE(2, read, ro, image, ) \ + SAMPLED_AND_IMAGETY_DEFINE(3, read, ro, image, ) #define IMAGETY_WRITE_3_DIM_IMAGE \ IMAGETY_DEFINE(1, write, wo, image, ) \ @@ -215,8 +238,8 @@ struct opencl_image_type { IMAGETY_DEFINE(3, discard_write, wo, image, ) #define IMAGETY_READ_2_DIM_IARRAY \ - IMAGETY_DEFINE(1, read, ro, image_array, array_) \ - IMAGETY_DEFINE(2, read, ro, image_array, array_) + SAMPLED_AND_IMAGETY_DEFINE(1, read, ro, image_array, array_) \ + SAMPLED_AND_IMAGETY_DEFINE(2, read, ro, image_array, array_) #define IMAGETY_WRITE_2_DIM_IARRAY \ IMAGETY_DEFINE(1, write, wo, image_array, array_) \ From 96ab6948053bea7d082f3ab877c70d44afa63184 Mon Sep 17 00:00:00 2001 From: Victor Lomuller Date: Fri, 3 Jul 2020 15:20:26 +0100 Subject: [PATCH 2/2] [SYCL-PTX] Address review comments Signed-off-by: Victor Lomuller --- clang/lib/Sema/SemaType.cpp | 11 +++++------ sycl/include/CL/__spirv/spirv_types.hpp | 1 + 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp index 6f2679a446e48..b5750511d8334 100644 --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -4952,12 +4952,11 @@ static TypeSourceInfo *GetFullTypeForDeclarator(TypeProcessingState &state, if (LangOpts.OpenCL || LangOpts.SYCLIsDevice) { // OpenCL v2.0 s6.12.5 - A block cannot be the return value of a // function. - if (T->isBlockPointerType() || T->isImageType() || T->isSamplerT()) { - if (!T->isSampledImageType()) { - S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return) - << T << 1 /*hint off*/; - D.setInvalidType(true); - } + if (!T->isSampledImageType() && + (T->isBlockPointerType() || T->isImageType() || T->isSamplerT())) { + S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return) + << T << 1 /*hint off*/; + D.setInvalidType(true); } // OpenCL doesn't support variadic functions and blocks // (s6.9.e and s6.12.5 OpenCL v2.0) except for printf. diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp index 386ec2245d64f..1b723575410d7 100644 --- a/sycl/include/CL/__spirv/spirv_types.hpp +++ b/sycl/include/CL/__spirv/spirv_types.hpp @@ -113,4 +113,5 @@ __SYCL_SPV_IMAGE_TYPE(image1d_array_wo); __SYCL_SPV_IMAGE_TYPE(image2d_array_wo); #undef __SYCL_SPV_IMAGE_TYPE +#undef __SYCL_SPV_SAMPLED_AND_IMAGE_TYPE #endif