forked from OSchip/llvm-project
Clean up AMDGCN tests
Differential Revision: https://reviews.llvm.org/D43340 llvm-svn: 325279
This commit is contained in:
Yaxun Liu
parent
5bb02f3c02
commit
f8ad59d99d
|
|
@ -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, *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;
|
||||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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 x = {1, 2, 3, 4};
|
|||
|
||||
// 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 x = {1, 2, 3, 4};
|
|||
// 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 x = {1, 2, 3, 4};
|
|||
// 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}}};
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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(Mat3X3 in) {
|
|||
// 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_large(Mat32X32 in) {
|
|||
// 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);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
|
|
|
|||
|
|
@ -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
|
||||
|
|
|
|||
|
|
@ -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 helper_no_markers() {
|
|||
|
||||
void lifetime_test() {
|
||||
// CHECK: @llvm.lifetime.start.p0i
|
||||
// AMDGIZ: @llvm.lifetime.start.p5i
|
||||
// AMDGCN: @llvm.lifetime.start.p5i
|
||||
helper_no_markers();
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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];
|
||||
}
|
||||
|
|
|
|||
|
|
@ -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
|
||||
}
|
||||
|
|
|
|||
Loading…
Reference in New Issue