[AMDGPU] Add an additional Code Object V3 assembler example

[scott.linder]

Summary

Document the intended use of the .amdgcn.next_free_{s,v}gpr in the context of multiple kernels and functions.

[kzhuravl]

LGTM

docs/AMDGPUUsage.rst

 .amdgcn.gfx_generation_stepping
 +++++++++++++++++++++++++++++++
 
 Set to the GFX stepping generation number of the target being assembled for.
 For example, when assembling for a "GFX704" target this will be set to the
 integer value "4". The possible GFX stepping generation numbers are presented
 in :ref:`amdgpu-processors`.
 
+.. _amdgpu-amdhsa-assembler-symbol-next_free_vgpr:
+
 .amdgcn.next_free_vgpr
 ++++++++++++++++++++++
 
 Set to zero before assembly begins. At each instruction, if the current value
 of this symbol is less than or equal to the maximum VGPR number explicitly
 referenced within that instruction then the symbol value is updated to equal
 that VGPR number plus one.
 
 May be used to set the `.amdhsa_next_free_vpgr` directive in
 :ref:`amdhsa-kernel-directives-table`.
 
 May be set at any time, e.g. manually set to zero at the start of each kernel.
 
+.. _amdgpu-amdhsa-assembler-symbol-next_free_sgpr:
+
 .amdgcn.next_free_sgpr
 ++++++++++++++++++++++
 
 Set to zero before assembly begins. At each instruction, if the current value
 of this symbol is less than or equal the maximum SGPR number explicitly
 referenced within that instruction then the symbol value is updated to equal
 that SGPR number plus one.
 
- .name: hello_world
       .kernarg_segment_align: 4
       .wavefront_size: 64
       .sgpr_count: 2
       .vgpr_count: 3
       .max_flat_workgroup_size: 256
   ...
   .end_amdgpu_metadata
 
+If an assembly source file contains multiple kernels and/or functions, the
+:ref:`amdgpu-amdhsa-assembler-symbol-next_free_vgpr` and
+:ref:`amdgpu-amdhsa-assembler-symbol-next_free_sgpr` symbols may be reset using
+the ``.set <symbol>, <expression>`` directive. For example, in the case of two
+kernels, where ``function1`` is only called from ``kernel1`` it is sufficient
+to group the function with the kernel that calls it and reset the symbols
+between the two connected components:
+
+.. code-block:: none
+
+  .amdgcn_target "amdgcn-amd-amdhsa--gfx900+xnack" // optional
+
+  // gpr tracking symbols are implicitly set to zero
+
+  .text
+  .globl kern0
+  .p2align 8
+  .type kern0,@function
+  kern0:
+    // ...
+    s_endpgm
+  .Lkern0_end:
+    .size   kern0, .Lkern0_end-kern0
+
+  .rodata
+  .p2align 6
+  .amdhsa_kernel kern0
+    // ...
+    .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
+    .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
+  .end_amdhsa_kernel
+
+  // reset symbols to begin tracking usage in func1 and kern1
+  .set .amdgcn.next_free_vgpr, 0
+  .set .amdgcn.next_free_sgpr, 0
+
+  .text
+  .hidden func1
+  .global func1
+  .p2align 2
+  .type func1,@function
+  func1:
+    // ...
+    s_setpc_b64 s[30:31]
+  .Lfunc1_end:
+  .size func1, .Lfunc1_end-func1
+
+  .globl kern1
+  .p2align 8
+  .type kern1,@function
+  kern1:
+    // ...
+    s_getpc_b64 s[4:5]
+    s_add_u32 s4, s4, func1@rel32@lo+4
+    s_addc_u32 s5, s5, func1@rel32@lo+4
+    s_swappc_b64 s[30:31], s[4:5]
+    // ...
+    s_endpgm
+  .Lkern1_end:
+    .size   kern1, .Lkern1_end-kern1
+
+  .rodata
+  .p2align 6
+  .amdhsa_kernel kern1
+    // ...
+    .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
+    .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
+  .end_amdhsa_kernel
+
+These symbols cannot identify connected components in order to automatically
+track the usage for each kernel. However, in some cases careful organization of
+the kernels and functions in the source file means there is minimal additional
+effort required to accurately calculate GPR usage.
+
 Additional Documentation
 ========================
 
 .. [AMD-RADEON-HD-2000-3000] `AMD R6xx shader ISA <http://developer.amd.com/wordpress/media/2012/10/R600_Instruction_Set_Architecture.pdf>`__
 .. [AMD-RADEON-HD-4000] `AMD R7xx shader ISA <http://developer.amd.com/wordpress/media/2012/10/R700-Family_Instruction_Set_Architecture.pdf>`__
 .. [AMD-RADEON-HD-5000] `AMD Evergreen shader ISA <http://developer.amd.com/wordpress/media/2012/10/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf>`__
 .. [AMD-RADEON-HD-6000] `AMD Cayman/Trinity shader ISA <http://developer.amd.com/wordpress/media/2012/10/AMD_HD_6900_Series_Instruction_Set_Architecture.pdf>`__
 .. [AMD-GCN-GFX6] `AMD Southern Islands Series ISA <http://developer.amd.com/wordpress/media/2012/12/AMD_Southern_Islands_Instruction_Set_Architecture.pdf>`__