[AMDGPU] Reserve extra SGPR blocks wth XNACK "any" TID Setting

[kerbowa]

Summary

ASMPrinter was relying on feature bits to setup extra SGRPs in the knerel
descriptor for the xnack_mask. This was broken for the dynamic XNACK "any" TID
setting which could cause user SGPRs to be clobbered if the number of SGPRs
reserved was near a granulated block boundary.

When XNACK was enabled this worked correctly in the ASMParser which meant some
kernels were only failing without "-save-temps".

Fixes: SWDEV-382764

[foad]

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll

4

Could you pipe the binary into llvm-readelf --notes - instead of llvm-objdump, and then match text instead of hex dumps?

[kerbowa]

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll

4

The notes section contains metadata that may not match the kernel descriptor (KD) in .rodata. In a failed test, the metadata was the same in both passing and failing cases, but the KD was different. I can add an extra runline though.

[kerbowa]

Add readelf run-line to test.

[kzhuravl]

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll

4

What was the differfence in the kernel descriptor between passing and failing case?

[kerbowa]

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll

4

The difference was the granulated SGPR field. It was 1 in the failing case, 2 in the passing version. I can't really replicate it here since the only reason it was passing with save-temps, is that the ASM parser was correctly checking the dynamic target ID setting and encoding the correct kernel descriptor. We already have MC tests for that. The metadata and what the ASM printer is doing has always been incorrect.

[kzhuravl]

LGTM, but please add a todo or a fixme (see comment below), thanks

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll

4

There are several internal discussions related to asm directives and I think we are going to end up adding the assembler directive for granulated sgpr count. Once that directive is added we can change this test from checking hex to checking that directory. Would you mind adding a TODO or a FIXME saying this test needs to be updated to use the directive instead of hex once the directive is available?

[kerbowa]

Actually, this is breaking tests with non-HSA. Is TargetID relevant for pal/graphics/ect @foad, or should the default there be XNACK- in the absence of any explicit subtarget features being added?

[kerbowa]

Update tests.

[kerbowa]

Added AMDGPU group to reviewers.

Is there any objection to changing the defaults for subtargets that support XNACK to always reserve extra SGPRs unless -xnack is explicitly requested? This would impact graphics as well. The old defaults were doing the opposite and only reserving the extra SGPRs with +xnack, meaning the default in the absence of either +/-xnack will be changing.

clang/test/Frontend/amdgcn-machine-analysis-remarks.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx908 -Rpass-analysis=kernel-resource-usage -S -O0 -verify %s -o /dev/null
 
 // expected-remark@+9 {{Function Name: foo}}
-// expected-remark@+8 {{    SGPRs: 9}}
+// expected-remark@+8 {{    SGPRs: 13}}
 // expected-remark@+7 {{    VGPRs: 10}}
 // expected-remark@+6 {{    AGPRs: 12}}
 // expected-remark@+5 {{    ScratchSize [bytes/lane]: 0}}
 // expected-remark@+4 {{    Occupancy [waves/SIMD]: 10}}
 // expected-remark@+3 {{    SGPRs Spill: 0}}
 // expected-remark@+2 {{    VGPRs Spill: 0}}
 // expected-remark@+1 {{    LDS Size [bytes/block]: 0}}
 __kernel void foo() {
   __asm volatile ("; clobber s8" :::"s8");
   __asm volatile ("; clobber v9" :::"v9");
   __asm volatile ("; clobber a11" :::"a11");
 }

llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp

void AMDGPUAsmPrinter::emitFunctionBodyEnd() {
   const GCNSubtarget &STM = MF->getSubtarget<GCNSubtarget>();
 
   SmallString<128> KernelName;
   getNameWithPrefix(KernelName, &MF->getFunction());
   getTargetStreamer()->EmitAmdhsaKernelDescriptor(
       STM, KernelName, getAmdhsaKernelDescriptor(*MF, CurrentProgramInfo),
       CurrentProgramInfo.NumVGPRsForWavesPerEU,
       CurrentProgramInfo.NumSGPRsForWavesPerEU -
-          IsaInfo::getNumExtraSGPRs(&STM,
-                                    CurrentProgramInfo.VCCUsed,
-                                    CurrentProgramInfo.FlatUsed),
+          IsaInfo::getNumExtraSGPRs(
+              &STM, CurrentProgramInfo.VCCUsed, CurrentProgramInfo.FlatUsed,
+              getTargetStreamer()->getTargetID()->isXnackOnOrAny()),
       CurrentProgramInfo.VCCUsed, CurrentProgramInfo.FlatUsed,
       CodeObjectVersion);
 
   Streamer.popSection();
 }
 
 void AMDGPUAsmPrinter::emitFunctionEntryLabel() {
   if (TM.getTargetTriple().getOS() == Triple::AMDHSA &&
void AMDGPUAsmPrinter::getSIProgramInfo(SIProgramInfo &ProgInfo,
   }
 
   const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();
 
   // The calculations related to SGPR/VGPR blocks are
   // duplicated in part in AMDGPUAsmParser::calculateGPRBlocks, and could be
   // unified.
   unsigned ExtraSGPRs = IsaInfo::getNumExtraSGPRs(
-      &STM, ProgInfo.VCCUsed, ProgInfo.FlatUsed);
+      &STM, ProgInfo.VCCUsed, ProgInfo.FlatUsed,
+      getTargetStreamer()->getTargetID()->isXnackOnOrAny());
 
   // Check the addressable register limit before we add ExtraSGPRs.
   if (STM.getGeneration() >= AMDGPUSubtarget::VOLCANIC_ISLANDS &&
       !STM.hasSGPRInitBug()) {
     unsigned MaxAddressableNumSGPRs = STM.getAddressableNumSGPRs();
     if (ProgInfo.NumSGPR > MaxAddressableNumSGPRs) {
       // This can happen due to a compiler bug or when using inline asm.
       LLVMContext &Ctx = MF.getFunction().getContext();

llvm/test/CodeGen/AMDGPU/GlobalISel/extractelement.ll

 ; GPRIDX-NEXT:     is_dynamic_callstack = 0
 ; GPRIDX-NEXT:     is_debug_enabled = 0
 ; GPRIDX-NEXT:     is_xnack_enabled = 1
 ; GPRIDX-NEXT:     workitem_private_segment_byte_size = 0
 ; GPRIDX-NEXT:     workgroup_group_segment_byte_size = 0
 ; GPRIDX-NEXT:     gds_segment_byte_size = 0
 ; GPRIDX-NEXT:     kernarg_segment_byte_size = 12
 ; GPRIDX-NEXT:     workgroup_fbarrier_count = 0
-; GPRIDX-NEXT:     wavefront_sgpr_count = 9
+; GPRIDX-NEXT:     wavefront_sgpr_count = 13
 ; GPRIDX-NEXT:     workitem_vgpr_count = 3
 ; GPRIDX-NEXT:     reserved_vgpr_first = 0
 ; GPRIDX-NEXT:     reserved_vgpr_count = 0
 ; GPRIDX-NEXT:     reserved_sgpr_first = 0
 ; GPRIDX-NEXT:     reserved_sgpr_count = 0
 ; GPRIDX-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; GPRIDX-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; GPRIDX-NEXT:     kernarg_segment_alignment = 4

 ; GPRIDX-NEXT:     amd_code_version_minor = 2
 ; GPRIDX-NEXT:     amd_machine_kind = 1
 ; GPRIDX-NEXT:     amd_machine_version_major = 9
 ; GPRIDX-NEXT:     amd_machine_version_minor = 0
 ; GPRIDX-NEXT:     amd_machine_version_stepping = 0
 ; GPRIDX-NEXT:     kernel_code_entry_byte_offset = 256
 ; GPRIDX-NEXT:     kernel_code_prefetch_byte_size = 0
 ; GPRIDX-NEXT:     granulated_workitem_vgpr_count = 0
-; GPRIDX-NEXT:     granulated_wavefront_sgpr_count = 0
+; GPRIDX-NEXT:     granulated_wavefront_sgpr_count = 1
 ; GPRIDX-NEXT:     priority = 0
 ; GPRIDX-NEXT:     float_mode = 240
 ; GPRIDX-NEXT:     priv = 0
 ; GPRIDX-NEXT:     enable_dx10_clamp = 1
 ; GPRIDX-NEXT:     debug_mode = 0
 ; GPRIDX-NEXT:     enable_ieee_mode = 1
 ; GPRIDX-NEXT:     enable_wgp_mode = 0
 ; GPRIDX-NEXT:     enable_mem_ordered = 0

 ; GPRIDX-NEXT:     is_dynamic_callstack = 0
 ; GPRIDX-NEXT:     is_debug_enabled = 0
 ; GPRIDX-NEXT:     is_xnack_enabled = 1
 ; GPRIDX-NEXT:     workitem_private_segment_byte_size = 0
 ; GPRIDX-NEXT:     workgroup_group_segment_byte_size = 0
 ; GPRIDX-NEXT:     gds_segment_byte_size = 0
 ; GPRIDX-NEXT:     kernarg_segment_byte_size = 12
 ; GPRIDX-NEXT:     workgroup_fbarrier_count = 0
-; GPRIDX-NEXT:     wavefront_sgpr_count = 6
+; GPRIDX-NEXT:     wavefront_sgpr_count = 10
 ; GPRIDX-NEXT:     workitem_vgpr_count = 2
 ; GPRIDX-NEXT:     reserved_vgpr_first = 0
 ; GPRIDX-NEXT:     reserved_vgpr_count = 0
 ; GPRIDX-NEXT:     reserved_sgpr_first = 0
 ; GPRIDX-NEXT:     reserved_sgpr_count = 0
 ; GPRIDX-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; GPRIDX-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; GPRIDX-NEXT:     kernarg_segment_alignment = 4

 ; GPRIDX-NEXT:     amd_code_version_minor = 2
 ; GPRIDX-NEXT:     amd_machine_kind = 1
 ; GPRIDX-NEXT:     amd_machine_version_major = 9
 ; GPRIDX-NEXT:     amd_machine_version_minor = 0
 ; GPRIDX-NEXT:     amd_machine_version_stepping = 0
 ; GPRIDX-NEXT:     kernel_code_entry_byte_offset = 256
 ; GPRIDX-NEXT:     kernel_code_prefetch_byte_size = 0
 ; GPRIDX-NEXT:     granulated_workitem_vgpr_count = 0
-; GPRIDX-NEXT:     granulated_wavefront_sgpr_count = 0
+; GPRIDX-NEXT:     granulated_wavefront_sgpr_count = 1
 ; GPRIDX-NEXT:     priority = 0
 ; GPRIDX-NEXT:     float_mode = 240
 ; GPRIDX-NEXT:     priv = 0
 ; GPRIDX-NEXT:     enable_dx10_clamp = 1
 ; GPRIDX-NEXT:     debug_mode = 0
 ; GPRIDX-NEXT:     enable_ieee_mode = 1
 ; GPRIDX-NEXT:     enable_wgp_mode = 0
 ; GPRIDX-NEXT:     enable_mem_ordered = 0

 ; GPRIDX-NEXT:     is_dynamic_callstack = 0
 ; GPRIDX-NEXT:     is_debug_enabled = 0
 ; GPRIDX-NEXT:     is_xnack_enabled = 1
 ; GPRIDX-NEXT:     workitem_private_segment_byte_size = 0
 ; GPRIDX-NEXT:     workgroup_group_segment_byte_size = 0
 ; GPRIDX-NEXT:     gds_segment_byte_size = 0
 ; GPRIDX-NEXT:     kernarg_segment_byte_size = 12
 ; GPRIDX-NEXT:     workgroup_fbarrier_count = 0
-; GPRIDX-NEXT:     wavefront_sgpr_count = 7
+; GPRIDX-NEXT:     wavefront_sgpr_count = 11
 ; GPRIDX-NEXT:     workitem_vgpr_count = 3
 ; GPRIDX-NEXT:     reserved_vgpr_first = 0
 ; GPRIDX-NEXT:     reserved_vgpr_count = 0
 ; GPRIDX-NEXT:     reserved_sgpr_first = 0
 ; GPRIDX-NEXT:     reserved_sgpr_count = 0
 ; GPRIDX-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; GPRIDX-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; GPRIDX-NEXT:     kernarg_segment_alignment = 4

llvm/test/CodeGen/AMDGPU/amdpal-callable.ll

-; RUN: llc -mtriple=amdgcn--amdpal -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,SDAG,GFX8 -enable-var-scope %s
-; RUN: llc -mtriple=amdgcn--amdpal -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,SDAG,GFX9 -enable-var-scope %s
-; RUN: llc -global-isel -mtriple=amdgcn--amdpal -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GISEL,GFX9 -enable-var-scope %s
+; RUN: llc -mtriple=amdgcn--amdpal -mattr=-xnack -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,SDAG,GFX8 -enable-var-scope %s
+; RUN: llc -mtriple=amdgcn--amdpal -mcpu=gfx900 -mattr=-xnack -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,SDAG,GFX9 -enable-var-scope %s
+; RUN: llc -global-isel -mtriple=amdgcn--amdpal -mattr=-xnack -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefixes=GCN,GISEL,GFX9 -enable-var-scope %s
 
 declare amdgpu_gfx float @extern_func(float) #0
 declare amdgpu_gfx float @extern_func_many_args(<64 x float>) #0
 
 @funcptr = external hidden unnamed_addr addrspace(4) constant ptr, align 4
 
 define amdgpu_gfx float @no_stack(float %arg0) #0 {
   %add = fadd float %arg0, 1.0

llvm/test/CodeGen/AMDGPU/occupancy-levels.ll

-; RUN: llc -march=amdgcn -mcpu=gfx900 < %s | FileCheck --check-prefixes=GCN,GFX9 %s
-; RUN: llc -march=amdgcn -mcpu=gfx1010 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1010,GFX1010W32 %s
-; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1010,GFX1010W64 %s
+; RUN: llc -march=amdgcn -mcpu=gfx900 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX9 %s
+; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=-xnack < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1010,GFX1010W32 %s
+; RUN: llc -march=amdgcn -mcpu=gfx1010 -mattr=-xnack -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1010,GFX1010W64 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1030 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1030,GFX1030W32 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1030 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1030,GFX1030W64 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1100 < %s | FileCheck --check-prefixes=GCN,GFX1100,GFX1100W32 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1100 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX1100,GFX1100W64 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1101 < %s | FileCheck --check-prefixes=GCN,GFX1100,GFX1100W32 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1101 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX1100,GFX1100W64 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1102 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W32,GFX1030,GFX1030W32 %s
 ; RUN: llc -march=amdgcn -mcpu=gfx1102 -mattr=+wavefrontsize64 < %s | FileCheck --check-prefixes=GCN,GFX10,GFX10W64,GFX1030,GFX1030W64 %s

llvm/test/CodeGen/AMDGPU/resource-optimization-remarks.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx908 -pass-remarks-output=%t -pass-remarks-analysis=kernel-resource-usage -filetype=obj -o /dev/null %s 2>&1 | FileCheck -check-prefix=STDERR %s
 ; RUN: FileCheck -check-prefix=REMARK %s < %t
 
 ; STDERR: remark: foo.cl:27:0: Function Name: test_kernel
-; STDERR-NEXT: remark: foo.cl:27:0:     SGPRs: 24
+; STDERR-NEXT: remark: foo.cl:27:0:     SGPRs: 28
 ; STDERR-NEXT: remark: foo.cl:27:0:     VGPRs: 9
 ; STDERR-NEXT: remark: foo.cl:27:0:     AGPRs: 43
 ; STDERR-NEXT: remark: foo.cl:27:0:     ScratchSize [bytes/lane]: 0
 ; STDERR-NEXT: remark: foo.cl:27:0:     Occupancy [waves/SIMD]: 5
 ; STDERR-NEXT: remark: foo.cl:27:0:     SGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:27:0:     VGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:27:0:     LDS Size [bytes/block]: 512
 
 ; REMARK-LABEL: --- !Analysis
 ; REMARK: Pass:            kernel-resource-usage
 ; REMARK-NEXT: Name:            FunctionName
 ; REMARK-NEXT: DebugLoc:        { File: foo.cl, Line: 27, Column: 0 }
 ; REMARK-NEXT: Function:        test_kernel
 ; REMARK-NEXT: Args:
 ; REMARK-NEXT:   - String:          'Function Name: '
 ; REMARK-NEXT:   - FunctionName:      test_kernel
 ; REMARK-NEXT: ...
 ; REMARK-NEXT: --- !Analysis
 ; REMARK-NEXT: Pass:            kernel-resource-usage
 ; REMARK-NEXT: Name:            NumSGPR
 ; REMARK-NEXT: DebugLoc:        { File: foo.cl, Line: 27, Column: 0 }
 ; REMARK-NEXT: Function:        test_kernel
 ; REMARK-NEXT: Args:
 ; REMARK-NEXT:   - String:          '    SGPRs: '
-; REMARK-NEXT:   - NumSGPR:         '24'
+; REMARK-NEXT:   - NumSGPR:         '28'
 ; REMARK-NEXT: ...
 ; REMARK-NEXT: --- !Analysis
 ; REMARK-NEXT: Pass:            kernel-resource-usage
 ; REMARK-NEXT: Name:            NumVGPR
 ; REMARK-NEXT: DebugLoc:        { File: foo.cl, Line: 27, Column: 0 }
 ; REMARK-NEXT: Function:        test_kernel
 ; REMARK-NEXT: Args:
 ; REMARK-NEXT:   - String:          '    VGPRs: '

 define void @test_func() !dbg !6 {
   call void asm sideeffect "; clobber v17", "~{v17}"()
   call void asm sideeffect "; clobber s11", "~{s11}"()
   call void asm sideeffect "; clobber a9", "~{a9}"()
   ret void
 }
 
 ; STDERR: remark: foo.cl:8:0: Function Name: empty_kernel
-; STDERR-NEXT: remark: foo.cl:8:0:     SGPRs: 0
+; STDERR-NEXT: remark: foo.cl:8:0:     SGPRs: 4
 ; STDERR-NEXT: remark: foo.cl:8:0:     VGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     AGPRs: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     ScratchSize [bytes/lane]: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     Occupancy [waves/SIMD]: 8
 ; STDERR-NEXT: remark: foo.cl:8:0:     SGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     VGPRs Spill: 0
 ; STDERR-NEXT: remark: foo.cl:8:0:     LDS Size [bytes/block]: 0
 define amdgpu_kernel void @empty_kernel() !dbg !7 {

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-any.ll

+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+
+; TODO: Update to check for granulated sgpr count directive once one is added.
+
+define amdgpu_kernel void @kern() {
+; ASM-LABEL: kern:
+; ASM: .amdhsa_next_free_sgpr 5
+; ASM: .amdhsa_reserve_xnack_mask 1
+
+; Verify that an extra SGPR block is reserved with XNACK "any" tid setting.
+; OBJ: Contents of section .rodata:
+; OBJ-NEXT: 0000 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0010 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0020 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0030 4000af00 88000000 01000000 00000000  @...............
+
+; ELF: AMDGPU Metadata
+; ELF: .sgpr_count:     9
+entry:
+  tail call void asm sideeffect "", "~{s[0:4]}"()
+  ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-off.ll

+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=-xnack --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+
+; TODO: Update to check for granulated sgpr count directive once one is added.
+
+define amdgpu_kernel void @kern() {
+; ASM-LABEL: kern:
+; ASM: .amdhsa_next_free_sgpr 5
+; ASM: .amdhsa_reserve_xnack_mask 0
+
+; Verify that an extra SGPR block is not reserved with XNACK "off" tid setting.
+; OBJ: Contents of section .rodata:
+; OBJ-NEXT: 0000 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0010 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0020 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0030 0000af00 88000000 01000000 00000000  ................
+
+; ELF: AMDGPU Metadata
+; ELF: .sgpr_count:     5
+entry:
+  tail call void asm sideeffect "", "~{s[0:4]}"()
+  ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/tid-kd-xnack-on.ll

+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack < %s | FileCheck --check-prefixes=ASM %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack --filetype=obj < %s | llvm-objdump -s -j .rodata - | FileCheck --check-prefixes=OBJ %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx90a -mattr=+xnack --filetype=obj < %s | llvm-readelf --notes - | FileCheck --check-prefixes=ELF %s
+
+; TODO: Update to check for granulated sgpr count directive once one is added.
+
+define amdgpu_kernel void @kern() {
+; ASM-LABEL: kern:
+; ASM: .amdhsa_next_free_sgpr 5
+; ASM: .amdhsa_reserve_xnack_mask 1
+
+; Verify that an extra SGPR block is reserved with XNACK "on" tid setting.
+; OBJ: Contents of section .rodata:
+; OBJ-NEXT: 0000 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0010 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0020 00000000 00000000 00000000 00000000  ................
+; OBJ-NEXT: 0030 4000af00 88000000 01000000 00000000  @...............
+
+; ELF: AMDGPU Metadata
+; ELF: .sgpr_count:     9
+entry:
+  tail call void asm sideeffect "", "~{s[0:4]}"()
+  ret void
+}
+
+!llvm.module.flags = !{!0}
+!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}

llvm/test/CodeGen/AMDGPU/trap-abis.ll

 ; RUN: sed 's/CODE_OBJECT_VERSION/300/g' %s | llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=-trap-handler -verify-machineinstrs | FileCheck --check-prefix=HSA-NOTRAP-GFX900-V3 %s
 ; RUN: sed 's/CODE_OBJECT_VERSION/400/g' %s | llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -mattr=-trap-handler -verify-machineinstrs | FileCheck --check-prefix=HSA-NOTRAP-GFX900-V4 %s
 
 declare void @llvm.trap() #0
 declare void @llvm.debugtrap() #1
 
 define amdgpu_kernel void @trap(ptr addrspace(1) nocapture readonly %arg0) {
 ; NOHSA-TRAP-GFX900-V2-LABEL: trap:
-; NOHSA-TRAP-GFX900-V2:       ; %bb.0:
+; NOHSA-TRAP-GFX900-V2:         .amd_kernel_code_t
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_code_version_major = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     priority = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     float_mode = 240
+; NOHSA-TRAP-GFX900-V2-NEXT:     priv = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_mode = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_fwd_progress = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_wave_byte_offset = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     user_sgpr_count = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_trap_handler = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_x = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_y = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_z = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_info = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_vgpr_workitem_id = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_exception_msb = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_lds_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_exception = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_buffer = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_dispatch_ptr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_queue_ptr = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_kernarg_segment_ptr = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_dispatch_id = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_flat_scratch_init = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_x = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_y = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_z = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_wavefront_size32 = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_ordered_append_gds = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     private_element_size = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_ptr64 = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 44
+; NOHSA-TRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 8
+; NOHSA-TRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 2
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     group_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     private_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     wavefront_size = 6
+; NOHSA-TRAP-GFX900-V2-NEXT:     call_convention = -1
+; NOHSA-TRAP-GFX900-V2-NEXT:     runtime_loader_kernel_symbol = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:    .end_amd_kernel_code_t
+; NOHSA-TRAP-GFX900-V2-NEXT:  ; %bb.0:
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_load_dwordx2 s[0:1], s[2:3], 0x24
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_mov_b32_e32 v0, 0
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_mov_b32_e32 v1, 1
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_waitcnt lgkmcnt(0)
 ; NOHSA-TRAP-GFX900-V2-NEXT:    global_store_dword v0, v1, s[0:1]
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_waitcnt vmcnt(0)
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_endpgm
 ;

 ; HSA-TRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
 ; HSA-TRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
-; HSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; HSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     priority = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     float_mode = 240
 ; HSA-TRAP-GFX900-V2-NEXT:     priv = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_mode = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0

 ; HSA-TRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 8
 ; HSA-TRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
-; HSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 8
+; HSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 12
 ; HSA-TRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 2
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4

 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
-; HSA-NOTRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; HSA-NOTRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     priority = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     float_mode = 240
 ; HSA-NOTRAP-GFX900-V2-NEXT:     priv = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_mode = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0

 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 8
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
-; HSA-NOTRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 8
+; HSA-NOTRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 12
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 2
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4
; HSA-NOTRAP-GFX900-V4-NEXT:    s_endpgm
   call void @llvm.trap()
   unreachable
   store volatile i32 2, ptr addrspace(1) %arg0
   ret void
 }
 
 define amdgpu_kernel void @non_entry_trap(ptr addrspace(1) nocapture readonly %arg0) local_unnamed_addr {
 ; NOHSA-TRAP-GFX900-V2-LABEL: non_entry_trap:
-; NOHSA-TRAP-GFX900-V2:       ; %bb.0: ; %entry
+; NOHSA-TRAP-GFX900-V2:         .amd_kernel_code_t
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_code_version_major = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     priority = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     float_mode = 240
+; NOHSA-TRAP-GFX900-V2-NEXT:     priv = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_mode = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_fwd_progress = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_wave_byte_offset = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     user_sgpr_count = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_trap_handler = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_x = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_y = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_z = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_info = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_vgpr_workitem_id = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_exception_msb = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_lds_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_exception = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_buffer = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_dispatch_ptr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_queue_ptr = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_kernarg_segment_ptr = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_dispatch_id = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_flat_scratch_init = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_x = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_y = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_z = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_wavefront_size32 = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_ordered_append_gds = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     private_element_size = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_ptr64 = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 44
+; NOHSA-TRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 8
+; NOHSA-TRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 2
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     group_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     private_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     wavefront_size = 6
+; NOHSA-TRAP-GFX900-V2-NEXT:     call_convention = -1
+; NOHSA-TRAP-GFX900-V2-NEXT:     runtime_loader_kernel_symbol = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:    .end_amd_kernel_code_t
+; NOHSA-TRAP-GFX900-V2-NEXT:  ; %bb.0: ; %entry
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_load_dwordx2 s[0:1], s[2:3], 0x24
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_mov_b32_e32 v0, 0
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_waitcnt lgkmcnt(0)
 ; NOHSA-TRAP-GFX900-V2-NEXT:    global_load_dword v1, v0, s[0:1] glc
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_waitcnt vmcnt(0)
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_cmp_eq_u32_e32 vcc, -1, v1
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_cbranch_vccz .LBB1_2
 ; NOHSA-TRAP-GFX900-V2-NEXT:  ; %bb.1: ; %ret

 ; HSA-TRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 8
 ; HSA-TRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
-; HSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 10
+; HSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 12
 ; HSA-TRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 2
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4

 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 8
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
-; HSA-NOTRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 10
+; HSA-NOTRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 12
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 2
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4

 
 ret:
   store volatile i32 3, ptr addrspace(1) %arg0
   ret void
 }
 
 define amdgpu_kernel void @debugtrap(ptr addrspace(1) nocapture readonly %arg0) {
 ; NOHSA-TRAP-GFX900-V2-LABEL: debugtrap:
-; NOHSA-TRAP-GFX900-V2:       ; %bb.0:
+; NOHSA-TRAP-GFX900-V2:         .amd_kernel_code_t
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_code_version_major = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     priority = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     float_mode = 240
+; NOHSA-TRAP-GFX900-V2-NEXT:     priv = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_mode = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_fwd_progress = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_wave_byte_offset = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     user_sgpr_count = 2
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_trap_handler = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_x = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_y = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_id_z = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_workgroup_info = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_vgpr_workitem_id = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_exception_msb = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     granulated_lds_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_exception = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_buffer = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_dispatch_ptr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_queue_ptr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_kernarg_segment_ptr = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_dispatch_id = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_flat_scratch_init = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_private_segment_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_x = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_y = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_sgpr_grid_workgroup_count_z = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_wavefront_size32 = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     enable_ordered_append_gds = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     private_element_size = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_ptr64 = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
+; NOHSA-TRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 44
+; NOHSA-TRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 6
+; NOHSA-TRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 3
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     group_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     private_segment_alignment = 4
+; NOHSA-TRAP-GFX900-V2-NEXT:     wavefront_size = 6
+; NOHSA-TRAP-GFX900-V2-NEXT:     call_convention = -1
+; NOHSA-TRAP-GFX900-V2-NEXT:     runtime_loader_kernel_symbol = 0
+; NOHSA-TRAP-GFX900-V2-NEXT:    .end_amd_kernel_code_t
+; NOHSA-TRAP-GFX900-V2-NEXT:  ; %bb.0:
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_load_dwordx2 s[0:1], s[0:1], 0x24
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_mov_b32_e32 v0, 0
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_mov_b32_e32 v1, 1
 ; NOHSA-TRAP-GFX900-V2-NEXT:    v_mov_b32_e32 v2, 2
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_waitcnt lgkmcnt(0)
 ; NOHSA-TRAP-GFX900-V2-NEXT:    global_store_dword v0, v1, s[0:1]
 ; NOHSA-TRAP-GFX900-V2-NEXT:    s_waitcnt vmcnt(0)
 ; NOHSA-TRAP-GFX900-V2-NEXT:    global_store_dword v0, v2, s[0:1]

 ; HSA-TRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
 ; HSA-TRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
-; HSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; HSA-TRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     priority = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     float_mode = 240
 ; HSA-TRAP-GFX900-V2-NEXT:     priv = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_mode = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0

 ; HSA-TRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
 ; HSA-TRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 8
 ; HSA-TRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
-; HSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 6
+; HSA-TRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 10
 ; HSA-TRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 3
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; HSA-TRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4

 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_code_version_minor = 2
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_kind = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_version_major = 9
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_version_minor = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     amd_machine_version_stepping = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernel_code_entry_byte_offset = 256
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernel_code_prefetch_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     granulated_workitem_vgpr_count = 0
-; HSA-NOTRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 0
+; HSA-NOTRAP-GFX900-V2-NEXT:     granulated_wavefront_sgpr_count = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     priority = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     float_mode = 240
 ; HSA-NOTRAP-GFX900-V2-NEXT:     priv = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_dx10_clamp = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_mode = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_ieee_mode = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_wgp_mode = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     enable_mem_ordered = 0

 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_dynamic_callstack = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_debug_enabled = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     is_xnack_enabled = 1
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workitem_private_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workgroup_group_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     gds_segment_byte_size = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernarg_segment_byte_size = 8
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workgroup_fbarrier_count = 0
-; HSA-NOTRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 6
+; HSA-NOTRAP-GFX900-V2-NEXT:     wavefront_sgpr_count = 10
 ; HSA-NOTRAP-GFX900-V2-NEXT:     workitem_vgpr_count = 3
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_vgpr_first = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_vgpr_count = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_sgpr_first = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     reserved_sgpr_count = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_wavefront_private_segment_offset_sgpr = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     debug_private_segment_buffer_sgpr = 0
 ; HSA-NOTRAP-GFX900-V2-NEXT:     kernarg_segment_alignment = 4