AMDGPU: Emit user sgpr count directives in text asm

[arsenm]

Summary

We were emitting these in the object file but not printing them.

[rampitec]

LGTM

[arsenm]

e6564f39c787d6aa06ad5573e55db3439ed3a263

llvm/docs/AMDGPUUsage.rst

0       1 bit   ENABLE_PRIVATE_SEGMENT          * Enable the setup of the
                                                        the setup of the
                                                        FLAT_SCRATCH register
                                                        pair (see
                                                        :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
 
                                                      Used by CP to set up
                                                      ``COMPUTE_PGM_RSRC2.SCRATCH_EN``.
      5:1     5 bits  USER_SGPR_COUNT                 The total number of SGPR
-                                                     user data registers
-                                                     requested. This number must
-                                                     match the number of user
+                                                     user data
+                                                     registers requested. This
+                                                     number must be greater than
+                                                     or equal to the number of user
                                                      data registers enabled.
 
                                                      Used by CP to set up
                                                      ``COMPUTE_PGM_RSRC2.USER_SGPR``.
      6       1 bit   ENABLE_TRAP_HANDLER             Must be 0.
 
                                                      This bit represents
                                                      ``COMPUTE_PGM_RSRC2.TRAP_PRESENT``,
.. table:: AMDHSA Kernel Assembler Directives
      Directive                                                Default             Supported On Description
      ======================================================== =================== ============ ===================
      ``.amdhsa_group_segment_fixed_size``                     0                   GFX6-GFX10   Controls GROUP_SEGMENT_FIXED_SIZE in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_private_segment_fixed_size``                   0                   GFX6-GFX10   Controls PRIVATE_SEGMENT_FIXED_SIZE in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_kernarg_size``                                 0                   GFX6-GFX10   Controls KERNARG_SIZE in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
+     ``.amdhsa_user_sgpr_count``                              0                   GFX6-GFX10   Controls USER_SGPR_COUNT in COMPUTE_PGM_RSRC2
+                                                                                               :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`
      ``.amdhsa_user_sgpr_private_segment_buffer``             0                   GFX6-GFX10   Controls ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_user_sgpr_dispatch_ptr``                       0                   GFX6-GFX10   Controls ENABLE_SGPR_DISPATCH_PTR in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_user_sgpr_queue_ptr``                          0                   GFX6-GFX10   Controls ENABLE_SGPR_QUEUE_PTR in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.
      ``.amdhsa_user_sgpr_kernarg_segment_ptr``                0                   GFX6-GFX10   Controls ENABLE_SGPR_KERNARG_SEGMENT_PTR in
                                                                                                :ref:`amdgpu-amdhsa-kernel-descriptor-v3-table`.

llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

bool AMDGPUAsmParser::ParseDirectiveAMDHSAKernel() {
 
   IsaVersion IVersion = getIsaVersion(getSTI().getCPU());
 
   SMRange VGPRRange;
   uint64_t NextFreeVGPR = 0;
   uint64_t AccumOffset = 0;
   SMRange SGPRRange;
   uint64_t NextFreeSGPR = 0;
-  unsigned UserSGPRCount = 0;
+
+  // Count the number of user SGPRs implied from the enabled feature bits.
+  unsigned ImpliedUserSGPRCount = 0;
+
+  // Track if the asm explicitly contains the directive for the user SGPR
+  // count.
+  Optional<unsigned> ExplicitUserSGPRCount;
   bool ReserveVCC = true;
   bool ReserveFlatScr = true;
   Optional<bool> EnableWavefrontSize32;
 
   while (true) {
     while (trySkipToken(AsmToken::EndOfStatement));
 
     StringRef ID;
AMDHSA_BITS_SET(FIELD, ENTRY, VALUE);
     } else if (ID == ".amdhsa_private_segment_fixed_size") {
       if (!isUInt<sizeof(KD.private_segment_fixed_size) * CHAR_BIT>(Val))
         return OutOfRangeError(ValRange);
       KD.private_segment_fixed_size = Val;
     } else if (ID == ".amdhsa_kernarg_size") {
       if (!isUInt<sizeof(KD.kernarg_size) * CHAR_BIT>(Val))
         return OutOfRangeError(ValRange);
       KD.kernarg_size = Val;
+    } else if (ID == ".amdhsa_user_sgpr_count") {
+      ExplicitUserSGPRCount = Val;
     } else if (ID == ".amdhsa_user_sgpr_private_segment_buffer") {
       if (hasArchitectedFlatScratch())
         return Error(IDRange.Start,
                      "directive is not supported with architected flat scratch",
                      IDRange);
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER,
                        Val, ValRange);
       if (Val)
-        UserSGPRCount += 4;
+        ImpliedUserSGPRCount += 4;
     } else if (ID == ".amdhsa_user_sgpr_dispatch_ptr") {
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR, Val,
                        ValRange);
       if (Val)
-        UserSGPRCount += 2;
+        ImpliedUserSGPRCount += 2;
     } else if (ID == ".amdhsa_user_sgpr_queue_ptr") {
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_QUEUE_PTR, Val,
                        ValRange);
       if (Val)
-        UserSGPRCount += 2;
+        ImpliedUserSGPRCount += 2;
     } else if (ID == ".amdhsa_user_sgpr_kernarg_segment_ptr") {
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_KERNARG_SEGMENT_PTR,
                        Val, ValRange);
       if (Val)
-        UserSGPRCount += 2;
+        ImpliedUserSGPRCount += 2;
     } else if (ID == ".amdhsa_user_sgpr_dispatch_id") {
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_ID, Val,
                        ValRange);
       if (Val)
-        UserSGPRCount += 2;
+        ImpliedUserSGPRCount += 2;
     } else if (ID == ".amdhsa_user_sgpr_flat_scratch_init") {
       if (hasArchitectedFlatScratch())
         return Error(IDRange.Start,
                      "directive is not supported with architected flat scratch",
                      IDRange);
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_FLAT_SCRATCH_INIT, Val,
                        ValRange);
       if (Val)
-        UserSGPRCount += 2;
+        ImpliedUserSGPRCount += 2;
     } else if (ID == ".amdhsa_user_sgpr_private_segment_size") {
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_SIZE,
                        Val, ValRange);
       if (Val)
-        UserSGPRCount += 1;
+        ImpliedUserSGPRCount += 1;
     } else if (ID == ".amdhsa_wavefront_size32") {
       if (IVersion.Major < 10)
         return Error(IDRange.Start, "directive requires gfx10+", IDRange);
       EnableWavefrontSize32 = Val;
       PARSE_BITS_ENTRY(KD.kernel_code_properties,
                        KERNEL_CODE_PROPERTY_ENABLE_WAVEFRONT_SIZE32,
                        Val, ValRange);
     } else if (ID == ".amdhsa_system_sgpr_private_segment_wavefront_offset") {
#undef PARSE_BITS_ENTRY
 
   if (!isUInt<COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT_WIDTH>(
           SGPRBlocks))
     return OutOfRangeError(SGPRRange);
   AMDHSA_BITS_SET(KD.compute_pgm_rsrc1,
                   COMPUTE_PGM_RSRC1_GRANULATED_WAVEFRONT_SGPR_COUNT,
                   SGPRBlocks);
 
+  if (ExplicitUserSGPRCount && ImpliedUserSGPRCount > *ExplicitUserSGPRCount)
+    return TokError("amdgpu_user_sgpr_count smaller than than implied by "
+                    "enabled user SGPRs");
+
+  unsigned UserSGPRCount =
+      ExplicitUserSGPRCount ? *ExplicitUserSGPRCount : ImpliedUserSGPRCount;
+
   if (!isUInt<COMPUTE_PGM_RSRC2_USER_SGPR_COUNT_WIDTH>(UserSGPRCount))
     return TokError("too many user SGPRs enabled");
   AMDHSA_BITS_SET(KD.compute_pgm_rsrc2, COMPUTE_PGM_RSRC2_USER_SGPR_COUNT,
                   UserSGPRCount);
 
   if (isGFX90A()) {
     if (Seen.find(".amdhsa_accum_offset") == Seen.end())
       return TokError(".amdhsa_accum_offset directive is required");

llvm/lib/Target/AMDGPU/MCTargetDesc/AMDGPUTargetStreamer.cpp

STREAM << "\t\t" << DIRECTIVE << " "                                         \
          << AMDHSA_BITS_GET(KERNEL_DESC.MEMBER_NAME, FIELD_NAME) << '\n';
 
   OS << "\t\t.amdhsa_group_segment_fixed_size " << KD.group_segment_fixed_size
      << '\n';
   OS << "\t\t.amdhsa_private_segment_fixed_size "
      << KD.private_segment_fixed_size << '\n';
   OS << "\t\t.amdhsa_kernarg_size " << KD.kernarg_size << '\n';
 
+  PRINT_FIELD(OS, ".amdhsa_user_sgpr_count", KD,
+              compute_pgm_rsrc2,
+              amdhsa::COMPUTE_PGM_RSRC2_USER_SGPR_COUNT);
+
   if (!hasArchitectedFlatScratch(STI))
     PRINT_FIELD(
         OS, ".amdhsa_user_sgpr_private_segment_buffer", KD,
         kernel_code_properties,
         amdhsa::KERNEL_CODE_PROPERTY_ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER);
   PRINT_FIELD(OS, ".amdhsa_user_sgpr_dispatch_ptr", KD,
               kernel_code_properties,
               amdhsa::KERNEL_CODE_PROPERTY_ENABLE_SGPR_DISPATCH_PTR);

llvm/test/CodeGen/AMDGPU/code-object-v3.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 < %s | FileCheck --check-prefixes=ALL-ASM,OSABI-AMDHSA-ASM %s
 ; RUN: llc -filetype=obj -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 < %s | llvm-readelf -S -r -s --notes - | FileCheck --check-prefix=OSABI-AMDHSA-ELF %s
 
 ; ALL-ASM-LABEL: {{^}}fadd:
 
 ; OSABI-AMDHSA-ASM-NOT: .hsa_code_object_version
 ; OSABI-AMDHSA-ASM-NOT: .hsa_code_object_isa
 ; OSABI-AMDHSA-ASM-NOT: .amdgpu_hsa_kernel
 ; OSABI-AMDHSA-ASM-NOT: .amd_kernel_code_t
 
 ; OSABI-AMDHSA-ASM: s_endpgm
 ; OSABI-AMDHSA-ASM: .section .rodata,#alloc
 ; OSABI-AMDHSA-ASM: .p2align 6
 ; OSABI-AMDHSA-ASM: .amdhsa_kernel fadd
+; OSABI-AMDHSA-ASM:     .amdhsa_user_sgpr_count 6
 ; OSABI-AMDHSA-ASM:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; OSABI-AMDHSA-ASM:     .amdhsa_user_sgpr_kernarg_segment_ptr 1
 ; OSABI-AMDHSA-ASM:     .amdhsa_next_free_vgpr 3
 ; OSABI-AMDHSA-ASM:     .amdhsa_next_free_sgpr 8
 ; OSABI-AMDHSA-ASM:     .amdhsa_reserve_vcc 0
 ; OSABI-AMDHSA-ASM:     .amdhsa_reserve_flat_scratch 0
 ; OSABI-AMDHSA-ASM: .end_amdhsa_kernel
 ; OSABI-AMDHSA-ASM: .text
 
 ; ALL-ASM-LABEL: {{^}}fsub:
 
 ; OSABI-AMDHSA-ASM-NOT: .amdgpu_hsa_kernel
 ; OSABI-AMDHSA-ASM-NOT: .amd_kernel_code_t
 
 ; OSABI-AMDHSA-ASM: s_endpgm
 ; OSABI-AMDHSA-ASM: .section .rodata,#alloc
 ; OSABI-AMDHSA-ASM: .p2align 6
 ; OSABI-AMDHSA-ASM: .amdhsa_kernel fsub
+; OSABI-AMDHSA-ASM:     .amdhsa_user_sgpr_count 6
 ; OSABI-AMDHSA-ASM:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; OSABI-AMDHSA-ASM:     .amdhsa_user_sgpr_kernarg_segment_ptr 1
 ; OSABI-AMDHSA-ASM:     .amdhsa_next_free_vgpr 3
 ; OSABI-AMDHSA-ASM:     .amdhsa_next_free_sgpr 8
 ; OSABI-AMDHSA-ASM:     .amdhsa_reserve_vcc 0
 ; OSABI-AMDHSA-ASM:     .amdhsa_reserve_flat_scratch 0
 ; OSABI-AMDHSA-ASM: .end_amdhsa_kernel
 ; OSABI-AMDHSA-ASM: .text

llvm/test/CodeGen/AMDGPU/indirect-call-known-callees.ll

 ; CHECK-NEXT:    ; implicit-def: $vgpr31
 ; CHECK-NEXT:    ; implicit-def: $vgpr1
 ; CHECK-NEXT:    s_xor_b64 exec, exec, s[34:35]
 ; CHECK-NEXT:    s_cbranch_execnz .LBB0_1
 ; CHECK-NEXT:  ; %bb.2:
 ; CHECK-NEXT:    s_endpgm
 
 ; CHECK: .amdhsa_kernarg_size 0
+; CHECK-NEXT: .amdhsa_user_sgpr_count 6
 ; CHECK-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1
 ; CHECK-NEXT: .amdhsa_user_sgpr_dispatch_ptr 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_queue_ptr 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_dispatch_id 0
 ; CHECK-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
 ; CHECK-NEXT: .amdhsa_user_sgpr_private_segment_size 0
 ; CHECK-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1

llvm/test/CodeGen/AMDGPU/kernarg-size.ll

 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 < %s | FileCheck --check-prefix=HSA %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=4 < %s | FileCheck --check-prefix=HSA %s
 ; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 < %s | FileCheck --check-prefix=HSA %s
 
 declare void @llvm.trap() #0
 declare void @llvm.debugtrap() #1
 
 ; HSA:      .amdhsa_kernel trap
 ; HSA-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; HSA-NEXT:     .amdhsa_private_segment_fixed_size 0
 ; HSA-NEXT:     .amdhsa_kernarg_size 8
+; HSA-NEXT:     .amdhsa_user_sgpr_count 8
 ; HSA-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; HSA:      .end_amdhsa_kernel
 
 define amdgpu_kernel void @trap(i32 addrspace(1)* nocapture readonly %arg0) {
   store volatile i32 1, i32 addrspace(1)* %arg0
   call void @llvm.trap()
   unreachable
   store volatile i32 2, i32 addrspace(1)* %arg0
   ret void
 }

llvm/test/CodeGen/AMDGPU/stack-realign-kernel.ll

 ; VI-NEXT:    s_waitcnt vmcnt(0)
 ; VI-NEXT:    s_endpgm
 ; VI-NEXT:    .section .rodata,#alloc
 ; VI-NEXT:    .p2align 6
 ; VI-NEXT:    .amdhsa_kernel max_alignment_128
 ; VI-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; VI-NEXT:     .amdhsa_private_segment_fixed_size 256
 ; VI-NEXT:     .amdhsa_kernarg_size 0
+; VI-NEXT:     .amdhsa_user_sgpr_count 6
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_id 0
 ; VI-NEXT:     .amdhsa_user_sgpr_flat_scratch_init 1
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_size 0
 ; VI-NEXT:     .amdhsa_system_sgpr_private_segment_wavefront_offset 1

 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    s_endpgm
 ; GFX9-NEXT:    .section .rodata,#alloc
 ; GFX9-NEXT:    .p2align 6
 ; GFX9-NEXT:    .amdhsa_kernel max_alignment_128
 ; GFX9-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; GFX9-NEXT:     .amdhsa_private_segment_fixed_size 256
 ; GFX9-NEXT:     .amdhsa_kernarg_size 0
+; GFX9-NEXT:     .amdhsa_user_sgpr_count 6
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_id 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_flat_scratch_init 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_size 0
 ; GFX9-NEXT:     .amdhsa_system_sgpr_private_segment_wavefront_offset 1

 ; VI-NEXT:    s_waitcnt vmcnt(0)
 ; VI-NEXT:    s_endpgm
 ; VI-NEXT:    .section .rodata,#alloc
 ; VI-NEXT:    .p2align 6
 ; VI-NEXT:    .amdhsa_kernel stackrealign_attr
 ; VI-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; VI-NEXT:     .amdhsa_private_segment_fixed_size 8
 ; VI-NEXT:     .amdhsa_kernarg_size 0
+; VI-NEXT:     .amdhsa_user_sgpr_count 6
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_id 0
 ; VI-NEXT:     .amdhsa_user_sgpr_flat_scratch_init 1
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_size 0
 ; VI-NEXT:     .amdhsa_system_sgpr_private_segment_wavefront_offset 1

 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    s_endpgm
 ; GFX9-NEXT:    .section .rodata,#alloc
 ; GFX9-NEXT:    .p2align 6
 ; GFX9-NEXT:    .amdhsa_kernel stackrealign_attr
 ; GFX9-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; GFX9-NEXT:     .amdhsa_private_segment_fixed_size 8
 ; GFX9-NEXT:     .amdhsa_kernarg_size 0
+; GFX9-NEXT:     .amdhsa_user_sgpr_count 6
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_id 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_flat_scratch_init 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_size 0
 ; GFX9-NEXT:     .amdhsa_system_sgpr_private_segment_wavefront_offset 1

 ; VI-NEXT:    s_waitcnt vmcnt(0)
 ; VI-NEXT:    s_endpgm
 ; VI-NEXT:    .section .rodata,#alloc
 ; VI-NEXT:    .p2align 6
 ; VI-NEXT:    .amdhsa_kernel alignstack_attr
 ; VI-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; VI-NEXT:     .amdhsa_private_segment_fixed_size 128
 ; VI-NEXT:     .amdhsa_kernarg_size 0
+; VI-NEXT:     .amdhsa_user_sgpr_count 6
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; VI-NEXT:     .amdhsa_user_sgpr_dispatch_id 0
 ; VI-NEXT:     .amdhsa_user_sgpr_flat_scratch_init 1
 ; VI-NEXT:     .amdhsa_user_sgpr_private_segment_size 0
 ; VI-NEXT:     .amdhsa_system_sgpr_private_segment_wavefront_offset 1

 ; GFX9-NEXT:    s_waitcnt vmcnt(0)
 ; GFX9-NEXT:    s_endpgm
 ; GFX9-NEXT:    .section .rodata,#alloc
 ; GFX9-NEXT:    .p2align 6
 ; GFX9-NEXT:    .amdhsa_kernel alignstack_attr
 ; GFX9-NEXT:     .amdhsa_group_segment_fixed_size 0
 ; GFX9-NEXT:     .amdhsa_private_segment_fixed_size 128
 ; GFX9-NEXT:     .amdhsa_kernarg_size 0
+; GFX9-NEXT:     .amdhsa_user_sgpr_count 6
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_buffer 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_queue_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_kernarg_segment_ptr 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_dispatch_id 0
 ; GFX9-NEXT:     .amdhsa_user_sgpr_flat_scratch_init 1
 ; GFX9-NEXT:     .amdhsa_user_sgpr_private_segment_size 0
 ; GFX9-NEXT:     .amdhsa_system_sgpr_private_segment_wavefront_offset 1

llvm/test/MC/AMDGPU/hsa-gfx10-v3.s

.amdhsa_kernel complete
   .amdhsa_exception_fp_ieee_inexact 1
   .amdhsa_exception_int_div_zero 1
 .end_amdhsa_kernel
 
 // ASM: .amdhsa_kernel complete
 // ASM-NEXT: .amdhsa_group_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_private_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_kernarg_size 8
+// ASM-NEXT: .amdhsa_user_sgpr_count 15
 // ASM-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1
 // ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
 // ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
 // ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
 // ASM-NEXT: .amdhsa_wavefront_size32 1

llvm/test/MC/AMDGPU/hsa-v3.s

.amdhsa_kernel complete
   .amdhsa_exception_fp_ieee_inexact 1
   .amdhsa_exception_int_div_zero 1
 .end_amdhsa_kernel
 
 // ASM: .amdhsa_kernel complete
 // ASM-NEXT: .amdhsa_group_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_private_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_kernarg_size 8
+// ASM-NEXT: .amdhsa_user_sgpr_count 15
 // ASM-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1
 // ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
 // ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
 // ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
 // ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
- .name:       amd_kernel_code_t_minimal
       .kernarg_segment_align: 64
       .wavefront_size: 128
       .sgpr_count: 14
       .vgpr_count: 40
       .max_flat_workgroup_size: 256
 .end_amdgpu_metadata
 
 // ASM:       .amdgpu_metadata
-// ASM:      amdhsa.kernels:  
+// ASM:      amdhsa.kernels:
 // ASM:        - .group_segment_fixed_size: 16
 // ASM:          .kernarg_segment_align: 64
 // ASM:          .kernarg_segment_size: 8
 // ASM:          .max_flat_workgroup_size: 256
 // ASM:          .name:           amd_kernel_code_t_test_all
 // ASM:          .private_segment_fixed_size: 32
 // ASM:          .sgpr_count:     14
 // ASM:          .symbol:         'amd_kernel_code_t_test_all@kd'
 // ASM:          .vgpr_count:     40
 // ASM:          .wavefront_size: 128
 // ASM:        - .group_segment_fixed_size: 16
 // ASM:          .kernarg_segment_align: 64
 // ASM:          .kernarg_segment_size: 8
 // ASM:          .max_flat_workgroup_size: 256
 // ASM:          .name:           amd_kernel_code_t_minimal
 // ASM:          .private_segment_fixed_size: 32
 // ASM:          .sgpr_count:     14
 // ASM:          .symbol:         'amd_kernel_code_t_minimal@kd'
 // ASM:          .vgpr_count:     40
 // ASM:          .wavefront_size: 128
-// ASM:      amdhsa.version:  
+// ASM:      amdhsa.version:
 // ASM-NEXT:   - 3
 // ASM-NEXT:   - 0
 // ASM:       .end_amdgpu_metadata

llvm/test/MC/AMDGPU/hsa-v4.s

 // ASM: .end_amdhsa_kernel
 
 // Test that we can specify all available directives with non-default values.
 .p2align 6
 .amdhsa_kernel complete
   .amdhsa_group_segment_fixed_size 1
   .amdhsa_private_segment_fixed_size 1
   .amdhsa_kernarg_size 8
+  .amdhsa_user_sgpr_count 15
   .amdhsa_user_sgpr_private_segment_buffer 1
   .amdhsa_user_sgpr_dispatch_ptr 1
   .amdhsa_user_sgpr_queue_ptr 1
   .amdhsa_user_sgpr_kernarg_segment_ptr 1
   .amdhsa_user_sgpr_dispatch_id 1
   .amdhsa_user_sgpr_flat_scratch_init 1
   .amdhsa_user_sgpr_private_segment_size 1
   .amdhsa_system_sgpr_private_segment_wavefront_offset 1
.amdhsa_kernel complete
   .amdhsa_exception_fp_ieee_inexact 1
   .amdhsa_exception_int_div_zero 1
 .end_amdhsa_kernel
 
 // ASM: .amdhsa_kernel complete
 // ASM-NEXT: .amdhsa_group_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_private_segment_fixed_size 1
 // ASM-NEXT: .amdhsa_kernarg_size 8
+// ASM-NEXT: .amdhsa_user_sgpr_count 15
 // ASM-NEXT: .amdhsa_user_sgpr_private_segment_buffer 1
 // ASM-NEXT: .amdhsa_user_sgpr_dispatch_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_queue_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_kernarg_segment_ptr 1
 // ASM-NEXT: .amdhsa_user_sgpr_dispatch_id 1
 // ASM-NEXT: .amdhsa_user_sgpr_flat_scratch_init 1
 // ASM-NEXT: .amdhsa_user_sgpr_private_segment_size 1
 // ASM-NEXT: .amdhsa_system_sgpr_private_segment_wavefront_offset 1
- .name:       amd_kernel_code_t_minimal
       .kernarg_segment_align: 64
       .wavefront_size: 128
       .sgpr_count: 14
       .vgpr_count: 40
       .max_flat_workgroup_size: 256
 .end_amdgpu_metadata
 
 // ASM:       .amdgpu_metadata
-// ASM:      amdhsa.kernels:  
+// ASM:      amdhsa.kernels:
 // ASM:        - .group_segment_fixed_size: 16
 // ASM:          .kernarg_segment_align: 64
 // ASM:          .kernarg_segment_size: 8
 // ASM:          .max_flat_workgroup_size: 256
 // ASM:          .name:           amd_kernel_code_t_test_all
 // ASM:          .private_segment_fixed_size: 32
 // ASM:          .sgpr_count:     14
 // ASM:          .symbol:         'amd_kernel_code_t_test_all@kd'
 // ASM:          .vgpr_count:     40
 // ASM:          .wavefront_size: 128
 // ASM:        - .group_segment_fixed_size: 16
 // ASM:          .kernarg_segment_align: 64
 // ASM:          .kernarg_segment_size: 8
 // ASM:          .max_flat_workgroup_size: 256
 // ASM:          .name:           amd_kernel_code_t_minimal
 // ASM:          .private_segment_fixed_size: 32
 // ASM:          .sgpr_count:     14
 // ASM:          .symbol:         'amd_kernel_code_t_minimal@kd'
 // ASM:          .vgpr_count:     40
 // ASM:          .wavefront_size: 128
-// ASM:      amdhsa.version:  
+// ASM:      amdhsa.version:
 // ASM-NEXT:   - 3
 // ASM-NEXT:   - 0
 // ASM:       .end_amdgpu_metadata

llvm/test/MC/AMDGPU/user-sgpr-count-diag.s

+// RUN: not llvm-mc --amdhsa-code-object-version=3 -triple amdgcn-amd-amdhsa -mcpu=gfx810 %s 2>&1 >/dev/null | FileCheck -check-prefix=ERR %s
+
+.amdhsa_kernel implied_count_too_low_0
+  .amdhsa_user_sgpr_count 0
+  .amdhsa_user_sgpr_queue_ptr 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+// ERR: [[@LINE+1]]:19: error: amdgpu_user_sgpr_count smaller than than implied by enabled user SGPRs
+.end_amdhsa_kernel
+
+.amdhsa_kernel implied_count_too_low_1
+  .amdhsa_user_sgpr_count 1
+  .amdhsa_user_sgpr_queue_ptr 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+// ERR: [[@LINE+1]]:19: error: amdgpu_user_sgpr_count smaller than than implied by enabled user SGPRs
+.end_amdhsa_kernel

llvm/test/MC/AMDGPU/user-sgpr-count.s

+// RUN: llvm-mc -triple amdgcn-amd-amdhsa -mcpu=gfx904 --amdhsa-code-object-version=3 -mattr=+xnack < %s | FileCheck --check-prefix=ASM %s
+
+.text
+// ASM: .text
+
+.amdgcn_target "amdgcn-amd-amdhsa--gfx904+xnack"
+// ASM: .amdgcn_target "amdgcn-amd-amdhsa--gfx904+xnack"
+
+
+// ASM-LABEL: .amdhsa_kernel user_sgprs_implied_count
+// ASM: .amdhsa_user_sgpr_count 15
+.amdhsa_kernel user_sgprs_implied_count_all
+  .amdhsa_user_sgpr_private_segment_buffer 1
+  .amdhsa_user_sgpr_dispatch_ptr 1
+  .amdhsa_user_sgpr_queue_ptr 1
+  .amdhsa_user_sgpr_kernarg_segment_ptr 1
+  .amdhsa_user_sgpr_dispatch_id 1
+  .amdhsa_user_sgpr_flat_scratch_init 1
+  .amdhsa_user_sgpr_private_segment_size 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+// ASM-LABEL: .amdhsa_kernel user_sgprs_implied_count_0
+// ASM: .amdhsa_user_sgpr_count 7
+.amdhsa_kernel user_sgprs_implied_count_0
+  .amdhsa_user_sgpr_queue_ptr 1
+  .amdhsa_user_sgpr_kernarg_segment_ptr 1
+  .amdhsa_user_sgpr_flat_scratch_init 1
+  .amdhsa_user_sgpr_private_segment_size 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+// ASM-LABEL: .amdhsa_kernel user_sgprs_implied_count_1
+// ASM: .amdhsa_user_sgpr_count 9
+.amdhsa_kernel user_sgprs_implied_count_1
+  .amdhsa_user_sgpr_private_segment_buffer 1
+  .amdhsa_user_sgpr_queue_ptr 1
+  .amdhsa_user_sgpr_kernarg_segment_ptr 1
+  .amdhsa_user_sgpr_private_segment_size 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+
+// ASM-LABEL: .amdhsa_kernel user_sgprs_implied_count_private_segment_buffer
+// ASM: .amdhsa_user_sgpr_count 4
+  .amdhsa_kernel user_sgprs_implied_count_private_segment_buffer
+  .amdhsa_user_sgpr_private_segment_buffer 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+
+// ASM-LABEL: .amdhsa_kernel explicit_user_sgpr_count_16
+.amdhsa_kernel explicit_user_sgpr_count_16
+  .amdhsa_user_sgpr_count 16
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+
+// ASM-LABEL: .amdhsa_kernel explicit_user_sgpr_count_0
+// ASM: .amdhsa_user_sgpr_count 0
+  .amdhsa_kernel explicit_user_sgpr_count_0
+  .amdhsa_user_sgpr_count 0
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+// ASM-LABEL: .amdhsa_kernel explicit_user_sgpr_count_1
+// ASM: .amdhsa_user_sgpr_count 1
+.amdhsa_kernel explicit_user_sgpr_count_1
+  .amdhsa_user_sgpr_count 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel
+
+.amdhsa_kernel explicit_user_sgpr_count_larger_than_implied
+  .amdhsa_user_sgpr_count 12
+  .amdhsa_user_sgpr_private_segment_buffer 1
+  .amdhsa_user_sgpr_queue_ptr 1
+  .amdhsa_user_sgpr_kernarg_segment_ptr 1
+  .amdhsa_next_free_vgpr 32
+  .amdhsa_next_free_sgpr 32
+.end_amdhsa_kernel