r325031 - [AMDGPU] Change constant addr space to 4

Yaxun Liu via cfe-commits cfe-commits at lists.llvm.org
Tue Feb 13 10:01:22 PST 2018


Author: yaxunl
Date: Tue Feb 13 10:01:21 2018
New Revision: 325031

URL: http://llvm.org/viewvc/llvm-project?rev=325031&view=rev
Log:
[AMDGPU] Change constant addr space to 4

Differential Revision: https://reviews.llvm.org/D43171

Added:
    cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl
Removed:
    cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl
Modified:
    cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
    cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
    cfe/trunk/test/CodeGen/target-data.c
    cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl
    cfe/trunk/test/CodeGenOpenCL/address-spaces.cl
    cfe/trunk/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl
    cfe/trunk/test/CodeGenOpenCL/amdgpu-nullptr.cl
    cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
    cfe/trunk/test/CodeGenOpenCL/cast_image.cl
    cfe/trunk/test/CodeGenOpenCL/opencl_types.cl
    cfe/trunk/test/CodeGenOpenCL/private-array-initialization.cl
    cfe/trunk/test/CodeGenOpenCL/size_t.cl
    cfe/trunk/test/CodeGenOpenCL/vla.cl

Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Tue Feb 13 10:01:21 2018
@@ -21,9 +21,9 @@
 // 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")

Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.cpp?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets/AMDGPU.cpp (original)
+++ cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Tue Feb 13 10:01:21 2018
@@ -38,7 +38,7 @@ static const char *const DataLayoutStrin
     "-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-p6:32:32"
+    "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6: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";
 
@@ -46,11 +46,11 @@ static const LangASMap AMDGPUPrivIsZeroD
     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
 };
 
@@ -58,11 +58,11 @@ static const LangASMap AMDGPUGenIsZeroDe
     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
 };
 
@@ -70,11 +70,11 @@ static const LangASMap AMDGPUPrivIsZeroD
     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
 };
 
@@ -82,11 +82,11 @@ static const LangASMap AMDGPUGenIsZeroDe
     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

Modified: cfe/trunk/test/CodeGen/target-data.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/target-data.c?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/target-data.c (original)
+++ cfe/trunk/test/CodeGen/target-data.c Tue Feb 13 10:01:21 2018
@@ -132,12 +132,12 @@
 
 // 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-p6: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-p6: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-p6: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-p6: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

Modified: cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl Tue Feb 13 10:01:21 2018
@@ -1,6 +1,5 @@
-// 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;
@@ -13,8 +12,10 @@ 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

Modified: cfe/trunk/test/CodeGenOpenCL/address-spaces.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-spaces.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-spaces.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/address-spaces.cl Tue Feb 13 10:01:21 2018
@@ -1,7 +1,7 @@
 // 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
 
@@ -33,7 +33,8 @@ 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
@@ -46,7 +47,8 @@ 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

