diff --git a/include/clang/AST/ASTContext.h b/include/clang/AST/ASTContext.h
index 86b2c42..d750ba9 100644
--- a/include/clang/AST/ASTContext.h
+++ b/include/clang/AST/ASTContext.h
@@ -716,6 +716,9 @@
   CanQualType PseudoObjectTy, ARCUnbridgedCastTy;
   CanQualType ObjCBuiltinIdTy, ObjCBuiltinClassTy, ObjCBuiltinSelTy;
   CanQualType ObjCBuiltinBoolTy;
+  CanQualType OCLImage1dTy, OCLImage1dArrayTy, OCLImage1dBufferTy;
+  CanQualType OCLImage2dTy, OCLImage2dArrayTy;
+  CanQualType OCLImage3dTy;
 
   // Types for deductions in C++0x [stmt.ranged]'s desugaring. Built on demand.
   mutable QualType AutoDeductTy;     // Deduction against 'auto'.
diff --git a/include/clang/AST/BuiltinTypes.def b/include/clang/AST/BuiltinTypes.def
index ba322fb..cb7cfed 100644
--- a/include/clang/AST/BuiltinTypes.def
+++ b/include/clang/AST/BuiltinTypes.def
@@ -154,6 +154,14 @@
 // type is a typedef of a PointerType to this.
 BUILTIN_TYPE(ObjCSel, ObjCBuiltinSelTy)
 
+// OpenCL image types.
+BUILTIN_TYPE(OCLImage1d, OCLImage1dTy)
+BUILTIN_TYPE(OCLImage1dArray, OCLImage1dArrayTy)
+BUILTIN_TYPE(OCLImage1dBuffer, OCLImage1dBufferTy)
+BUILTIN_TYPE(OCLImage2d, OCLImage2dTy)
+BUILTIN_TYPE(OCLImage2dArray, OCLImage2dArrayTy)
+BUILTIN_TYPE(OCLImage3d, OCLImage3dTy)
+
 // This represents the type of an expression whose type is
 // totally unknown, e.g. 'T::foo'.  It is permitted for this to
 // appear in situations where the structure of the type is
diff --git a/include/clang/AST/Type.h b/include/clang/AST/Type.h
index 368be8d..cf3b565 100644
--- a/include/clang/AST/Type.h
+++ b/include/clang/AST/Type.h
@@ -1580,6 +1580,17 @@
   bool isNullPtrType() const;                   // C++0x nullptr_t
   bool isAtomicType() const;                    // C11 _Atomic()
 
+  bool isImage1dT() const;                      // OpenCL image1d_t
+  bool isImage1dArrayT() const;                 // OpenCL image1d_array_t
+  bool isImage1dBufferT() const;                // OpenCL image1d_buffer_t
+  bool isImage2dT() const;                      // OpenCL image2d_t
+  bool isImage2dArrayT() const;                 // OpenCL image2d_array_t
+  bool isImage3dT() const;                      // OpenCL image3d_t
+
+  bool isImageType() const;                     // Any OpenCL image type
+
+  bool isOpenCLSpecificType() const;            // Any OpenCL specific type
+
   /// Determines if this type, which must satisfy
   /// isObjCLifetimeType(), is implicitly __unsafe_unretained rather
   /// than implicitly __strong.
@@ -4883,6 +4894,40 @@
 inline bool Type::isObjCBuiltinType() const {
   return isObjCIdType() || isObjCClassType() || isObjCSelType();
 }
+
+inline bool Type::isImage1dT() const {
+  return isSpecificBuiltinType(BuiltinType::OCLImage1d);
+}
+
+inline bool Type::isImage1dArrayT() const {
+  return isSpecificBuiltinType(BuiltinType::OCLImage1dArray);
+}
+
+inline bool Type::isImage1dBufferT() const {
+  return isSpecificBuiltinType(BuiltinType::OCLImage1dBuffer);
+}
+
+inline bool Type::isImage2dT() const {
+  return isSpecificBuiltinType(BuiltinType::OCLImage2d);
+}
+
+inline bool Type::isImage2dArrayT() const {
+  return isSpecificBuiltinType(BuiltinType::OCLImage2dArray);
+}
+
+inline bool Type::isImage3dT() const {
+  return isSpecificBuiltinType(BuiltinType::OCLImage3d);
+}
+inline bool Type::isImageType() const {
+  return isImage3dT() ||
+         isImage2dT() || isImage2dArrayT() ||
+         isImage1dT() || isImage1dArrayT() || isImage1dBufferT();
+}
+
+inline bool Type::isOpenCLSpecificType() const {
+  return isImageType();
+}
+
 inline bool Type::isTemplateTypeParmType() const {
   return isa<TemplateTypeParmType>(CanonicalType);
 }
diff --git a/include/clang/Basic/Specifiers.h b/include/clang/Basic/Specifiers.h
index c82b8cb..fedb321 100644
--- a/include/clang/Basic/Specifiers.h
+++ b/include/clang/Basic/Specifiers.h
@@ -62,13 +62,19 @@
     TST_auto,         // C++0x auto
     TST_unknown_anytype, // __unknown_anytype extension
     TST_atomic,       // C11 _Atomic
