[OpenCL] Add LangAS::opencl_private to represent private address space in AST

Currently Clang uses default address space (0) to represent private address space for OpenCL
in AST. There are two issues with this:

Multiple address spaces including private address space cannot be diagnosed.
There is no mangling for default address space. For example, if private int* is emitted as
i32 addrspace(5)* in IR. It is supposed to be mangled as PUAS5i but it is mangled as
Pi instead.

This patch attempts to represent OpenCL private address space explicitly in AST. It adds
a new enum LangAS::opencl_private and adds it to the variable types which are implicitly
private:

automatic variables without address space qualifier

function parameter

pointee type without address space qualifier (OpenCL 1.2 and below)

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


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

5 files changed, 77 insertions, 22 deletions

diff --git a/test/SemaOpenCL/address-spaces.cl b/test/SemaOpenCL/address-spaces.cl
index 7026d1bc08..97dd1f4f90 100644
--- a/test/SemaOpenCL/address-spaces.cl
+++ b/test/SemaOpenCL/address-spaces.cl
@@ -1,4 +1,5 @@
 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 
 __constant int ci = 1;
 
@@ -7,9 +8,15 @@ __kernel void foo(__global int *gip) {
   __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
 
   int *ip;
+#if __OPENCL_C_VERSION__ < 200
   ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
   ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
   ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
+#else
+  ip = gip;
+  ip = &li;
+  ip = &ci; // expected-error {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
+#endif
 }
 
 void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
@@ -40,3 +47,19 @@ void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l
   l = (local int*) l2;
   p = (private int*) p2;
 }
+
+__private int func_return_priv(void);       //expected-error {{return value cannot be qualified with address space}}
+__global int func_return_global(void);      //expected-error {{return value cannot be qualified with address space}}
+__local int func_return_local(void);        //expected-error {{return value cannot be qualified with address space}}
+__constant int func_return_constant(void);  //expected-error {{return value cannot be qualified with address space}}
+#if __OPENCL_C_VERSION__ >= 200
+__generic int func_return_generic(void);    //expected-error {{return value cannot be qualified with address space}}
+#endif
+
+void func_multiple_addr(void) {
+  typedef __private int private_int_t;
+  __local __private int var1;   // expected-error {{multiple address spaces specified for type}}
+  __local __private int *var2;  // expected-error {{multiple address spaces specified for type}}
+  __local private_int_t var3;   // expected-error {{multiple address spaces specified for type}}
+  __local private_int_t *var4;  // expected-error {{multiple address spaces specified for type}}
+}
diff --git a/test/SemaOpenCL/cl20-device-side-enqueue.cl b/test/SemaOpenCL/cl20-device-side-enqueue.cl
index b4714dc290..207fe7d340 100644
--- a/test/SemaOpenCL/cl20-device-side-enqueue.cl
+++ b/test/SemaOpenCL/cl20-device-side-enqueue.cl
@@ -222,7 +222,7 @@ kernel void foo(global int *buf)
 
 kernel void bar(global int *buf)
 {
-  ndrange_t n;
+  __private ndrange_t n;
   buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){});
   buf[0] = get_kernel_sub_group_count_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected 'ndrange_t' argument type}}
   buf[0] = get_kernel_sub_group_count_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected block argument type}}
diff --git a/test/SemaOpenCL/extern.cl b/test/SemaOpenCL/extern.cl
deleted file mode 100644
index 66efd70058..0000000000
--- a/test/SemaOpenCL/extern.cl
+++ /dev/null
@@ -1,9 +0,0 @@
-// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
-// expected-no-diagnostics
-
-// CHECK: @foo = external addrspace(2) constant float
-extern constant float foo;
-
-kernel void test(global float* buf) {
-  buf[0] += foo;
-}
diff --git a/test/SemaOpenCL/storageclass-cl20.cl b/test/SemaOpenCL/storageclass-cl20.cl
index b12676fe74..581701d2a6 100644
--- a/test/SemaOpenCL/storageclass-cl20.cl
+++ b/test/SemaOpenCL/storageclass-cl20.cl
@@ -1,21 +1,41 @@
 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
 
