[OpenCL] Align fake address space map with the SPIR target maps.

[echuraev]

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.

[Anastasia]

I believe it's not that important what those numbers actually are, but having the consistency with SPIR might be better indeed at least even for documentation purposes!

@pekka.jaaskelainen, I think you are using x86 target in the POCL toolchain, do you envision any issues with this change?

[pekka.jaaskelainen]

In D28048#629569, @Anastasia wrote:

@pekka.jaaskelainen, I think you are using x86 target in the POCL toolchain, do you envision any issues with this change?

No, this should simplify things actually.

[Anastasia]

Excellent! Thanks!

lib/AST/ASTContext.cpp

 
 static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
                                              const LangOptions &LOpts) {
   if (LOpts.FakeAddressSpaceMap) {
     // The fake address space map must have a distinct entry for each
     // 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
       7  // cuda_shared
     };
     return &FakeAddrSpaceMap;
   } else {
     return &T.getAddressSpaceMap();

test/CodeGen/blocks-opencl.cl

 // RUN: %clang_cc1 -O0 %s -ffake-address-space-map -emit-llvm -o - -fblocks -triple x86_64-unknown-unknown | FileCheck %s
 // This used to crash due to trying to generate a bitcase from a cstring
 // in the constant address space to i8* in AS0.
 
 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()
 {
   float (^const X)(float) = ^(float x) {
     return x + 42.0f;
   };
   dummy(X);
 }
 

test/CodeGenOpenCL/address-space-constant-initializers.cl

 // RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
 
 typedef struct {
     int i;
     float f; // At non-zero offset.
 } ArrayStruct;
 
 __constant ArrayStruct constant_array_struct = { 0, 0.0f };
 
 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
 };
 

test/CodeGenOpenCL/address-spaces-mangling.cl

 void f(global int *arg) { }
 // ASMANG: @_Z1fPU3AS1i
 // NOASMANG: @_Z1fPU8CLglobali
 // OCL-20-DAG: @_Z1fPU3AS1i
 // OCL-12-DAG: @_Z1fPU3AS1i
 
 __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

test/CodeGenOpenCL/address-spaces.cl

 // RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s
 // RUN: %clang_cc1 %s -O0 -DCL20 -cl-std=CL2.0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefix=CL20
 
 // CHECK: i32* %arg
 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
 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
 int i;
 // CL20-DAG: @i = common addrspace(1) global i32 0
 int *ptr;
 // CL20-DAG: @ptr = common addrspace(1) global i32 addrspace(4)* null
 #endif

test/CodeGenOpenCL/cl20-device-side-enqueue.cl

kernel void device_side_enqueue(global int *a, global int *b, int i) {
   // COMMON: %event_wait_list = alloca %opencl.clk_event_t*
   clk_event_t event_wait_list;
   // COMMON: %event_wait_list2 = alloca [1 x %opencl.clk_event_t*]
   clk_event_t event_wait_list2[] = {clk_event};
 
   // 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,
                  ^(void) {
                    a[i] = b[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: [[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,
                  ^(void) {
                    a[i] = b[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;
                  },
                  256);
   char c;
   // 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: [[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;
                  },
                  c);
 
   // 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: [[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;
                  },
                  256);
 
   // 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: [[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: [[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;
                  },
                  c);
 
   long l;
   // 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: [[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;
                  },
                  l);
 
   // 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;
                  },
                  4294967296L);
 
 
   void (^const block_A)(void) = ^{
     return;
   };
   void (^const block_B)(local void *) = ^(local void *a) {
     return;
   };
 
   // COMMON: [[BL:%[0-9]+]] = load void ()*, void ()** %block_A
   // COMMON: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
   // COMMON: call i32 @__get_kernel_work_group_size_impl(i8* [[BL_I8]])
   unsigned size = get_kernel_work_group_size(block_A);
-  // COMMON: [[BL:%[0-9]+]] = load void (i8 addrspace(2)*)*, void (i8 addrspace(2)*)** %block_B
-  // COMMON: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(2)*)* [[BL]] to i8*
+  // COMMON: [[BL:%[0-9]+]] = load void (i8 addrspace(3)*)*, void (i8 addrspace(3)*)** %block_B
+  // COMMON: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(3)*)* [[BL]] to i8*
   // COMMON: call i32 @__get_kernel_work_group_size_impl(i8* [[BL_I8]])
   size = get_kernel_work_group_size(block_B);
   // COMMON: [[BL:%[0-9]+]] = load void ()*, void ()** %block_A
   // COMMON: [[BL_I8:%[0-9]+]] = bitcast void ()* [[BL]] to i8*
   // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* [[BL_I8]])
   size = get_kernel_preferred_work_group_size_multiple(block_A);
-  // COMMON: [[BL:%[0-9]+]] = load void (i8 addrspace(2)*)*, void (i8 addrspace(2)*)* addrspace(1)* @block_G
-  // COMMON: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(2)*)* [[BL]] to i8*
+  // COMMON: [[BL:%[0-9]+]] = load void (i8 addrspace(3)*)*, void (i8 addrspace(3)*)* addrspace(1)* @block_G
+  // COMMON: [[BL_I8:%[0-9]+]] = bitcast void (i8 addrspace(3)*)* [[BL]] to i8*
   // COMMON: call i32 @__get_kernel_preferred_work_group_multiple_impl(i8* [[BL_I8]])
   size = get_kernel_preferred_work_group_size_multiple(block_G);
 }

test/CodeGenOpenCL/const-str-array-decay.cl

 // RUN: %clang_cc1 %s -emit-llvm -o - -ffake-address-space-map | FileCheck %s
 
 int test_func(constant char* foo);
 
 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
 

test/CodeGenOpenCL/constant-addr-space-globals.cl

 // RUN: %clang_cc1 %s -triple "spir64-unknown-unknown" -cl-opt-disable -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];
 }
 
 // Test that we don't use directly initializers for const aggregates
 // but create a copy in the original address space (unless a variable itself is
 // 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};
   // CHECK: %arr2 = alloca [3 x i32]
   const int arr2[] = {4, 5, 6};
   // CHECK: %arr3 = alloca [3 x i32]
   int arr3[] = {7, 8, 9};
 
   foo(arr1, arr2, arr3);
 }

test/CodeGenOpenCL/local-initializer-undef.cl

 // RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s
 
 typedef struct Foo {
     int x;
     float y;
     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;
     __local int lds_int_arr[128];
     __local Foo lds_struct;
     __local Foo lds_struct_arr[64];
 
     lds_int = 1;
     lds_int_arr[0] = 1;
     lds_struct.x = 1;
     lds_struct_arr[0].x = 1;
 }

test/CodeGenOpenCL/local.cl

 // RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck %s
 
 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);
 }
 
 // CHECK-LABEL: define {{.*}}void @_Z3barPU7CLlocali
 __kernel void __attribute__((__overloadable__)) bar(local int *x) {
   *x = 5;
 }

test/CodeGenOpenCL/memcpy.cl

 // RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
 
 // 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);
 }

test/CodeGenOpenCL/str_literals.cl

 // RUN: %clang_cc1 %s -cl-opt-disable -emit-llvm -o - -ffake-address-space-map | FileCheck %s
 
 __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)*

test/SemaOpenCL/extern.cl

 // 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) {
   buf[0] += foo;
 }