+    TST_image1d_t,        // OpenCL image1d_t
+    TST_image1d_array_t,  // OpenCL image1d_array_t
+    TST_image1d_buffer_t, // OpenCL image1d_buffer_t
+    TST_image2d_t,        // OpenCL image2d_t
+    TST_image2d_array_t,  // OpenCL image2d_array_t
+    TST_image3d_t,        // OpenCL image3d_t
     TST_error         // erroneous type
   };
   
   /// \brief Structure that packs information about the type specifiers that
   /// were written in a particular type specifier sequence.
   struct WrittenBuiltinSpecs {
-    /*DeclSpec::TST*/ unsigned Type  : 5;
+    /*DeclSpec::TST*/ unsigned Type  : 6;
     /*DeclSpec::TSS*/ unsigned Sign  : 2;
     /*DeclSpec::TSW*/ unsigned Width : 2;
     bool ModeAttr : 1;
diff --git a/include/clang/Basic/TokenKinds.def b/include/clang/Basic/TokenKinds.def
index f02ba64..7bc489b 100644
--- a/include/clang/Basic/TokenKinds.def
+++ b/include/clang/Basic/TokenKinds.def
@@ -448,6 +448,12 @@
 ALIAS("write_only", __write_only    , KEYOPENCL)
 ALIAS("read_write", __read_write    , KEYOPENCL)
 KEYWORD(__builtin_astype            , KEYOPENCL)
+KEYWORD(image1d_t                   , KEYOPENCL)
+KEYWORD(image1d_array_t             , KEYOPENCL)
+KEYWORD(image1d_buffer_t            , KEYOPENCL)
+KEYWORD(image2d_t                   , KEYOPENCL)
+KEYWORD(image2d_array_t             , KEYOPENCL)
+KEYWORD(image3d_t                   , KEYOPENCL)
 
 // Borland Extensions.
 KEYWORD(__pascal                    , KEYALL)
diff --git a/include/clang/Sema/DeclSpec.h b/include/clang/Sema/DeclSpec.h
index d261810..bc8fc64 100644
--- a/include/clang/Sema/DeclSpec.h
+++ b/include/clang/Sema/DeclSpec.h
@@ -276,6 +276,12 @@
   static const TST TST_auto = clang::TST_auto;
   static const TST TST_unknown_anytype = clang::TST_unknown_anytype;
   static const TST TST_atomic = clang::TST_atomic;
+  static const TST TST_image1d_t = clang::TST_image1d_t;
+  static const TST TST_image1d_array_t = clang::TST_image1d_array_t;
+  static const TST TST_image1d_buffer_t = clang::TST_image1d_buffer_t;
+  static const TST TST_image2d_t = clang::TST_image2d_t;
+  static const TST TST_image2d_array_t = clang::TST_image2d_array_t;
+  static const TST TST_image3d_t = clang::TST_image3d_t;
   static const TST TST_error = clang::TST_error;
 
   // type-qualifiers
@@ -306,7 +312,7 @@
   /*TSW*/unsigned TypeSpecWidth : 2;
   /*TSC*/unsigned TypeSpecComplex : 2;
   /*TSS*/unsigned TypeSpecSign : 2;
-  /*TST*/unsigned TypeSpecType : 5;
+  /*TST*/unsigned TypeSpecType : 6;
   unsigned TypeAltiVecVector : 1;
   unsigned TypeAltiVecPixel : 1;
   unsigned TypeAltiVecBool : 1;
diff --git a/include/clang/Serialization/ASTBitCodes.h b/include/clang/Serialization/ASTBitCodes.h
index 4273e14..0810d15 100644
--- a/include/clang/Serialization/ASTBitCodes.h
+++ b/include/clang/Serialization/ASTBitCodes.h
@@ -701,7 +701,19 @@
       /// \brief The __va_list_tag placeholder type.
       PREDEF_TYPE_VA_LIST_TAG = 36,
       /// \brief The placeholder type for builtin functions.
-      PREDEF_TYPE_BUILTIN_FN = 37
+      PREDEF_TYPE_BUILTIN_FN = 37,
+      /// \brief OpenCL 1d image type.
+      PREDEF_TYPE_IMAGE1D_ID    = 38,
+      /// \brief OpenCL 1d image array type.
+      PREDEF_TYPE_IMAGE1D_ARR_ID = 39,
+      /// \brief OpenCL 1d image buffer type.
+      PREDEF_TYPE_IMAGE1D_BUFF_ID = 40,
+      /// \brief OpenCL 2d image type.
+      PREDEF_TYPE_IMAGE2D_ID    = 41,
+      /// \brief OpenCL 2d image array type.
+      PREDEF_TYPE_IMAGE2D_ARR_ID = 42,
+      /// \brief OpenCL 3d image type.
+      PREDEF_TYPE_IMAGE3D_ID    = 43
     };
 
     /// \brief The number of predefined type IDs that are reserved for
diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp
index f3fa135..b04b830 100644
--- a/lib/AST/ASTContext.cpp
+++ b/lib/AST/ASTContext.cpp
@@ -874,6 +874,15 @@
   InitBuiltinType(ObjCBuiltinIdTy, BuiltinType::ObjCId);
   InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass);
   InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel);
