Index: test/SemaOpenCL/kernel-errors.cl
===================================================================
--- test/SemaOpenCL/kernel-errors.cl	(revision 0)
+++ test/SemaOpenCL/kernel-errors.cl	(revision 0)
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 -fsyntax-only -pedantic -verify %s
+
+__kernel void foo(int *x, __global int* y) { // expected-error {{kernel pointer arguments must have a global, local, or constant address space qualifier}}
+}
+
+__kernel int bar()  { // expected-error {{a function using the __kernel qualifier must return type void}}
+  return 6;
+}
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td	(revision 161998)
+++ include/clang/Basic/DiagnosticSemaKinds.td	(working copy)
@@ -5856,7 +5856,10 @@
   "kernel functions cannot be declared static">;
 def err_static_function_scope : Error<
   "variables in function scope cannot be declared static">;
-
+def err_expected_kernel_void_return_type : Error<
+  "a function using the __kernel qualifier must return type void">;
+def err_kernel_arg_with_private_address_space : Error<
+  "kernel pointer arguments must have a global, local, or constant address space qualifier">;
 } // end of sema category
 
 let CategoryName = "Related Result Type Issue" in {
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp	(revision 161998)
+++ lib/Sema/SemaDecl.cpp	(working copy)
@@ -5781,14 +5781,35 @@
     }
   }
 
-  // OpenCL v1.2 s6.8 static is invalid for kernel functions.
-  if ((getLangOpts().OpenCLVersion >= 120)
-      && NewFD->hasAttr<OpenCLKernelAttr>()
-      && (SC == SC_Static)) {
-    Diag(D.getIdentifierLoc(), diag::err_static_kernel);
-    D.setInvalidType();
+  if (getLangOpts().OpenCL && NewFD->hasAttr<OpenCLKernelAttr>()) {
+    // OpenCL v1.2 s6.8 static is invalid for kernel functions.
+    if ((getLangOpts().OpenCLVersion >= 120)
+        && (SC == SC_Static)) {
+      Diag(D.getIdentifierLoc(), diag::err_static_kernel);
+      D.setInvalidType();
+    }
+    // OpenCL v1.1 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. "
+    for (unsigned p = 0, NumParams = NewFD->getNumParams();
+         p < NumParams; ++p) {
+      ParmVarDecl *Param = NewFD->getParamDecl(p);
+      QualType T = Param->getType();
+      const PointerType *PT = dyn_cast<PointerType>(T.getTypePtr());
+      if (PT
+          && PT->getPointeeType().getAddressSpace() == 0) {
+        Diag(Param->getLocation(),
+             diag::err_kernel_arg_with_private_address_space);
+        Param->setInvalidDecl();
+      }
+    }
+    // OpenCL v1.1 s6.8j Kernels may only have void return type.
+    if (!NewFD->getResultType()->isVoidType()) {
+      Diag(D.getIdentifierLoc(),
+           diag::err_expected_kernel_void_return_type);
+      D.setInvalidType();
+    }
   }
-
   MarkUnusedFileScopedDecl(NewFD);
 
   if (getLangOpts().CUDA)