Index: include/clang/AST/ASTContext.h
===================================================================
--- include/clang/AST/ASTContext.h	(revision 165097)
+++ include/clang/AST/ASTContext.h	(working copy)
@@ -699,6 +699,9 @@
   CanQualType PseudoObjectTy, ARCUnbridgedCastTy;
   CanQualType ObjCBuiltinIdTy, ObjCBuiltinClassTy, ObjCBuiltinSelTy;
   CanQualType ObjCBuiltinBoolTy;
+  CanQualType Image1dTTy, Image1dArrayTTy, Image1dBufferTTy;
+  CanQualType Image2dTTy, Image2dArrayTTy, Image3dTTy;
+  CanQualType SamplerTTy, EventTTy;
 
   // Types for deductions in C++0x [stmt.ranged]'s desugaring. Built on demand.
   mutable QualType AutoDeductTy;     // Deduction against 'auto'.
Index: include/clang/AST/BuiltinTypes.def
===================================================================
--- include/clang/AST/BuiltinTypes.def	(revision 165097)
+++ include/clang/AST/BuiltinTypes.def	(working copy)
@@ -154,6 +154,20 @@
 // type is a typedef of a PointerType to this.
 BUILTIN_TYPE(ObjCSel, ObjCBuiltinSelTy)
 
+// OpenCL image types
+BUILTIN_TYPE(Image1d_T, Image1dTTy)
+BUILTIN_TYPE(Image1dArray_T, Image1dArrayTTy)
+BUILTIN_TYPE(Image1dBuffer_T, Image1dBufferTTy)
+BUILTIN_TYPE(Image2d_T, Image2dTTy)
+BUILTIN_TYPE(Image2dArray_T, Image2dArrayTTy)
+BUILTIN_TYPE(Image3d_T, Image3dTTy)
+
+// OpenCL sampler_t
+BUILTIN_TYPE(Sampler_T, SamplerTTy)
+
+// OpenCL event_t
+BUILTIN_TYPE(Event_T, EventTTy)
+
 // 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
Index: include/clang/AST/Type.h
===================================================================
--- include/clang/AST/Type.h	(revision 165097)
+++ include/clang/AST/Type.h	(working copy)
@@ -1584,6 +1584,18 @@
   bool isNullPtrType() const;                   // C++0x nullptr_t
   bool isAtomicType() const;                    // C11 _Atomic()
 
+  bool isImage1dT() const;                      // OpenCL image1d_t
+  bool isImage1dArrayT() const;                 // OpenCL image1d_t
+  bool isImage1dBufferT() const;                // OpenCL image1d_t
+  bool isImage2dT() const;                      // OpenCL image2d_t
+  bool isImage2dArrayT() const;                 // OpenCL image2d_t
+  bool isImage3dT() const;                      // OpenCL image3d_t
+
+  bool isImageType() const;                     // Any OpenCL image type
+
+  bool isSamplerT() const;                      // OpenCL sampler_t
+  bool isEventT() const;                        // OpenCL event_t
+
   /// Determines if this type, which must satisfy
   /// isObjCLifetimeType(), is implicitly __unsafe_unretained rather
   /// than implicitly __strong.
@@ -4936,6 +4948,61 @@
 inline bool Type::isObjCBuiltinType() const {
   return isObjCIdType() || isObjCClassType() || isObjCSelType();
 }
+
+inline bool Type::isImage1dT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Image1d_T;
+  return false;
+}
+
+inline bool Type::isImage1dArrayT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Image1dArray_T;
+  return false;
+}
+
+inline bool Type::isImage1dBufferT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Image1dBuffer_T;
+  return false;
+}
+
+inline bool Type::isImage2dT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Image2d_T;
+  return false;
+}
+
+inline bool Type::isImage2dArrayT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Image2dArray_T;
+  return false;
+}
+
+inline bool Type::isImage3dT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Image3d_T;
+  return false;
+}
+
+inline bool Type::isSamplerT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Sampler_T;
+  return false;
+}
+
+inline bool Type::isEventT() const {
+  if (const BuiltinType *BT = dyn_cast<BuiltinType>(CanonicalType))
+    return BT->getKind() == BuiltinType::Event_T;
+  return false;
+}
+
+inline bool Type::isImageType() const {
+  return isImage1dT() || isImage1dArrayT() || isImage1dBufferT() ||
+         isImage2dT() || isImage2dArrayT() ||
+         isImage3dT();
+}
+
 inline bool Type::isTemplateTypeParmType() const {
   return isa<TemplateTypeParmType>(CanonicalType);
 }
