[OpenCL] Improve address space diagnostics.

Reviewers: Anastasia

Subscribers: bader, yaxunl, cfe-commits

Differential Revision: https://reviews.llvm.org/D27671

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@289536 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/Basic/DiagnosticSemaKinds.td

@@ -8076,8 +8076,9 @@ def err_static_kernel : Error<
   "kernel functions cannot be declared static">;
 def err_opencl_ptrptr_kernel_param : Error<
   "kernel parameter cannot be declared as a pointer to a pointer">;
-def err_opencl_private_ptr_kernel_param : Error<
-  "kernel parameter cannot be declared as a pointer to the __private address space">;
+def err_kernel_arg_address_space : Error<
+  "pointer arguments to kernel functions must reside in '__global', "
+  "'__constant' or '__local' address space">;
 def err_opencl_function_variable : Error<
   "%select{non-kernel function|function scope}0 variable cannot be declared in %1 address space">;
 def err_static_function_scope : Error<

lib/Sema/SemaDecl.cpp

@@ -7586,7 +7586,7 @@ enum OpenCLParamType {
   ValidKernelParam,
   PtrPtrKernelParam,
   PtrKernelParam,
-  PrivatePtrKernelParam,
+  InvalidAddrSpacePtrKernelParam,
   InvalidKernelParam,
   RecordKernelParam
 };
@@ -7596,8 +7596,10 @@ static OpenCLParamType getOpenCLKernelParameterType(Sema &S, QualType PT) {
     QualType PointeeType = PT->getPointeeType();
     if (PointeeType->isPointerType())
       return PtrPtrKernelParam;
-    return PointeeType.getAddressSpace() == 0 ? PrivatePtrKernelParam
-                                              : PtrKernelParam;
+    if (PointeeType.getAddressSpace() == LangAS::opencl_generic ||
+        PointeeType.getAddressSpace() == 0)
+      return InvalidAddrSpacePtrKernelParam;
+    return PtrKernelParam;
   }
 
   // TODO: Forbid the other integer types (size_t, ptrdiff_t...) when they can
@@ -7645,11 +7647,12 @@ static void checkIsValidOpenCLKernelParameter(
     D.setInvalidType();
     return;
 
-  case PrivatePtrKernelParam:
-    // OpenCL v1.2 s6.9.a:
-    // A kernel function argument cannot be declared as a
-    // pointer to the private address space.
-    S.Diag(Param->getLocation(), diag::err_opencl_private_ptr_kernel_param);
+  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);
     D.setInvalidType();
     return;
 
@@ -7736,7 +7739,7 @@ static void checkIsValidOpenCLKernelParameter(
       // do not allow OpenCL objects to be passed as elements of the struct or
       // union.
       if (ParamType == PtrKernelParam || ParamType == PtrPtrKernelParam ||
-          ParamType == PrivatePtrKernelParam) {
+          ParamType == InvalidAddrSpacePtrKernelParam) {
         S.Diag(Param->getLocation(),
                diag::err_record_with_pointers_kernel_param)
           << PT->isUnionType()

test/SemaOpenCL/invalid-kernel.cl

@@ -1,10 +1,11 @@
 // RUN: %clang_cc1 -verify %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -verify %s
 
-kernel void no_ptrptr(global int **i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}}
+kernel void no_ptrptr(global int * global *i) { } // expected-error{{kernel parameter cannot be declared as a pointer to a pointer}}
 
-__kernel void no_privateptr(__private int *i) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}}
+__kernel void no_privateptr(__private int *i) { } // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 
-__kernel void no_privatearray(__private int i[]) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}}
+__kernel void no_privatearray(__private int i[]) { } // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
 
 kernel int bar()  { // expected-error {{kernel must have void return type}}
   return 6;
@@ -29,3 +30,6 @@ int* local x(int* x) { // expected-error {{return value cannot be qualified with
 int* constant x(int* x) { // expected-error {{return value cannot be qualified with address space}}
   return x + 1;
 }
+
+__kernel void testKernel(int *ptr) { // expected-error {{pointer arguments to kernel functions must reside in '__global', '__constant' or '__local' address space}}
+}