Index: lib/AST/ASTContext.cpp =================================================================== --- lib/AST/ASTContext.cpp +++ lib/AST/ASTContext.cpp @@ -704,8 +704,8 @@ // 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 Index: test/CodeGen/blocks-opencl.cl =================================================================== --- test/CodeGen/blocks-opencl.cl +++ 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() { Index: test/CodeGenOpenCL/address-space-constant-initializers.cl =================================================================== --- test/CodeGenOpenCL/address-space-constant-initializers.cl +++ test/CodeGenOpenCL/address-space-constant-initializers.cl @@ -11,8 +11,8 @@ __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 Index: test/CodeGenOpenCL/address-spaces-mangling.cl =================================================================== --- test/CodeGenOpenCL/address-spaces-mangling.cl +++ test/CodeGenOpenCL/address-spaces-mangling.cl @@ -31,14 +31,14 @@ __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 Index: test/CodeGenOpenCL/address-spaces.cl =================================================================== --- test/CodeGenOpenCL/address-spaces.cl +++ test/CodeGenOpenCL/address-spaces.cl @@ -7,10 +7,10 @@ // 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 @@ // 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 Index: test/CodeGenOpenCL/cl20-device-side-enqueue.cl =================================================================== --- test/CodeGenOpenCL/cl20-device-side-enqueue.cl +++ test/CodeGenOpenCL/cl20-device-side-enqueue.cl @@ -22,7 +22,7 @@ // 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, @@ -35,7 +35,7 @@ // 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, @@ -46,8 +46,8 @@ // 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; @@ -58,9 +58,9 @@ // 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; @@ -73,8 +73,8 @@ // 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; @@ -88,9 +88,9 @@ // 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; @@ -102,9 +102,9 @@ // 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; @@ -114,8 +114,8 @@ // 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; @@ -134,16 +134,16 @@ // 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); } Index: test/CodeGenOpenCL/const-str-array-decay.cl =================================================================== --- test/CodeGenOpenCL/const-str-array-decay.cl +++ test/CodeGenOpenCL/const-str-array-decay.cl @@ -6,6 +6,6 @@ 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 Index: test/CodeGenOpenCL/constant-addr-space-globals.cl =================================================================== --- test/CodeGenOpenCL/constant-addr-space-globals.cl +++ test/CodeGenOpenCL/constant-addr-space-globals.cl @@ -12,9 +12,9 @@ // 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}; Index: test/CodeGenOpenCL/local-initializer-undef.cl =================================================================== --- test/CodeGenOpenCL/local-initializer-undef.cl +++ test/CodeGenOpenCL/local-initializer-undef.cl @@ -6,10 +6,10 @@ 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; Index: test/CodeGenOpenCL/local.cl =================================================================== --- test/CodeGenOpenCL/local.cl +++ 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); } Index: test/CodeGenOpenCL/memcpy.cl =================================================================== --- test/CodeGenOpenCL/memcpy.cl +++ 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); } Index: test/CodeGenOpenCL/str_literals.cl =================================================================== --- test/CodeGenOpenCL/str_literals.cl +++ 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)* Index: test/SemaOpenCL/extern.cl =================================================================== --- test/SemaOpenCL/extern.cl +++ 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) {