From 520113df8784cf844b14b26f6d8271ee5d465ad5 Mon Sep 17 00:00:00 2001 From: Tom Stellard Date: Thu, 16 Oct 2014 15:29:19 +0000 Subject: [PATCH] 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 --- lib/AST/Type.cpp | 2 +- test/CodeGenOpenCL/address-space-constant-initializers.cl | 2 +- test/CodeGenOpenCL/constant-addr-space-globals.cl | 8 ++++++++ test/CodeGenOpenCL/opencl_types.cl | 2 +- test/CodeGenOpenCL/str_literals.cl | 4 ++-- test/SemaOpenCL/address-spaces.cl | 5 ++++- test/SemaOpenCL/extern.cl | 2 +- 7 files changed, 18 insertions(+), 7 deletions(-) create mode 100644 test/CodeGenOpenCL/constant-addr-space-globals.cl diff --git a/lib/AST/Type.cpp b/lib/AST/Type.cpp index 35676da641..bfd726c165 100644 --- 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 ae8cedc1ca..079b0706f4 100644 --- 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 index 0000000000..92fb9790b5 --- /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 7e99fc548e..ed2bf6dd11 100644 --- 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 43c90f83f6..092b6372a4 100644 --- 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 b188ea4576..7026d1bc08 100644 --- 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 5b88d70793..b2e4857e28 100644 --- 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) { -- 2.40.0