diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -9749,8 +9749,8 @@
 def err_opencl_ptrptr_kernel_param : Error<
   "kernel parameter cannot be declared as a pointer to a pointer">;
 def err_kernel_arg_address_space : Error<
-  "pointer arguments to kernel functions must reside in '__global', "
-  "'__constant' or '__local' address space">;
+  "%select{|nested }0pointer arguments to kernel functions must reside in "
+  "'__global', '__constant' or '__local' address space">;
 def err_opencl_ext_vector_component_invalid_length : Error<
   "vector component access has invalid length %0.  Supported: 1,2,3,4,8,16.">;
 def err_opencl_function_variable : Error<
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -8590,14 +8590,18 @@
   return false;
 }
 
+static bool hasOpenCLValidPtrParameterAddrSpace(QualType Ty) {
+  LangAS AS = Ty.getAddressSpace();
+  return AS != LangAS::opencl_generic && AS != LangAS::opencl_private &&
+         AS != LangAS::Default;
+}
+
 static OpenCLParamType getOpenCLKernelParameterType(Sema &S, QualType PT) {
   if (PT->isPointerType()) {
     QualType PointeeType = PT->getPointeeType();
     if (PointeeType->isPointerType())
       return PtrPtrKernelParam;
-    if (PointeeType.getAddressSpace() == LangAS::opencl_generic ||
-        PointeeType.getAddressSpace() == LangAS::opencl_private ||
-        PointeeType.getAddressSpace() == LangAS::Default)
+    if (!hasOpenCLValidPtrParameterAddrSpace(PointeeType))
       return InvalidAddrSpacePtrKernelParam;
     return PtrKernelParam;
   }
@@ -8650,20 +8654,38 @@
     return;
 
   switch (getOpenCLKernelParameterType(S, PT)) {
-  case PtrPtrKernelParam:
-    // OpenCL v1.2 s6.9.a:
-    // A kernel function argument cannot be declared as a
-    // pointer to a pointer type.
-    S.Diag(Param->getLocation(), diag::err_opencl_ptrptr_kernel_param);
-    D.setInvalidType();
+  case PtrPtrKernelParam: {
+    // OpenCL v3.0 s6.11.a:
+    // A kernel function argument cannot be declared as a pointer to a pointer
+    // type. [...] This restriction only applies to OpenCL C 1.2 or below.
+    if (S.getLangOpts().OpenCLVersion < 120 &&
+        !S.getLangOpts().OpenCLCPlusPlus) {
+      S.Diag(Param->getLocation(), diag::err_opencl_ptrptr_kernel_param);
+      D.setInvalidType();
+      return;
+    }
+
+    // Loop through inner pointer types to ensure address spaces are valid.
+    QualType NestedPT = PT->getPointeeType();
+    do {
+      if (!hasOpenCLValidPtrParameterAddrSpace(NestedPT->getPointeeType())) {
+        S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space) << 1;
+        D.setInvalidType();
+        return;
+      }
+      NestedPT = NestedPT->getPointeeType();
+    } while (NestedPT->isPointerType());
+
+    ValidTypes.insert(PT.getTypePtr());
     return;
+  }
 
   case InvalidAddrSpacePtrKernelParam:
     // OpenCL v1.0 s6.5:
     // __kernel function arguments declared to be a pointer of a type can point
     // to one of the following address spaces only : __global, __local or
     // __constant.
-    S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space);
+    S.Diag(Param->getLocation(), diag::err_kernel_arg_address_space) << 0;
     D.setInvalidType();
     return;
 
diff --git a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
--- a/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
+++ b/clang/test/SemaOpenCL/invalid-kernel-parameters.cl
@@ -5,8 +5,12 @@
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-// expected-error@+1{{kernel parameter cannot be declared as a pointer to a pointer}}
+#if __OPENCL_C_VERSION__ < CL_VERSION_2_0
+// expected-error@+3{{kernel parameter cannot be declared as a pointer to a pointer}}
+// expected-error@+3{{kernel parameter cannot be declared as a pointer to a pointer}}
+#endif
 kernel void no_ptrptr(global int * global *i) { }
+kernel void no_ptrptrptr(global int * global * global *i) { }
 
 // expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 __kernel void no_privateptr(__private int *i) { }
@@ -17,6 +21,15 @@
 // expected-error@+1{{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 __kernel void no_addrsp_ptr(int *ptr) { }
 
+#if __OPENCL_C_VERSION__ >= CL_VERSION_2_0
+kernel void no_ptr_private_ptr(private int * global *i) { }
+// expected-error@-1{{nested pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+kernel void no_ptr_ptr_private_ptr(private int * global * global *i) { }
+// expected-error@-1{{nested pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+kernel void no_ptr_private_ptr_ptr(global int * private * global *i) { }
+// expected-error@-1{{nested pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+#endif
+
 // Disallowed: parameters with type
 // bool, half, size_t, ptrdiff_t, intptr_t, and uintptr_t
 // or a struct / union with any of these types in them