[AMDGPU] Change constant addr space to 4 for clang

[t-tye]

LGTM. Other clean up can be done as a separate patch.

test/CodeGenOpenCL/address-spaces.cl

37

Suggest using the same name across all the OpenCL tests as some are using GIZ, some AMD and some AMDGCN. AMDGCN seems the clearer now that only have a single address space mapping? Similar comment for other places GIZ/AMD is being mentioned.

test/CodeGenOpenCL/amdgpu-env-amdgiz.cl

4

Should this test be renamed to amdgpu-env-amdgcn.cl now there is a single address space mapping (and delete any relating to the old mapping if they exist)?

[yaxunl]

test/CodeGenOpenCL/address-spaces.cl

37

Will have a separate patch to cleanup the tests.

test/CodeGenOpenCL/amdgpu-env-amdgiz.cl

4

will rename it when committing.

include/clang/Basic/BuiltinsAMDGPU.def

 
 #if defined(BUILTIN) && !defined(TARGET_BUILTIN)
 #   define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
 #endif
 //===----------------------------------------------------------------------===//
 // SI+ only builtins.
 //===----------------------------------------------------------------------===//
 
-BUILTIN(__builtin_amdgcn_dispatch_ptr, "Uc*2", "nc")
-BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*2", "nc")
-BUILTIN(__builtin_amdgcn_implicitarg_ptr, "Uc*2", "nc")
+BUILTIN(__builtin_amdgcn_dispatch_ptr, "Uc*4", "nc")
+BUILTIN(__builtin_amdgcn_kernarg_segment_ptr, "Uc*4", "nc")
+BUILTIN(__builtin_amdgcn_implicitarg_ptr, "Uc*4", "nc")
 
 BUILTIN(__builtin_amdgcn_workgroup_id_x, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workgroup_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workgroup_id_z, "Ui", "nc")
 
 BUILTIN(__builtin_amdgcn_workitem_id_x, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_y, "Ui", "nc")
 BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc")

lib/Basic/Targets/AMDGPU.cpp

static const char *const DataLayoutStringR600 =
     "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5";
 
 static const char *const DataLayoutStringSIPrivateIsZero =
     "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32"
     "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
     "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
 
 static const char *const DataLayoutStringSIGenericIsZero =
-    "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32"
+    "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32"
     "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
     "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5";
 
 static const LangASMap AMDGPUPrivIsZeroDefIsGenMap = {
     4, // Default
     1, // opencl_global
     3, // opencl_local
-    2, // opencl_constant
+    4, // opencl_constant
     0, // opencl_private
     4, // opencl_generic
     1, // cuda_device
-    2, // cuda_constant
+    4, // cuda_constant
     3  // cuda_shared
 };
 
 static const LangASMap AMDGPUGenIsZeroDefIsGenMap = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
-    2, // opencl_constant
+    4, // opencl_constant
     5, // opencl_private
     0, // opencl_generic
     1, // cuda_device
