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...
>From 3c6a668a2cfc5c5d8e092b508d56edcb58b4dcf7 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 | 3 ++- 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 | 1 + test/SemaOpenCL/extern.cl | 2 +- 7 files changed, 16 insertions(+), 6 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..2b98eb6 100644 --- a/lib/AST/Type.cpp +++ b/lib/AST/Type.cpp @@ -26,6 +26,7 @@ #include "llvm/ADT/StringExtras.h" #include "llvm/Support/raw_ostream.h" #include <algorithm> +#include "llvm/Support/Debug.h" using namespace clang; bool Qualifiers::isStrictSupersetOf(Qualifiers Other) const { @@ -70,7 +71,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..9b81c36 100644 --- a/test/SemaOpenCL/address-spaces.cl +++ b/test/SemaOpenCL/address-spaces.cl @@ -1,6 +1,7 @@ // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only __constant int ci = 1; +const __constant int cci = 1; __kernel void foo(__global int *gip) { __local int li; 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
