Index: include/clang/AST/DeclOpenMP.h
===================================================================
--- include/clang/AST/DeclOpenMP.h
+++ include/clang/AST/DeclOpenMP.h
@@ -189,8 +189,9 @@
   void anchor() override;
 
   OMPCapturedExprDecl(ASTContext &C, DeclContext *DC, IdentifierInfo *Id,
-                      QualType Type, SourceLocation StartLoc)
-      : VarDecl(OMPCapturedExpr, C, DC, StartLoc, StartLoc, Id, Type, nullptr,
+                      QualType Type, TypeSourceInfo *TInfo,
+                      SourceLocation StartLoc)
+      : VarDecl(OMPCapturedExpr, C, DC, StartLoc, StartLoc, Id, Type, TInfo,
                 SC_None) {
     setImplicit();
   }
Index: lib/AST/DeclOpenMP.cpp
===================================================================
--- lib/AST/DeclOpenMP.cpp
+++ lib/AST/DeclOpenMP.cpp
@@ -92,13 +92,14 @@
 OMPCapturedExprDecl *OMPCapturedExprDecl::Create(ASTContext &C, DeclContext *DC,
                                                  IdentifierInfo *Id, QualType T,
                                                  SourceLocation StartLoc) {
-  return new (C, DC) OMPCapturedExprDecl(C, DC, Id, T, StartLoc);
+  return new (C, DC) OMPCapturedExprDecl(
+      C, DC, Id, T, C.getTrivialTypeSourceInfo(T), StartLoc);
 }
 
 OMPCapturedExprDecl *OMPCapturedExprDecl::CreateDeserialized(ASTContext &C,
                                                              unsigned ID) {
-  return new (C, ID)
-      OMPCapturedExprDecl(C, nullptr, nullptr, QualType(), SourceLocation());
+  return new (C, ID) OMPCapturedExprDecl(C, nullptr, nullptr, QualType(),
+                                         /*TInfo=*/nullptr, SourceLocation());
 }
 
 SourceRange OMPCapturedExprDecl::getSourceRange() const {
Index: test/SemaOpenCL/with_openmp.cl
===================================================================
--- /dev/null
+++ test/SemaOpenCL/with_openmp.cl
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -verify -fopenmp -ast-dump -x cl %s 2>&1 | FileCheck %s
+// expected-no-diagnostics
+
+// CHECK: FunctionDecl{{.*}}getVal
+// CHECK: ParmVarDecl{{.*}}v
+// CHECK: ParmVarDecl{{.*}}m
+// CHECK: OMPDeclareSimdDeclAttr
+#pragma omp declare simd
+inline float getVal(const float v, const float m)
+{
+  return v * m + 0.5F;
+}
+
+// CHECK: FunctionDecl{{.*}}KFunc
+// CHECK-NEXT: ParmVarDecl{{.*}}data
+// CHECK-NEXT: ParmVarDecl{{.*}}size
+__kernel void KFunc(__global int *data, int size) {
+  float foo[128];
+// CHECK: OMPSimdDirective
+// CHECK: OMPCapturedExprDecl
+// CHECK-NEXT: ImplicitCastExpr{{.*}}LValueToRValue
+// CHECK-NEXT: DeclRefExpr {{.*}}ParmVar{{.*}}size
+  #pragma omp simd simdlen(16)
+  for (int i = 0; i < size; ++i) {
+    foo[i] = getVal(*data, 1);
+  }
+
+// CHECK: OMPParallelForDirective
+// CHECK: OMPCapturedExprDecl
+// CHECK-NEXT: ImplicitCastExpr{{.*}}LValueToRValue
+// CHECK-NEXT: DeclRefExpr {{.*}}ParmVar{{.*}}size
+  #pragma omp parallel for
+  for (int i = 0; i < size; ++i) {
+    foo[i] = *data + 1;
+  }
+// CHECK: OMPTaskLoopDirective
+// CHECK: OMPCapturedExprDecl
+// CHECK-NEXT: ImplicitCastExpr{{.*}}LValueToRValue
+// CHECK-NEXT: DeclRefExpr {{.*}}ParmVar{{.*}}size
+  #pragma omp taskloop
+  for (int i = 0; i < size; ++i) {
+    foo[i] = *data + 1;
+  }
+
+// CHECK: OMPParallelDirective
+  #pragma omp parallel
+  {
+// CHECK: OMPForDirective
+// CHECK: OMPCapturedExprDecl
+// CHECK-NEXT: ImplicitCastExpr{{.*}}LValueToRValue
+// CHECK-NEXT: DeclRefExpr {{.*}}ParmVar{{.*}}size
+    #pragma omp for
+    for (int i = 0; i < size; ++i) {
+      foo[i] = *data + 1;
+    }
+  }
+}