+
+  if (LangOpts.OpenCL) { 
+    InitBuiltinType(OCLImage1dTy, BuiltinType::OCLImage1d);
+    InitBuiltinType(OCLImage1dArrayTy, BuiltinType::OCLImage1dArray);
+    InitBuiltinType(OCLImage1dBufferTy, BuiltinType::OCLImage1dBuffer);
+    InitBuiltinType(OCLImage2dTy, BuiltinType::OCLImage2d);
+    InitBuiltinType(OCLImage2dArrayTy, BuiltinType::OCLImage2dArray);
+    InitBuiltinType(OCLImage3dTy, BuiltinType::OCLImage3d);
+  }
   
   // Builtin type for __objc_yes and __objc_no
   ObjCBuiltinBoolTy = (Target.useSignedCharForObjCBool() ?
@@ -1412,6 +1421,16 @@
       Width = Target->getPointerWidth(0); 
       Align = Target->getPointerAlign(0);
       break;
+    case BuiltinType::OCLImage1d:
+    case BuiltinType::OCLImage1dArray:
+    case BuiltinType::OCLImage1dBuffer:
+    case BuiltinType::OCLImage2d:
+    case BuiltinType::OCLImage2dArray:
+    case BuiltinType::OCLImage3d:
+      // Currently these types are pointers to opaque types.
+      Width = Target->getPointerWidth(0);
+      Align = Target->getPointerAlign(0);
+      break;
     }
     break;
   case Type::ObjCObjectPointer:
diff --git a/lib/AST/ItaniumMangle.cpp b/lib/AST/ItaniumMangle.cpp
index 566a389..b6077ec 100644
--- a/lib/AST/ItaniumMangle.cpp
+++ b/lib/AST/ItaniumMangle.cpp
@@ -1880,6 +1880,12 @@
   case BuiltinType::ObjCId: Out << "11objc_object"; break;
   case BuiltinType::ObjCClass: Out << "10objc_class"; break;
   case BuiltinType::ObjCSel: Out << "13objc_selector"; break;
+  case BuiltinType::OCLImage1d: Out << "11ocl_image1d"; break;
+  case BuiltinType::OCLImage1dArray: Out << "16ocl_image1darray"; break;
+  case BuiltinType::OCLImage1dBuffer: Out << "17ocl_image1dbuffer"; break;
+  case BuiltinType::OCLImage2d: Out << "11ocl_image2d"; break;
+  case BuiltinType::OCLImage2dArray: Out << "16ocl_image2darray"; break;
+  case BuiltinType::OCLImage3d: Out << "11ocl_image3d"; break;
   }
 }
 
diff --git a/lib/AST/MicrosoftMangle.cpp b/lib/AST/MicrosoftMangle.cpp
index 0da7f51..bd0125d 100644
--- a/lib/AST/MicrosoftMangle.cpp
+++ b/lib/AST/MicrosoftMangle.cpp
@@ -1053,6 +1053,13 @@
   case BuiltinType::ObjCId: Out << "PAUobjc_object@@"; break;
   case BuiltinType::ObjCClass: Out << "PAUobjc_class@@"; break;
   case BuiltinType::ObjCSel: Out << "PAUobjc_selector@@"; break;
+
+  case BuiltinType::OCLImage1d: Out << "PAUocl_image1d@@"; break;
+  case BuiltinType::OCLImage1dArray: Out << "PAUocl_image1darray@@"; break;
+  case BuiltinType::OCLImage1dBuffer: Out << "PAUocl_image1dbuffer@@"; break;
+  case BuiltinType::OCLImage2d: Out << "PAUocl_image2d@@"; break;
+  case BuiltinType::OCLImage2dArray: Out << "PAUocl_image2darray@@"; break;
+  case BuiltinType::OCLImage3d: Out << "PAUocl_image3d@@"; break;
  
   case BuiltinType::NullPtr: Out << "$$T"; break;
 
diff --git a/lib/AST/NSAPI.cpp b/lib/AST/NSAPI.cpp
index 0837509..b92b08c 100644
--- a/lib/AST/NSAPI.cpp
+++ b/lib/AST/NSAPI.cpp
@@ -337,6 +337,12 @@
   case BuiltinType::ObjCClass:
   case BuiltinType::ObjCId:
   case BuiltinType::ObjCSel:
+  case BuiltinType::OCLImage1d:
+  case BuiltinType::OCLImage1dArray:
+  case BuiltinType::OCLImage1dBuffer:
+  case BuiltinType::OCLImage2d:
+  case BuiltinType::OCLImage2dArray:
+  case BuiltinType::OCLImage3d:
   case BuiltinType::BoundMember:
   case BuiltinType::Dependent:
   case BuiltinType::Overload:
diff --git a/lib/AST/Type.cpp b/lib/AST/Type.cpp
index 97448ee..26eee2d 100644
--- a/lib/AST/Type.cpp
+++ b/lib/AST/Type.cpp
@@ -1512,6 +1512,12 @@
   case ObjCId:            return "id";
   case ObjCClass:         return "Class";
   case ObjCSel:           return "SEL";
+  case OCLImage1d:        return "image1d_t";
+  case OCLImage1dArray:   return "image1d_array_t";
+  case OCLImage1dBuffer:  return "image1d_buffer_t";
+  case OCLImage2d:        return "image2d_t";
+  case OCLImage2dArray:   return "image2d_array_t";
+  case OCLImage3d:        return "image3d_t";
   }
   
   llvm_unreachable("Invalid builtin type.");
