Skip to content

Commit 0a8d17e

Browse files
saiislamsaiislam
authored
[AMDGPU] Make default AMDHSA Code Object Version to be 5 (#65410)
Also update LIT tests and docs. For more details, see https://llvm.org/docs/AMDGPUUsage.html#code-object-v5-metadata Reviewed By: arsenm, jhuber6 Github PR: #65410 Differential Revision: https://reviews.llvm.org/D129818
1 parent 2318bc8 commit 0a8d17e

File tree

109 files changed

+9194
-9722
lines changed

Some content is hidden

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

109 files changed

+9194
-9722
lines changed

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

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

46184618
def mcode_object_version_EQ : Joined<["-"], "mcode-object-version=">, Group<m_Group>,
4619-
HelpText<"Specify code object ABI version. Defaults to 4. (AMDGPU only)">,
4619+
HelpText<"Specify code object ABI version. Defaults to 5. (AMDGPU only)">,
46204620
Visibility<[ClangOption, CC1Option]>,
46214621
Values<"none,2,3,4,5">,
46224622
NormalizedValuesScope<"TargetOptions">,
46234623
NormalizedValues<["COV_None", "COV_2", "COV_3", "COV_4", "COV_5"]>,
4624-
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_4">;
4624+
MarshallingInfoEnum<TargetOpts<"CodeObjectVersion">, "COV_5">;
46254625

46264626
defm cumode : SimpleMFlag<"cumode",
46274627
"Specify CU wavefront", "Specify WGP wavefront",

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

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

23422342
unsigned tools::getAMDGPUCodeObjectVersion(const Driver &D,
23432343
const llvm::opt::ArgList &Args) {
2344-
unsigned CodeObjVer = 4; // default
2344+
unsigned CodeObjVer = 5; // default
23452345
if (auto *CodeObjArg = getAMDGPUCodeObjectArgument(D, Args))
23462346
StringRef(CodeObjArg->getValue()).getAsInteger(0, CodeObjVer);
23472347
return CodeObjVer;

‎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=V4
4+
// RUN: -o - %s | FileCheck %s -check-prefix=V5
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: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,10 @@
11
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
2-
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
2+
// RUN: -fcuda-is-device -mcode-object-version=4 -emit-llvm -o - -x hip %s \
33
// RUN: | FileCheck -check-prefix=PRECOV5 %s
44

55

66
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
7-
// RUN: -fcuda-is-device -mcode-object-version=5 -emit-llvm -o - -x hip %s \
7+
// RUN: -fcuda-is-device -emit-llvm -o - -x hip %s \
88
// RUN: | FileCheck -check-prefix=COV5 %s
99

1010
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \

‎clang/test/CodeGenHIP/default-attributes.hip

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -46,11 +46,11 @@ __global__ void kernel() {
4646
// OPT: attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
4747
// OPT: attributes #1 = { mustprogress nofree norecurse nosync nounwind willreturn memory(none) "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" }
4848
//.
49-
// OPTNONE: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
49+
// OPTNONE: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
5050
// OPTNONE: !1 = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
5151
// OPTNONE: !2 = !{i32 1, !"wchar_size", i32 4}
5252
//.
53-
// OPT: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
53+
// OPT: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
5454
// OPT: !1 = !{i32 1, !"amdgpu_printf_kind", !"hostcall"}
5555
// OPT: !2 = !{i32 1, !"wchar_size", i32 4}
5656
//.

‎clang/test/CodeGenOpenCL/amdgpu-enqueue-kernel.cl

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -703,7 +703,7 @@ kernel void test_target_features_kernel(global int *i) {
703703
// GFX900: attributes #8 = { nounwind }
704704
// GFX900: attributes #9 = { convergent nounwind }
705705
//.
706-
// NOCPU: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
706+
// NOCPU: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
707707
// NOCPU: !1 = !{i32 1, !"wchar_size", i32 4}
708708
// NOCPU: !2 = !{i32 2, i32 0}
709709
// NOCPU: !3 = !{i32 1, i32 0, i32 1, i32 0}
@@ -721,7 +721,7 @@ kernel void test_target_features_kernel(global int *i) {
721721
// NOCPU: !15 = !{i32 1}
722722
// NOCPU: !16 = !{!"int*"}
723723
//.
724-
// GFX900: !0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
724+
// GFX900: !0 = !{i32 1, !"amdgpu_code_object_version", i32 500}
725725
// GFX900: !1 = !{i32 1, !"wchar_size", i32 4}
726726
// GFX900: !2 = !{i32 2, i32 0}
727727
// GFX900: !3 = !{!4, !4, i64 0}

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

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

601601
// CHECK-LABEL: @test_get_workgroup_size(
602-
// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
603-
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 4
602+
// CHECK: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
603+
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 12
604604
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
605-
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 6
605+
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 14
606606
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
607-
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 8
608-
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
607+
// CHECK: getelementptr inbounds i8, ptr addrspace(4) %{{.*}}, i64 16
608+
// CHECK: load i16, ptr addrspace(4) %{{.*}}, align 8, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
609609
void test_get_workgroup_size(int d, global int *out)
610610
{
611611
switch (d) {

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

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

165-
// Test default code object version with old device library without abi_version_400.bc
166-
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
165+
// Test default code object version with old device library without abi_version_500.bc
166+
// RUN: not %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
167167
// RUN: --hip-device-lib-path=%S/Inputs/rocm/amdgcn/bitcode-no-abi-ver \
168168
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
169-
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI4
169+
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI5
170170

171171
// Test -mcode-object-version=3
172172
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
@@ -193,12 +193,12 @@
193193
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
194194
// RUN: 2>&1 | FileCheck %s --check-prefixes=ABI5
195195

196-
// Test -mcode-object-version=5 with old device library without abi_version_400.bc
197-
// RUN: not %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
198-
// RUN: -mcode-object-version=5 \
196+
// Test -mcode-object-version=4 with old device library without abi_version_400.bc
197+
// RUN: %clang -### --target=x86_64-linux-gnu --offload-arch=gfx900 \
198+
// RUN: -mcode-object-version=4 \
199199
// RUN: --hip-device-lib-path=%S/Inputs/rocm/amdgcn/bitcode-no-abi-ver \
200200
// RUN: --rocm-path=%S/Inputs/rocm %S/Inputs/hip_multiple_inputs/b.hip \
201-
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI5
201+
// RUN: 2>&1 | FileCheck %s --check-prefixes=NOABI4
202202

203203
// ALL-NOT: error:
204204
// ALL: {{"[^"]*clang[^"]*"}}

‎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: 2
16+
# CHECK-NEXT: ABIVersion: 3
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: 2
18+
; AMDHSA: ABIVersion: 3
1919

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

‎llvm/docs/AMDGPUUsage.rst

Lines changed: 7 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1418,12 +1418,12 @@ The AMDGPU backend uses the following ELF header:
14181418

14191419
* ``ELFABIVERSION_AMDGPU_HSA_V4`` is used to specify the version of AMD HSA
14201420
runtime ABI for code object V4. Specify using the Clang option
1421-
``-mcode-object-version=4``. This is the default code object
1422-
version if not specified.
1421+
``-mcode-object-version=4``.
14231422

14241423
* ``ELFABIVERSION_AMDGPU_HSA_V5`` is used to specify the version of AMD HSA
14251424
runtime ABI for code object V5. Specify using the Clang option
1426-
``-mcode-object-version=5``.
1425+
``-mcode-object-version=5``. This is the default code object
1426+
version if not specified.
14271427

14281428
* ``ELFABIVERSION_AMDGPU_PAL`` is used to specify the version of AMD PAL
14291429
runtime ABI.
@@ -3852,6 +3852,10 @@ same *vendor-name*.
38523852
Code Object V4 Metadata
38533853
+++++++++++++++++++++++
38543854

3855+
. warning::
3856+
Code object V4 is not the default code object version emitted by this version
3857+
of LLVM.
3858+
38553859
Code object V4 metadata is the same as
38563860
:ref:`amdgpu-amdhsa-code-object-metadata-v3` with the changes and additions
38573861
defined in table :ref:`amdgpu-amdhsa-code-object-metadata-map-table-v4`.
@@ -3882,11 +3886,6 @@ defined in table :ref:`amdgpu-amdhsa-code-object-metadata-map-table-v4`.
38823886
Code Object V5 Metadata
38833887
+++++++++++++++++++++++
38843888

3885-
.. warning::
3886-
Code object V5 is not the default code object version emitted by this version
3887-
of LLVM.
3888-
3889-
38903889
Code object V5 metadata is the same as
38913890
:ref:`amdgpu-amdhsa-code-object-metadata-v4` with the changes defined in table
38923891
:ref:`amdgpu-amdhsa-code-object-metadata-map-table-v5`, table

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@
3434
static llvm::cl::opt<unsigned>
3535
AmdhsaCodeObjectVersion("amdhsa-code-object-version", llvm::cl::Hidden,
3636
llvm::cl::desc("AMDHSA Code Object Version"),
37-
llvm::cl::init(4));
37+
llvm::cl::init(5));
3838

3939
namespace {
4040

@@ -177,7 +177,7 @@ unsigned getCodeObjectVersion(const Module &M) {
177177
}
178178

179179
// Default code object version.
180-
return AMDHSA_COV4;
180+
return AMDHSA_COV5;
181181
}
182182

183183
unsigned getMultigridSyncArgImplicitArgPosition(unsigned CodeObjectVersion) {

‎llvm/test/CodeGen/AMDGPU/GlobalISel/crash-stack-address-O0.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
define amdgpu_kernel void @stack_write_fi() {
88
; CHECK-LABEL: stack_write_fi:
99
; CHECK: ; %bb.0: ; %entry
10-
; CHECK-NEXT: s_add_u32 s0, s0, s17
10+
; CHECK-NEXT: s_add_u32 s0, s0, s15
1111
; CHECK-NEXT: s_addc_u32 s1, s1, 0
1212
; CHECK-NEXT: s_mov_b32 s5, 0
1313
; CHECK-NEXT: s_mov_b32 s4, 0

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

Lines changed: 28 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -7,43 +7,42 @@ 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: $sgpr14, $sgpr15, $sgpr16, $vgpr0, $vgpr1, $vgpr2, $sgpr4_sgpr5, $sgpr6_sgpr7, $sgpr8_sgpr9, $sgpr10_sgpr11
10+
; CHECK-NEXT: liveins: $sgpr12, $sgpr13, $sgpr14, $vgpr0, $vgpr1, $vgpr2, $sgpr4_sgpr5, $sgpr6_sgpr7, $sgpr8_sgpr9
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 $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
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
2221
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc, 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
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
3029
; CHECK-NEXT: [[S_MOV_B32_:%[0-9]+]]:sreg_32 = S_MOV_B32 10, debug-location !6
31-
; CHECK-NEXT: [[COPY16:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_]], debug-location !6
32-
; CHECK-NEXT: [[V_LSHLREV_B32_e64_:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY16]], [[COPY1]], implicit $exec, debug-location !6
30+
; CHECK-NEXT: [[COPY14:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_]], debug-location !6
31+
; CHECK-NEXT: [[V_LSHLREV_B32_e64_:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY14]], [[COPY1]], implicit $exec, debug-location !6
3332
; CHECK-NEXT: [[S_MOV_B32_1:%[0-9]+]]:sreg_32 = S_MOV_B32 20, debug-location !6
34-
; CHECK-NEXT: [[COPY17:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_1]], debug-location !6
35-
; CHECK-NEXT: [[V_LSHLREV_B32_e64_1:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY17]], [[COPY]], implicit $exec, debug-location !6
33+
; CHECK-NEXT: [[COPY15:%[0-9]+]]:vgpr_32 = COPY [[S_MOV_B32_1]], debug-location !6
34+
; CHECK-NEXT: [[V_LSHLREV_B32_e64_1:%[0-9]+]]:vgpr_32 = V_LSHLREV_B32_e64 [[COPY15]], [[COPY]], implicit $exec, debug-location !6
3635
; 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
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
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
4746
; CHECK-NEXT: $vgpr31 = COPY [[V_OR3_B32_e64_]], debug-location !6
4847
; 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
4948
; 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 = COPY $sgpr6_sgpr7
35+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = 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]]
41+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
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 = COPY $sgpr6_sgpr7
84+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = 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]]
90+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
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 = COPY $sgpr6_sgpr7
134+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = 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]]
140+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
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 = COPY $sgpr6_sgpr7
184+
; CHECK-NEXT: [[COPY7:%[0-9]+]]:sgpr_64(p4) = COPY $sgpr6_sgpr7
185185
; CHECK-NEXT: [[COPY8:%[0-9]+]]:sgpr_64 = COPY $sgpr4_sgpr5
186186
; CHECK-NEXT: [[GV:%[0-9]+]]:ccr_sgpr_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]]
188+
; CHECK-NEXT: [[COPY10:%[0-9]+]]:_(p4) = COPY [[COPY7]](p4)
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)