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<ASTContext> {
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..b5750511d8334 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -4952,7 +4952,8 @@ 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() &&
+ (T->isBlockPointerType() || T->isImageType() || T->isSamplerT())) {
S.Diag(D.getIdentifierLoc(), diag::err_opencl_invalid_return)
<< T << 1 /*hint off*/;
D.setInvalidType(true);
diff --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp
index 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..1b723575410d7 100644
--- a/sycl/include/CL/__spirv/spirv_types.hpp
+++ b/sycl/include/CL/__spirv/spirv_types.hpp
@@ -95,14 +95,23 @@ 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
+#undef __SYCL_SPV_SAMPLED_AND_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 <CL/__spirv/spirv_ops.hpp>
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// Type trait to get the associated sampled image type for a given image type.
+template <typename ImageType> struct sampled_opencl_image_type;
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+
#define INVOKE_SPIRV_CALL_ARG1(call) \
template <typename R, typename T1> inline R __invoke_##call(T1 ParT1) { \
using Ret = cl::sycl::detail::ConvertToOpenCLType_t<R>; \
@@ -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<RetType>;
using TempArgT = cl::sycl::detail::ConvertToOpenCLType_t<CoordT>;
- using SampledT = void*;
+ using SampledT =
+ typename cl::sycl::detail::sampled_opencl_image_type<ImageT>::type;
TempArgT TmpCoords =
cl::sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
@@ -191,6 +203,9 @@ struct opencl_image_type<Dimensions, AccessMode, access::target::host_image> {
using type =
opencl_image_type<Dimensions, AccessMode, access::target::host_image> *;
};
+template <typename T> struct sampled_opencl_image_type<T *> {
+ using type = void *;
+};
#define IMAGETY_DEFINE(Dim, AccessMode, AMSuffix, Target, Ifarray_) \
template <> \
@@ -198,11 +213,19 @@ struct opencl_image_type<Dimensions, AccessMode, access::target::host_image> {
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<typename opencl_image_type< \
+ Dim, access::mode::AccessMode, access::target::Target>::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<Dimensions, AccessMode, access::target::host_image> {
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_) \