Index: include/clang/Basic/Specifiers.h
===================================================================
--- include/clang/Basic/Specifiers.h	(revision 165097)
+++ include/clang/Basic/Specifiers.h	(working copy)
@@ -62,6 +62,14 @@
     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_sampler_t,        // OpenCL sampler_t
+    TST_event_t,          // OpenCL event_t
     TST_error         // erroneous type
   };
   
Index: include/clang/Basic/TokenKinds.def
===================================================================
--- include/clang/Basic/TokenKinds.def	(revision 165097)
+++ include/clang/Basic/TokenKinds.def	(working copy)
@@ -448,6 +448,14 @@
 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)
+KEYWORD(sampler_t                   , KEYOPENCL)
+KEYWORD(event_t                     , KEYOPENCL)
 
 // Borland Extensions.
 KEYWORD(__pascal                    , KEYALL)
Index: include/clang/Sema/DeclSpec.h
===================================================================
--- include/clang/Sema/DeclSpec.h	(revision 165097)
+++ include/clang/Sema/DeclSpec.h	(working copy)
@@ -276,6 +276,14 @@
   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_sampler_t = clang::TST_sampler_t;
+  static const TST TST_event_t = clang::TST_event_t;
   static const TST TST_error = clang::TST_error;
 
   // type-qualifiers
@@ -306,7 +314,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;
Index: include/clang/Serialization/ASTBitCodes.h
===================================================================
--- include/clang/Serialization/ASTBitCodes.h	(revision 165097)
+++ include/clang/Serialization/ASTBitCodes.h	(working copy)
@@ -644,7 +644,23 @@
       /// \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 OpenCL sampler type.
+      PREDEF_TYPE_SAMPLER_ID    = 44,
+      /// \brief OpenCL event type.
+      PREDEF_TYPE_EVENT_ID      = 45
     };
 
     /// \brief The number of predefined type IDs that are reserved for
Index: lib/AST/ASTContext.cpp
===================================================================
--- lib/AST/ASTContext.cpp	(revision 165097)
+++ lib/AST/ASTContext.cpp	(working copy)
@@ -818,6 +818,19 @@
   InitBuiltinType(ObjCBuiltinIdTy, BuiltinType::ObjCId);
   InitBuiltinType(ObjCBuiltinClassTy, BuiltinType::ObjCClass);
   InitBuiltinType(ObjCBuiltinSelTy, BuiltinType::ObjCSel);
+
+  if (LangOpts.OpenCL) { 
+    InitBuiltinType(Image1dTTy, BuiltinType::Image1d_T);
+    InitBuiltinType(Image1dArrayTTy, BuiltinType::Image1dArray_T);
+    InitBuiltinType(Image1dBufferTTy, BuiltinType::Image1dBuffer_T);
+    InitBuiltinType(Image2dTTy, BuiltinType::Image2d_T);
+    InitBuiltinType(Image2dArrayTTy, BuiltinType::Image2dArray_T);
+    InitBuiltinType(Image3dTTy, BuiltinType::Image3d_T);
+
+    InitBuiltinType(SamplerTTy, BuiltinType::Sampler_T);
+    InitBuiltinType(EventTTy, BuiltinType::Event_T);
+  }
+
   
   // Builtin type for __objc_yes and __objc_no
   ObjCBuiltinBoolTy = (Target.useSignedCharForObjCBool() ?
@@ -1335,6 +1348,17 @@
       Width = Target->getPointerWidth(0); 
       Align = Target->getPointerAlign(0);
       break;
+    case BuiltinType::Sampler_T:
+    case BuiltinType::Event_T:
+    case BuiltinType::Image1d_T:
+    case BuiltinType::Image1dArray_T:
+    case BuiltinType::Image1dBuffer_T:
+    case BuiltinType::Image2d_T:
+    case BuiltinType::Image2dArray_T:
+    case BuiltinType::Image3d_T:
+      Width = Target->getPointerWidth(0); // Currently these types are pointers
+      Align = Target->getPointerAlign(0); // to opaque types
+      break;
     }
     break;
   case Type::ObjCObjectPointer:
Index: lib/AST/Type.cpp
===================================================================
--- lib/AST/Type.cpp	(revision 165097)
+++ lib/AST/Type.cpp	(working copy)
@@ -1509,6 +1509,14 @@
   case ObjCId:            return "id";
   case ObjCClass:         return "Class";
   case ObjCSel:           return "SEL";
+  case Image1d_T:         return "image1d_t";
+  case Image1dArray_T:    return "image1d_array_t";
+  case Image1dBuffer_T:   return "image1d_buffer_t";
+  case Image2d_T:         return "image2d_t";
+  case Image2dArray_T:    return "image2d_array_t";
+  case Image3d_T:         return "image3d_t";
+  case Sampler_T:         return "sampler_t";
+  case Event_T:           return "event_t";
   }
   
   llvm_unreachable("Invalid builtin type.");
Index: lib/AST/TypeLoc.cpp
===================================================================
--- lib/AST/TypeLoc.cpp	(revision 165097)
+++ lib/AST/TypeLoc.cpp	(working copy)
@@ -241,6 +241,14 @@
   case BuiltinType::ObjCId:
   case BuiltinType::ObjCClass:
   case BuiltinType::ObjCSel:
+  case BuiltinType::Image1d_T:
+  case BuiltinType::Image1dArray_T:
+  case BuiltinType::Image1dBuffer_T:
+  case BuiltinType::Image2d_T:
+  case BuiltinType::Image2dArray_T:
+  case BuiltinType::Image3d_T:
+  case BuiltinType::Sampler_T:
+  case BuiltinType::Event_T:
   case BuiltinType::BuiltinFn:
     return TST_unspecified;
   }
Index: lib/CodeGen/CGDebugInfo.cpp
===================================================================
--- lib/CodeGen/CGDebugInfo.cpp	(revision 165097)
+++ lib/CodeGen/CGDebugInfo.cpp	(working copy)
@@ -402,6 +402,16 @@
                                  0);
     return SelTy;
   }
+
+  case BuiltinType::Image1d_T:
+  case BuiltinType::Image1dArray_T:
+  case BuiltinType::Image1dBuffer_T:
+  case BuiltinType::Image2d_T:
+  case BuiltinType::Image2dArray_T:
+  case BuiltinType::Image3d_T:
+  case BuiltinType::Sampler_T:
+  case BuiltinType::Event_T: Encoding = llvm::dwarf::DW_ATE_unsigned; break; 
+
   case BuiltinType::UChar:
   case BuiltinType::Char_U: Encoding = llvm::dwarf::DW_ATE_unsigned_char; break;
   case BuiltinType::Char_S:
Index: lib/CodeGen/CGRTTI.cpp
===================================================================
--- lib/CodeGen/CGRTTI.cpp	(revision 165097)
+++ lib/CodeGen/CGRTTI.cpp	(working copy)
@@ -191,6 +191,14 @@
     case BuiltinType::Char32:
     case BuiltinType::Int128:
     case BuiltinType::UInt128:
+    case BuiltinType::Image1d_T:
+    case BuiltinType::Image1dArray_T:
+    case BuiltinType::Image1dBuffer_T:
+    case BuiltinType::Image2d_T:
+    case BuiltinType::Image2dArray_T:
+    case BuiltinType::Image3d_T:
+    case BuiltinType::Sampler_T:
+    case BuiltinType::Event_T:
       return true;
       
     case BuiltinType::Dependent:
Index: lib/CodeGen/CodeGenTypes.cpp
===================================================================
--- lib/CodeGen/CodeGenTypes.cpp	(revision 165097)
+++ lib/CodeGen/CodeGenTypes.cpp	(working copy)
@@ -366,6 +366,46 @@
     case BuiltinType::Int128:
       ResultType = llvm::IntegerType::get(getLLVMContext(), 128);
       break;