diff --git a/lib/AST/TypeLoc.cpp b/lib/AST/TypeLoc.cpp
index c021cf8..a5baf70 100644
--- a/lib/AST/TypeLoc.cpp
+++ b/lib/AST/TypeLoc.cpp
@@ -262,6 +262,12 @@
   case BuiltinType::ObjCId:
   case BuiltinType::ObjCClass:
   case BuiltinType::ObjCSel:
+  case BuiltinType::OCLImage1d:
+  case BuiltinType::OCLImage1dArray:
+  case BuiltinType::OCLImage1dBuffer:
+  case BuiltinType::OCLImage2d:
+  case BuiltinType::OCLImage2dArray:
+  case BuiltinType::OCLImage3d:
   case BuiltinType::BuiltinFn:
     return TST_unspecified;
   }
diff --git a/lib/CodeGen/CGDebugInfo.cpp b/lib/CodeGen/CGDebugInfo.cpp
index 844514b..276c922 100644
--- a/lib/CodeGen/CGDebugInfo.cpp
+++ b/lib/CodeGen/CGDebugInfo.cpp
@@ -407,6 +407,26 @@
                                  0);
     return SelTy;
   }
+
+  case BuiltinType::OCLImage1d:
+    return getOrCreateStructPtrType("opencl_image1d_t",
+                                    OCLImage1dDITy);
+  case BuiltinType::OCLImage1dArray:
+    return getOrCreateStructPtrType("opencl_image1d_array_t", 
+                                    OCLImage1dArrayDITy);
+  case BuiltinType::OCLImage1dBuffer:
+    return getOrCreateStructPtrType("opencl_image1d_buffer_t",
+                                    OCLImage1dBufferDITy);
+  case BuiltinType::OCLImage2d:
+    return getOrCreateStructPtrType("opencl_image2d_t",
+                                    OCLImage2dDITy);
+  case BuiltinType::OCLImage2dArray:
+    return getOrCreateStructPtrType("opencl_image2d_array_t",
+                                    OCLImage2dArrayDITy);
+  case BuiltinType::OCLImage3d:
+    return getOrCreateStructPtrType("opencl_image3d_t",
+                                    OCLImage3dDITy);
+
   case BuiltinType::UChar:
   case BuiltinType::Char_U: Encoding = llvm::dwarf::DW_ATE_unsigned_char; break;
   case BuiltinType::Char_S:
@@ -613,6 +633,18 @@
                                     Size, Align);
 }
 
+llvm::DIType CGDebugInfo::getOrCreateStructPtrType(StringRef Name, llvm::DIType &Cache) {
+    if (Cache.Verify())
+      return Cache;
+    Cache =
+      DBuilder.createForwardDecl(llvm::dwarf::DW_TAG_structure_type,
+                                 Name, TheCU, getOrCreateMainFile(),
+                                 0);
+    unsigned Size = CGM.getContext().getTypeSize(CGM.getContext().VoidPtrTy);
+    Cache = DBuilder.createPointerType(Cache, Size);
+    return Cache;
+}
+
 llvm::DIType CGDebugInfo::CreateType(const BlockPointerType *Ty,
                                      llvm::DIFile Unit) {
   if (BlockLiteralGenericSet)
diff --git a/lib/CodeGen/CGDebugInfo.h b/lib/CodeGen/CGDebugInfo.h
index cb9ce8b..b80dcae 100644
--- a/lib/CodeGen/CGDebugInfo.h
+++ b/lib/CodeGen/CGDebugInfo.h
@@ -52,6 +52,9 @@
   llvm::DIType ClassTy;
   llvm::DIType ObjTy;
   llvm::DIType SelTy;
+  llvm::DIType OCLImage1dDITy, OCLImage1dArrayDITy, OCLImage1dBufferDITy;
+  llvm::DIType OCLImage2dDITy, OCLImage2dArrayDITy;
+  llvm::DIType OCLImage3dDITy;
   
   /// TypeCache - Cache of previously constructed Types.
   llvm::DenseMap<void *, llvm::WeakVH> TypeCache;
@@ -116,7 +119,9 @@
   llvm::DIType CreatePointerLikeType(unsigned Tag,
                                      const Type *Ty, QualType PointeeTy,
                                      llvm::DIFile F);
-  
+
+  llvm::DIType getOrCreateStructPtrType(StringRef Name, llvm::DIType &Cache);
+
   llvm::DISubprogram CreateCXXMemberFunction(const CXXMethodDecl *Method,
                                              llvm::DIFile F,
                                              llvm::DIType RecordTy);
diff --git a/lib/CodeGen/CGOpenCLRuntime.cpp b/lib/CodeGen/CGOpenCLRuntime.cpp
index 3a0e116..6052c51 100644
--- a/lib/CodeGen/CGOpenCLRuntime.cpp
+++ b/lib/CodeGen/CGOpenCLRuntime.cpp
@@ -16,6 +16,8 @@
 #include "CGOpenCLRuntime.h"
 #include "CodeGenFunction.h"
 #include "llvm/GlobalValue.h"
+#include "llvm/DerivedTypes.h"
+#include <assert.h>
 
 using namespace clang;
 using namespace CodeGen;
@@ -26,3 +28,32 @@
                                                 const VarDecl &D) {
   return CGF.EmitStaticVarDecl(D, llvm::GlobalValue::InternalLinkage);
 }
