r219929 - OpenCL: Emit global variables in the constant addr space as constant globals
Tom Stellard
thomas.stellard at amd.com
Thu Oct 16 08:29:19 PDT 2014
Author: tstellar
Date: Thu Oct 16 10:29:19 2014
New Revision: 219929
URL: http://llvm.org/viewvc/llvm-project?rev=219929&view=rev
Log:
OpenCL: Emit global variables in the constant addr space as constant globals
Added:
cfe/trunk/test/CodeGenOpenCL/constant-addr-space-globals.cl
Modified:
cfe/trunk/lib/AST/Type.cpp
cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl
cfe/trunk/test/CodeGenOpenCL/opencl_types.cl
cfe/trunk/test/CodeGenOpenCL/str_literals.cl
cfe/trunk/test/SemaOpenCL/address-spaces.cl
cfe/trunk/test/SemaOpenCL/extern.cl
Modified: cfe/trunk/lib/AST/Type.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Type.cpp?rev=219929&r1=219928&r2=219929&view=diff
==============================================================================
--- cfe/trunk/lib/AST/Type.cpp (original)
+++ cfe/trunk/lib/AST/Type.cpp Thu Oct 16 10:29:19 2014
@@ -70,7 +70,7 @@ bool QualType::isConstant(QualType T, AS
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,
Modified: cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl?rev=219929&r1=219928&r2=219929&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl Thu Oct 16 10:29:19 2014
@@ -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
Added: cfe/trunk/test/CodeGenOpenCL/constant-addr-space-globals.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/constant-addr-space-globals.cl?rev=219929&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/constant-addr-space-globals.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/constant-addr-space-globals.cl Thu Oct 16 10:29:19 2014
@@ -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];
+}
Modified: cfe/trunk/test/CodeGenOpenCL/opencl_types.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/opencl_types.cl?rev=219929&r1=219928&r2=219929&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/opencl_types.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/opencl_types.cl Thu Oct 16 10:29:19 2014
@@ -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*
Modified: cfe/trunk/test/CodeGenOpenCL/str_literals.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/str_literals.cl?rev=219929&r1=219928&r2=219929&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/str_literals.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/str_literals.cl Thu Oct 16 10:29:19 2014
@@ -5,5 +5,5 @@ __constant char * __constant y = "hello
// 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)*
Modified: cfe/trunk/test/SemaOpenCL/address-spaces.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/address-spaces.cl?rev=219929&r1=219928&r2=219929&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/address-spaces.cl (original)
+++ cfe/trunk/test/SemaOpenCL/address-spaces.cl Thu Oct 16 10:29:19 2014
@@ -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
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)
Modified: cfe/trunk/test/SemaOpenCL/extern.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/extern.cl?rev=219929&r1=219928&r2=219929&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/extern.cl (original)
+++ cfe/trunk/test/SemaOpenCL/extern.cl Thu Oct 16 10:29:19 2014
@@ -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) {
More information about the cfe-commits
mailing list