Clean up AMDGCN tests

[yaxunl]

Summary

Remove amdgiz and use AMDGCN as check prefix.

[t-tye]

LGTM Plus I think there was a test with giz in its name. Should that be renamed?

[yaxunl]

In D43340#1008980, @t-tye wrote:

LGTM Plus I think there was a test with giz in its name. Should that be renamed?

That's test/CodeGenOpenCL/amdgpu-env-amdgiz.cl. I already renamed it in my last commit.

[b-sumner]

Looks good to me.

test/CodeGen/address-space.c

-// RUN: %clang_cc1 -triple x86_64-apple-darwin -emit-llvm < %s | FileCheck -check-prefixes=CHECK,X86,GIZ %s
-// RUN: %clang_cc1 -triple amdgcn---amdgiz -emit-llvm < %s | FileCheck -check-prefixes=CHECK,AMDGIZ,GIZ %s
+// RUN: %clang_cc1 -triple x86_64-apple-darwin -emit-llvm < %s | FileCheck -check-prefixes=CHECK,X86 %s
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm < %s | FileCheck -check-prefixes=CHECK,AMDGCN %s
 
 // CHECK: @foo = common addrspace(1) global
 int foo __attribute__((address_space(1)));
 
 // CHECK: @ban = common addrspace(1) global
 int ban[10] __attribute__((address_space(1)));
 
 // CHECK: @a = common global
 int a __attribute__((address_space(0)));
 
 // CHECK-LABEL: define i32 @test1() 
 // CHECK: load i32, i32 addrspace(1)* @foo
 int test1() { return foo; }
 
 // CHECK-LABEL: define i32 @test2(i32 %i) 
 // CHECK: load i32, i32 addrspace(1)*
 // CHECK-NEXT: ret i32
 int test2(int i) { return ban[i]; }
 
 // Both A and B point into addrspace(2).
 __attribute__((address_space(2))) int *A, *B;
 
 // CHECK-LABEL: define void @test3()
 // X86: load i32 addrspace(2)*, i32 addrspace(2)** @B
-// AMDGIZ: load i32 addrspace(2)*, i32 addrspace(2)** addrspacecast (i32 addrspace(2)* addrspace(1)* @B to i32 addrspace(2)**)
+// AMDGCN: load i32 addrspace(2)*, i32 addrspace(2)** addrspacecast (i32 addrspace(2)* addrspace(1)* @B to i32 addrspace(2)**)
 // CHECK: load i32, i32 addrspace(2)*
 // X86: load i32 addrspace(2)*, i32 addrspace(2)** @A
-// AMDGIZ: load i32 addrspace(2)*, i32 addrspace(2)** addrspacecast (i32 addrspace(2)* addrspace(1)* @A to i32 addrspace(2)**)
+// AMDGCN: load i32 addrspace(2)*, i32 addrspace(2)** addrspacecast (i32 addrspace(2)* addrspace(1)* @A to i32 addrspace(2)**)
 // CHECK: store i32 {{.*}}, i32 addrspace(2)*
 void test3() {
   *A = *B;
 }
 
 // PR7437
 typedef struct {
   float aData[1];
 } MyStruct;
 
 // CHECK-LABEL: define void @test4(
-// GIZ: call void @llvm.memcpy.p0i8.p2i8
-// GIZ: call void @llvm.memcpy.p2i8.p0i8
+// CHECK: call void @llvm.memcpy.p0i8.p2i8
+// CHECK: call void @llvm.memcpy.p2i8.p0i8
 void test4(MyStruct __attribute__((address_space(2))) *pPtr) {
   MyStruct s = pPtr[0];
   pPtr[0] = s;
 }

test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp

 // RUN: %clang_cc1 -std=c++11 -triple x86_64-none-linux-gnu -emit-llvm -o - %s | FileCheck -check-prefixes=X86,CHECK %s
-// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa-amdgiz -DNO_TLS -emit-llvm -o - %s | FileCheck -check-prefixes=AMD,CHECK %s
+// RUN: %clang_cc1 -std=c++11 -triple amdgcn-amd-amdhsa -DNO_TLS -emit-llvm -o - %s | FileCheck -check-prefixes=AMDGCN,CHECK %s
 
 namespace std {
   typedef decltype(sizeof(int)) size_t;
 
   // libc++'s implementation
   template <class _E>
   class initializer_list
   {
struct witharg1 {
   ~witharg1();
 };
 struct wantslist1 {
   wantslist1(std::initializer_list<destroyme1>);
   ~wantslist1();
 };
 // X86: @_ZGR15globalInitList1_ = internal constant [3 x i32] [i32 1, i32 2, i32 3]
 // X86: @globalInitList1 = global %{{[^ ]+}} { i32* getelementptr inbounds ([3 x i32], [3 x i32]* @_ZGR15globalInitList1_, i32 0, i32 0), i{{32|64}} 3 }
-// AMD: @_ZGR15globalInitList1_ = internal addrspace(1) constant [3 x i32] [i32 1, i32 2, i32 3]
-// AMD: @globalInitList1 = addrspace(1) global %{{[^ ]+}} { i32* addrspacecast (i32 addrspace(1)* getelementptr inbounds ([3 x i32], [3 x i32] addrspace(1)* @_ZGR15globalInitList1_, i32 0, i32 0) to i32*), i{{32|64}} 3 }
+// AMDGCN: @_ZGR15globalInitList1_ = internal addrspace(1) constant [3 x i32] [i32 1, i32 2, i32 3]
+// AMDGCN: @globalInitList1 = addrspace(1) global %{{[^ ]+}} { i32* addrspacecast (i32 addrspace(1)* getelementptr inbounds ([3 x i32], [3 x i32] addrspace(1)* @_ZGR15globalInitList1_, i32 0, i32 0) to i32*), i{{32|64}} 3 }
 std::initializer_list<int> globalInitList1 = {1, 2, 3};
 
 #ifndef NO_TLS
 namespace thread_local_global_array {
 // FIXME: We should be able to constant-evaluate this even though the
 // initializer is not a constant expression (pointers to thread_local
 // objects aren't really a problem).
 //
 // X86: @_ZN25thread_local_global_array1xE = thread_local global
 // X86: @_ZGRN25thread_local_global_array1xE_ = internal thread_local constant [4 x i32] [i32 1, i32 2, i32 3, i32 4]
 std::initializer_list<int> thread_local x = {1, 2, 3, 4};
 }
 #endif
 
 // X86: @globalInitList2 = global %{{[^ ]+}} zeroinitializer
 // X86: @_ZGR15globalInitList2_ = internal global [2 x %[[WITHARG:[^ ]*]]] zeroinitializer
-// AMD: @globalInitList2 = addrspace(1) global %{{[^ ]+}} zeroinitializer
-// AMD: @_ZGR15globalInitList2_ = internal addrspace(1) global [2 x %[[WITHARG:[^ ]*]]] zeroinitializer
+// AMDGCN: @globalInitList2 = addrspace(1) global %{{[^ ]+}} zeroinitializer
+// AMDGCN: @_ZGR15globalInitList2_ = internal addrspace(1) global [2 x %[[WITHARG:[^ ]*]]] zeroinitializer
 
 // X86: @_ZN15partly_constant1kE = global i32 0, align 4
 // X86: @_ZN15partly_constant2ilE = global {{.*}} null, align 8
 // X86: @[[PARTLY_CONSTANT_OUTER:_ZGRN15partly_constant2ilE.*]] = internal global {{.*}} zeroinitializer, align 8
 // X86: @[[PARTLY_CONSTANT_INNER:_ZGRN15partly_constant2ilE.*]] = internal global [3 x {{.*}}] zeroinitializer, align 8
 // X86: @[[PARTLY_CONSTANT_FIRST:_ZGRN15partly_constant2ilE.*]] = internal constant [3 x i32] [i32 1, i32 2, i32 3], align 4
 // X86: @[[PARTLY_CONSTANT_SECOND:_ZGRN15partly_constant2ilE.*]] = internal global [2 x i32] zeroinitializer, align 4
 // X86: @[[PARTLY_CONSTANT_THIRD:_ZGRN15partly_constant2ilE.*]] = internal constant [4 x i32] [i32 5, i32 6, i32 7, i32 8], align 4
-// AMD: @_ZN15partly_constant1kE = addrspace(1) global i32 0, align 4
-// AMD: @_ZN15partly_constant2ilE = addrspace(2) global {{.*}} null, align 8
-// AMD: @[[PARTLY_CONSTANT_OUTER:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) global {{.*}} zeroinitializer, align 8
-// AMD: @[[PARTLY_CONSTANT_INNER:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) global [3 x {{.*}}] zeroinitializer, align 8
-// AMD: @[[PARTLY_CONSTANT_FIRST:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4
-// AMD: @[[PARTLY_CONSTANT_SECOND:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) global [2 x i32] zeroinitializer, align 4
-// AMD: @[[PARTLY_CONSTANT_THIRD:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) constant [4 x i32] [i32 5, i32 6, i32 7, i32 8], align 4
+// AMDGCN: @_ZN15partly_constant1kE = addrspace(1) global i32 0, align 4
+// AMDGCN: @_ZN15partly_constant2ilE = addrspace(2) global {{.*}} null, align 8
+// AMDGCN: @[[PARTLY_CONSTANT_OUTER:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) global {{.*}} zeroinitializer, align 8
+// AMDGCN: @[[PARTLY_CONSTANT_INNER:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) global [3 x {{.*}}] zeroinitializer, align 8
+// AMDGCN: @[[PARTLY_CONSTANT_FIRST:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) constant [3 x i32] [i32 1, i32 2, i32 3], align 4
+// AMDGCN: @[[PARTLY_CONSTANT_SECOND:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) global [2 x i32] zeroinitializer, align 4
+// AMDGCN: @[[PARTLY_CONSTANT_THIRD:_ZGRN15partly_constant2ilE.*]] = internal addrspace(2) constant [4 x i32] [i32 5, i32 6, i32 7, i32 8], align 4
 
 // X86: @[[REFTMP1:.*]] = private constant [2 x i32] [i32 42, i32 43], align 4
 // X86: @[[REFTMP2:.*]] = private constant [3 x %{{.*}}] [%{{.*}} { i32 1 }, %{{.*}} { i32 2 }, %{{.*}} { i32 3 }], align 4
-// AMD: @[[REFTMP1:.*]] = private addrspace(2) constant [2 x i32] [i32 42, i32 43], align 4
-// AMD: @[[REFTMP2:.*]] = private addrspace(2) constant [3 x %{{.*}}] [%{{.*}} { i32 1 }, %{{.*}} { i32 2 }, %{{.*}} { i32 3 }], align 4
+// AMDGCN: @[[REFTMP1:.*]] = private addrspace(2) constant [2 x i32] [i32 42, i32 43], align 4
+// AMDGCN: @[[REFTMP2:.*]] = private addrspace(2) constant [3 x %{{.*}}] [%{{.*}} { i32 1 }, %{{.*}} { i32 2 }, %{{.*}} { i32 3 }], align 4
 
 // CHECK: appending global
 
 // thread_local initializer:
 // X86-LABEL: define internal void @__cxx_global_var_init
 // X86: store i32* getelementptr inbounds ([4 x i32], [4 x i32]* @_ZGRN25thread_local_global_array1xE_, i64 0, i64 0),
 // X86:       i32** getelementptr inbounds ({{.*}}, {{.*}}* @_ZN25thread_local_global_array1xE, i32 0, i32 0), align 8
 // X86: store i64 4, i64* getelementptr inbounds ({{.*}}, {{.*}}* @_ZN25thread_local_global_array1xE, i32 0, i32 1), align 8
 
 // CHECK-LABEL: define internal void @__cxx_global_var_init
 // X86: call void @_ZN8witharg1C1ERK10destroyme1(%[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* @_ZGR15globalInitList2_, i{{32|64}} 0, i{{32|64}} 0
 // X86: call void @_ZN8witharg1C1ERK10destroyme1(%[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* @_ZGR15globalInitList2_, i{{32|64}} 0, i{{32|64}} 1
-// AMD: call void @_ZN8witharg1C1ERK10destroyme1(%[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* addrspacecast ({{[^@]+}} @_ZGR15globalInitList2_ {{[^)]+}}), i{{32|64}} 0, i{{32|64}} 0
-// AMD: call void @_ZN8witharg1C1ERK10destroyme1(%[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* addrspacecast ({{[^@]+}} @_ZGR15globalInitList2_ {{[^)]+}}), i{{32|64}} 0, i{{32|64}} 1
+// AMDGCN: call void @_ZN8witharg1C1ERK10destroyme1(%[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* addrspacecast ({{[^@]+}} @_ZGR15globalInitList2_ {{[^)]+}}), i{{32|64}} 0, i{{32|64}} 0
+// AMDGCN: call void @_ZN8witharg1C1ERK10destroyme1(%[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* addrspacecast ({{[^@]+}} @_ZGR15globalInitList2_ {{[^)]+}}), i{{32|64}} 0, i{{32|64}} 1
 // CHECK: call i32 @__cxa_atexit
 // X86: store %[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* @_ZGR15globalInitList2_, i64 0, i64 0),
 // X86:       %[[WITHARG]]** getelementptr inbounds (%{{.*}}, %{{.*}}* @globalInitList2, i32 0, i32 0), align 8
 // X86: store i64 2, i64* getelementptr inbounds (%{{.*}}, %{{.*}}* @globalInitList2, i32 0, i32 1), align 8
-// AMD: store %[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* addrspacecast ({{[^@]+}} @_ZGR15globalInitList2_ {{[^)]+}}), i64 0, i64 0),
-// AMD:       %[[WITHARG]]** getelementptr inbounds (%{{.*}}, %{{.*}}* addrspacecast ({{[^@]+}} @globalInitList2 {{[^)]+}}), i32 0, i32 0), align 8
-// AMD: store i64 2, i64* getelementptr inbounds (%{{.*}}, %{{.*}}* addrspacecast ({{[^@]+}} @globalInitList2 {{[^)]+}}), i32 0, i32 1), align 8
+// AMDGCN: store %[[WITHARG]]* getelementptr inbounds ([2 x %[[WITHARG]]], [2 x %[[WITHARG]]]* addrspacecast ({{[^@]+}} @_ZGR15globalInitList2_ {{[^)]+}}), i64 0, i64 0),
+// AMDGCN:       %[[WITHARG]]** getelementptr inbounds (%{{.*}}, %{{.*}}* addrspacecast ({{[^@]+}} @globalInitList2 {{[^)]+}}), i32 0, i32 0), align 8
+// AMDGCN: store i64 2, i64* getelementptr inbounds (%{{.*}}, %{{.*}}* addrspacecast ({{[^@]+}} @globalInitList2 {{[^)]+}}), i32 0, i32 1), align 8
 // CHECK: call void @_ZN10destroyme1D1Ev
 // CHECK-NEXT: call void @_ZN10destroyme1D1Ev
 // CHECK-NEXT: ret void
 std::initializer_list<witharg1> globalInitList2 = {
   witharg1(destroyme1()), witharg1(destroyme1())
 };
 
 void fn1(int i) {
   // CHECK-LABEL: define void @_Z3fn1i
   // temporary array
   // X86: [[array:%[^ ]+]] = alloca [3 x i32]
-  // AMD: [[alloca:%[^ ]+]] = alloca [3 x i32], align 4, addrspace(5)
-  // AMD: [[array:%[^ ]+]] = addrspacecast [3 x i32] addrspace(5)* [[alloca]] to [3 x i32]*
+  // AMDGCN: [[alloca:%[^ ]+]] = alloca [3 x i32], align 4, addrspace(5)
+  // AMDGCN: [[array:%[^ ]+]] = addrspacecast [3 x i32] addrspace(5)* [[alloca]] to [3 x i32]*
   // CHECK: getelementptr inbounds [3 x i32], [3 x i32]* [[array]], i{{32|64}} 0
   // CHECK-NEXT: store i32 1, i32*
   // CHECK-NEXT: getelementptr
   // CHECK-NEXT: store
   // CHECK-NEXT: getelementptr
   // CHECK-NEXT: load
   // CHECK-NEXT: store
   // init the list
namespace B19773010 {
   enum E { ENUM_CONSTANT };
   struct testcase {
     testcase(std::initializer_list<pair<const char *, E>>);
   };
   void f1() {
     // CHECK-LABEL: @_ZN9B197730102f1Ev
     testcase a{{"", ENUM_CONSTANT}};
     // X86: store %"struct.B19773010::pair"* getelementptr inbounds ([1 x %"struct.B19773010::pair"], [1 x %"struct.B19773010::pair"]* bitcast ([1 x { i8*, i32 }]* @.ref.tmp{{.*}} to [1 x %"struct.B19773010::pair"]*), i64 0, i64 0), %"struct.B19773010::pair"** %{{.*}}, align 8
-    // AMD: store %"struct.B19773010::pair"* getelementptr inbounds ([1 x %"struct.B19773010::pair"], [1 x %"struct.B19773010::pair"]* addrspacecast{{.*}} bitcast ([1 x { i8*, i32 }] addrspace(2)* @.ref.tmp{{.*}} to [1 x %"struct.B19773010::pair"] addrspace(2)*){{.*}}, i64 0, i64 0), %"struct.B19773010::pair"** %{{.*}}, align 8
+    // AMDGCN: store %"struct.B19773010::pair"* getelementptr inbounds ([1 x %"struct.B19773010::pair"], [1 x %"struct.B19773010::pair"]* addrspacecast{{.*}} bitcast ([1 x { i8*, i32 }] addrspace(2)* @.ref.tmp{{.*}} to [1 x %"struct.B19773010::pair"] addrspace(2)*){{.*}}, i64 0, i64 0), %"struct.B19773010::pair"** %{{.*}}, align 8
   }
   void f2() {
     // CHECK-LABEL: @_ZN9B197730102f2Ev
     // X86: store %"struct.B19773010::pair"* getelementptr inbounds ([1 x %"struct.B19773010::pair"], [1 x %"struct.B19773010::pair"]* bitcast ([1 x { i8*, i32 }]* @_ZGRZN9B197730102f2EvE1p_ to [1 x %"struct.B19773010::pair"]*), i64 0, i64 0), %"struct.B19773010::pair"** getelementptr inbounds ([2 x %"class.std::initializer_list.10"], [2 x %"class.std::initializer_list.10"]* @_ZZN9B197730102f2EvE1p, i64 0, i64 1, i32 0), align 16
-    // AMD: store %"struct.B19773010::pair"* getelementptr inbounds ([1 x %"struct.B19773010::pair"], [1 x %"struct.B19773010::pair"]* addrspacecast{{.*}} bitcast ([1 x { i8*, i32 }] addrspace(1)* @_ZGRZN9B197730102f2EvE1p_ to [1 x %"struct.B19773010::pair"] addrspace(1)*){{.*}}, i64 0, i64 0), %"struct.B19773010::pair"** getelementptr inbounds ([2 x %"class.std::initializer_list.10"], [2 x %"class.std::initializer_list.10"]* addrspacecast{{.*}}@_ZZN9B197730102f2EvE1p{{.*}}, i64 0, i64 1, i32 0), align 8
+    // AMDGCN: store %"struct.B19773010::pair"* getelementptr inbounds ([1 x %"struct.B19773010::pair"], [1 x %"struct.B19773010::pair"]* addrspacecast{{.*}} bitcast ([1 x { i8*, i32 }] addrspace(1)* @_ZGRZN9B197730102f2EvE1p_ to [1 x %"struct.B19773010::pair"] addrspace(1)*){{.*}}, i64 0, i64 0), %"struct.B19773010::pair"** getelementptr inbounds ([2 x %"class.std::initializer_list.10"], [2 x %"class.std::initializer_list.10"]* addrspacecast{{.*}}@_ZZN9B197730102f2EvE1p{{.*}}, i64 0, i64 1, i32 0), align 8
     static std::initializer_list<pair<const char *, E>> a, p[2] =
         {a, {{"", ENUM_CONSTANT}}};
   }
 
   void PR22940_helper(const pair<void*, int>&) { }
   void PR22940() {
     // CHECK-LABEL: @_ZN9B197730107PR22940Ev
     // CHECK: call {{.*}} @_ZN9B197730104pairIPviEC{{.}}Ev(
     // CHECK: call {{.*}} @_ZN9B1977301014PR22940_helperERKNS_4pairIPviEE(
     PR22940_helper(pair<void*, int>());
   }
 }

test/CodeGenCXX/vla.cpp

 // RUN: %clang_cc1 -std=c++11 -triple x86_64-apple-darwin %s -emit-llvm -o - | FileCheck -check-prefixes=X64,CHECK %s
-// RUN: %clang_cc1 -std=c++11 -triple amdgcn---amdgiz %s -emit-llvm -o - | FileCheck -check-prefixes=AMD,CHECK %s
+// RUN: %clang_cc1 -std=c++11 -triple amdgcn %s -emit-llvm -o - | FileCheck -check-prefixes=AMDGCN,CHECK %s
 
 template<typename T>
 struct S {
   static int n;
 };
 template<typename T> int S<T>::n = 5;
 
 int f() {
   // Make sure that the reference here is enough to trigger the instantiation of
   // the static data member.
   // CHECK: @_ZN1SIiE1nE = linkonce_odr{{.*}} global i32 5
   int a[S<int>::n];
   return sizeof a;
 }
 
 // rdar://problem/9506377
 void test0(void *array, int n) {
   // CHECK-LABEL: define void @_Z5test0Pvi(
   // X64:        [[ARRAY:%.*]] = alloca i8*, align 8
-  // AMD:        [[ARRAY0:%.*]] = alloca i8*, align 8, addrspace(5)
-  // AMD-NEXT:   [[ARRAY:%.*]] = addrspacecast i8* addrspace(5)* [[ARRAY0]] to i8**
+  // AMDGCN:        [[ARRAY0:%.*]] = alloca i8*, align 8, addrspace(5)
+  // AMDGCN-NEXT:   [[ARRAY:%.*]] = addrspacecast i8* addrspace(5)* [[ARRAY0]] to i8**
   // X64-NEXT:   [[N:%.*]] = alloca i32, align 4
-  // AMD:        [[N0:%.*]] = alloca i32, align 4, addrspace(5)
-  // AMD-NEXT:   [[N:%.*]] = addrspacecast i32 addrspace(5)* [[N0]] to i32*
+  // AMDGCN:        [[N0:%.*]] = alloca i32, align 4, addrspace(5)
+  // AMDGCN-NEXT:   [[N:%.*]] = addrspacecast i32 addrspace(5)* [[N0]] to i32*
   // X64-NEXT:   [[REF:%.*]] = alloca i16*, align 8
-  // AMD:        [[REF0:%.*]] = alloca i16*, align 8, addrspace(5)
-  // AMD-NEXT:   [[REF:%.*]] = addrspacecast i16* addrspace(5)* [[REF0]] to i16**
+  // AMDGCN:        [[REF0:%.*]] = alloca i16*, align 8, addrspace(5)
+  // AMDGCN-NEXT:   [[REF:%.*]] = addrspacecast i16* addrspace(5)* [[REF0]] to i16**
   // X64-NEXT:   [[S:%.*]] = alloca i16, align 2
-  // AMD:        [[S0:%.*]] = alloca i16, align 2, addrspace(5)
-  // AMD-NEXT:   [[S:%.*]] = addrspacecast i16 addrspace(5)* [[S0]] to i16*
+  // AMDGCN:        [[S0:%.*]] = alloca i16, align 2, addrspace(5)
+  // AMDGCN-NEXT:   [[S:%.*]] = addrspacecast i16 addrspace(5)* [[S0]] to i16*
   // CHECK-NEXT: store i8* 
   // CHECK-NEXT: store i32
 
   // Capture the bounds.
   // CHECK-NEXT: [[T0:%.*]] = load i32, i32* [[N]], align 4
   // CHECK-NEXT: [[DIM0:%.*]] = zext i32 [[T0]] to i64
   // CHECK-NEXT: [[T0:%.*]] = load i32, i32* [[N]], align 4
   // CHECK-NEXT: [[T1:%.*]] = add nsw i32 [[T0]], 1
void test0(void *array, int n) {
 
   // CHECK-NEXT: ret void
 }
 
 
 void test2(int b) {
   // CHECK-LABEL: define void {{.*}}test2{{.*}}(i32 %b)
   int varr[b];
-  // AMD: %__end1 = alloca i32*, align 8, addrspace(5)
-  // AMD: [[END:%.*]] = addrspacecast i32* addrspace(5)* %__end1 to i32**
+  // AMDGCN: %__end1 = alloca i32*, align 8, addrspace(5)
+  // AMDGCN: [[END:%.*]] = addrspacecast i32* addrspace(5)* %__end1 to i32**
   // get the address of %b by checking the first store that stores it 
   //CHECK: store i32 %b, i32* [[PTR_B:%.*]]
 
   // get the size of the VLA by getting the first load of the PTR_B
   //CHECK: [[VLA_NUM_ELEMENTS_PREZEXT:%.*]] = load i32, i32* [[PTR_B]]
   //CHECK-NEXT: [[VLA_NUM_ELEMENTS_PRE:%.*]] = zext i32 [[VLA_NUM_ELEMENTS_PREZEXT]]
   
   b = 15;
   //CHECK: store i32 15, i32* [[PTR_B]]
   
   // Now get the sizeof, and then divide by the element size
   
   
   //CHECK: [[VLA_SIZEOF:%.*]] = mul nuw i64 4, [[VLA_NUM_ELEMENTS_PRE]]
   //CHECK-NEXT: [[VLA_NUM_ELEMENTS_POST:%.*]] = udiv i64 [[VLA_SIZEOF]], 4
   //CHECK-NEXT: [[VLA_END_PTR:%.*]] = getelementptr inbounds i32, i32* {{%.*}}, i64 [[VLA_NUM_ELEMENTS_POST]]
   //X64-NEXT: store i32* [[VLA_END_PTR]], i32** %__end1
-  //AMD-NEXT: store i32* [[VLA_END_PTR]], i32** [[END]]
+  //AMDGCN-NEXT: store i32* [[VLA_END_PTR]], i32** [[END]]
   for (int d : varr) 0;
 }
 
 void test3(int b, int c) {
   // CHECK-LABEL: define void {{.*}}test3{{.*}}(i32 %b, i32 %c)
   int varr[b][c];
-  // AMD: %__end1 = alloca i32*, align 8, addrspace(5)
-  // AMD: [[END:%.*]] = addrspacecast i32* addrspace(5)* %__end1 to i32**
+  // AMDGCN: %__end1 = alloca i32*, align 8, addrspace(5)
+  // AMDGCN: [[END:%.*]] = addrspacecast i32* addrspace(5)* %__end1 to i32**
   // get the address of %b by checking the first store that stores it 
   //CHECK: store i32 %b, i32* [[PTR_B:%.*]]
   //CHECK-NEXT: store i32 %c, i32* [[PTR_C:%.*]]
   
   // get the size of the VLA by getting the first load of the PTR_B
   //CHECK: [[VLA_DIM1_PREZEXT:%.*]] = load i32, i32* [[PTR_B]]
   //CHECK-NEXT: [[VLA_DIM1_PRE:%.*]] = zext i32 [[VLA_DIM1_PREZEXT]]
   //CHECK: [[VLA_DIM2_PREZEXT:%.*]] = load i32, i32* [[PTR_C]]
   //CHECK-NEXT: [[VLA_DIM2_PRE:%.*]] = zext i32 [[VLA_DIM2_PREZEXT]]
   
   b = 15;
   c = 15;
   //CHECK: store i32 15, i32* [[PTR_B]]
   //CHECK: store i32 15, i32* [[PTR_C]]
   // Now get the sizeof, and then divide by the element size
   
   // multiply the two dimensions, then by the element type and then divide by the sizeof dim2
   //CHECK: [[VLA_DIM1_X_DIM2:%.*]] = mul nuw i64 [[VLA_DIM1_PRE]], [[VLA_DIM2_PRE]]
   //CHECK-NEXT: [[VLA_SIZEOF:%.*]] = mul nuw i64 4, [[VLA_DIM1_X_DIM2]]
   //CHECK-NEXT: [[VLA_SIZEOF_DIM2:%.*]] = mul nuw i64 4, [[VLA_DIM2_PRE]]
   //CHECK-NEXT: [[VLA_NUM_ELEMENTS:%.*]] = udiv i64 [[VLA_SIZEOF]], [[VLA_SIZEOF_DIM2]]
   //CHECK-NEXT: [[VLA_END_INDEX:%.*]] = mul nsw i64 [[VLA_NUM_ELEMENTS]], [[VLA_DIM2_PRE]]
   //CHECK-NEXT: [[VLA_END_PTR:%.*]] = getelementptr inbounds i32, i32* {{%.*}}, i64 [[VLA_END_INDEX]]
   //X64-NEXT: store i32* [[VLA_END_PTR]], i32** %__end
-  //AMD-NEXT: store i32* [[VLA_END_PTR]], i32** [[END]]
+  //AMDGCN-NEXT: store i32* [[VLA_END_PTR]], i32** [[END]]
  
   for (auto &d : varr) 0;
 }
 
 

test/CodeGenOpenCL/addr-space-struct-arg.cl

 // RUN: %clang_cc1 %s -emit-llvm -o - -O0 -finclude-default-header -ffake-address-space-map -triple i686-pc-darwin | FileCheck -enable-var-scope -check-prefixes=COM,X86 %s
-// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -finclude-default-header -triple amdgcn-amdhsa-amd-amdgizcl | FileCheck -enable-var-scope -check-prefixes=COM,AMD %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -finclude-default-header -triple amdgcn-amdhsa-amd | FileCheck -enable-var-scope -check-prefixes=COM,AMDGCN %s
 
 typedef struct {
   int cells[9];
 } Mat3X3;
 
 typedef struct {
   int cells[16];
 } Mat4X4;

 
 struct LargeStructTwoMember {
   int2 x[40];
   int2 y[20];
 };
 
 
 // X86-LABEL: define void @foo(%struct.Mat4X4* noalias sret %agg.result, %struct.Mat3X3* byval align 4 %in)
-// AMD-LABEL: define %struct.Mat4X4 @foo([9 x i32] %in.coerce)
+// AMDGCN-LABEL: define %struct.Mat4X4 @foo([9 x i32] %in.coerce)
 Mat4X4 __attribute__((noinline)) foo(Mat3X3 in) {
   Mat4X4 out;
   return out;
 }
 
 // COM-LABEL: define {{.*}} void @ker
 // Expect two mem copies: one for the argument "in", and one for
 // the return value.
 // X86: call void @llvm.memcpy.p0i8.p1i8.i32(i8*
 // X86: call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)*
 
-// AMD: load [9 x i32], [9 x i32] addrspace(1)*
-// AMD: call %struct.Mat4X4 @foo([9 x i32]
-// AMD: call void @llvm.memcpy.p1i8.p5i8.i64(i8 addrspace(1)*
+// AMDGCN: load [9 x i32], [9 x i32] addrspace(1)*
+// AMDGCN: call %struct.Mat4X4 @foo([9 x i32]
+// AMDGCN: call void @llvm.memcpy.p1i8.p5i8.i64(i8 addrspace(1)*
 kernel void ker(global Mat3X3 *in, global Mat4X4 *out) {
   out[0] = foo(in[1]);
 }
 
 // X86-LABEL: define void @foo_large(%struct.Mat64X64* noalias sret %agg.result, %struct.Mat32X32* byval align 4 %in)
-// AMD-LABEL: define void @foo_large(%struct.Mat64X64 addrspace(5)* noalias sret %agg.result, %struct.Mat32X32 addrspace(5)* byval align 4 %in)
+// AMDGCN-LABEL: define void @foo_large(%struct.Mat64X64 addrspace(5)* noalias sret %agg.result, %struct.Mat32X32 addrspace(5)* byval align 4 %in)
 Mat64X64 __attribute__((noinline)) foo_large(Mat32X32 in) {
   Mat64X64 out;
   return out;
 }
 
 // COM-LABEL: define {{.*}} void @ker_large
 // Expect two mem copies: one for the argument "in", and one for
 // the return value.
 // X86: call void @llvm.memcpy.p0i8.p1i8.i32(i8*
 // X86: call void @llvm.memcpy.p1i8.p0i8.i32(i8 addrspace(1)*
-// AMD: call void @llvm.memcpy.p5i8.p1i8.i64(i8 addrspace(5)*
-// AMD: call void @llvm.memcpy.p1i8.p5i8.i64(i8 addrspace(1)*
+// AMDGCN: call void @llvm.memcpy.p5i8.p1i8.i64(i8 addrspace(5)*
+// AMDGCN: call void @llvm.memcpy.p1i8.p5i8.i64(i8 addrspace(1)*
 kernel void ker_large(global Mat32X32 *in, global Mat64X64 *out) {
   out[0] = foo_large(in[1]);
 }
 
-// AMD-LABEL: define void @FuncOneMember(<2 x i32> %u.coerce)
+// AMDGCN-LABEL: define void @FuncOneMember(<2 x i32> %u.coerce)
 void FuncOneMember(struct StructOneMember u) {
   u.x = (int2)(0, 0);
 }
 
-// AMD-LABEL: define void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* byval align 8 %u)
+// AMDGCN-LABEL: define void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* byval align 8 %u)
 void FuncOneLargeMember(struct LargeStructOneMember u) {
   u.x[0] = (int2)(0, 0);
 }
 
-// AMD-LABEL: define amdgpu_kernel void @KernelOneMember
-// AMD-SAME:  (<2 x i32> %[[u_coerce:.*]])
-// AMD:  %[[u:.*]] = alloca %struct.StructOneMember, align 8, addrspace(5)
-// AMD:  %[[coerce_dive:.*]] = getelementptr inbounds %struct.StructOneMember, %struct.StructOneMember addrspace(5)* %[[u]], i32 0, i32 0
-// AMD:  store <2 x i32> %[[u_coerce]], <2 x i32> addrspace(5)* %[[coerce_dive]]
-// AMD:  call void @FuncOneMember(<2 x i32>
+// AMDGCN-LABEL: define amdgpu_kernel void @KernelOneMember
+// AMDGCN-SAME:  (<2 x i32> %[[u_coerce:.*]])
+// AMDGCN:  %[[u:.*]] = alloca %struct.StructOneMember, align 8, addrspace(5)
+// AMDGCN:  %[[coerce_dive:.*]] = getelementptr inbounds %struct.StructOneMember, %struct.StructOneMember addrspace(5)* %[[u]], i32 0, i32 0
+// AMDGCN:  store <2 x i32> %[[u_coerce]], <2 x i32> addrspace(5)* %[[coerce_dive]]
+// AMDGCN:  call void @FuncOneMember(<2 x i32>
 kernel void KernelOneMember(struct StructOneMember u) {
   FuncOneMember(u);
 }
 
-// AMD-LABEL: define amdgpu_kernel void @KernelLargeOneMember(
-// AMD:  %[[U:.*]] = alloca %struct.LargeStructOneMember, align 8, addrspace(5)
-// AMD:  store %struct.LargeStructOneMember %u.coerce, %struct.LargeStructOneMember addrspace(5)* %[[U]], align 8
-// AMD:  call void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* byval align 8 %[[U]])
+// AMDGCN-LABEL: define amdgpu_kernel void @KernelLargeOneMember(
+// AMDGCN:  %[[U:.*]] = alloca %struct.LargeStructOneMember, align 8, addrspace(5)
+// AMDGCN:  store %struct.LargeStructOneMember %u.coerce, %struct.LargeStructOneMember addrspace(5)* %[[U]], align 8
+// AMDGCN:  call void @FuncOneLargeMember(%struct.LargeStructOneMember addrspace(5)* byval align 8 %[[U]])
 kernel void KernelLargeOneMember(struct LargeStructOneMember u) {
   FuncOneLargeMember(u);
 }
 
-// AMD-LABEL: define void @FuncTwoMember(<2 x i32> %u.coerce0, <2 x i32> %u.coerce1)
+// AMDGCN-LABEL: define void @FuncTwoMember(<2 x i32> %u.coerce0, <2 x i32> %u.coerce1)
 void FuncTwoMember(struct StructTwoMember u) {
   u.y = (int2)(0, 0);
 }
 
-// AMD-LABEL: define void @FuncLargeTwoMember(%struct.LargeStructTwoMember addrspace(5)* byval align 8 %u)
+// AMDGCN-LABEL: define void @FuncLargeTwoMember(%struct.LargeStructTwoMember addrspace(5)* byval align 8 %u)
 void FuncLargeTwoMember(struct LargeStructTwoMember u) {
   u.y[0] = (int2)(0, 0);
 }
 
 
-// AMD-LABEL: define amdgpu_kernel void @KernelTwoMember
-// AMD-SAME:  (%struct.StructTwoMember %[[u_coerce:.*]])
-// AMD:  %[[u:.*]] = alloca %struct.StructTwoMember, align 8, addrspace(5)
-// AMD: %[[LD0:.*]] = load <2 x i32>, <2 x i32> addrspace(5)*
-// AMD: %[[LD1:.*]] = load <2 x i32>, <2 x i32> addrspace(5)*
-// AMD: call void @FuncTwoMember(<2 x i32> %[[LD0]], <2 x i32> %[[LD1]])
+// AMDGCN-LABEL: define amdgpu_kernel void @KernelTwoMember
+// AMDGCN-SAME:  (%struct.StructTwoMember %[[u_coerce:.*]])
+// AMDGCN:  %[[u:.*]] = alloca %struct.StructTwoMember, align 8, addrspace(5)
+// AMDGCN: %[[LD0:.*]] = load <2 x i32>, <2 x i32> addrspace(5)*
+// AMDGCN: %[[LD1:.*]] = load <2 x i32>, <2 x i32> addrspace(5)*
+// AMDGCN: call void @FuncTwoMember(<2 x i32> %[[LD0]], <2 x i32> %[[LD1]])
 kernel void KernelTwoMember(struct StructTwoMember u) {
   FuncTwoMember(u);
 }
 
-// AMD-LABEL: define amdgpu_kernel void @KernelLargeTwoMember
-// AMD-SAME:  (%struct.LargeStructTwoMember %[[u_coerce:.*]])
-// AMD:  %[[u:.*]] = alloca %struct.LargeStructTwoMember, align 8, addrspace(5)
-// AMD:  store %struct.LargeStructTwoMember %[[u_coerce]], %struct.LargeStructTwoMember addrspace(5)* %[[u]]
-// AMD:  call void @FuncLargeTwoMember(%struct.LargeStructTwoMember addrspace(5)* byval align 8 %[[u]])
+// AMDGCN-LABEL: define amdgpu_kernel void @KernelLargeTwoMember
+// AMDGCN-SAME:  (%struct.LargeStructTwoMember %[[u_coerce:.*]])
+// AMDGCN:  %[[u:.*]] = alloca %struct.LargeStructTwoMember, align 8, addrspace(5)
+// AMDGCN:  store %struct.LargeStructTwoMember %[[u_coerce]], %struct.LargeStructTwoMember addrspace(5)* %[[u]]
+// AMDGCN:  call void @FuncLargeTwoMember(%struct.LargeStructTwoMember addrspace(5)* byval align 8 %[[u]])
 kernel void KernelLargeTwoMember(struct LargeStructTwoMember u) {
   FuncLargeTwoMember(u);
 }

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

 // 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
+// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck -check-prefix=AMDGCN %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;
 
 // 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)*) }
+// AMDGCN: %struct.ConstantArrayPointerStruct = type { float addrspace(4)* }
+// AMDGCN: 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
 };
 
 __kernel void initializer_cast_is_valid_crash()
 {
   unsigned char v512[64] = {
       0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,
       0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,
       0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,
       0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00
   };
 
 }

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 -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
+// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s
+// RUN: %clang_cc1 %s -O0 -triple amdgcn-amd-amdhsa -DCL20 -cl-std=CL2.0 -emit-llvm -o - | FileCheck %s --check-prefixes=CL20,CL20AMDGCN
+// RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s
+// RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,AMDGCN %s
 
 // SPIR: %struct.S = type { i32, i32, i32* }
 // CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* }
 struct S {
   int x;
   int y;
   int *z;
 };
 
 // CL20-DAG: @g_extern_var = external addrspace(1) global float
 // CL20-DAG: @l_extern_var = external addrspace(1) global float
 // CL20-DAG: @test_static.l_static_var = internal addrspace(1) global float 0.000000e+00
 // CL20-DAG: @g_static_var = internal addrspace(1) global float 0.000000e+00
 
 #ifdef CL20
 // CL20-DAG: @g_s = common addrspace(1) global %struct.S zeroinitializer
 struct S g_s;
 #endif
 
 // SPIR: i32* %arg
-// GIZ: i32 addrspace(5)* %arg
+// AMDGCN: i32 addrspace(5)* %arg
 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) {}
 
 // SPIR: i32 addrspace(2)* %arg
-// GIZ: i32 addrspace(4)* %arg
+// AMDGCN: i32 addrspace(4)* %arg
 void f__c(__constant int *arg) {}
 
 // SPIR: i32* %arg
-// GIZ: i32 addrspace(5)* %arg
+// AMDGCN: 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) {}
 
 // SPIR: i32 addrspace(2)* %arg
-// GIZ: i32 addrspace(4)* %arg
+// AMDGCN: 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
+// CL20AMDGCN-DAG: @ptr = common addrspace(1) global i32* null
 #endif
 
 // SPIR: i32* %arg
-// GIZ: i32 addrspace(5)* %arg
+// AMDGCN: i32 addrspace(5)* %arg
 // CL20SPIR-DAG: i32 addrspace(4)* %arg
-// CL20GIZ-DAG: i32* %arg
+// CL20AMDGCN-DAG: i32* %arg
 void f(int *arg) {
 
   int i;
 // SPIR: %i = alloca i32,
-// GIZ: %i = alloca i32{{.*}}addrspace(5)
+// AMDGCN: %i = alloca i32{{.*}}addrspace(5)
 // CL20SPIR-DAG: %i = alloca i32,
-// CL20GIZ-DAG: %i = alloca i32{{.*}}addrspace(5)
+// CL20AMDGCN-DAG: %i = alloca i32{{.*}}addrspace(5)
 
 #ifdef CL20
   static int ii;
 // CL20-DAG: @f.ii = internal addrspace(1) global i32 0
 #endif
 }
 
 typedef int int_td;

test/CodeGenOpenCL/blocks.cl

 // RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck -check-prefixes=COMMON,SPIR %s
-// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa-opencl | FileCheck -check-prefixes=COMMON,AMD %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -emit-llvm -o - -O0 -triple amdgcn-amd-amdhsa | FileCheck -check-prefixes=COMMON,AMDGCN %s
 
 // SPIR: %struct.__opencl_block_literal_generic = type { i32, i32, i8 addrspace(4)* }
-// AMD: %struct.__opencl_block_literal_generic = type { i32, i32, i8* }
+// AMDGCN: %struct.__opencl_block_literal_generic = type { i32, i32, i8* }
 // SPIR: @__block_literal_global = internal addrspace(1) constant { i32, i32, i8 addrspace(4)* } { i32 12, i32 4, i8 addrspace(4)* addrspacecast (i8* bitcast (void (i8 addrspace(4)*, i8 addrspace(3)*)* @block_A_block_invoke to i8*) to i8 addrspace(4)*) }
-// AMD: @__block_literal_global = internal addrspace(1) constant { i32, i32, i8* } { i32 16, i32 8, i8* bitcast (void (i8*, i8 addrspace(3)*)* @block_A_block_invoke to i8*) }
+// AMDGCN: @__block_literal_global = internal addrspace(1) constant { i32, i32, i8* } { i32 16, i32 8, i8* bitcast (void (i8*, i8 addrspace(3)*)* @block_A_block_invoke to i8*) }
 // COMMON-NOT: .str
 
 // SPIR-LABEL: define internal {{.*}}void @block_A_block_invoke(i8 addrspace(4)* %.block_descriptor, i8 addrspace(3)* %a)
-// AMD-LABEL: define internal {{.*}}void @block_A_block_invoke(i8* %.block_descriptor, i8 addrspace(3)* %a)
+// AMDGCN-LABEL: define internal {{.*}}void @block_A_block_invoke(i8* %.block_descriptor, i8 addrspace(3)* %a)
 void (^block_A)(local void *) = ^(local void *a) {
   return;
 };
 
 // COMMON-LABEL: define {{.*}}void @foo()
 void foo(){
   int i;
   // COMMON-NOT: %block.isa
   // COMMON-NOT: %block.flags
   // COMMON-NOT: %block.reserved
   // COMMON-NOT: %block.descriptor
   // SPIR: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %block, i32 0, i32 0
-  // AMD: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %block, i32 0, i32 0
+  // AMDGCN: %[[block_size:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %block, i32 0, i32 0
   // SPIR: store i32 16, i32* %[[block_size]]
-  // AMD: store i32 20, i32 addrspace(5)* %[[block_size]]
+  // AMDGCN: store i32 20, i32 addrspace(5)* %[[block_size]]
   // SPIR: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %block, i32 0, i32 1
-  // AMD: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %block, i32 0, i32 1
+  // AMDGCN: %[[block_align:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %block, i32 0, i32 1
   // SPIR: store i32 4, i32* %[[block_align]]
-  // AMD: store i32 8, i32 addrspace(5)* %[[block_align]]
+  // AMDGCN: store i32 8, i32 addrspace(5)* %[[block_align]]
   // SPIR: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %[[block:.*]], i32 0, i32 2
   // SPIR: store i8 addrspace(4)* addrspacecast (i8* bitcast (i32 (i8 addrspace(4)*)* @__foo_block_invoke to i8*) to i8 addrspace(4)*), i8 addrspace(4)** %[[block_invoke]]
   // SPIR: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }>* %[[block]], i32 0, i32 3
   // SPIR: %[[i_value:.*]] = load i32, i32* %i
   // SPIR: store i32 %[[i_value]], i32* %[[block_captured]],
   // SPIR: %[[blk_ptr:.*]] = bitcast <{ i32, i32, i8 addrspace(4)*, i32 }>* %[[block]] to i32 ()*
   // SPIR: %[[blk_gen_ptr:.*]] = addrspacecast i32 ()* %[[blk_ptr]] to i32 () addrspace(4)*
   // SPIR: store i32 () addrspace(4)* %[[blk_gen_ptr]], i32 () addrspace(4)** %[[block_B:.*]],
   // SPIR: %[[blk_gen_ptr:.*]] = load i32 () addrspace(4)*, i32 () addrspace(4)** %[[block_B]]
   // SPIR: %[[block_literal:.*]] = bitcast i32 () addrspace(4)* %[[blk_gen_ptr]] to %struct.__opencl_block_literal_generic addrspace(4)*
   // SPIR: %[[invoke_addr:.*]] = getelementptr inbounds %struct.__opencl_block_literal_generic, %struct.__opencl_block_literal_generic addrspace(4)* %[[block_literal]], i32 0, i32 2
   // SPIR: %[[blk_gen_ptr:.*]] = bitcast %struct.__opencl_block_literal_generic addrspace(4)* %[[block_literal]] to i8 addrspace(4)*
   // SPIR: %[[invoke_func_ptr:.*]] = load i8 addrspace(4)*, i8 addrspace(4)* addrspace(4)* %[[invoke_addr]]
   // SPIR: %[[invoke_func:.*]] = addrspacecast i8 addrspace(4)* %[[invoke_func_ptr]] to i32 (i8 addrspace(4)*)*
   // SPIR: call {{.*}}i32 %[[invoke_func]](i8 addrspace(4)* %[[blk_gen_ptr]])
-  // AMD: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block:.*]], i32 0, i32 2
-  // AMD: store i8* bitcast (i32 (i8*)* @__foo_block_invoke to i8*), i8* addrspace(5)* %[[block_invoke]]
-  // AMD: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block]], i32 0, i32 3
-  // AMD: %[[i_value:.*]] = load i32, i32 addrspace(5)* %i
-  // AMD: store i32 %[[i_value]], i32 addrspace(5)* %[[block_captured]],
-  // AMD: %[[blk_ptr:.*]] = bitcast <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block]] to i32 () addrspace(5)*
-  // AMD: %[[blk_gen_ptr:.*]] = addrspacecast i32 () addrspace(5)* %[[blk_ptr]] to i32 ()*
-  // AMD: store i32 ()* %[[blk_gen_ptr]], i32 ()* addrspace(5)* %[[block_B:.*]],
-  // AMD: %[[blk_gen_ptr:.*]] = load i32 ()*, i32 ()* addrspace(5)* %[[block_B]]
-  // AMD: %[[block_literal:.*]] = bitcast i32 ()* %[[blk_gen_ptr]] to %struct.__opencl_block_literal_generic*
-  // AMD: %[[invoke_addr:.*]] = getelementptr inbounds %struct.__opencl_block_literal_generic, %struct.__opencl_block_literal_generic* %[[block_literal]], i32 0, i32 2
-  // AMD: %[[blk_gen_ptr:.*]] = bitcast %struct.__opencl_block_literal_generic* %[[block_literal]] to i8*
-  // AMD: %[[invoke_func_ptr:.*]] = load i8*, i8** %[[invoke_addr]]
-  // AMD: %[[invoke_func:.*]] = bitcast i8* %[[invoke_func_ptr]] to i32 (i8*)*
-  // AMD: call {{.*}}i32 %[[invoke_func]](i8* %[[blk_gen_ptr]])
+  // AMDGCN: %[[block_invoke:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block:.*]], i32 0, i32 2
+  // AMDGCN: store i8* bitcast (i32 (i8*)* @__foo_block_invoke to i8*), i8* addrspace(5)* %[[block_invoke]]
+  // AMDGCN: %[[block_captured:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block]], i32 0, i32 3
+  // AMDGCN: %[[i_value:.*]] = load i32, i32 addrspace(5)* %i
+  // AMDGCN: store i32 %[[i_value]], i32 addrspace(5)* %[[block_captured]],
+  // AMDGCN: %[[blk_ptr:.*]] = bitcast <{ i32, i32, i8*, i32 }> addrspace(5)* %[[block]] to i32 () addrspace(5)*
+  // AMDGCN: %[[blk_gen_ptr:.*]] = addrspacecast i32 () addrspace(5)* %[[blk_ptr]] to i32 ()*
+  // AMDGCN: store i32 ()* %[[blk_gen_ptr]], i32 ()* addrspace(5)* %[[block_B:.*]],
+  // AMDGCN: %[[blk_gen_ptr:.*]] = load i32 ()*, i32 ()* addrspace(5)* %[[block_B]]
+  // AMDGCN: %[[block_literal:.*]] = bitcast i32 ()* %[[blk_gen_ptr]] to %struct.__opencl_block_literal_generic*
+  // AMDGCN: %[[invoke_addr:.*]] = getelementptr inbounds %struct.__opencl_block_literal_generic, %struct.__opencl_block_literal_generic* %[[block_literal]], i32 0, i32 2
+  // AMDGCN: %[[blk_gen_ptr:.*]] = bitcast %struct.__opencl_block_literal_generic* %[[block_literal]] to i8*
+  // AMDGCN: %[[invoke_func_ptr:.*]] = load i8*, i8** %[[invoke_addr]]
+  // AMDGCN: %[[invoke_func:.*]] = bitcast i8* %[[invoke_func_ptr]] to i32 (i8*)*
+  // AMDGCN: call {{.*}}i32 %[[invoke_func]](i8* %[[blk_gen_ptr]])
 
   int (^ block_B)(void) = ^{
     return i;
   };
   block_B();
 }
 
 // SPIR-LABEL: define internal {{.*}}i32 @__foo_block_invoke(i8 addrspace(4)* %.block_descriptor)
 // SPIR:  %[[block:.*]] = bitcast i8 addrspace(4)* %.block_descriptor to <{ i32, i32, i8 addrspace(4)*, i32 }> addrspace(4)*
 // SPIR:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, i8 addrspace(4)*, i32 }>, <{ i32, i32, i8 addrspace(4)*, i32 }> addrspace(4)* %[[block]], i32 0, i32 3
 // SPIR:  %[[block_capture:.*]] = load i32, i32 addrspace(4)* %[[block_capture_addr]]
-// AMD-LABEL: define internal {{.*}}i32 @__foo_block_invoke(i8* %.block_descriptor)
-// AMD:  %[[block:.*]] = bitcast i8* %.block_descriptor to <{ i32, i32, i8*, i32 }>*
-// AMD:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }>* %[[block]], i32 0, i32 3
-// AMD:  %[[block_capture:.*]] = load i32, i32* %[[block_capture_addr]]
+// AMDGCN-LABEL: define internal {{.*}}i32 @__foo_block_invoke(i8* %.block_descriptor)
+// AMDGCN:  %[[block:.*]] = bitcast i8* %.block_descriptor to <{ i32, i32, i8*, i32 }>*
+// AMDGCN:  %[[block_capture_addr:.*]] = getelementptr inbounds <{ i32, i32, i8*, i32 }>, <{ i32, i32, i8*, i32 }>* %[[block]], i32 0, i32 3
+// AMDGCN:  %[[block_capture:.*]] = load i32, i32* %[[block_capture_addr]]
 
 // COMMON-NOT: define{{.*}}@__foo_block_invoke_kernel

test/CodeGenOpenCL/lifetime.cl

 // RUN: %clang_cc1 -emit-llvm -o - %s | FileCheck %s
-// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn---amdgizcl %s | FileCheck %s -check-prefix=AMDGIZ
+// RUN: %clang_cc1 -emit-llvm -o - -triple amdgcn %s | FileCheck %s -check-prefix=AMDGCN
 
 void use(char *a);
 
 __attribute__((always_inline)) void helper_no_markers() {
   char a;
   use(&a);
 }
 
 void lifetime_test() {
 // CHECK: @llvm.lifetime.start.p0i
-// AMDGIZ: @llvm.lifetime.start.p5i
+// AMDGCN: @llvm.lifetime.start.p5i
   helper_no_markers();
 }

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 -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,AMD %s
+// RUN: %clang_cc1 -emit-llvm -triple amdgcn-amd-amdhsa -O0 -cl-std=CL2.0 -o - %s | FileCheck -check-prefixes=CHECK,AMDGCN %s
 
 constant int sz0 = 5;
 // SPIR: @sz0 = addrspace(2) constant i32 5
-// AMD: @sz0 = addrspace(4) constant i32 5
+// AMDGCN: @sz0 = addrspace(4) constant i32 5
 const global int sz1 = 16;
 // CHECK: @sz1 = addrspace(1) constant i32 16
 const constant int sz2 = 8;
 // SPIR: @sz2 = addrspace(2) constant i32 8
-// AMD: @sz2 = addrspace(4) constant i32 8
+// AMDGCN: @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)
+// AMDGCN: %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)
+// AMDGCN: %vla1 = alloca [16 x i8]{{.*}}addrspace(5)
   local short vla2[sz2];
 }

test/Index/pipe-size.cl

 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple x86_64-unknown-linux-gnu %s -o - | FileCheck %s --check-prefix=X86
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR
 // RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPIR64
-// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa-amdgizcl %s -o - | FileCheck %s --check-prefix=AMD
+// RUN: %clang_cc1 -x cl -O0 -cl-std=CL2.0 -emit-llvm -triple amdgcn-amd-amdhsa %s -o - | FileCheck %s --check-prefix=AMDGCN
 __kernel void testPipe( pipe int test )
 {
     int s = sizeof(test);
     // X86: store %opencl.pipe_t* %test, %opencl.pipe_t** %test.addr, align 8
     // X86: store i32 8, i32* %s, align 4
     // SPIR: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 4
     // SPIR: store i32 4, i32* %s, align 4
     // SPIR64: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)** %test.addr, align 8
     // SPIR64: store i32 8, i32* %s, align 4
-    // AMD: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 8
-    // AMD: store i32 8, i32 addrspace(5)* %s, align 4
+    // AMDGCN: store %opencl.pipe_t addrspace(1)* %test, %opencl.pipe_t addrspace(1)* addrspace(5)* %test.addr, align 8
+    // AMDGCN: store i32 8, i32 addrspace(5)* %s, align 4
 }