-static constant int G1 = 0;
 int G2 = 0;
 global int G3 = 0;
 local int G4 = 0;              // expected-error{{program scope variable must reside in global or constant address space}}
 
-void kernel foo() {
-  static int S1 = 5;
-  static global int S2 = 5;
-  static private int S3 = 5;   // expected-error{{static local variable must reside in global or constant address space}}
+static float g_implicit_static_var = 0;
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;
+static local float g_local_static_var = 0;     // expected-error {{program scope variable must reside in global or constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+static generic float g_generic_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+
+extern float g_implicit_extern_var;
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;
+extern local float g_local_extern_var;     // expected-error {{extern variable must reside in global or constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+extern generic float g_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
 
+void kernel foo() {
   constant int L1 = 0;
   local int L2;
   global int L3;                              // expected-error{{function scope variable cannot be declared in global address space}}
   generic int L4;                             // expected-error{{automatic variable qualified with an invalid address space}}
   __attribute__((address_space(100))) int L5; // expected-error{{automatic variable qualified with an invalid address space}}
 
-  extern global int G5;
-  extern int G6; // expected-error{{extern variable must reside in global or constant address space}}
+  static float l_implicit_static_var = 0;
+  static constant float l_constant_static_var = 0;
+  static global float l_global_static_var = 0;
+  static local float l_local_static_var = 0;     // expected-error {{static local variable must reside in global or constant address space}}
+  static private float l_private_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+  static generic float l_generic_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+
+  extern float l_implicit_extern_var;
+  extern constant float l_constant_extern_var;
+  extern global float l_global_extern_var;
+  extern local float l_local_extern_var;     // expected-error {{extern variable must reside in global or constant address space}}
+  extern private float l_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+  extern generic float l_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
 }
diff --git a/test/SemaOpenCL/storageclass.cl b/test/SemaOpenCL/storageclass.cl
index 9a461068f2..cbcb94509b 100644
--- a/test/SemaOpenCL/storageclass.cl
+++ b/test/SemaOpenCL/storageclass.cl
@@ -5,6 +5,20 @@ constant int G2 = 0;
 int G3 = 0;        // expected-error{{program scope variable must reside in constant address space}}
 global int G4 = 0; // expected-error{{program scope variable must reside in constant address space}}
 
+static float g_implicit_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;   // expected-error {{program scope variable must reside in constant address space}}
+static local float g_local_static_var = 0;     // expected-error {{program scope variable must reside in constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static generic float g_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{program scope variable must reside in constant address space}}
+
+extern float g_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;   // expected-error {{extern variable must reside in constant address space}}
+extern local float g_local_extern_var;     // expected-error {{extern variable must reside in constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern generic float g_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
+
 void kernel foo(int x) {
   // static is not allowed at local scope before CL2.0
   static int S1 = 5;          // expected-error{{variables in function scope cannot be declared static}}
@@ -45,10 +59,17 @@ void f() {
     __attribute__((address_space(100))) int L4; // expected-error{{automatic variable qualified with an invalid address space}}
   }
 
+  static float l_implicit_static_var = 0;          // expected-error {{variables in function scope cannot be declared static}}
+  static constant float l_constant_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+  static global float l_global_static_var = 0;     // expected-error {{variables in function scope cannot be declared static}}
+  static local float l_local_static_var = 0;       // expected-error {{variables in function scope cannot be declared static}}
+  static private float l_private_static_var = 0;   // expected-error {{variables in function scope cannot be declared static}}
+  static generic float l_generic_static_var = 0;   // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{variables in function scope cannot be declared static}}
 
-  extern constant float L5;
-  extern local float L6; // expected-error{{extern variable must reside in constant address space}}
-
-  static int L7 = 0;     // expected-error{{variables in function scope cannot be declared static}}
-  static int L8;         // expected-error{{variables in function scope cannot be declared static}}
+  extern float l_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+  extern constant float l_constant_extern_var;
+  extern global float l_global_extern_var;   // expected-error {{extern variable must reside in constant address space}}
+  extern local float l_local_extern_var;     // expected-error {{extern variable must reside in constant address space}}
+  extern private float l_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+  extern generic float l_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
 }