+
+    case BuiltinType::Image1d_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.image1d_t"), 0);
+      break;
+
+    case BuiltinType::Image1dArray_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.image1d_array_t"), 0);
+      break;
+
+    case BuiltinType::Image1dBuffer_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.image1d_buffer_t"), 0);
+      break;
+
+    case BuiltinType::Image2d_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.image2d_t"), 0);
+      break;
+
+    case BuiltinType::Image2dArray_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.image2d_array_t"), 0);
+      break;
+
+    case BuiltinType::Image3d_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.image3d_t"), 0);
+      break;
+
+    case BuiltinType::Sampler_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.sampler_t"), 0);
+      break;
+
+    case BuiltinType::Event_T:
+      ResultType = llvm::PointerType::get(llvm::StructType::create(
+                             getLLVMContext(),"opencl.event_t"), 0);
+      break;
     
     case BuiltinType::Dependent:
 #define BUILTIN_TYPE(Id, SingletonId)
Index: lib/Parse/ParseDecl.cpp
===================================================================
--- lib/Parse/ParseDecl.cpp	(revision 165097)
+++ lib/Parse/ParseDecl.cpp	(working copy)
@@ -2713,6 +2713,38 @@
     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_sampler_t:
+      isInvalid = DS.SetTypeSpecType(DeclSpec::TST_sampler_t, Loc,
+                                      PrevSpec, DiagID);
+      break;
+    case tok::kw_event_t:
+      isInvalid = DS.SetTypeSpecType(DeclSpec::TST_event_t, Loc, 
+                                      PrevSpec, DiagID);
+      break;
     case tok::kw___unknown_anytype:
       isInvalid = DS.SetTypeSpecType(TST_unknown_anytype, Loc,
                                      PrevSpec, DiagID);
@@ -3541,6 +3573,16 @@
   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:
+  case tok::kw_sampler_t:
+  case tok::kw_event_t:
+
     // struct-or-union-specifier (C99) or class-specifier (C++)
   case tok::kw_class:
   case tok::kw_struct:
@@ -3613,6 +3655,16 @@
   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:
+  case tok::kw_sampler_t:
+  case tok::kw_event_t:
+
     // struct-or-union-specifier (C99) or class-specifier (C++)
   case tok::kw_class:
   case tok::kw_struct:
@@ -3751,6 +3803,16 @@
   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:
+  case tok::kw_sampler_t:
+  case tok::kw_event_t:
+
     // struct-or-union-specifier (C99) or class-specifier (C++)
   case tok::kw_class:
   case tok::kw_struct:
Index: lib/Parse/ParseExpr.cpp
===================================================================
--- lib/Parse/ParseExpr.cpp	(revision 165097)
+++ lib/Parse/ParseExpr.cpp	(working copy)
@@ -1078,7 +1078,15 @@
   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:
+  case tok::kw_sampler_t:
+  case tok::kw_event_t: {
     if (!getLangOpts().CPlusPlus) {
       Diag(Tok, diag::err_expected_expression);
       return ExprError();
Index: lib/Parse/ParseTentative.cpp
===================================================================
--- lib/Parse/ParseTentative.cpp	(revision 165097)
+++ lib/Parse/ParseTentative.cpp	(working copy)
@@ -837,6 +837,14 @@
   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_sampler_t:
+  case tok::kw_event_t:
     return TPResult::False();
 
   default:
Index: lib/Sema/DeclSpec.cpp
===================================================================
--- lib/Sema/DeclSpec.cpp	(revision 165097)
+++ lib/Sema/DeclSpec.cpp	(working copy)
@@ -276,6 +276,14 @@
     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:
+    case TST_sampler_t:
+    case TST_event_t:
       return false;
 
     case TST_decltype:
@@ -410,6 +418,14 @@
   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_sampler_t:   return "sampler_t";
+  case DeclSpec::TST_event_t:     return "event_t";
   case DeclSpec::TST_error:       return "(error)";
   }
   llvm_unreachable("Unknown typespec!");
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp	(revision 165097)
+++ lib/Sema/SemaType.cpp	(working copy)
@@ -891,6 +891,38 @@
     }
     break;
 