-    2, // cuda_constant
+    4, // cuda_constant
     3  // cuda_shared
 };
 
 static const LangASMap AMDGPUPrivIsZeroDefIsPrivMap = {
     0, // Default
     1, // opencl_global
     3, // opencl_local
-    2, // opencl_constant
+    4, // opencl_constant
     0, // opencl_private
     4, // opencl_generic
     1, // cuda_device
-    2, // cuda_constant
+    4, // cuda_constant
     3  // cuda_shared
 };
 
 static const LangASMap AMDGPUGenIsZeroDefIsPrivMap = {
     5, // Default
     1, // opencl_global
     3, // opencl_local
-    2, // opencl_constant
+    4, // opencl_constant
     5, // opencl_private
     0, // opencl_generic
     1, // cuda_device
-    2, // cuda_constant
+    4, // cuda_constant
     3  // cuda_shared
 };
 } // namespace targets
 } // namespace clang
 
 const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                                               \
   {#ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr},

test/CodeGen/target-data.c

 // R600: target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
 
 // RUN: %clang_cc1 -triple r600-unknown -target-cpu cayman -o - -emit-llvm %s \
 // RUN: | FileCheck %s -check-prefix=R600D
 // R600D: target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
 
 // RUN: %clang_cc1 -triple amdgcn-unknown -target-cpu hawaii -o - -emit-llvm %s \
 // RUN: | FileCheck %s -check-prefix=R600SI
-// R600SI: target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
+// R600SI: target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
 
 // Test default -target-cpu
 // RUN: %clang_cc1 -triple amdgcn-unknown -o - -emit-llvm %s \
 // RUN: | FileCheck %s -check-prefix=R600SIDefault
-// R600SIDefault: target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
+// R600SIDefault: target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
 
 // RUN: %clang_cc1 -triple arm64-unknown -o - -emit-llvm %s | \
 // RUN: FileCheck %s -check-prefix=AARCH64
 // AARCH64: target datalayout = "e-m:e-i8:8:32-i16:16:32-i64:64-i128:128-n32:64-S128"
 
 // RUN: %clang_cc1 -triple thumb-unknown-gnueabi -o - -emit-llvm %s | \
 // RUN: FileCheck %s -check-prefix=THUMB
 // THUMB: target datalayout = "e-m:e-p:32:32-i64:64-v128:64:128-a:0:32-n32-S64"

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

-// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-opencl -emit-llvm -o - | FileCheck %s
-// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa-amdgizcl -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck -check-prefix=FAKE %s
+// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck -check-prefix=AMD %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(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)*) }
+// FAKE: %struct.ConstantArrayPointerStruct = type { float addrspace(2)* }
+// FAKE: 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)*) }
+// AMD: %struct.ConstantArrayPointerStruct = type { float addrspace(4)* }
+// AMD: addrspace(4) constant %struct.ConstantArrayPointerStruct { float addrspace(4)* bitcast (i8 addrspace(4)* getelementptr (i8, i8 addrspace(4)* bitcast (%struct.ArrayStruct addrspace(4)* @constant_array_struct to i8 addrspace(4)*), i64 4) to float addrspace(4)*) }
 // Bug  18567
 __constant ConstantArrayPointerStruct constant_array_pointer_struct = {
     &constant_array_struct.f
 };
 

test/CodeGenOpenCL/address-spaces.cl

 // RUN: %clang_cc1 %s -O0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,SPIR
 // RUN: %clang_cc1 %s -O0 -DCL20 -cl-std=CL2.0 -ffake-address-space-map -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20SPIR
-// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-opencl -emit-llvm -o - | FileCheck --check-prefixes=CHECK,GIZ %s
-// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa-opencl -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20GIZ
+// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck --check-prefixes=CHECK,GIZ %s
+// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20GIZ
 // RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,GIZ %s
 // RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,GIZ %s
 
 // SPIR: %struct.S = type { i32, i32, i32* }
 // CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* }
 struct S {
   int x;
   int y;

 void f__p(__private int *arg) {}
 
 // CHECK: i32 addrspace(1)* %arg
 void f__g(__global int *arg) {}
 
 // CHECK: i32 addrspace(3)* %arg
 void f__l(__local int *arg) {}
 
-// CHECK: i32 addrspace(2)* %arg
+// SPIR: i32 addrspace(2)* %arg
+// GIZ: i32 addrspace(4)* %arg
 void f__c(__constant int *arg) {}
 
 // SPIR: i32* %arg
 // GIZ: i32 addrspace(5)* %arg
 void fp(private int *arg) {}
 
 // CHECK: i32 addrspace(1)* %arg
 void fg(global int *arg) {}
 
 // CHECK: i32 addrspace(3)* %arg
 void fl(local int *arg) {}
 
-// CHECK: i32 addrspace(2)* %arg
+// SPIR: i32 addrspace(2)* %arg
+// GIZ: i32 addrspace(4)* %arg
 void fc(constant int *arg) {}
 
 #ifdef CL20
 int i;
 // CL20-DAG: @i = common addrspace(1) global i32 0
 int *ptr;
 // CL20SPIR-DAG: @ptr = common addrspace(1) global i32 addrspace(4)* null
 // CL20GIZ-DAG: @ptr = common addrspace(1) global i32* null

test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl

 // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FILEVAR14]], expr: !DIExpression())
 int *constant FileVar14 = 0;
 
 kernel void kernel1(
     // CHECK-DAG: ![[KERNELARG0:[0-9]+]] = !DILocalVariable(name: "KernelArg0", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
     // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)* addrspace(5)* {{.*}}, metadata ![[KERNELARG0]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
     global int *KernelArg0,
     // CHECK-DAG: ![[KERNELARG1:[0-9]+]] = !DILocalVariable(name: "KernelArg1", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
-    // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)* addrspace(5)* {{.*}}, metadata ![[KERNELARG1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
+    // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(4)* addrspace(5)* {{.*}}, metadata ![[KERNELARG1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
     constant int *KernelArg1,
     // CHECK-DAG: ![[KERNELARG2:[0-9]+]] = !DILocalVariable(name: "KernelArg2", arg: {{[0-9]+}}, scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
     // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)* addrspace(5)* {{.*}}, metadata ![[KERNELARG2]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
     local int *KernelArg2) {
   private int *Tmp0;
   int *Tmp1;
 
   // CHECK-DAG: ![[FUNCVAR0:[0-9]+]] = !DILocalVariable(name: "FuncVar0", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR0]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   global int *FuncVar0 = KernelArg0;
   // CHECK-DAG: ![[FUNCVAR1:[0-9]+]] = !DILocalVariable(name: "FuncVar1", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
-  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
+  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(4)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR1]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   constant int *FuncVar1 = KernelArg1;
   // CHECK-DAG: ![[FUNCVAR2:[0-9]+]] = !DILocalVariable(name: "FuncVar2", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR2]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   local int *FuncVar2 = KernelArg2;
   // CHECK-DAG: ![[FUNCVAR3:[0-9]+]] = !DILocalVariable(name: "FuncVar3", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(5)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR3]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   private int *FuncVar3 = Tmp0;
   // CHECK-DAG: ![[FUNCVAR4:[0-9]+]] = !DILocalVariable(name: "FuncVar4", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
kernel void kernel1(
   // CHECK-DAG: ![[FUNCVAR14:[0-9]+]] = distinct !DIGlobalVariable(name: "FuncVar14", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}}, isLocal: true, isDefinition: true)
   // CHECK-DAG: !DIGlobalVariableExpression(var: ![[FUNCVAR14]], expr: !DIExpression(DW_OP_constu, 2, DW_OP_swap, DW_OP_xderef))
   int *local FuncVar14; FuncVar14 = Tmp1;
 
   // CHECK-DAG: ![[FUNCVAR15:[0-9]+]] = !DILocalVariable(name: "FuncVar15", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(1)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR15]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   global int *private FuncVar15 = KernelArg0;
   // CHECK-DAG: ![[FUNCVAR16:[0-9]+]] = !DILocalVariable(name: "FuncVar16", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
-  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(2)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR16]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
+  // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(4)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR16]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   constant int *private FuncVar16 = KernelArg1;
   // CHECK-DAG: ![[FUNCVAR17:[0-9]+]] = !DILocalVariable(name: "FuncVar17", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(3)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR17]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   local int *private FuncVar17 = KernelArg2;
   // CHECK-DAG: ![[FUNCVAR18:[0-9]+]] = !DILocalVariable(name: "FuncVar18", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32 addrspace(5)* addrspace(5)* {{.*}}, metadata ![[FUNCVAR18]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   private int *private FuncVar18 = Tmp0;
   // CHECK-DAG: ![[FUNCVAR19:[0-9]+]] = !DILocalVariable(name: "FuncVar19", scope: !{{[0-9]+}}, file: !{{[0-9]+}}, line: {{[0-9]+}}, type: !{{[0-9]+}})
   // CHECK-DAG: call void @llvm.dbg.declare(metadata i32* addrspace(5)* {{.*}}, metadata ![[FUNCVAR19]], metadata !DIExpression(DW_OP_constu, 1, DW_OP_swap, DW_OP_xderef)), !dbg !{{[0-9]+}}
   int *private FuncVar19 = Tmp1;
 }

test/CodeGenOpenCL/amdgpu-env-amdgiz.cl

 // RUN: %clang_cc1 %s -O0 -triple amdgcn -emit-llvm -o - | FileCheck %s
 // RUN: %clang_cc1 %s -O0 -triple amdgcn---opencl -emit-llvm -o - | FileCheck %s
 
-// CHECK: target datalayout = "e-p:64:64-p1:64:64-p2:64:64-p3:32:32-p4:32:32-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
+// CHECK: target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-A5"
 void foo(void) {}

test/CodeGenOpenCL/amdgpu-nullptr.cl

 private char *private_p = 0;
 
 // CHECK: @local_p = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 local char *local_p = 0;
 
 // CHECK: @global_p = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
 global char *global_p = 0;
 
-// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
+// CHECK: @constant_p = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
 constant char *constant_p = 0;
 
 // CHECK: @generic_p = local_unnamed_addr addrspace(1) global i8* null, align 8
 generic char *generic_p = 0;
 
 // Test NULL as initializer.
 
 // CHECK: @private_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(5)* null, align 4
 private char *private_p_NULL = NULL;
 
 // CHECK: @local_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 local char *local_p_NULL = NULL;
 
 // CHECK: @global_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
 global char *global_p_NULL = NULL;
 
-// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
+// CHECK: @constant_p_NULL = local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
 constant char *constant_p_NULL = NULL;
 
 // CHECK: @generic_p_NULL = local_unnamed_addr addrspace(1) global i8* null, align 8
 generic char *generic_p_NULL = NULL;
 
 // Test constant folding of null pointer.
 // A null pointer should be folded to a null pointer in the target address space.
 

 
 // Test static variable initialization.
 
 // NOOPT: @test_static_var_private.sp1 = internal addrspace(1) global i8 addrspace(5)* null, align 4
 // NOOPT: @test_static_var_private.sp2 = internal addrspace(1) global i8 addrspace(5)* null, align 4
 // NOOPT: @test_static_var_private.sp3 = internal addrspace(1) global i8 addrspace(5)* null, align 4
 // NOOPT: @test_static_var_private.sp4 = internal addrspace(1) global i8 addrspace(5)* null, align 4
 // NOOPT: @test_static_var_private.sp5 = internal addrspace(1) global i8 addrspace(5)* null, align 4
-// NOOPT: @test_static_var_private.SS1 = internal addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8* null }, align 8
+// NOOPT: @test_static_var_private.SS1 = internal addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }, align 8
 // NOOPT: @test_static_var_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
 
 void test_static_var_private(void) {
   static private char *sp1 = 0;
   static private char *sp2 = NULL;
   static private char *sp3;
   static private char *sp4 = (private char*)((void)0, 0);
   const int x = 0;
   static private char *sp5 = (private char*)x;
   static StructTy1 SS1;
   static StructTy2 SS2;
 }
 
 // NOOPT: @test_static_var_local.sp1 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 // NOOPT: @test_static_var_local.sp2 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 // NOOPT: @test_static_var_local.sp3 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 // NOOPT: @test_static_var_local.sp4 = internal addrspace(1) global i8 addrspace(3)* null, align 4
 // NOOPT: @test_static_var_local.sp5 = internal addrspace(1) global i8 addrspace(3)* null, align 4
-// NOOPT: @test_static_var_local.SS1 = internal addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8* null }, align 8
+// NOOPT: @test_static_var_local.SS1 = internal addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }, align 8
 // NOOPT: @test_static_var_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
 void test_static_var_local(void) {
   static local char *sp1 = 0;
   static local char *sp2 = NULL;
   static local char *sp3;
   static local char *sp4 = (local char*)((void)0, 0);
   const int x = 0;
   static local char *sp5 = (local char*)x;
   static StructTy1 SS1;
   static StructTy2 SS2;
 }
 
 // Test function-scope variable initialization.
 // NOOPT-LABEL: @test_func_scope_var_private(
 // NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp1, align 4
 // NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp2, align 4
 // NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp3, align 4
 // NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp4, align 4
 // NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1 addrspace(5)* %SS1 to i8 addrspace(5)*
-// NOOPT: call void @llvm.memcpy.p5i8.p2i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(2)* align 8 bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_private.SS1 to i8 addrspace(2)*), i64 32, i1 false)
+// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @test_func_scope_var_private.SS1 to i8 addrspace(4)*), i64 32, i1 false)
 // NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2 addrspace(5)* %SS2 to i8 addrspace(5)*
 // NOOPT: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 8 %[[SS2]], i8 0, i64 24, i1 false)
 void test_func_scope_var_private(void) {
   private char *sp1 = 0;
   private char *sp2 = NULL;
   private char *sp3 = (private char*)((void)0, 0);
   const int x = 0;
   private char *sp4 = (private char*)x;
   StructTy1 SS1 = {0, 0, 0, 0, 0};
   StructTy2 SS2 = {0, 0, 0};
 }
 
 // Test function-scope variable initialization.
 // NOOPT-LABEL: @test_func_scope_var_local(
 // NOOPT: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(5)* %sp1, align 4
 // NOOPT: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(5)* %sp2, align 4
 // NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)* addrspace(5)* %sp3, align 4
 // NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)* addrspace(5)* %sp4, align 4
 // NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1 addrspace(5)* %SS1 to i8 addrspace(5)*
-// NOOPT: call void @llvm.memcpy.p5i8.p2i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(2)* align 8 bitcast (%struct.StructTy1 addrspace(2)* @test_func_scope_var_local.SS1 to i8 addrspace(2)*), i64 32, i1 false)
+// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @test_func_scope_var_local.SS1 to i8 addrspace(4)*), i64 32, i1 false)
 // NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2 addrspace(5)* %SS2 to i8 addrspace(5)*
 // NOOPT: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 8 %[[SS2]], i8 0, i64 24, i1 false)
 void test_func_scope_var_local(void) {
   local char *sp1 = 0;
   local char *sp2 = NULL;
   local char *sp3 = (local char*)((void)0, 0);
   const int x = 0;
   local char *sp4 = (local char*)x;

 // and does not have explicit section.
 
 // CHECK: @p1 = common local_unnamed_addr addrspace(1) global i8 addrspace(5)* null, align 4
 private char *p1;
 
 // CHECK: @p2 = weak local_unnamed_addr addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 local char *p2;
 
-// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(2)* null, align 8
+// CHECK: @p3 = common local_unnamed_addr addrspace(1) global i8 addrspace(4)* null, align 8
 constant char *p3;
 
 // CHECK: @p4 = common local_unnamed_addr addrspace(1) global i8 addrspace(1)* null, align 8
 global char *p4;
 
 // CHECK: @p5 = common local_unnamed_addr addrspace(1) global i8* null, align 8
 generic char *p5;
 
 // Test default initialization of sturcture.
 
-// CHECK: @S1 = weak local_unnamed_addr addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8* null }, align 8
+// CHECK: @S1 = weak local_unnamed_addr addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }, align 8
 StructTy1 S1;
 
 // CHECK: @S2 = common local_unnamed_addr addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
 StructTy2 S2;
 
 // Test default initialization of array.
-// CHECK: @A1 = weak local_unnamed_addr addrspace(1) global [2 x %struct.StructTy1] [%struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8* null }, %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(2)* null, i8 addrspace(1)* null, i8* null }], align 8
+// CHECK: @A1 = weak local_unnamed_addr addrspace(1) global [2 x %struct.StructTy1] [%struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }, %struct.StructTy1 { i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }], align 8
 StructTy1 A1[2];
 
 // CHECK: @A2 = common local_unnamed_addr addrspace(1) global [2 x %struct.StructTy2] zeroinitializer, align 8
 StructTy2 A2[2];
 
 // Test comparison with 0.
 
 // CHECK-LABEL: cmp_private

 // CHECK-LABEL: cmp_global
 // CHECK: icmp eq i8 addrspace(1)* %p, null
 void cmp_global(global char* p) {
   if (p != 0)
     *p = 0;
 }
 
 // CHECK-LABEL: cmp_constant
-// CHECK: icmp eq i8 addrspace(2)* %p, null
+// CHECK: icmp eq i8 addrspace(4)* %p, null
 char cmp_constant(constant char* p) {
   if (p != 0)
     return *p;
   else
     return 0;
 }
 
 // CHECK-LABEL: cmp_generic

 // CHECK-LABEL: cmp_NULL_global
 // CHECK: icmp eq i8 addrspace(1)* %p, null
 void cmp_NULL_global(global char* p) {
   if (p != NULL)
     *p = 0;
 }
 
 // CHECK-LABEL: cmp_NULL_constant
-// CHECK: icmp eq i8 addrspace(2)* %p, null
+// CHECK: icmp eq i8 addrspace(4)* %p, null
 char cmp_NULL_constant(constant char* p) {
   if (p != NULL)
     return *p;
   else
     return 0;
 }
 
 // CHECK-LABEL: cmp_NULL_generic
 // CHECK: icmp eq i8* %p, null
 void cmp_NULL_generic(generic char* p) {
   if (p != NULL)
     *p = 0;
 }
 
 // Test storage 0 as null pointer.
 // CHECK-LABEL: test_storage_null_pointer
 // CHECK: store i8 addrspace(5)* null, i8 addrspace(5)** %arg_private
 // CHECK: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)** %arg_local
 // CHECK: store i8 addrspace(1)* null, i8 addrspace(1)** %arg_global
-// CHECK: store i8 addrspace(2)* null, i8 addrspace(2)** %arg_constant
+// CHECK: store i8 addrspace(4)* null, i8 addrspace(4)** %arg_constant
 // CHECK: store i8* null, i8** %arg_generic
 void test_storage_null_pointer(private char** arg_private,
                                local char** arg_local,
                                global char** arg_global,
                                constant char** arg_constant,
                                generic char** arg_generic) {
    *arg_private = 0;
    *arg_local = 0;
    *arg_global = 0;
    *arg_constant = 0;
    *arg_generic = 0;
 }
 
 // Test storage NULL as null pointer.
 // CHECK-LABEL: test_storage_null_pointer_NULL
 // CHECK: store i8 addrspace(5)* null, i8 addrspace(5)** %arg_private
 // CHECK: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)** %arg_local
 // CHECK: store i8 addrspace(1)* null, i8 addrspace(1)** %arg_global
-// CHECK: store i8 addrspace(2)* null, i8 addrspace(2)** %arg_constant
+// CHECK: store i8 addrspace(4)* null, i8 addrspace(4)** %arg_constant
 // CHECK: store i8* null, i8** %arg_generic
 void test_storage_null_pointer_NULL(private char** arg_private,
                                     local char** arg_local,
                                     global char** arg_global,
                                     constant char** arg_constant,
                                     generic char** arg_generic) {
    *arg_private = NULL;
    *arg_local = NULL;
    *arg_global = NULL;
    *arg_constant = NULL;
    *arg_generic = NULL;
 }
 
 // Test pass null pointer to function as argument.
 void test_pass_null_pointer_arg_calee(private char* arg_private,
                                       local char* arg_local,
                                       global char* arg_global,
                                       constant char* arg_constant,
                                       generic char* arg_generic);
 
 // CHECK-LABEL: test_pass_null_pointer_arg
-// CHECK: call void @test_pass_null_pointer_arg_calee(i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(1)* null, i8 addrspace(2)* null, i8* null)
-// CHECK: call void @test_pass_null_pointer_arg_calee(i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(1)* null, i8 addrspace(2)* null, i8* null)
+// CHECK: call void @test_pass_null_pointer_arg_calee(i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(1)* null, i8 addrspace(4)* null, i8* null)
+// CHECK: call void @test_pass_null_pointer_arg_calee(i8 addrspace(5)* null, i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(1)* null, i8 addrspace(4)* null, i8* null)
 void test_pass_null_pointer_arg(void) {
   test_pass_null_pointer_arg_calee(0, 0, 0, 0, 0);
   test_pass_null_pointer_arg_calee(NULL, NULL, NULL, NULL, NULL);
 }
 
 // Test cast null pointer to size_t.
 void test_cast_null_pointer_to_sizet_calee(size_t arg_private,
                                            size_t arg_local,

 // CHECK-LABEL: cast_bool_global
 // CHECK: icmp eq i8 addrspace(1)* %p, null
 void cast_bool_global(global char* p) {
   if (p)
     *p = 0;
 }
 
 // CHECK-LABEL: cast_bool_constant
-// CHECK: icmp eq i8 addrspace(2)* %p, null
+// CHECK: icmp eq i8 addrspace(4)* %p, null
 char cast_bool_constant(constant char* p) {
   if (p)
     return *p;
   else
     return 0;
 }
 
 // CHECK-LABEL: cast_bool_generic

test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_read_exec_hi(
 // CHECK: call i32 @llvm.read_register.i32(metadata ![[EXEC_HI:[0-9]+]]) #[[READ_EXEC_ATTRS]]
 void test_read_exec_hi(global uint* out) {
   *out = __builtin_amdgcn_read_exec_hi();
 }
 
 // CHECK-LABEL: @test_dispatch_ptr
-// CHECK: call i8 addrspace(2)* @llvm.amdgcn.dispatch.ptr()
-void test_dispatch_ptr(__attribute__((address_space(2))) unsigned char ** out)
+// CHECK: call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+void test_dispatch_ptr(__attribute__((address_space(4))) unsigned char ** out)
 {
   *out = __builtin_amdgcn_dispatch_ptr();
 }
 
 // CHECK-LABEL: @test_kernarg_segment_ptr
-// CHECK: call i8 addrspace(2)* @llvm.amdgcn.kernarg.segment.ptr()
-void test_kernarg_segment_ptr(__attribute__((address_space(2))) unsigned char ** out)
+// CHECK: call i8 addrspace(4)* @llvm.amdgcn.kernarg.segment.ptr()
+void test_kernarg_segment_ptr(__attribute__((address_space(4))) unsigned char ** out)
 {
   *out = __builtin_amdgcn_kernarg_segment_ptr();
 }
 
 // CHECK-LABEL: @test_implicitarg_ptr
-// CHECK: call i8 addrspace(2)* @llvm.amdgcn.implicitarg.ptr()
-void test_implicitarg_ptr(__attribute__((address_space(2))) unsigned char ** out)
+// CHECK: call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+void test_implicitarg_ptr(__attribute__((address_space(4))) unsigned char ** out)
 {
   *out = __builtin_amdgcn_implicitarg_ptr();
 }
 
 // CHECK-LABEL: @test_get_group_id(
 // CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
 // CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
 // CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()

test/CodeGenOpenCL/cast_image.cl

 // RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn--amdhsa %s | FileCheck --check-prefix=AMDGCN %s
 // RUN: %clang_cc1 -emit-llvm -o - -triple spir-unknown-unknown %s | FileCheck --check-prefix=SPIR %s
 
 #ifdef __AMDGCN__
 
 constant int* convert(image2d_t img) {
-  // AMDGCN: bitcast %opencl.image2d_ro_t addrspace(2)* %img to i32 addrspace(2)*
+  // AMDGCN: bitcast %opencl.image2d_ro_t addrspace(4)* %img to i32 addrspace(4)*
   return __builtin_astype(img, constant int*);
 }
 
 #else
 
 global int* convert(image2d_t img) {
   // SPIR: bitcast %opencl.image2d_ro_t addrspace(1)* %img to i32 addrspace(1)*
   return __builtin_astype(img, global int*);
 }
 
 #endif

test/CodeGenOpenCL/opencl_types.cl

 // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "spir-unknown-unknown" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-SPIR
 // RUN: %clang_cc1 -cl-std=CL2.0 %s -triple "amdgcn--amdhsa" -emit-llvm -o - -O0 | FileCheck %s --check-prefixes=CHECK-COM,CHECK-AMDGCN
 
 #define CLK_ADDRESS_CLAMP_TO_EDGE       2
 #define CLK_NORMALIZED_COORDS_TRUE      1
 #define CLK_FILTER_NEAREST              0x10
 #define CLK_FILTER_LINEAR               0x20
 
 constant sampler_t glb_smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_NEAREST;
 // CHECK-COM-NOT: constant i32
 
 void fnc1(image1d_t img) {}
 // CHECK-SPIR: @fnc1(%opencl.image1d_ro_t addrspace(1)*
-// CHECK-AMDGCN: @fnc1(%opencl.image1d_ro_t addrspace(2)*
+// CHECK-AMDGCN: @fnc1(%opencl.image1d_ro_t addrspace(4)*
 
 void fnc1arr(image1d_array_t img) {}
 // CHECK-SPIR: @fnc1arr(%opencl.image1d_array_ro_t addrspace(1)*
-// CHECK-AMDGCN: @fnc1arr(%opencl.image1d_array_ro_t addrspace(2)*
+// CHECK-AMDGCN: @fnc1arr(%opencl.image1d_array_ro_t addrspace(4)*
 
 void fnc1buff(image1d_buffer_t img) {}
 // CHECK-SPIR: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(1)*
-// CHECK-AMDGCN: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(2)*
+// CHECK-AMDGCN: @fnc1buff(%opencl.image1d_buffer_ro_t addrspace(4)*
 
 void fnc2(image2d_t img) {}
 // CHECK-SPIR: @fnc2(%opencl.image2d_ro_t addrspace(1)*
-// CHECK-AMDGCN: @fnc2(%opencl.image2d_ro_t addrspace(2)*
+// CHECK-AMDGCN: @fnc2(%opencl.image2d_ro_t addrspace(4)*
 
 void fnc2arr(image2d_array_t img) {}
 // CHECK-SPIR: @fnc2arr(%opencl.image2d_array_ro_t addrspace(1)*
-// CHECK-AMDGCN: @fnc2arr(%opencl.image2d_array_ro_t addrspace(2)*
+// CHECK-AMDGCN: @fnc2arr(%opencl.image2d_array_ro_t addrspace(4)*
 
 void fnc3(image3d_t img) {}
 // CHECK-SPIR: @fnc3(%opencl.image3d_ro_t addrspace(1)*
-// CHECK-AMDGCN: @fnc3(%opencl.image3d_ro_t addrspace(2)*
+// CHECK-AMDGCN: @fnc3(%opencl.image3d_ro_t addrspace(4)*
 
 void fnc4smp(sampler_t s) {}
 // CHECK-SPIR-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
-// CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
+// CHECK-AMDGCN-LABEL: define {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)*
 
 kernel void foo(image1d_t img) {
   sampler_t smp = CLK_ADDRESS_CLAMP_TO_EDGE|CLK_NORMALIZED_COORDS_TRUE|CLK_FILTER_LINEAR;
-  // CHECK-COM: alloca %opencl.sampler_t addrspace(2)*
+  // CHECK-SPIR: alloca %opencl.sampler_t addrspace(2)*
+  // CHECK-AMDGCN: alloca %opencl.sampler_t addrspace(4)*
   event_t evt;
   // CHECK-SPIR: alloca %opencl.event_t*
   // CHECK-AMDGCN: alloca %opencl.event_t addrspace(5)*
   clk_event_t clk_evt;
   // CHECK-SPIR: alloca %opencl.clk_event_t*
   // CHECK-AMDGCN: alloca %opencl.clk_event_t addrspace(1)*
   queue_t queue;
   // CHECK-SPIR: alloca %opencl.queue_t*
   // CHECK-AMDGCN: alloca %opencl.queue_t addrspace(1)*
   reserve_id_t rid;
   // CHECK-SPIR: alloca %opencl.reserve_id_t*
   // CHECK-AMDGCN: alloca %opencl.reserve_id_t addrspace(1)*
-  // CHECK-COM: store %opencl.sampler_t addrspace(2)*
+  // CHECK-SPIR: store %opencl.sampler_t addrspace(2)*
+  // CHECK-AMDGCN: store %opencl.sampler_t addrspace(4)*
   fnc4smp(smp);
-  // CHECK-COM: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
+  // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
+  // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)*
   fnc4smp(glb_smp);
-  // CHECK-COM: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
+  // CHECK-SPIR: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(2)*
+  // CHECK-AMDGCN: call {{.*}}void @fnc4smp(%opencl.sampler_t addrspace(4)*
 }
 
 kernel void foo_pipe(read_only pipe int p) {}
 // CHECK-SPIR: @foo_pipe(%opencl.pipe_t addrspace(1)* %p)
 // CHECK_AMDGCN: @foo_pipe(%opencl.pipe_t addrspace(1)* %p)
 
 void __attribute__((overloadable)) bad1(image1d_t b, image2d_t c, image2d_t d) {}
 // CHECK-SPIR-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}
-// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(%opencl.image1d_ro_t addrspace(2)*{{.*}}%opencl.image2d_ro_t addrspace(2)*{{.*}}%opencl.image2d_ro_t addrspace(2)*{{.*}})
+// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1@Z"}}(%opencl.image1d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}})

test/CodeGenOpenCL/size_t.cl

 //SZ64: define{{.*}} i64 @test_ptrtoint_global(i8 addrspace(1)* %x)
 //SZ64: ptrtoint i8 addrspace(1)* %{{.*}} to i64
 intptr_t test_ptrtoint_global(global char* x) {
   return (intptr_t)x;
 }
 
 //SZ32: define{{.*}} i32 @test_ptrtoint_constant(i8 addrspace(2)* %x)
 //SZ32: ptrtoint i8 addrspace(2)* %{{.*}} to i32
-//SZ64: define{{.*}} i64 @test_ptrtoint_constant(i8 addrspace(2)* %x)
-//SZ64: ptrtoint i8 addrspace(2)* %{{.*}} to i64
+//SZ64ONLY: define{{.*}} i64 @test_ptrtoint_constant(i8 addrspace(2)* %x)
+//SZ64ONLY: ptrtoint i8 addrspace(2)* %{{.*}} to i64
+//AMDGCN: define{{.*}} i64 @test_ptrtoint_constant(i8 addrspace(4)* %x)
+//AMDGCN: ptrtoint i8 addrspace(4)* %{{.*}} to i64
 uintptr_t test_ptrtoint_constant(constant char* x) {
   return (uintptr_t)x;
 }
 
 //SZ32: define{{.*}} i32 @test_ptrtoint_local(i8 addrspace(3)* %x)
 //SZ32: ptrtoint i8 addrspace(3)* %{{.*}} to i32
 //SZ64: define{{.*}} i64 @test_ptrtoint_local(i8 addrspace(3)* %x)
 //SZ64: ptrtoint i8 addrspace(3)* %{{.*}} to i64

test/CodeGenOpenCL/vla.cl

 // RUN: %clang_cc1 -emit-llvm -triple "spir-unknown-unknown" -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,SPIR %s
-// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa-opencl -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,SPIR %s
-// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa-amdgizcl -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,GIZ %s
+// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,AMD %s
 
 constant int sz0 = 5;
-// CHECK: @sz0 = addrspace(2) constant i32 5
+// SPIR: @sz0 = addrspace(2) constant i32 5
+// AMD: @sz0 = addrspace(4) constant i32 5
 const global int sz1 = 16;
 // CHECK: @sz1 = addrspace(1) constant i32 16
 const constant int sz2 = 8;
-// CHECK: @sz2 = addrspace(2) constant i32 8
+// SPIR: @sz2 = addrspace(2) constant i32 8
+// AMD: @sz2 = addrspace(4) constant i32 8
 // CHECK: @testvla.vla2 = internal addrspace(3) global [8 x i16] undef
 
 kernel void testvla()
 {
   int vla0[sz0];
 // SPIR: %vla0 = alloca [5 x i32]
 // SPIR-NOT: %vla0 = alloca [5 x i32]{{.*}}addrspace
 // GIZ: %vla0 = alloca [5 x i32]{{.*}}addrspace(5)
   char vla1[sz1];
 // SPIR: %vla1 = alloca [16 x i8]
 // SPIR-NOT: %vla1 = alloca [16 x i8]{{.*}}addrspace
 // GIZ: %vla1 = alloca [16 x i8]{{.*}}addrspace(5)
   local short vla2[sz2];
 }