+
+llvm::Type *CGOpenCLRuntime::convertOpenCLSpecificType(const Type *T) {
+  assert(T->isOpenCLSpecificType() &&
+         "Not an OpenCL specific type!");
+
+  switch (cast<BuiltinType>(T)->getKind()) {
+  default: 
+    llvm_unreachable("Unexpected opencl builtin type!");
+    return 0;
+  case BuiltinType::OCLImage1d:
+    return llvm::PointerType::get(llvm::StructType::create(
+                           CGM.getLLVMContext(), "opencl.image1d_t"), 0);
+  case BuiltinType::OCLImage1dArray:
+    return llvm::PointerType::get(llvm::StructType::create(
+                           CGM.getLLVMContext(), "opencl.image1d_array_t"), 0);
+  case BuiltinType::OCLImage1dBuffer:
+    return llvm::PointerType::get(llvm::StructType::create(
+                           CGM.getLLVMContext(), "opencl.image1d_buffer_t"), 0);
+  case BuiltinType::OCLImage2d:
+    return llvm::PointerType::get(llvm::StructType::create(
+                           CGM.getLLVMContext(), "opencl.image2d_t"), 0);
+  case BuiltinType::OCLImage2dArray:
+    return llvm::PointerType::get(llvm::StructType::create(
+                           CGM.getLLVMContext(), "opencl.image2d_array_t"), 0);
+  case BuiltinType::OCLImage3d:
+    return llvm::PointerType::get(llvm::StructType::create(
+                           CGM.getLLVMContext(), "opencl.image3d_t"), 0);
+  }
+}
diff --git a/lib/CodeGen/CGOpenCLRuntime.h b/lib/CodeGen/CGOpenCLRuntime.h
index 9a8430f..ba7ea9b 100644
--- a/lib/CodeGen/CGOpenCLRuntime.h
+++ b/lib/CodeGen/CGOpenCLRuntime.h
@@ -16,6 +16,10 @@
 #ifndef CLANG_CODEGEN_OPENCLRUNTIME_H
 #define CLANG_CODEGEN_OPENCLRUNTIME_H
 
+#include "clang/AST/Type.h"
+#include "llvm/Type.h"
+#include "llvm/Value.h"
+
 namespace clang {
 
 class VarDecl;
@@ -38,6 +42,8 @@
   /// CodeGenFunction::EmitStaticVarDecl to emit an internal global for D.
   virtual void EmitWorkGroupLocalVarDecl(CodeGenFunction &CGF,
                                          const VarDecl &D);
+
+  virtual llvm::Type *convertOpenCLSpecificType(const Type *T);
 };
 
 }
diff --git a/lib/CodeGen/CGRTTI.cpp b/lib/CodeGen/CGRTTI.cpp
index 53716a0..41b73e2 100644
--- a/lib/CodeGen/CGRTTI.cpp
+++ b/lib/CodeGen/CGRTTI.cpp
@@ -191,6 +191,12 @@
     case BuiltinType::Char32:
     case BuiltinType::Int128:
     case BuiltinType::UInt128:
+    case BuiltinType::OCLImage1d:
+    case BuiltinType::OCLImage1dArray:
+    case BuiltinType::OCLImage1dBuffer:
+    case BuiltinType::OCLImage2d:
+    case BuiltinType::OCLImage2dArray:
+    case BuiltinType::OCLImage3d:
       return true;
       
     case BuiltinType::Dependent:
diff --git a/lib/CodeGen/CodeGenTypes.cpp b/lib/CodeGen/CodeGenTypes.cpp
index a0effa8..8cb6dd0 100644
--- a/lib/CodeGen/CodeGenTypes.cpp
+++ b/lib/CodeGen/CodeGenTypes.cpp
@@ -14,6 +14,7 @@
 #include "CodeGenTypes.h"
 #include "CGCXXABI.h"
 #include "CGCall.h"
+#include "CGOpenCLRuntime.h"
 #include "CGRecordLayout.h"
 #include "TargetInfo.h"
 #include "clang/AST/ASTContext.h"
@@ -366,6 +367,15 @@
     case BuiltinType::Int128:
       ResultType = llvm::IntegerType::get(getLLVMContext(), 128);
       break;
+
+    case BuiltinType::OCLImage1d:
+    case BuiltinType::OCLImage1dArray:
+    case BuiltinType::OCLImage1dBuffer:
+    case BuiltinType::OCLImage2d:
+    case BuiltinType::OCLImage2dArray:
+    case BuiltinType::OCLImage3d:
+      ResultType = CGM.getOpenCLRuntime().convertOpenCLSpecificType(Ty);
+      break;
     
     case BuiltinType::Dependent:
 #define BUILTIN_TYPE(Id, SingletonId)
diff --git a/lib/Parse/ParseDecl.cpp b/lib/Parse/ParseDecl.cpp
index 518c8a07..390fd34 100644
--- a/lib/Parse/ParseDecl.cpp
+++ b/lib/Parse/ParseDecl.cpp
@@ -2753,6 +2753,30 @@
     case tok::kw___pixel:
       isInvalid = DS.SetTypeAltiVecPixel(true, Loc, PrevSpec, DiagID);
       break;
