Index: test/CodeGen/address-space.c =================================================================== --- test/CodeGen/address-space.c +++ test/CodeGen/address-space.c @@ -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 @@ // 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 @@ } 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; Index: test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp =================================================================== --- test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp +++ test/CodeGenCXX/cxx0x-initializer-stdinitializerlist.cpp @@ -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 @@ }; // 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 globalInitList1 = {1, 2, 3}; #ifndef NO_TLS @@ -67,8 +67,8 @@ // 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 @@ // 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 @@ // 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 @@ // 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 @@ // 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> a, p[2] = {a, {{"", ENUM_CONSTANT}}}; } Index: test/CodeGenCXX/vla.cpp =================================================================== --- test/CodeGenCXX/vla.cpp +++ test/CodeGenCXX/vla.cpp @@ -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 struct S { @@ -19,17 +19,17 @@ 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 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 @@ //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 @@ //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; } Index: test/CodeGenOpenCL/addr-space-struct-arg.cl =================================================================== --- test/CodeGenOpenCL/addr-space-struct-arg.cl +++ test/CodeGenOpenCL/addr-space-struct-arg.cl @@ -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 @@ // 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 @@ // 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 @@ // 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); } Index: test/CodeGenOpenCL/address-space-constant-initializers.cl =================================================================== --- test/CodeGenOpenCL/address-space-constant-initializers.cl +++ test/CodeGenOpenCL/address-space-constant-initializers.cl @@ -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 @@ // 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 Index: test/CodeGenOpenCL/address-spaces.cl =================================================================== --- test/CodeGenOpenCL/address-spaces.cl +++ test/CodeGenOpenCL/address-spaces.cl @@ -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 @@ #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__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 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 @@ // 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; Index: test/CodeGenOpenCL/blocks.cl =================================================================== --- test/CodeGenOpenCL/blocks.cl +++ test/CodeGenOpenCL/blocks.cl @@ -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 @@ // 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 @@ // 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 @@ // 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 Index: test/CodeGenOpenCL/lifetime.cl =================================================================== --- test/CodeGenOpenCL/lifetime.cl +++ test/CodeGenOpenCL/lifetime.cl @@ -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 @@ void lifetime_test() { // CHECK: @llvm.lifetime.start.p0i -// AMDGIZ: @llvm.lifetime.start.p5i +// AMDGCN: @llvm.lifetime.start.p5i helper_no_markers(); } Index: test/CodeGenOpenCL/vla.cl =================================================================== --- test/CodeGenOpenCL/vla.cl +++ test/CodeGenOpenCL/vla.cl @@ -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 @@ 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]; } Index: test/Index/pipe-size.cl =================================================================== --- test/Index/pipe-size.cl +++ test/Index/pipe-size.cl @@ -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 @@ // 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 }