Modified: cfe/trunk/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-debug-info-variable-expression.cl Tue Feb 13 10:01:21 2018
@@ -54,7 +54,7 @@ kernel void kernel1(
     // 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]+}}
@@ -66,7 +66,7 @@ kernel void kernel1(
   // 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]+}}
@@ -114,7 +114,7 @@ kernel void kernel1(
   // 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]+}}

Added: cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl?rev=325031&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgcn.cl Tue Feb 13 10:01:21 2018
@@ -0,0 +1,5 @@
+// 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:32:32-p3:32:32-p4:64:64-p5:32:32-p6: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) {}

Removed: cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl?rev=325030&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-env-amdgiz.cl (removed)
@@ -1,5 +0,0 @@
-// 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-p6: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) {}

Modified: cfe/trunk/test/CodeGenOpenCL/amdgpu-nullptr.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-nullptr.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-nullptr.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-nullptr.cl Tue Feb 13 10:01:21 2018
@@ -30,7 +30,7 @@ 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
@@ -47,7 +47,7 @@ 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
@@ -104,7 +104,7 @@ int fold_int5_local = (int) &((local Str
 // 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) {
@@ -123,7 +123,7 @@ void test_static_var_private(void) {
 // 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;
@@ -143,7 +143,7 @@ void test_static_var_local(void) {
 // 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) {
@@ -163,7 +163,7 @@ void test_func_scope_var_private(void) {
 // 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) {
@@ -189,7 +189,7 @@ 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
@@ -200,14 +200,14 @@ 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
@@ -237,7 +237,7 @@ void cmp_global(global char* p) {
 }
 
 // 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;
@@ -276,7 +276,7 @@ void cmp_NULL_global(global char* p) {
 }
 
 // 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;
@@ -296,7 +296,7 @@ void cmp_NULL_generic(generic char* p) {
 // 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,
@@ -315,7 +315,7 @@ void test_storage_null_pointer(private c
 // 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,
@@ -337,8 +337,8 @@ void test_pass_null_pointer_arg_calee(pr
                                       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);
@@ -486,7 +486,7 @@ void cast_bool_global(global char* p) {
 }
 
 // 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;

Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Tue Feb 13 10:01:21 2018
@@ -434,22 +434,22 @@ void test_read_exec_hi(global uint* out)
 }
 
 // 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();
 }

Modified: cfe/trunk/test/CodeGenOpenCL/cast_image.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/cast_image.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/cast_image.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/cast_image.cl Tue Feb 13 10:01:21 2018
@@ -4,7 +4,7 @@
 #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*);
 }
 

Modified: cfe/trunk/test/CodeGenOpenCL/opencl_types.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/opencl_types.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/opencl_types.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/opencl_types.cl Tue Feb 13 10:01:21 2018
@@ -11,35 +11,36 @@ constant sampler_t glb_smp = CLK_ADDRESS
 
 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)*
@@ -52,11 +53,14 @@ kernel void foo(image1d_t img) {
   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) {}
@@ -65,4 +69,4 @@ kernel void foo_pipe(read_only pipe int
 
 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 at Z"}}
-// CHECK-AMDGCN-LABEL: @{{_Z4bad114ocl_image1d_ro14ocl_image2d_roS0_|"\\01\?bad1@@\$\$J0YAXPAUocl_image1d_ro@@PAUocl_image2d_ro@@1 at 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 at Z"}}(%opencl.image1d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}}%opencl.image2d_ro_t addrspace(4)*{{.*}})

Modified: cfe/trunk/test/CodeGenOpenCL/private-array-initialization.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/private-array-initialization.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/private-array-initialization.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/private-array-initialization.cl Tue Feb 13 10:01:21 2018
@@ -10,7 +10,7 @@ void test() {
 
 // PRIVATE5: %arr = alloca [3 x i32], align 4, addrspace(5)
 // PRIVATE5: %0 = bitcast [3 x i32] addrspace(5)* %arr to i8 addrspace(5)*
-// PRIVATE5: call void @llvm.memcpy.p5i8.p2i8.i64(i8 addrspace(5)* align 4 %0, i8 addrspace(2)* align 4 bitcast ([3 x i32] addrspace(2)* @test.arr to i8 addrspace(2)*), i64 12, i1 false)
+// PRIVATE5: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 4 %0, i8 addrspace(4)* align 4 bitcast ([3 x i32] addrspace(4)* @test.arr to i8 addrspace(4)*), i64 12, i1 false)
 }
 
 __kernel void initializer_cast_is_valid_crash() {

Modified: cfe/trunk/test/CodeGenOpenCL/size_t.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/size_t.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/size_t.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/size_t.cl Tue Feb 13 10:01:21 2018
@@ -23,8 +23,10 @@ intptr_t test_ptrtoint_global(global cha
 
 //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;
 }

Modified: cfe/trunk/test/CodeGenOpenCL/vla.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/vla.cl?rev=325031&r1=325030&r2=325031&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/vla.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/vla.cl Tue Feb 13 10:01:21 2018
@@ -1,13 +1,14 @@
 // 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()




More information about the cfe-commits mailing list