+    case tok::kw_image1d_t:
+       isInvalid = DS.SetTypeSpecType(DeclSpec::TST_image1d_t, Loc,
+                                      PrevSpec, DiagID);
+      break;
+    case tok::kw_image1d_array_t:
+       isInvalid = DS.SetTypeSpecType(DeclSpec::TST_image1d_array_t, Loc,
+                                      PrevSpec, DiagID);
+      break;
+    case tok::kw_image1d_buffer_t:
+       isInvalid = DS.SetTypeSpecType(DeclSpec::TST_image1d_buffer_t, Loc,
+                                      PrevSpec, DiagID);
+      break;
+    case tok::kw_image2d_t:
+       isInvalid = DS.SetTypeSpecType(DeclSpec::TST_image2d_t, Loc,
+                                      PrevSpec, DiagID);
+      break;
+    case tok::kw_image2d_array_t:
+       isInvalid = DS.SetTypeSpecType(DeclSpec::TST_image2d_array_t, Loc,
+                                      PrevSpec, DiagID);
+      break;
+    case tok::kw_image3d_t:
+      isInvalid = DS.SetTypeSpecType(DeclSpec::TST_image3d_t, Loc,
+                                     PrevSpec, DiagID);
+      break;
     case tok::kw___unknown_anytype:
       isInvalid = DS.SetTypeSpecType(TST_unknown_anytype, Loc,
                                      PrevSpec, DiagID);
@@ -3596,6 +3620,14 @@
   case tok::kw__Decimal128:
   case tok::kw___vector:
 
+    // OpenCL specific types:
+  case tok::kw_image1d_t:
+  case tok::kw_image1d_array_t:
+  case tok::kw_image1d_buffer_t:
+  case tok::kw_image2d_t:
+  case tok::kw_image2d_array_t:
+  case tok::kw_image3d_t:
+
     // struct-or-union-specifier (C99) or class-specifier (C++)
   case tok::kw_class:
   case tok::kw_struct:
@@ -3668,6 +3700,14 @@
   case tok::kw__Decimal128:
   case tok::kw___vector:
 
+    // OpenCL specific types:
+  case tok::kw_image1d_t:
+  case tok::kw_image1d_array_t:
+  case tok::kw_image1d_buffer_t:
+  case tok::kw_image2d_t:
+  case tok::kw_image2d_array_t:
+  case tok::kw_image3d_t:
+
     // struct-or-union-specifier (C99) or class-specifier (C++)
   case tok::kw_class:
   case tok::kw_struct:
@@ -3812,6 +3852,14 @@
   case tok::kw__Decimal128:
   case tok::kw___vector:
 
+    // OpenCL specific types:
+  case tok::kw_image1d_t:
+  case tok::kw_image1d_array_t:
+  case tok::kw_image1d_buffer_t:
+  case tok::kw_image2d_t:
+  case tok::kw_image2d_array_t:
+  case tok::kw_image3d_t:
+
     // struct-or-union-specifier (C99) or class-specifier (C++)
   case tok::kw_class:
   case tok::kw_struct:
diff --git a/lib/Parse/ParseExpr.cpp b/lib/Parse/ParseExpr.cpp
index 14980ee..b7705f8 100644
--- a/lib/Parse/ParseExpr.cpp
+++ b/lib/Parse/ParseExpr.cpp
@@ -1078,7 +1078,13 @@
   case tok::kw_void:
   case tok::kw_typename:
   case tok::kw_typeof:
-  case tok::kw___vector: {
+  case tok::kw___vector:
+  case tok::kw_image1d_t:
+  case tok::kw_image1d_array_t:
+  case tok::kw_image1d_buffer_t:
+  case tok::kw_image2d_t:
+  case tok::kw_image2d_array_t:
+  case tok::kw_image3d_t: {
     if (!getLangOpts().CPlusPlus) {
       Diag(Tok, diag::err_expected_expression);
       return ExprError();
diff --git a/lib/Parse/ParseTentative.cpp b/lib/Parse/ParseTentative.cpp
index b26181f..01863b4 100644
--- a/lib/Parse/ParseTentative.cpp
+++ b/lib/Parse/ParseTentative.cpp
@@ -837,6 +837,12 @@
   case tok::kw___vector:
   case tok::kw___pixel:
   case tok::kw__Atomic:
+  case tok::kw_image1d_t:
+  case tok::kw_image1d_array_t:
+  case tok::kw_image1d_buffer_t:
+  case tok::kw_image2d_t:
+  case tok::kw_image2d_array_t:
+  case tok::kw_image3d_t:
   case tok::kw___unknown_anytype:
     return TPResult::False();
 
diff --git a/lib/Sema/DeclSpec.cpp b/lib/Sema/DeclSpec.cpp
index 040e638..b34b953 100644
--- a/lib/Sema/DeclSpec.cpp
+++ b/lib/Sema/DeclSpec.cpp
@@ -280,6 +280,12 @@
     case TST_unspecified:
     case TST_void:
     case TST_wchar:
+    case TST_image1d_t:
+    case TST_image1d_array_t:
+    case TST_image1d_buffer_t:
+    case TST_image2d_t:
+    case TST_image2d_array_t:
+    case TST_image3d_t:
       return false;
 
     case TST_decltype:
@@ -414,6 +420,12 @@
   case DeclSpec::TST_underlyingType: return "__underlying_type";
   case DeclSpec::TST_unknown_anytype: return "__unknown_anytype";
   case DeclSpec::TST_atomic: return "_Atomic";
+  case DeclSpec::TST_image1d_t:   return "image1d_t";
+  case DeclSpec::TST_image1d_array_t: return "image1d_array_t";
+  case DeclSpec::TST_image1d_buffer_t: return "image1d_buffer_t";
+  case DeclSpec::TST_image2d_t:   return "image2d_t";
+  case DeclSpec::TST_image2d_array_t: return "image2d_array_t";
+  case DeclSpec::TST_image3d_t:   return "image3d_t";
   case DeclSpec::TST_error:       return "(error)";
   }
   llvm_unreachable("Unknown typespec!");
diff --git a/lib/Sema/SemaTemplateVariadic.cpp b/lib/Sema/SemaTemplateVariadic.cpp
index f797825..6c70320 100644
--- a/lib/Sema/SemaTemplateVariadic.cpp
+++ b/lib/Sema/SemaTemplateVariadic.cpp
@@ -731,6 +731,12 @@
   case TST_class:
   case TST_auto:
   case TST_unknown_anytype:
+  case TST_image1d_t:
+  case TST_image1d_array_t:
+  case TST_image1d_buffer_t:
+  case TST_image2d_t:
+  case TST_image2d_array_t:
+  case TST_image3d_t:
   case TST_error:
     break;
   }
diff --git a/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp
index 2c7b7c1..1afc8c7 100644
--- a/lib/Sema/SemaType.cpp
+++ b/lib/Sema/SemaType.cpp
@@ -903,6 +903,30 @@
     }
     break;
 
