Skip to content

Commit ca856ff

Browse files
ronliebronlieb
committed
Revert "enable code-object-version=5"
very sorry wrong repo. This reverts commit d882ba7.
1 parent b09a5e5 commit ca856ff

File tree

58 files changed

+5474
-5407
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

58 files changed

+5474
-5407
lines changed

‎clang/include/clang/Driver/Options.td

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3690,12 +3690,12 @@ defm amdgpu_ieee : BoolOption<"m", "amdgpu-ieee",
36903690
NegFlag<SetFalse, [CC1Option]>>, Group<m_Group>;
36913691

36923692
def mcode_object_version_EQ : Joined<["-"], "mcode-object-version=">, Group<m_Group>,
3693-
HelpText<"Specify code object ABI version. Defaults to 5. (AMDGPU only)">,
3693+
HelpText<"Specify code object ABI version. Defaults to 4. (AMDGPU only)">,
36943694
Flags<[CC1Option]>,
36953695
Values<"none,2,3,4,5">,
36963696
NormalizedValuesScope<"TargetOptions">,
36973697
NormalizedValues<["COV_None", "COV_2", "COV_3", "COV_4", "COV_5"]>,
3698-
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_5">;
3698+
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_4">;
36993699

37003700
defm code_object_v3_legacy : SimpleMFlag<"code-object-v3",
37013701
"Legacy option to specify code object ABI V3",

‎clang/lib/Driver/ToolChains/CommonArgs.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2157,7 +2157,7 @@ void tools::checkAMDGPUCodeObjectVersion(const Driver &D,
21572157

21582158
unsigned tools::getAMDGPUCodeObjectVersion(const Driver &D,
21592159
const llvm::opt::ArgList &Args) {
2160-
unsigned CodeObjVer = 5; // default
2160+
unsigned CodeObjVer = 4; // default
21612161
if (auto *CodeObjArg = getAMDGPUCodeObjectArgument(D, Args)) {
21622162
if (CodeObjArg->getOption().getID() ==
21632163
options::OPT_mno_code_object_v3_legacy) {

‎clang/test/CodeGenCUDA/amdgpu-code-object-version.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
// Create module flag for code object version.
22

33
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
4-
// RUN: -o - %s | FileCheck %s -check-prefix=V5
4+
// RUN: -o - %s | FileCheck %s -check-prefix=V4
55

66
// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm \
77
// RUN: -mcode-object-version=2 -o - %s | FileCheck -check-prefix=V2 %s

‎clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa \
2-
// RUN: -fcuda-is-device -mcode-object-version=4 -emit-llvm -o - -x hip %s \
2+
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
33
// RUN: | FileCheck -check-prefix=PRECOV5 %s
44

55

‎clang/test/CodeGenOpenCL/builtins-amdgcn.cl

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -583,13 +583,13 @@ void test_get_local_id(int d, global int *out)
583583
}
584584

585585
// CHECK-LABEL: @test_get_workgroup_size(
586-
// CHECK: call align 8 dereferenceable(256) i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
587-
// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 12
586+
// CHECK: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
587+
// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 4
588588
// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load
589-
// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 14
589+
// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 6
590590
// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
591-
// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 16
592-
// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 8, !range [[$WS_RANGE:![0-9]*]], !invariant.load
591+
// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i64 8
592+
// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load
593593
void test_get_workgroup_size(int d, global int *out)
594594
{
595595
switch (d) {

‎clang/test/Driver/hip-device-libs.hip

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -139,13 +139,13 @@
139139
// Test default code object version.
140140
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
141141
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
142-
// RUN: 2>&1 | FileCheck %s --check-prefixes=ABI5
142+
// RUN: 2>&1 | FileCheck %s --check-prefixes=ABI4
143143

144144
// Test default code object version with old device library without abi_version_400.bc
145145
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
146146
// RUN: --hip-device-lib-path=%S/Inputs/rocm/amdgcn/bitcode-no-abi-ver \
147147
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
148-
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI5
148+
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI4
149149

150150
// Test -mcode-object-version=3
151151
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \

‎lld/test/ELF/emulation-amdgpu.s

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@
1313
# CHECK-NEXT: DataEncoding: LittleEndian (0x1)
1414
# CHECK-NEXT: FileVersion: 1
1515
# CHECK-NEXT: OS/ABI: AMDGPU_HSA (0x40)
16-
# CHECK-NEXT: ABIVersion: 3
16+
# CHECK-NEXT: ABIVersion: 2
1717
# CHECK-NEXT: Unused: (00 00 00 00 00 00 00)
1818
# CHECK-NEXT: }
1919
# CHECK-NEXT: Type: Executable (0x2)

‎lld/test/ELF/lto/amdgcn-oses.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -15,7 +15,7 @@
1515
; RUN: llvm-readobj --file-headers %t/mesa3d.so | FileCheck %s --check-prefixes=GCN,NON-AMDHSA,MESA3D
1616

1717
; AMDHSA: OS/ABI: AMDGPU_HSA (0x40)
18-
; AMDHSA: ABIVersion: 3
18+
; AMDHSA: ABIVersion: 2
1919

2020
; AMDPAL: OS/ABI: AMDGPU_PAL (0x41)
2121
; MESA3D: OS/ABI: AMDGPU_MESA3D (0x42)

‎llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@
3232
static llvm::cl::opt<unsigned>
3333
AmdhsaCodeObjectVersion("amdhsa-code-object-version", llvm::cl::Hidden,
3434
llvm::cl::desc("AMDHSA Code Object Version"),
35-
llvm::cl::init(5));
35+
llvm::cl::init(4));
3636

3737
namespace {
3838

‎llvm/test/CodeGen/AMDGPU/GlobalISel/dropped_debug_info_assert.ll

Lines changed: 29 additions & 28 deletions
Original file line numberDiff line numberDiff line change
@@ -7,42 +7,43 @@ declare void @callee()
77
define amdgpu_kernel void @call_debug_loc() {
88
; CHECK-LABEL: name: call_debug_loc
99
; CHECK: bb.1.entry:
10-
; CHECK-NEXT: liveins: $sgpr12, $sgpr13, $sgpr14, $vgpr0, $vgpr1, $vgpr2, $sgpr4_sgpr5, $sgpr6_sgpr7, $sgpr8_sgpr9
10+
; CHECK-NEXT: liveins: $sgpr14, $sgpr15, $sgpr16, $vgpr0, $vgpr1, $vgpr2, $sgpr4_sgpr5, $sgpr6_sgpr7, $sgpr8_sgpr9, $sgpr10_sgpr11
1111
; CHECK-NEXT: {{ $}}
1212
; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr_32 = COPY $vgpr2, debug-location !6
1313
; CHECK-NEXT: [[COPY1:%[0-9]+]]:vgpr_32 = COPY $vgpr1, debug-location !6
1414
; CHECK-NEXT: [[COPY2:%[0-9]+]]:vgpr_32 = COPY $vgpr0, debug-location !6
15-
; CHECK-NEXT: [[COPY3:%[0-9]+]]:sgpr_32 = COPY $sgpr14, debug-location !6
16-
; CHECK-NEXT: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr13, debug-location !6
17-
; CHECK-NEXT: [[COPY5:%[0-9]+]]:sgpr_32 = COPY $sgpr12, debug-location !6
18-
; CHECK-NEXT: [[COPY6:%[0-9]+]]:sgpr_64 = COPY $sgpr8_sgpr9, debug-location !6
19-
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5, debug-location !6
20-
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sreg_64 = COPY $sgpr6_sgpr7
15+
; CHECK-NEXT: [[COPY3:%[0-9]+]]:sgpr_32 = COPY $sgpr16, debug-location !6
16+
; CHECK-NEXT: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr15, debug-location !6
17+
; CHECK-NEXT: [[COPY5:%[0-9]+]]:sgpr_32 = COPY $sgpr14, debug-location !6
18+
; CHECK-NEXT: [[COPY6:%[0-9]+]]:sgpr_64 = COPY $sgpr10_sgpr11, debug-location !6
19+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64 = COPY $sgpr6_sgpr7, debug-location !6
20+
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5, debug-location !6
21+
; CHECK-NEXT: [[COPY9:%[0-9]+]]:sreg_64 = COPY $sgpr8_sgpr9
2122
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc, debug-location !6
22-
; CHECK-NEXT: [[COPY9:%[0-9]+]]:sreg_64 = COPY [[COPY7]], debug-location !6
23-
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64 = IMPLICIT_DEF debug-location !6
24-
; CHECK-NEXT: [[COPY10:%[0-9]+]]:sreg_64 = COPY [[COPY6]], debug-location !6
25-
; CHECK-NEXT: [[COPY11:%[0-9]+]]:sreg_32 = COPY [[COPY5]], debug-location !6
26-
; CHECK-NEXT: [[COPY12:%[0-9]+]]:sreg_32 = COPY [[COPY4]], debug-location !6
27-
; CHECK-NEXT: [[COPY13:%[0-9]+]]:sreg_32 = COPY [[COPY3]], debug-location !6
28-
; CHECK-NEXT: [[DEF1:%[0-9]+]]:sreg_32 = IMPLICIT_DEF debug-location !6
23+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:sreg_64 = COPY [[COPY8]], debug-location !6
24+
; CHECK-NEXT: [[COPY11:%[0-9]+]]:sreg_64 = COPY [[COPY7]], debug-location !6
25+
; CHECK-NEXT: [[COPY12:%[0-9]+]]:sreg_64 = COPY [[COPY6]], debug-location !6
26+
; CHECK-NEXT: [[COPY13:%[0-9]+]]:sreg_32 = COPY [[COPY5]], debug-location !6
27+
; CHECK-NEXT: [[COPY14:%[0-9]+]]:sreg_32 = COPY [[COPY4]], debug-location !6
28+
; CHECK-NEXT: [[COPY15:%[0-9]+]]:sreg_32 = COPY [[COPY3]], debug-location !6
29+
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_32 = IMPLICIT_DEF debug-location !6
2930
; CHECK-NEXT: [[S_MOV_B32_:%[0-9]+]]:sreg_32 = S_MOV_B32 10
30-
; CHECK-NEXT: [[COPY14:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_]]
31-
; CHECK-NEXT: [[V_LSHLREV_B32_e64_:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY14]], [[COPY1]], implicit $exec, debug-location !6
31+
; CHECK-NEXT: [[COPY16:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_]]
32+
; CHECK-NEXT: [[V_LSHLREV_B32_e64_:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY16]], [[COPY1]], implicit $exec, debug-location !6
3233
; CHECK-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_MOV_B32 20
33-
; CHECK-NEXT: [[COPY15:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_1]]
34-
; CHECK-NEXT: [[V_LSHLREV_B32_e64_1:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY15]], [[COPY]], implicit $exec, debug-location !6
34+
; CHECK-NEXT: [[COPY17:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_1]]
35+
; CHECK-NEXT: [[V_LSHLREV_B32_e64_1:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY17]], [[COPY]], implicit $exec, debug-location !6
3536
; CHECK-NEXT: [[V_OR3_B32_e64_:%[0-9]+]]:vgpr_32 = V_OR3_B32_e64 [[COPY2]], [[V_LSHLREV_B32_e64_]], [[V_LSHLREV_B32_e64_1]], implicit $exec, debug-location !6
36-
; CHECK-NEXT: [[COPY16:%[0-9]+]]:sgpr_128 = COPY $sgpr0_sgpr1_sgpr2_sgpr3, debug-location !6
37-
; CHECK-NEXT: $sgpr0_sgpr1_sgpr2_sgpr3 = COPY [[COPY16]], debug-location !6
38-
; CHECK-NEXT: $sgpr4_sgpr5 = COPY [[COPY9]], debug-location !6
39-
; CHECK-NEXT: $sgpr6_sgpr7 = COPY [[DEF]], debug-location !6
40-
; CHECK-NEXT: $sgpr8_sgpr9 = COPY [[COPY8]], debug-location !6
41-
; CHECK-NEXT: $sgpr10_sgpr11 = COPY [[COPY10]], debug-location !6
42-
; CHECK-NEXT: $sgpr12 = COPY [[COPY11]], debug-location !6
43-
; CHECK-NEXT: $sgpr13 = COPY [[COPY12]], debug-location !6
44-
; CHECK-NEXT: $sgpr14 = COPY [[COPY13]], debug-location !6
45-
; CHECK-NEXT: $sgpr15 = COPY [[DEF1]], debug-location !6
37+
; CHECK-NEXT: [[COPY18:%[0-9]+]]:sgpr_128 = COPY $sgpr0_sgpr1_sgpr2_sgpr3, debug-location !6
38+
; CHECK-NEXT: $sgpr0_sgpr1_sgpr2_sgpr3 = COPY [[COPY18]], debug-location !6
39+
; CHECK-NEXT: $sgpr4_sgpr5 = COPY [[COPY10]], debug-location !6
40+
; CHECK-NEXT: $sgpr6_sgpr7 = COPY [[COPY11]], debug-location !6
41+
; CHECK-NEXT: $sgpr8_sgpr9 = COPY [[COPY9]], debug-location !6
42+
; CHECK-NEXT: $sgpr10_sgpr11 = COPY [[COPY12]], debug-location !6
43+
; CHECK-NEXT: $sgpr12 = COPY [[COPY13]], debug-location !6
44+
; CHECK-NEXT: $sgpr13 = COPY [[COPY14]], debug-location !6
45+
; CHECK-NEXT: $sgpr14 = COPY [[COPY15]], debug-location !6
46+
; CHECK-NEXT: $sgpr15 = COPY [[DEF]], debug-location !6
4647
; CHECK-NEXT: $vgpr31 = COPY [[V_OR3_B32_e64_]], debug-location !6
4748
; CHECK-NEXT: [[SI_PC_ADD_REL_OFFSET:%[0-9]+]]:sreg_64 = SI_PC_ADD_REL_OFFSET target-flags(amdgpu-gotprel32-lo) @callee + 4, target-flags(amdgpu-gotprel32-hi) @callee + 12, implicit-def $scc, debug-location !6
4849
; CHECK-NEXT: [[S_LOAD_DWORDX2_IMM:%[0-9]+]]:sreg_64_xexec = S_LOAD_DWORDX2_IMM [[SI_PC_ADD_REL_OFFSET]], 0, 0, debug-location !6 :: (dereferenceable invariant load (p0) from got, addrspace 4)

‎llvm/test/CodeGen/AMDGPU/GlobalISel/irtranslator-assert-align.ll

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -32,13 +32,13 @@ define void @call_result_align_1() {
3232
; CHECK-NEXT: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr12
3333
; CHECK-NEXT: [[COPY5:%[0-9]+]]:sgpr_64 = COPY $sgpr10_sgpr11
3434
; CHECK-NEXT: [[COPY6:%[0-9]+]]:sgpr_64 = COPY $sgpr8_sgpr9
35-
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr6_sgpr7
35+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64 = COPY $sgpr6_sgpr7
3636
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5
3737
; CHECK-NEXT: [[C:%[0-9]+]]:_(s8) = G_CONSTANT i8 0
3838
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
3939
; CHECK-NEXT: [[GV:%[0-9]+]]:_(p0) = G_GLOBAL_VALUE @returns_ptr
4040
; CHECK-NEXT: [[COPY9:%[0-9]+]]:_(p4) = COPY [[COPY8]]
41-
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
41+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]]
4242
; CHECK-NEXT: [[COPY11:%[0-9]+]]:_(p4) = COPY [[COPY6]]
4343
; CHECK-NEXT: [[COPY12:%[0-9]+]]:_(s64) = COPY [[COPY5]]
4444
; CHECK-NEXT: [[COPY13:%[0-9]+]]:_(s32) = COPY [[COPY4]]
@@ -81,13 +81,13 @@ define void @call_result_align_8() {
8181
; CHECK-NEXT: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr12
8282
; CHECK-NEXT: [[COPY5:%[0-9]+]]:sgpr_64 = COPY $sgpr10_sgpr11
8383
; CHECK-NEXT: [[COPY6:%[0-9]+]]:sgpr_64 = COPY $sgpr8_sgpr9
84-
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr6_sgpr7
84+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64 = COPY $sgpr6_sgpr7
8585
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5
8686
; CHECK-NEXT: [[C:%[0-9]+]]:_(s8) = G_CONSTANT i8 0
8787
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
8888
; CHECK-NEXT: [[GV:%[0-9]+]]:_(p0) = G_GLOBAL_VALUE @returns_ptr
8989
; CHECK-NEXT: [[COPY9:%[0-9]+]]:_(p4) = COPY [[COPY8]]
90-
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
90+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]]
9191
; CHECK-NEXT: [[COPY11:%[0-9]+]]:_(p4) = COPY [[COPY6]]
9292
; CHECK-NEXT: [[COPY12:%[0-9]+]]:_(s64) = COPY [[COPY5]]
9393
; CHECK-NEXT: [[COPY13:%[0-9]+]]:_(s32) = COPY [[COPY4]]
@@ -131,13 +131,13 @@ define void @declaration_result_align_8() {
131131
; CHECK-NEXT: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr12
132132
; CHECK-NEXT: [[COPY5:%[0-9]+]]:sgpr_64 = COPY $sgpr10_sgpr11
133133
; CHECK-NEXT: [[COPY6:%[0-9]+]]:sgpr_64 = COPY $sgpr8_sgpr9
134-
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr6_sgpr7
134+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64 = COPY $sgpr6_sgpr7
135135
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5
136136
; CHECK-NEXT: [[C:%[0-9]+]]:_(s8) = G_CONSTANT i8 0
137137
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
138138
; CHECK-NEXT: [[GV:%[0-9]+]]:_(p0) = G_GLOBAL_VALUE @returns_ptr_align8
139139
; CHECK-NEXT: [[COPY9:%[0-9]+]]:_(p4) = COPY [[COPY8]]
140-
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
140+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]]
141141
; CHECK-NEXT: [[COPY11:%[0-9]+]]:_(p4) = COPY [[COPY6]]
142142
; CHECK-NEXT: [[COPY12:%[0-9]+]]:_(s64) = COPY [[COPY5]]
143143
; CHECK-NEXT: [[COPY13:%[0-9]+]]:_(s32) = COPY [[COPY4]]
@@ -181,11 +181,11 @@ define ptr addrspace(1) @tail_call_assert_align() {
181181
; CHECK-NEXT: [[COPY4:%[0-9]+]]:sgpr_32 = COPY $sgpr12
182182
; CHECK-NEXT: [[COPY5:%[0-9]+]]:sgpr_64 = COPY $sgpr10_sgpr11
183183
; CHECK-NEXT: [[COPY6:%[0-9]+]]:sgpr_64 = COPY $sgpr8_sgpr9
184-
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr6_sgpr7
184+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64 = COPY $sgpr6_sgpr7
185185
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5
186186
; CHECK-NEXT: [[GV:%[0-9]+]]:sreg_64(p0) = G_GLOBAL_VALUE @returns_ptr_align8
187187
; CHECK-NEXT: [[COPY9:%[0-9]+]]:_(p4) = COPY [[COPY8]]
188-
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
188+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]]
189189
; CHECK-NEXT: [[COPY11:%[0-9]+]]:_(p4) = COPY [[COPY6]]
190190
; CHECK-NEXT: [[COPY12:%[0-9]+]]:_(s64) = COPY [[COPY5]]
191191
; CHECK-NEXT: [[COPY13:%[0-9]+]]:_(s32) = COPY [[COPY4]]

0 commit comments

Comments
 (0)