r325279 - Clean up AMDGCN tests
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Thu Feb 15 11:12:41 PST 2018
Author: yaxunl
Date: Thu Feb 15 11:12:41 2018
New Revision: 325279
URL: http://llvm.org/viewvc/llvm-project?rev=325279&view=rev
Log:
Clean up AMDGCN tests
Differential Revision: https://reviews.llvm.org/D43340
Modified:
cfe/trunk/test/CodeGen/address-space.c
cfe/trunk/test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp
cfe/trunk/test/CodeGenCXX/vla.cpp
cfe/trunk/test/CodeGenOpenCL/addr-space-struct-arg.cl
cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl
cfe/trunk/test/CodeGenOpenCL/address-spaces.cl
cfe/trunk/test/CodeGenOpenCL/blocks.cl
cfe/trunk/test/CodeGenOpenCL/lifetime.cl
cfe/trunk/test/CodeGenOpenCL/vla.cl
cfe/trunk/test/Index/pipe-size.cl
Modified: cfe/trunk/test/CodeGen/address-space.c
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGen/address-space.c?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGen/address-space.c (original)
+++ cfe/trunk/test/CodeGen/address-space.c Thu Feb 15 11:12:41 2018
@@ -1,5 +1,5 @@
-// 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)));
@@ -24,10 +24,10 @@ __attribute__((address_space(2))) int *A
// 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;
@@ -39,8 +39,8 @@ typedef struct {
} 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;
Modified: cfe/trunk/test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp (original)
+++ cfe/trunk/test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp Thu Feb 15 11:12:41 2018
@@ -1,5 +1,5 @@
// 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;
@@ -49,8 +49,8 @@ struct 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
@@ -67,8 +67,8 @@ std::initializer_list<int> thread_local
// 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
@@ -77,18 +77,18 @@ std::initializer_list<int> thread_local
// 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
@@ -101,15 +101,15 @@ std::initializer_list<int> thread_local
// 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
@@ -121,8 +121,8 @@ 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
@@ -518,12 +518,12 @@ namespace B19773010 {
// 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}}};
}
Modified: cfe/trunk/test/CodeGenCXX/vla.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCXX/vla.cpp?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCXX/vla.cpp (original)
+++ cfe/trunk/test/CodeGenCXX/vla.cpp Thu Feb 15 11:12:41 2018
@@ -1,5 +1,5 @@
// 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 {
@@ -19,17 +19,17 @@ int f() {
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
@@ -68,8 +68,8 @@ void test0(void *array, int n) {
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:%.*]]
@@ -87,15 +87,15 @@ void test2(int b) {
//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:%.*]]
@@ -120,7 +120,7 @@ void test3(int b, int c) {
//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;
}
Modified: cfe/trunk/test/CodeGenOpenCL/addr-space-struct-arg.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/addr-space-struct-arg.cl?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/addr-space-struct-arg.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/addr-space-struct-arg.cl Thu Feb 15 11:12:41 2018
@@ -1,5 +1,5 @@
// 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];
@@ -37,7 +37,7 @@ struct LargeStructTwoMember {
// 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;
@@ -49,15 +49,15 @@ Mat4X4 __attribute__((noinline)) foo(Mat
// 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;
@@ -68,66 +68,66 @@ Mat64X64 __attribute__((noinline)) foo_l
// 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);
}
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=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/address-space-constant-initializers.cl Thu Feb 15 11:12:41 2018
@@ -1,5 +1,5 @@
// 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;
@@ -14,8 +14,8 @@ typedef struct {
// 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
Modified: cfe/trunk/test/CodeGenOpenCL/address-spaces.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-spaces.cl?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-spaces.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/address-spaces.cl Thu Feb 15 11:12:41 2018
@@ -1,9 +1,9 @@
// 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)* }
@@ -24,7 +24,7 @@ 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
@@ -34,11 +34,11 @@ void f__g(__global int *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
@@ -48,7 +48,7 @@ void fg(global int *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
@@ -56,20 +56,20 @@ 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;
Modified: cfe/trunk/test/CodeGenOpenCL/blocks.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/blocks.cl?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/blocks.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/blocks.cl Thu Feb 15 11:12:41 2018
@@ -1,14 +1,14 @@
// 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;
};
@@ -21,13 +21,13 @@ void foo(){
// 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
@@ -43,21 +43,21 @@ void foo(){
// 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;
@@ -69,9 +69,9 @@ void foo(){
// 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
Modified: cfe/trunk/test/CodeGenOpenCL/lifetime.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/lifetime.cl?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/lifetime.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/lifetime.cl Thu Feb 15 11:12:41 2018
@@ -1,5 +1,5 @@
// 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);
@@ -10,6 +10,6 @@ __attribute__((always_inline)) void help
void lifetime_test() {
// CHECK: @llvm.lifetime.start.p0i
-// AMDGIZ: @llvm.lifetime.start.p5i
+// AMDGCN: @llvm.lifetime.start.p5i
helper_no_markers();
}
Modified: cfe/trunk/test/CodeGenOpenCL/vla.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/vla.cl?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/vla.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/vla.cl Thu Feb 15 11:12:41 2018
@@ -1,14 +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 -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()
@@ -16,10 +16,10 @@ 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];
}
Modified: cfe/trunk/test/Index/pipe-size.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/Index/pipe-size.cl?rev=325279&r1=325278&r2=325279&view=diff
==============================================================================
--- cfe/trunk/test/Index/pipe-size.cl (original)
+++ cfe/trunk/test/Index/pipe-size.cl Thu Feb 15 11:12:41 2018
@@ -1,7 +1,7 @@
// 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);
@@ -11,6 +11,6 @@ __kernel void testPipe( pipe int test )
// 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
}
More information about the cfe-commits
mailing list