+  case DeclSpec::TST_image1d_t:
+    Result = Context.OCLImage1dTy;
+    break;
+
+  case DeclSpec::TST_image1d_array_t:
+    Result = Context.OCLImage1dArrayTy;
+    break;
+
+  case DeclSpec::TST_image1d_buffer_t:
+    Result = Context.OCLImage1dBufferTy;
+    break;
+
+  case DeclSpec::TST_image2d_t:
+    Result = Context.OCLImage2dTy;
+    break;
+
+  case DeclSpec::TST_image2d_array_t:
+    Result = Context.OCLImage2dArrayTy;
+    break;
+
+  case DeclSpec::TST_image3d_t:
+    Result = Context.OCLImage3dTy;
+    break;
+
   case DeclSpec::TST_error:
     Result = Context.IntTy;
     declarator.setInvalidType(true);
diff --git a/lib/Serialization/ASTCommon.cpp b/lib/Serialization/ASTCommon.cpp
index aec8a81..bcd8e87 100644
--- a/lib/Serialization/ASTCommon.cpp
+++ b/lib/Serialization/ASTCommon.cpp
@@ -60,6 +60,12 @@
   case BuiltinType::ObjCId:     ID = PREDEF_TYPE_OBJC_ID;       break;
   case BuiltinType::ObjCClass:  ID = PREDEF_TYPE_OBJC_CLASS;    break;
   case BuiltinType::ObjCSel:    ID = PREDEF_TYPE_OBJC_SEL;      break;
+  case BuiltinType::OCLImage1d:       ID = PREDEF_TYPE_IMAGE1D_ID;      break;
+  case BuiltinType::OCLImage1dArray:  ID = PREDEF_TYPE_IMAGE1D_ARR_ID;  break;
+  case BuiltinType::OCLImage1dBuffer: ID = PREDEF_TYPE_IMAGE1D_BUFF_ID; break;
+  case BuiltinType::OCLImage2d:       ID = PREDEF_TYPE_IMAGE2D_ID;      break;
+  case BuiltinType::OCLImage2dArray:  ID = PREDEF_TYPE_IMAGE2D_ARR_ID;  break;
+  case BuiltinType::OCLImage3d:       ID = PREDEF_TYPE_IMAGE3D_ID;      break;
   case BuiltinType::BuiltinFn:
                                 ID = PREDEF_TYPE_BUILTIN_FN; break;
 
diff --git a/lib/Serialization/ASTReader.cpp b/lib/Serialization/ASTReader.cpp
index c0976ee..08bc6b4 100644
--- a/lib/Serialization/ASTReader.cpp
+++ b/lib/Serialization/ASTReader.cpp
@@ -4883,6 +4883,12 @@
     case PREDEF_TYPE_OBJC_ID:       T = Context.ObjCBuiltinIdTy;    break;
     case PREDEF_TYPE_OBJC_CLASS:    T = Context.ObjCBuiltinClassTy; break;
     case PREDEF_TYPE_OBJC_SEL:      T = Context.ObjCBuiltinSelTy;   break;
+    case PREDEF_TYPE_IMAGE1D_ID:    T = Context.OCLImage1dTy;       break;
+    case PREDEF_TYPE_IMAGE1D_ARR_ID: T = Context.OCLImage1dArrayTy; break;
+    case PREDEF_TYPE_IMAGE1D_BUFF_ID: T = Context.OCLImage1dBufferTy; break;
+    case PREDEF_TYPE_IMAGE2D_ID:    T = Context.OCLImage2dTy;       break;
+    case PREDEF_TYPE_IMAGE2D_ARR_ID: T = Context.OCLImage2dArrayTy; break;
+    case PREDEF_TYPE_IMAGE3D_ID:    T = Context.OCLImage3dTy;       break;
     case PREDEF_TYPE_AUTO_DEDUCT:   T = Context.getAutoDeductType(); break;
         
     case PREDEF_TYPE_AUTO_RREF_DEDUCT: 
diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl
new file mode 100644
index 0000000..35444ed
--- /dev/null
+++ b/test/CodeGenOpenCL/opencl_types.cl
@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s
+
+void fnc1(image1d_t img) {}
+// CHECK: @fnc1(%opencl.image1d_t*
+
+void fnc1arr(image1d_array_t img) {}
+// CHECK: @fnc1arr(%opencl.image1d_array_t*
+
+void fnc1buff(image1d_buffer_t img) {}
+// CHECK: @fnc1buff(%opencl.image1d_buffer_t*
+
+void fnc2(image2d_t img) {}
+// CHECK: @fnc2(%opencl.image2d_t*
+
+void fnc2arr(image2d_array_t img) {}
+// CHECK: @fnc2arr(%opencl.image2d_array_t*
+
+void fnc3(image3d_t img) {}
+// CHECK: @fnc3(%opencl.image3d_t*
+
+kernel void foo(image1d_t img) {
+}
diff --git a/test/PCH/ocl_types.cl b/test/PCH/ocl_types.cl
new file mode 100644
index 0000000..972853b
--- /dev/null
+++ b/test/PCH/ocl_types.cl
@@ -0,0 +1,18 @@
+// Test this without pch.
+// RUN: %clang_cc1 -include %S/ocl_types.h -fsyntax-only %s
+
+// Test with pch.
+// RUN: %clang_cc1 -x cl -emit-pch -o %t %S/ocl_types.h
+// RUN: %clang_cc1 -include-pch %t -fsyntax-only %s -ast-print
+
+void foo1(img1d_t img);
+
+void foo2(img1darr_t img);
+
+void foo3(img1dbuff_t img);
+
+void foo4(img2d_t img);
+
+void foo5(img2darr_t img);
+
+void foo6(img3d_t img);
diff --git a/test/PCH/ocl_types.h b/test/PCH/ocl_types.h
new file mode 100644
index 0000000..bec065a
--- /dev/null
+++ b/test/PCH/ocl_types.h
@@ -0,0 +1,19 @@
+/* Used with the ocl_types.cl test */
+
+// image1d_t
+typedef image1d_t img1d_t;
+
+// image1d_array_t
+typedef image1d_array_t img1darr_t;
+
+// image1d_buffer_t
+typedef image1d_buffer_t img1dbuff_t;
+
+// image2d_t
+typedef image2d_t img2d_t;
+
+// image2d_array_t
+typedef image2d_array_t img2darr_t;
+
+// image3d_t
+typedef image3d_t img3d_t;
diff --git a/test/Parser/opencl-image-access.cl b/test/Parser/opencl-image-access.cl
index 313587c..e08d129 100644
--- a/test/Parser/opencl-image-access.cl
+++ b/test/Parser/opencl-image-access.cl
@@ -1,7 +1,5 @@
 // RUN: %clang_cc1 %s -fsyntax-only
 
-typedef void* image2d_t;
-
 __kernel void f__ro(__read_only image2d_t a) { }
 
 __kernel void f__wo(__write_only image2d_t a) { }
diff --git a/tools/libclang/CIndex.cpp b/tools/libclang/CIndex.cpp
index 1fb9a70..3532abd 100644
--- a/tools/libclang/CIndex.cpp
+++ b/tools/libclang/CIndex.cpp
@@ -1367,6 +1367,12 @@
   case BuiltinType::Void:
   case BuiltinType::NullPtr:
   case BuiltinType::Dependent:
+  case BuiltinType::OCLImage1d:
+  case BuiltinType::OCLImage1dArray:
+  case BuiltinType::OCLImage1dBuffer:
+  case BuiltinType::OCLImage2d:
+  case BuiltinType::OCLImage2dArray:
+  case BuiltinType::OCLImage3d:
 #define BUILTIN_TYPE(Id, SingletonId)
 #define SIGNED_TYPE(Id, SingletonId) case BuiltinType::Id:
 #define UNSIGNED_TYPE(Id, SingletonId) case BuiltinType::Id:
diff --git a/tools/libclang/CIndexUSRs.cpp b/tools/libclang/CIndexUSRs.cpp
index b76363e..1cfb6a4 100644
--- a/tools/libclang/CIndexUSRs.cpp
+++ b/tools/libclang/CIndexUSRs.cpp
@@ -587,6 +587,12 @@
 #define PLACEHOLDER_TYPE(Id, SingletonId) case BuiltinType::Id:
 #include "clang/AST/BuiltinTypes.def"
         case BuiltinType::Dependent:
+        case BuiltinType::OCLImage1d:
+        case BuiltinType::OCLImage1dArray:
+        case BuiltinType::OCLImage1dBuffer:
+        case BuiltinType::OCLImage2d:
+        case BuiltinType::OCLImage2dArray:
+        case BuiltinType::OCLImage3d:
           IgnoreResults = true;
           return;
         case BuiltinType::ObjCId: