On Fri, Oct 03, 2014 at 11:35:00AM -0700, Tom Stellard wrote: > On Fri, Oct 03, 2014 at 11:23:24AM -0700, Tom Stellard wrote: > > On Wed, Oct 01, 2014 at 10:30:19AM +0300, Pekka Jääskeläinen wrote: > > > Hi Tom, > > > > > > This patch seemed familiar to me, and by googling I found this > > > patch for the constant string literals: > > > > > > http://reviews.llvm.org/D894 > > > > > > "I updated this patch to remove the 'const' part as I now believe it > > > is wrong, and I changed IsModifiable as Tanya suggested.", writes > > > Joey, > > > referring to this comment from Tanya: > > > http://lists.cs.uiuc.edu/pipermail/cfe-commits/Week-of-Mon-20130603/081103.html > > > > > > The comment might apply to this patch as well. Perhaps adding a > > > test for const __constant is a good idea? > > > > > > As you didn't change this test: > > > > > > > +++ test/SemaOpenCL/address-spaces.cl > > > > @@ -9,5 +9,5 @@ > > > > int *ip; > > > > 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}} > > > > + ip = &ci; // expected-error {{assigning 'const __constant int > > > *' to 'int *' changes address space of pointer}} > > > > > > I assume the other part of the comment (confusing error message) > > > is not a problem anymore. > > > > Hi Pekka, > > > > > > From what I can tell this patch is slightly different from the > > other one, because instead of actually adding the const modifier > > to __constant address space types, it updates the query that determines > > whether or not a type is constant. So there shouldn't be any issues > > with mixing const and __constant. > > > > Here is an updated patch with a test case for const __constant, > > though it's not quite the same as the one you posted above. > > In the thread Eli mentioned that const __constant was illegal in OpenCL > > C, I couldn't find where in the spec where it says this. It would > > be good to get clarification on this, but I think that fix is unrelated > > and maybe the new test case is too. > > > > What do you think? > > > > Forgot the patch...
Hi Pekka, I understand now why you recommended the const __constant test cases. I have added these to the patch. I also removed an unnecessary include I added to Type.cpp -Tom
>From 4e4596def13a92d3565285e377d52e86a084bab2 Mon Sep 17 00:00:00 2001 From: Tom Stellard <[email protected]> Date: Tue, 30 Sep 2014 20:41:14 -0400 Subject: [PATCH] OpenCL: Emit global variables in the constant addr space as constant globals --- 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 35676da..bfd726c 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 ae8cedc..079b070 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 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 7e99fc5..ed2bf6d 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 43c90f8..092b637 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 b188ea4..7026d1b 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 ee5e072..764e724 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 %s -o - -verify | FileCheck %s // expected-no-diagnostics -// CHECK: @foo = external global float +// CHECK: @foo = external constant float extern constant float foo; kernel void test(global float* buf) { -- 1.8.5.5
_______________________________________________ cfe-commits mailing list [email protected] http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