+  case DeclSpec::TST_image1d_t:
+    Result = Context.Image1dTTy;
+    break;
+
+  case DeclSpec::TST_image1d_array_t:
+    Result = Context.Image1dArrayTTy;
+    break;
+
+  case DeclSpec::TST_image1d_buffer_t:
+    Result = Context.Image1dBufferTTy;
+    break;
+
+  case DeclSpec::TST_image2d_t:
+    Result = Context.Image2dTTy;
+    break;
+
+  case DeclSpec::TST_image2d_array_t:
+    Result = Context.Image2dArrayTTy;
+    break;
+
+  case DeclSpec::TST_image3d_t:
+    Result = Context.Image3dTTy;
+    break;
+
+  case DeclSpec::TST_sampler_t:
+    Result = Context.SamplerTTy;
+    break;
+
+  case DeclSpec::TST_event_t:
+    Result = Context.EventTTy;
+    break;
+
   case DeclSpec::TST_error:
     Result = Context.IntTy;
     declarator.setInvalidType(true);
Index: lib/Serialization/ASTCommon.cpp
===================================================================
--- lib/Serialization/ASTCommon.cpp	(revision 165097)
+++ lib/Serialization/ASTCommon.cpp	(working copy)
@@ -60,6 +60,14 @@
   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::Image1d_T:       ID = PREDEF_TYPE_IMAGE1D_ID;      break;
+  case BuiltinType::Image1dArray_T:  ID = PREDEF_TYPE_IMAGE1D_ARR_ID;  break;
+  case BuiltinType::Image1dBuffer_T: ID = PREDEF_TYPE_IMAGE1D_BUFF_ID; break;
+  case BuiltinType::Image2d_T:       ID = PREDEF_TYPE_IMAGE2D_ID;      break;
+  case BuiltinType::Image2dArray_T:  ID = PREDEF_TYPE_IMAGE2D_ARR_ID;  break;
+  case BuiltinType::Image3d_T:       ID = PREDEF_TYPE_IMAGE3D_ID;      break;
+  case BuiltinType::Sampler_T:       ID = PREDEF_TYPE_SAMPLER_ID;      break;
+  case BuiltinType::Event_T:         ID = PREDEF_TYPE_EVENT_ID;        break;
   case BuiltinType::BuiltinFn:
                                 ID = PREDEF_TYPE_BUILTIN_FN; break;
 
Index: lib/Serialization/ASTReader.cpp
===================================================================
--- lib/Serialization/ASTReader.cpp	(revision 165097)
+++ lib/Serialization/ASTReader.cpp	(working copy)
@@ -4489,6 +4489,14 @@
     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.Image1dTTy;         break;
+    case PREDEF_TYPE_IMAGE1D_ARR_ID: T = Context.Image1dArrayTTy;   break;
+    case PREDEF_TYPE_IMAGE1D_BUFF_ID: T = Context.Image1dBufferTTy; break;
+    case PREDEF_TYPE_IMAGE2D_ID:    T = Context.Image2dTTy;         break;
+    case PREDEF_TYPE_IMAGE2D_ARR_ID: T = Context.Image2dArrayTTy;   break;
+    case PREDEF_TYPE_IMAGE3D_ID:    T = Context.Image3dTTy;         break;
+    case PREDEF_TYPE_SAMPLER_ID:    T = Context.SamplerTTy;         break;
+    case PREDEF_TYPE_EVENT_ID:      T = Context.EventTTy;           break;
     case PREDEF_TYPE_AUTO_DEDUCT:   T = Context.getAutoDeductType(); break;
         
     case PREDEF_TYPE_AUTO_RREF_DEDUCT: 
Index: test/CodeGenOpenCL/opencl_types.cl
===================================================================
--- test/CodeGenOpenCL/opencl_types.cl	(revision 0)
+++ test/CodeGenOpenCL/opencl_types.cl	(revision 0)
@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s
+
+sampler_t glb_smp;
+// CHECK: global %opencl.sampler_t*
+
+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) {
+	sampler_t smp;
+// CHECK: alloca %opencl.sampler_t*
+	event_t evt;
+// CHECK: alloca %opencl.event_t*
+}
Index: test/Parser/opencl-image-access.cl
===================================================================
--- test/Parser/opencl-image-access.cl	(revision 165097)
+++ test/Parser/opencl-image-access.cl	(working copy)
@@ -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) { }
