summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorEgor Churaev <egor.churaev@gmail.com>2016-12-23 16:11:25 +0000
committerEgor Churaev <egor.churaev@gmail.com>2016-12-23 16:11:25 +0000
commit28f00aab73730b8f74894fce1e82b8545c729ffb (patch)
tree045cfd5ebb73689af4b0767849c3933f73028d8f
parent206a510e5406fcab172b15a11cea64b65c953823 (diff)
downloadbcm5719-llvm-28f00aab73730b8f74894fce1e82b8545c729ffb.tar.gz
bcm5719-llvm-28f00aab73730b8f74894fce1e82b8545c729ffb.zip
[OpenCL] Align fake address space map with the SPIR target maps.
Summary: We compile user opencl kernel code with spir triple. But built-ins are written in OpenCL and we compile it with triple x86_64 to be able to use x86 intrinsics. And we need address spaces to match in both cases. So, we change fake address space map in OpenCL for matching with spir. On CPU address spaces are not really important but we'd like to preserve address space information in order to perform optimizations relying on this info like enhanced alias analysis. Reviewers: pekka.jaaskelainen, Anastasia Subscribers: pekka.jaaskelainen, yaxunl, bader, cfe-commits Differential Revision: https://reviews.llvm.org/D28048 llvm-svn: 290436
-rw-r--r--clang/lib/AST/ASTContext.cpp4
-rw-r--r--clang/test/CodeGen/blocks-opencl.cl2
-rw-r--r--clang/test/CodeGenOpenCL/address-space-constant-initializers.cl4
-rw-r--r--clang/test/CodeGenOpenCL/address-spaces-mangling.cl12
-rw-r--r--clang/test/CodeGenOpenCL/address-spaces.cl8
-rw-r--r--clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl30
-rw-r--r--clang/test/CodeGenOpenCL/const-str-array-decay.cl2
-rw-r--r--clang/test/CodeGenOpenCL/constant-addr-space-globals.cl6
-rw-r--r--clang/test/CodeGenOpenCL/local-initializer-undef.cl8
-rw-r--r--clang/test/CodeGenOpenCL/local.cl2
-rw-r--r--clang/test/CodeGenOpenCL/memcpy.cl2
-rw-r--r--clang/test/CodeGenOpenCL/str_literals.cl8
-rw-r--r--clang/test/SemaOpenCL/extern.cl2
13 files changed, 45 insertions, 45 deletions
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 9a25c6d6605..1b5988d0198 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -704,8 +704,8 @@ static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
// language-specific address space.
static const unsigned FakeAddrSpaceMap[] = {
1, // opencl_global
- 2, // opencl_local
- 3, // opencl_constant
+ 3, // opencl_local
+ 2, // opencl_constant
4, // opencl_generic
5, // cuda_device
6, // cuda_constant
diff --git a/clang/test/CodeGen/blocks-opencl.cl b/clang/test/CodeGen/blocks-opencl.cl
index 61c479b7b98..12513331e9d 100644
--- a/clang/test/CodeGen/blocks-opencl.cl
+++ b/clang/test/CodeGen/blocks-opencl.cl
@@ -5,7 +5,7 @@
void dummy(float (^const op)(float)) {
}
-// CHECK: i8 addrspace(3)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(3)* @.str, i32 0, i32 0)
+// CHECK: i8 addrspace(2)* getelementptr inbounds ([9 x i8], [9 x i8] addrspace(2)* @.str, i32 0, i32 0)
kernel void test_block()
{
diff --git a/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl b/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl
index 44c4ddf8780..45c64c2f716 100644
--- a/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl
+++ b/clang/test/CodeGenOpenCL/address-space-constant-initializers.cl
@@ -11,8 +11,8 @@ typedef struct {
__constant float* constant_float_ptr;
} ConstantArrayPointerStruct;
-// CHECK: %struct.ConstantArrayPointerStruct = type { float addrspace(3)* }
-// CHECK: addrspace(3) constant %struct.ConstantArrayPointerStruct { float addrspace(3)* bitcast (i8 addrspace(3)* getelementptr (i8, i8 addrspace(3)* bitcast (%struct.ArrayStruct addrspace(3)* @constant_array_struct to i8 addrspace(3)*), i64 4) to float addrspace(3)*) }
+// CHECK: %struct.ConstantArrayPointerStruct = type { float addrspace(2)* }
+// CHECK: addrspace(2) constant %struct.ConstantArrayPointerStruct { float addrspace(2)* bitcast (i8 addrspace(2)* getelementptr (i8, i8 addrspace(2)* bitcast (%struct.ArrayStruct addrspace(2)* @constant_array_struct to i8 addrspace(2)*), i64 4) to float addrspace(2)*) }
// Bug 18567
__constant ConstantArrayPointerStruct constant_array_pointer_struct = {
&constant_array_struct.f
diff --git a/clang/test/CodeGenOpenCL/address-spaces-mangling.cl b/clang/test/CodeGenOpenCL/address-spaces-mangling.cl
index 57fd1329642..3c74c718c2a 100644
--- a/clang/test/CodeGenOpenCL/address-spaces-mangling.cl
+++ b/clang/test/CodeGenOpenCL/address-spaces-mangling.cl
@@ -31,14 +31,14 @@ void f(global int *arg) { }
__attribute__((overloadable))
void f(local int *arg) { }
-// ASMANG: @_Z1fPU3AS2i
+// ASMANG: @_Z1fPU3AS3i
// NOASMANG: @_Z1fPU7CLlocali
-// OCL-20-DAG: @_Z1fPU3AS2i
-// OCL-12-DAG: @_Z1fPU3AS2i
+// OCL-20-DAG: @_Z1fPU3AS3i
+// OCL-12-DAG: @_Z1fPU3AS3i
__attribute__((overloadable))
void f(constant int *arg) { }
-// ASMANG: @_Z1fPU3AS3i
+// ASMANG: @_Z1fPU3AS2i
// NOASMANG: @_Z1fPU10CLconstanti
-// OCL-20-DAG: @_Z1fPU3AS3i
-// OCL-12-DAG: @_Z1fPU3AS3i
+// OCL-20-DAG: @_Z1fPU3AS2i
+// OCL-12-DAG: @_Z1fPU3AS2i
diff --git a/clang/test/CodeGenOpenCL/address-spaces.cl b/clang/test/CodeGenOpenCL/address-spaces.cl
index 68fa02dadb2..3b34986d893 100644
--- a/clang/test/CodeGenOpenCL/address-spaces.cl
+++ b/clang/test/CodeGenOpenCL/address-spaces.cl
@@ -7,10 +7,10 @@ void f__p(__private int *arg) {}
// CHECK: i32 addrspace(1)* %arg
void f__g(__global int *arg) {}
-// CHECK: i32 addrspace(2)* %arg
+// CHECK: i32 addrspace(3)* %arg
void f__l(__local int *arg) {}
-// CHECK: i32 addrspace(3)* %arg
+// CHECK: i32 addrspace(2)* %arg
void f__c(__constant int *arg) {}
// CHECK: i32* %arg
@@ -19,10 +19,10 @@ void fp(private int *arg) {}
// CHECK: i32 addrspace(1)* %arg
void fg(global int *arg) {}
-// CHECK: i32 addrspace(2)* %arg
+// CHECK: i32 addrspace(3)* %arg
void fl(local int *arg) {}
-// CHECK: i32 addrspace(3)* %arg
+// CHECK: i32 addrspace(2)* %arg
void fc(constant int *arg) {}
#ifdef CL20
diff --git a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
index 9986f579957..add13b9e3fc 100644
--- a/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
+++ b/clang/test/CodeGenOpenCL/cl20-device-side-enqueue.cl
@@ -24,7 +24,7 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
- // COMMON: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(3)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block to void ()*
+ // COMMON: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(2)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block to void ()*
// COMMON: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
// COMMON: call i32 @__enqueue_kernel_basic(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange,
@@ -37,7 +37,7 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
// COMMON: [[WAIT_EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %event_wait_list to %opencl.clk_event_t{{.*}}* addrspace(4)*
// COMMON: [[EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)*
- // COMMON: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(3)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block3 to void ()*
+ // COMMON: [[BL:%[0-9]+]] = bitcast <{ i8*, i32, i32, i8*, %struct.__block_descriptor addrspace(2)*, i32{{.*}}, i32{{.*}}, i32{{.*}} }>* %block3 to void ()*
// COMMON: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
// COMMON: call i32 @__enqueue_kernel_basic_events(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* [[BL_I8]])
enqueue_kernel(default_queue, flags, ndrange, 2, &event_wait_list, &clk_event,
@@ -48,8 +48,8 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t{{.*}}*, %opencl.queue_t{{.*}}** %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
- // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256)
- // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 256)
+ // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256)
+ // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 256)
enqueue_kernel(default_queue, flags, ndrange,
^(local void *p) {
return;
@@ -60,9 +60,9 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
// B32: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32
- // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
+ // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
// B64: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i64
- // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]])
+ // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]])
enqueue_kernel(default_queue, flags, ndrange,
^(local void *p) {
return;
@@ -75,8 +75,8 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[AD:%arraydecay[0-9]*]] = getelementptr inbounds [1 x %opencl.clk_event_t*], [1 x %opencl.clk_event_t*]* %event_wait_list2, i32 0, i32 0
// COMMON: [[WAIT_EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** [[AD]] to %opencl.clk_event_t{{.*}}* addrspace(4)*
// COMMON: [[EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)*
- // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256)
- // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 256)
+ // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 256)
+ // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}} [[WAIT_EVNT]], %opencl.clk_event_t{{.*}} [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 256)
enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
^(local void *p) {
return;
@@ -90,9 +90,9 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[WAIT_EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** [[AD]] to %opencl.clk_event_t{{.*}}* addrspace(4)*
// COMMON: [[EVNT:%[0-9]+]] = addrspacecast %opencl.clk_event_t{{.*}}** %clk_event to %opencl.clk_event_t{{.*}}* addrspace(4)*
// B32: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i32
- // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
+ // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
// B64: [[SIZE:%[0-9]+]] = zext i8 {{%[0-9]+}} to i64
- // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]])
+ // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i32, %opencl.clk_event_t{{.*}}* addrspace(4)*, %opencl.clk_event_t{{.*}}* addrspace(4)*, i8*, i32, ...) @__enqueue_kernel_events_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* {{.*}}, i32 2, %opencl.clk_event_t{{.*}}* addrspace(4)* [[WAIT_EVNT]], %opencl.clk_event_t{{.*}}* addrspace(4)* [[EVNT]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]])
enqueue_kernel(default_queue, flags, ndrange, 2, event_wait_list2, &clk_event,
^(local void *p) {
return;
@@ -104,9 +104,9 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
// B32: [[SIZE:%[0-9]+]] = trunc i64 {{%[0-9]+}} to i32
- // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
+ // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 [[SIZE]])
// B64: [[SIZE:%[0-9]+]] = load i64, i64* %l
- // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]])
+ // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 [[SIZE]])
enqueue_kernel(default_queue, flags, ndrange,
^(local void *p) {
return;
@@ -116,8 +116,8 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
// COMMON: [[DEF_Q:%[0-9]+]] = load %opencl.queue_t*, %opencl.queue_t** %default_queue
// COMMON: [[FLAGS:%[0-9]+]] = load i32, i32* %flags
// COMMON: [[NDR:%[0-9]+]] = load %opencl.ndrange_t*, %opencl.ndrange_t** %ndrange
- // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 0)
- // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(3)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 4294967296)
+ // B32: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i32 0)
+ // B64: call i32 (%opencl.queue_t{{.*}}*, i32, %opencl.ndrange_t*, i8*, i32, ...) @__enqueue_kernel_vaargs(%opencl.queue_t{{.*}}* [[DEF_Q]], i32 [[FLAGS]], %opencl.ndrange_t* [[NDR]], i8* bitcast ({ i8**, i32, i32, i8*, %struct.__block_descriptor addrspace(2)* }* @__block_literal_global{{(.[0-9]+)?}} to i8*), i32 1, i64 4294967296)
enqueue_kernel(default_queue, flags, ndrange,
^(local void *p) {
return;
@@ -131,7 +131,7 @@ kernel void device_side_enqueue(global int *a, global int *b, int i) {
return;
};
- // COMMON: store void (i8 addrspace(2)*)* bitcast ([[BL_B:[^@]+@__block_literal_global.[0-9]+]] to void (i8 addrspace(2)*)*), void (i8 addrspace(2)*)** %block_B
+ // COMMON: store void (i8 addrspace(3)*)* bitcast ([[BL_B:[^@]+@__block_literal_global.[0-9]+]] to void (i8 addrspace(3)*)*), void (i8 addrspace(3)*)** %block_B
void (^const block_B)(local void *) = ^(local void *a) {
return;
};
diff --git a/clang/test/CodeGenOpenCL/const-str-array-decay.cl b/clang/test/CodeGenOpenCL/const-str-array-decay.cl
index af8c35542f3..353aa3a9a08 100644
--- a/clang/test/CodeGenOpenCL/const-str-array-decay.cl
+++ b/clang/test/CodeGenOpenCL/const-str-array-decay.cl
@@ -6,6 +6,6 @@ kernel void str_array_decy() {
test_func("Test string literal");
}
-// CHECK: i8 addrspace(3)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(3)*
+// CHECK: i8 addrspace(2)* getelementptr inbounds ([20 x i8], [20 x i8] addrspace(2)*
// CHECK-NOT: addrspacecast
diff --git a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
index 4b889133240..4f0d1ea23e5 100644
--- a/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
+++ b/clang/test/CodeGenOpenCL/constant-addr-space-globals.cl
@@ -12,9 +12,9 @@ kernel void test(global float *out) {
// in the constant address space).
void foo(constant const int *p1, const int *p2, const int *p3);
-// CHECK: @k.arr1 = internal addrspace(3) constant [3 x i32] [i32 1, i32 2, i32 3]
-// CHECK: @k.arr2 = private unnamed_addr addrspace(3) constant [3 x i32] [i32 4, i32 5, i32 6]
-// CHECK: @k.arr3 = private unnamed_addr addrspace(3) constant [3 x i32] [i32 7, i32 8, i32 9]
+// CHECK: @k.arr1 = internal addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3]
+// CHECK: @k.arr2 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 4, i32 5, i32 6]
+// CHECK: @k.arr3 = private unnamed_addr addrspace(2) constant [3 x i32] [i32 7, i32 8, i32 9]
kernel void k(void) {
// CHECK-NOT: %arr1 = alloca [3 x i32]
constant const int arr1[] = {1, 2, 3};
diff --git a/clang/test/CodeGenOpenCL/local-initializer-undef.cl b/clang/test/CodeGenOpenCL/local-initializer-undef.cl
index 5d34f56b821..f209b1f6b44 100644
--- a/clang/test/CodeGenOpenCL/local-initializer-undef.cl
+++ b/clang/test/CodeGenOpenCL/local-initializer-undef.cl
@@ -6,10 +6,10 @@ typedef struct Foo {
float z;
} Foo;
-// CHECK-DAG: @test.lds_int = internal addrspace(2) global i32 undef
-// CHECK-DAG: @test.lds_int_arr = internal addrspace(2) global [128 x i32] undef
-// CHECK-DAG: @test.lds_struct = internal addrspace(2) global %struct.Foo undef
-// CHECK-DAG: @test.lds_struct_arr = internal addrspace(2) global [64 x %struct.Foo] undef
+// CHECK-DAG: @test.lds_int = internal addrspace(3) global i32 undef
+// CHECK-DAG: @test.lds_int_arr = internal addrspace(3) global [128 x i32] undef
+// CHECK-DAG: @test.lds_struct = internal addrspace(3) global %struct.Foo undef
+// CHECK-DAG: @test.lds_struct_arr = internal addrspace(3) global [64 x %struct.Foo] undef
__kernel void test()
{
__local int lds_int;
diff --git a/clang/test/CodeGenOpenCL/local.cl b/clang/test/CodeGenOpenCL/local.cl
index da371f8e464..6f44b684e94 100644
--- a/clang/test/CodeGenOpenCL/local.cl
+++ b/clang/test/CodeGenOpenCL/local.cl
@@ -3,7 +3,7 @@
void func(local int*);
__kernel void foo(void) {
- // CHECK: @foo.i = internal addrspace(2) global i32 undef
+ // CHECK: @foo.i = internal addrspace(3) global i32 undef
__local int i;
func(&i);
}
diff --git a/clang/test/CodeGenOpenCL/memcpy.cl b/clang/test/CodeGenOpenCL/memcpy.cl
index cb27803adf4..5911b5c5aa6 100644
--- a/clang/test/CodeGenOpenCL/memcpy.cl
+++ b/clang/test/CodeGenOpenCL/memcpy.cl
@@ -2,7 +2,7 @@
// CHECK-LABEL: @test
// CHECK-NOT: addrspacecast
-// CHECK: call void @llvm.memcpy.p1i8.p3i8
+// CHECK: call void @llvm.memcpy.p1i8.p2i8
kernel void test(global float *g, constant float *c) {
__builtin_memcpy(g, c, 32);
}
diff --git a/clang/test/CodeGenOpenCL/str_literals.cl b/clang/test/CodeGenOpenCL/str_literals.cl
index 1c0acd155da..8fe5cefc8d7 100644
--- a/clang/test/CodeGenOpenCL/str_literals.cl
+++ b/clang/test/CodeGenOpenCL/str_literals.cl
@@ -3,7 +3,7 @@
__constant char * __constant x = "hello world";
__constant char * __constant y = "hello world";
-// CHECK: unnamed_addr addrspace(3) constant
-// CHECK-NOT: addrspace(3) unnamed_addr constant
-// CHECK: @x = addrspace(3) constant i8 addrspace(3)*
-// CHECK: @y = addrspace(3) constant i8 addrspace(3)*
+// CHECK: unnamed_addr addrspace(2) constant
+// CHECK-NOT: addrspace(2) unnamed_addr constant
+// CHECK: @x = addrspace(2) constant i8 addrspace(2)*
+// CHECK: @y = addrspace(2) constant i8 addrspace(2)*
diff --git a/clang/test/SemaOpenCL/extern.cl b/clang/test/SemaOpenCL/extern.cl
index 5f1f9f80b0f..66efd700583 100644
--- a/clang/test/SemaOpenCL/extern.cl
+++ b/clang/test/SemaOpenCL/extern.cl
@@ -1,7 +1,7 @@
// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
// expected-no-diagnostics
-// CHECK: @foo = external addrspace(3) constant float
+// CHECK: @foo = external addrspace(2) constant float
extern constant float foo;
kernel void test(global float* buf) {
OpenPOWER on IntegriCloud