OpenCL: Emit global variables in the constant addr space as constant globals

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

diff --git a/lib/AST/Type.cpp b/lib/AST/Type.cpp
index 35676da641da026abe9e1377aeb18dcad2430326..bfd726c165e529b8425749ad2f153e215715442f 100644 (file)
--- a/lib/AST/Type.cpp
+++ b/lib/AST/Type.cpp
@@ -70,7 +70,7 @@ bool QualType::isConstant(QualType T, ASTContext &Ctx) {
   if (const ArrayType *AT = Ctx.getAsArrayType(T))
     return AT->getElementType().isConstant(Ctx);
 
-  return false;
+  return T.getAddressSpace() == LangAS::opencl_constant;
 }
 
 unsigned ConstantArrayType::getNumAddressingBits(ASTContext &Context,

diff --git a/test/CodeGenOpenCL/address-space-constant-initializers.cl b/test/CodeGenOpenCL/address-space-constant-initializers.cl
index ae8cedc1ca5b01863bd6983facf0bf559804e64f..079b0706f4982bd90c05fc6c36e45408e365c0ca 100644 (file)
--- a/test/CodeGenOpenCL/address-space-constant-initializers.cl
+++ b/test/CodeGenOpenCL/address-space-constant-initializers.cl
@@ -12,7 +12,7 @@ typedef struct {
 } ConstantArrayPointerStruct;
 
 // CHECK: %struct.ConstantArrayPointerStruct = type { float addrspace(3)* }
-// CHECK: addrspace(3) global %struct.ConstantArrayPointerStruct { float addrspace(3)* bitcast (i8 addrspace(3)* getelementptr (i8 addrspace(3)* bitcast (%struct.ArrayStruct addrspace(3)* @constant_array_struct to i8 addrspace(3)*), i64 4) to float addrspace(3)*) }
+// CHECK: addrspace(3) constant %struct.ConstantArrayPointerStruct { float addrspace(3)* bitcast (i8 addrspace(3)* getelementptr (i8 addrspace(3)* bitcast (%struct.ArrayStruct addrspace(3)* @constant_array_struct to i8 addrspace(3)*), i64 4) to float addrspace(3)*) }
 // Bug  18567
 __constant ConstantArrayPointerStruct constant_array_pointer_struct = {
     &constant_array_struct.f

diff --git a/test/CodeGenOpenCL/constant-addr-space-globals.cl b/test/CodeGenOpenCL/constant-addr-space-globals.cl
new file mode 100644 (file)
index 0000000..92fb979
--- /dev/null
+++ b/test/CodeGenOpenCL/constant-addr-space-globals.cl
@@ -0,0 +1,8 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+// CHECK: @array = addrspace({{[0-9]+}}) constant
+__constant float array[2] = {0.0f, 1.0f};
+
+kernel void test(global float *out) {
+  *out = array[0];
+}

diff --git a/test/CodeGenOpenCL/opencl_types.cl b/test/CodeGenOpenCL/opencl_types.cl
index 7e99fc548ec53d992c51e8521d670e511b5ed8ff..ed2bf6dd1145e5f7772f348b4f034457818d5a9d 100644 (file)
--- a/test/CodeGenOpenCL/opencl_types.cl
+++ b/test/CodeGenOpenCL/opencl_types.cl
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 %s -emit-llvm -o - -O0 | FileCheck %s
 
 constant sampler_t glb_smp = 7;
-// CHECK: global i32 7
+// CHECK: constant i32 7
 
 void fnc1(image1d_t img) {}
 // CHECK: @fnc1(%opencl.image1d_t*

diff --git a/test/CodeGenOpenCL/str_literals.cl b/test/CodeGenOpenCL/str_literals.cl
index 43c90f83f6e7aa9cff0baf1341057deb7c3f5e65..092b6372a414dc6e325f1944f8f0f6bd2b11d03e 100644 (file)
--- a/test/CodeGenOpenCL/str_literals.cl
+++ b/test/CodeGenOpenCL/str_literals.cl
@@ -5,5 +5,5 @@ __constant char * __constant y = "hello world";
 
 // CHECK: unnamed_addr addrspace(3) constant
 // CHECK-NOT: addrspace(3) unnamed_addr constant
-// CHECK: @x = addrspace(3) global i8 addrspace(3)*
-// CHECK: @y = addrspace(3) global i8 addrspace(3)*
+// CHECK: @x = addrspace(3) constant i8 addrspace(3)*
+// CHECK: @y = addrspace(3) constant i8 addrspace(3)*

diff --git a/test/SemaOpenCL/address-spaces.cl b/test/SemaOpenCL/address-spaces.cl
index b188ea4576f1f3e08af6fe635360b532237c34b8..7026d1bc086d6ac248ab5007f39211b3c239b03b 100644 (file)
--- a/test/SemaOpenCL/address-spaces.cl
+++ b/test/SemaOpenCL/address-spaces.cl
@@ -12,14 +12,16 @@ __kernel void foo(__global int *gip) {
   ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
 }
 
-void explicit_cast(global int* g, local int* l, constant int* c, private int* p)
+void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
 {
   g = (global int*) l;    // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
   g = (global int*) c;    // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
+  g = (global int*) cc;   // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
   g = (global int*) p;    // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
 
   l = (local int*) g;     // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
   l = (local int*) c;     // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
+  l = (local int*) cc;    // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
   l = (local int*) p;     // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
 
   c = (constant int*) g;  // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
@@ -29,6 +31,7 @@ void explicit_cast(global int* g, local int* l, constant int* c, private int* p)
   p = (private int*) g;   // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
   p = (private int*) l;   // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
   p = (private int*) c;   // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
+  p = (private int*) cc;  // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
 }
 
 void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l2, private int* p, private int* p2)

diff --git a/test/SemaOpenCL/extern.cl b/test/SemaOpenCL/extern.cl
index 5b88d70793f32cde8534f5772a2407ee734d878e..b2e4857e28daaf22acc25bf669e7a2267751b549 100644 (file)
--- a/test/SemaOpenCL/extern.cl
+++ b/test/SemaOpenCL/extern.cl
@@ -1,7 +1,7 @@
 // RUN: %clang_cc1 -x cl -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
 // expected-no-diagnostics
 
-// CHECK: @foo = external addrspace(3) global float
+// CHECK: @foo = external addrspace(3) constant float
 extern constant float foo;
 
 kernel void test(global float* buf) {