r289536 - [OpenCL] Improve address space diagnostics.

[Egor Churaev via cfe-commits]

Author: echuraev
Date: Tue Dec 13 08:07:23 2016
New Revision: 289536

URL: http://llvm.org/viewvc/llvm-project?rev=289536&view=rev
Log:
[OpenCL] Improve address space diagnostics.

Reviewers: Anastasia

Subscribers: bader, yaxunl, cfe-commits

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

Modified:
    cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
    cfe/trunk/lib/Sema/SemaDecl.cpp
    cfe/trunk/test/SemaOpenCL/invalid-kernel.cl

Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=289536&r1=289535&r2=289536&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Tue Dec 13 08:07:23 2016
@@ -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<

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=289536&r1=289535&r2=289536&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Dec 13 08:07:23 2016
@@ -7586,7 +7586,7 @@ enum OpenCLParamType {
   ValidKernelParam,
   PtrPtrKernelParam,
   PtrKernelParam,
-  PrivatePtrKernelParam,
+  InvalidAddrSpacePtrKernelParam,
   InvalidKernelParam,
   RecordKernelParam
 };
@@ -7596,8 +7596,10 @@ static OpenCLParamType getOpenCLKernelPa
     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 checkIsValidOpenCLKernelPara
     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 checkIsValidOpenCLKernelPara
       // 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()

Modified: cfe/trunk/test/SemaOpenCL/invalid-kernel.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/invalid-kernel.cl?rev=289536&r1=289535&r2=289536&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/invalid-kernel.cl (original)
+++ cfe/trunk/test/SemaOpenCL/invalid-kernel.cl Tue Dec 13 08:07:23 2016
@@ -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
 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}}
+}