[AMDGPU] AMDGPUUsage clarify address space information and other typo and formatting fixes

[t-tye]

Summary

• Clarify AMDGPU address spaces.
• Correct path to AMDGPU backend since now in the mono-repo.
• Fix numerous text style and typo issues.
• Correct reStructure text formatting warnings.
• Made reStructure directive usage more consistent.
• Add references for gfx10 ISA specification.

[foad]

llvm/docs/AMDGPUUsage.rst

331

What is the text "64-Bit Process Address Space" supposed to mean here? Is it some kind of column heading? Is it just a mistake?

[t-tye]

llvm/docs/AMDGPUUsage.rst

331

It looks like it is a heading for both the "Address Size" and "NULL Value" columns. Both those columns provide information about address spaces on AMDGPU hardware which uses a 64 bit virtual address space. Not sure if it really adds any extra meaning or could be dropped.

llvm/docs/AMDGPUUsage.rst

 =============================
 User Guide for AMDGPU Backend
 =============================
 
 .. contents::
    :local:
 
 Introduction
 ============
 
 The AMDGPU backend provides ISA code generation for AMD GPUs, starting with the
 R600 family up until the current GCN families. It lives in the
-``lib/Target/AMDGPU`` directory.
+``llvm/lib/Target/AMDGPU`` directory.
 
 LLVM
 ====
 
 .. _amdgpu-target-triples:
 
 Target Triples
 --------------
.. table:: AMDGPU Environments
      *<empty>*    Default.
      ============ ==============================================================
 
 .. _amdgpu-processors:
 
 Processors
 ----------
 
-Use the ``clang -mcpu <Processor>`` option to specify the AMD GPU processor. The
+Use the ``clang -mcpu <Processor>`` option to specify the AMDGPU processor. The
 names from both the *Processor* and *Alternative Processor* can be used.
 
   .. table:: AMDGPU Processors
      :name: amdgpu-processor-table
 
      =========== =============== ============ ===== ================= ======= ======================
      Processor   Alternative     Target       dGPU/ Target            ROCm    Example
                  Processor       Triple       APU   Features          Support Products
``gfx900``                  ``amdgcn``   dGPU  - xnack           ROCm    - Radeon Vega
                                                                               - Radeon RX Vega 64
                                                                               - Radeon RX Vega 64
                                                                                 Liquid
                                                                               - Radeon Instinct MI25
      ``gfx902``                  ``amdgcn``   APU   - xnack                   - Ryzen 3 2200G
                                                       [on]                    - Ryzen 5 2400G
      ``gfx904``                  ``amdgcn``   dGPU  - xnack                   *TBA*
                                                       [off]
-                                                                              .. TODO
+                                                                              .. TODO::
                                                                                  Add product
                                                                                  names.
      ``gfx906``                  ``amdgcn``   dGPU  - xnack                   - Radeon Instinct MI50
                                                       [off]                   - Radeon Instinct MI60
      ``gfx908``                  ``amdgcn``   dGPU  - xnack                   *TBA*
                                                       [off]
                                                       sram-ecc
                                                       [on]
      ``gfx909``                  ``amdgcn``   APU   - xnack                   *TBA* (Raven Ridge 2)
                                                       [on]
-                                                                              .. TODO
+                                                                              .. TODO::
                                                                                  Add product
                                                                                  names.
      **GCN GFX10** [AMD-GCN-GFX10]_
      -----------------------------------------------------------------------------------------------
      ``gfx1010``                 ``amdgcn``   dGPU  - xnack                   *TBA*
                                                       [off]
                                                     - wavefrontsize64
                                                       [off]
                                                     - cumode
                                                       [off]
-                                                                              .. TODO
+                                                                              .. TODO::
                                                                                  Add product
                                                                                  names.
      ``gfx1011``                 ``amdgcn``   dGPU  - xnack                   *TBA*
                                                       [off]
                                                     - wavefrontsize64
                                                       [off]
                                                     - cumode
                                                       [off]
-                                                                              .. TODO
+                                                                              .. TODO::
                                                                                  Add product
                                                                                  names.
      ``gfx1012``                 ``amdgcn``   dGPU  - xnack                   *TBA*
                                                       [off]
                                                     - wavefrontsize64
                                                       [off]
                                                     - cumode
                                                       [off]
-                                                                              .. TODO
+                                                                              .. TODO::
                                                                                  Add product
                                                                                  names.
      =========== =============== ============ ===== ================= ======= ======================
 
 .. _amdgpu-target-features:
 
 Target Features
 ---------------
 
 Target features control how code is generated to support certain
 processor specific features. Not all target features are supported by
 all processors. The runtime must ensure that the features supported by
 the device used to execute the code match the features enabled when
 generating the code. A mismatch of features may result in incorrect
 execution, or a reduction in performance.
 
 The target features supported by each processor, and the default value
 used if not specified explicitly, is listed in
 :ref:`amdgpu-processor-table`.
 
-Use the ``clang -m[no-]<TargetFeature>`` option to specify the AMD GPU
+Use the ``clang -m[no-]<TargetFeature>`` option to specify the AMDGPU
 target features.
 
 For example:
 
 ``-mxnack``
   Enable the ``xnack`` feature.
 ``-mno-xnack``
   Disable the ``xnack`` feature.
-m[no-]cumode          Control the default wavefront execution mode used
                             (see :ref:`amdgpu-amdhsa-memory-model`).
      ====================== ==================================================
 
 .. _amdgpu-address-spaces:
 
 Address Spaces
 --------------
 
-The AMDGPU backend uses the following address space mappings.
+The AMDGPU architecture supports a number of memory address spaces. The address
+space names use the OpenCL standard names, with some additions.
 
-The memory space names used in the table, aside from the region memory space, is
-from the OpenCL standard.
+The AMDGPU address spaces correspond to architecture-specific LLVM address
+space numbers used in LLVM IR.
 
-LLVM Address Space number is used throughout LLVM (for example, in LLVM IR).
+The AMDGPU address spaces are described in
+:ref:`amdgpu-address-spaces-table`. Only 64-bit process address spaces are
+supported for the ``amdgcn`` target.
+
+  .. table:: AMDGPU Address Spaces
+     :name: amdgpu-address-spaces-table
+
+     ================================= =============== =========== ================ ======= ============================
+     ..                                                                                     64-Bit Process Address Space
+     --------------------------------- --------------- ----------- ---------------- ------------------------------------
+     Address Space Name                LLVM IR Address HSA Segment Hardware         Address NULL Value
+                                       Space Number    Name        Name             Size
+     ================================= =============== =========== ================ ======= ============================
+     Generic                           0               flat        flat             64      0x0000000000000000
+     Global                            1               global      global           64      0x0000000000000000
+     Region                            2               N/A         GDS              32      *not implemented for AMDHSA*
+     Local                             3               group       LDS              32      0xFFFFFFFF
+     Constant                          4               constant    *same as global* 64      0x0000000000000000
+     Private                           5               private     scratch          32      0x00000000
+     Constant 32-bit                   6               *TODO*
+     Buffer Fat Pointer (experimental) 7               *TODO*
+     ================================= =============== =========== ================ ======= ============================
 
-  .. table:: Address Space Mapping
-     :name: amdgpu-address-space-mapping-table
-
-     ================== =================================
-     LLVM Address Space Memory Space
-     ================== =================================
-     0                  Generic (Flat)
-     1                  Global
-     2                  Region (GDS)
-     3                  Local (group/LDS)
-     4                  Constant
-     5                  Private (Scratch)
-     6                  Constant 32-bit
-     7                  Buffer Fat Pointer (experimental)
-     ================== =================================
+**Generic**
+  The generic address space uses the hardware flat address support available in
+  GFX7-GFX10. This uses two fixed ranges of virtual addresses (the private and
+  local apertures), that are outside the range of addressable global memory, to
+  map from a flat address to a private or local address.
+
+  FLAT instructions can take a flat address and access global, private
+  (scratch), and group (LDS) memory depending on if the address is within one
+  of the aperture ranges. Flat access to scratch requires hardware aperture
+  setup and setup in the kernel prologue (see
+  :ref:`amdgpu-amdhsa-flat-scratch`). Flat access to LDS requires hardware
+  aperture setup and M0 (GFX7-GFX8) register setup (see
+  :ref:`amdgpu-amdhsa-m0`).
+
+  To convert between a private or group address space address (termed a segment
+  address) and a flat address the base address of the corresponding aperture
+  can be used. For GFX7-GFX8 these are available in the
+  :ref:`amdgpu-amdhsa-hsa-aql-queue` the address of which can be obtained with
+  Queue Ptr SGPR (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). For
+  GFX9-GFX10 the aperture base addresses are directly available as inline
+  constant registers ``SRC_SHARED_BASE/LIMIT`` and ``SRC_PRIVATE_BASE/LIMIT``.
+  In 64-bit address mode the aperture sizes are 2^32 bytes and the base is
+  aligned to 2^32 which makes it easier to convert from flat to segment or
+  segment to flat.
+
+  A global address space address has the same value when used as a flat address
+  so no conversion is needed.
+
+**Global and Constant**
+  The global and constant address spaces both use global virtual addresses,
+  which are the same virtual address space used by the CPU. However, some
+  virtual addresses may only be accessible to the CPU, some only accessible
+  by the GPU, and some by both.
+
+  Using the constant address space indicates that the data will not change
+  during the execution of the kernel. This allows scalar read instructions to
+  be used. The vector and scalar L1 caches are invalidated of volatile data
+  before each kernel dispatch execution to allow constant memory to change
+  values between kernel dispatches.
+
+**Region**
+  The region address space uses the hardware Global Data Store (GDS). All
+  wavefronts executing on the same device will access the same memory for any
+  given region address. However, the same region address accessed by wavefronts
+  executing on different devices will access different memory. It is higher
+  performance than global memory. It is allocated by the runtime. The data
+  store (DS) instructions can be used to access it.
+
+**Local**
+  The local address space uses the hardware Local Data Store (LDS) which is
+  automatically allocated when the hardware creates the wavefronts of a
+  work-group, and freed when all the wavefronts of a work-group have
+  terminated. All wavefronts belonging to the same work-group will access the
+  same memory for any given local address. However, the same local address
+  accessed by wavefronts belonging to different work-groups will access
+  different memory. It is higher performance than global memory. The data store
+  (DS) instructions can be used to access it.
+
+**Private**
+  The private address space uses the hardware scratch memory support which
+  automatically allocates memory when it creates a wavefront, and frees it when
+  a wavefronts terminates. The memory accessed by a lane of a wavefront for any
+  given private address will be different to the memory accessed by another lane
+  of the same or different wavefront for the same private address.
+
+  If a kernel dispatch uses scratch, then the hardware allocates memory from a
+  pool of backing memory allocated by the runtime for each wavefront. The lanes
+  of the wavefront access this using dword (4 byte) interleaving. The mapping
+  used from private address to backing memory address is:
+
+    ``wavefront-scratch-base +
+    ((private-address / 4) * wavefront-size * 4) +
+    (wavefront-lane-id * 4) + (private-address % 4)``
+
+  If each lane of a wavefront accesses the same private address, the
+  interleaving results in adjacent dwords being accessed and hence requires
+  fewer cache lines to be fetched.
+
+  There are different ways that the wavefront scratch base address is
+  determined by a wavefront (see
+  :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
+
+  Scratch memory can be accessed in an interleaved manner using buffer
+  instructions with the scratch buffer descriptor and per wavefront scratch
+  offset, by the scratch instructions, or by flat instructions. Multi-dword
+  access is not supported except by flat and scratch instructions in
+  GFX9-GFX10.
+
+**Constant 32-bit**
+  *TODO*
 
+**Buffer Fat Pointer**
   The buffer fat pointer is an experimental address space that is currently
-unsupported in the backend. It exposes a non-integral pointer that is in future
-intended to support the modelling of 128-bit buffer descriptors + a 32-bit
-offset into the buffer descriptor (in total encapsulating a 160-bit 'pointer'),
-allowing us to use normal LLVM load/store/atomic operations to model the buffer
-descriptors used heavily in graphics workloads targeting the backend.
+  unsupported in the backend. It exposes a non-integral pointer that is in
+  the future intended to support the modelling of 128-bit buffer descriptors
+  plus a 32-bit offset into the buffer (in total encapsulating a 160-bit
+  *pointer*), allowing normal LLVM load/store/atomic operations to be used to
+  model the buffer descriptors used heavily in graphics workloads targeting
+  the backend.
 
 .. _amdgpu-memory-scopes:
 
 Memory Scopes
 -------------
 
 This section provides LLVM memory synchronization scopes supported by the AMDGPU
 backend memory model when the target triple OS is ``amdhsa`` (see
 :ref:`amdgpu-amdhsa-memory-model` and :ref:`amdgpu-target-triples`).
 
 The memory model supported is based on the HSA memory model [HSA]_ which is
 based in turn on HRF-indirect with scope inclusion [HRF]_. The happens-before
-relation is transitive over the synchonizes-with relation independent of scope,
-and synchonizes-with allows the memory scope instances to be inclusive (see
+relation is transitive over the synchronizes-with relation independent of scope,
+and synchronizes-with allows the memory scope instances to be inclusive (see
 table :ref:`amdgpu-amdhsa-llvm-sync-scopes-table`).
 
 This is different to the OpenCL [OpenCL]_ memory model which does not have scope
 inclusion and requires the memory scopes to exactly match. However, this
 is conservatively correct for OpenCL.
 
   .. table:: AMDHSA LLVM Sync Scopes
      :name: amdgpu-amdhsa-llvm-sync-scopes-table

 
 AMDGPU Intrinsics
 -----------------
 
 The AMDGPU backend implements the following LLVM IR intrinsics.
 
 *This section is WIP.*
 
-.. TODO
-   List AMDGPU intrinsics
+.. TODO::
+
+   List AMDGPU intrinsics.
 
 AMDGPU Attributes
 -----------------
 
 The AMDGPU backend supports the following LLVM IR attributes.
 
   .. table:: AMDGPU LLVM IR Attributes
      :name: amdgpu-llvm-ir-attributes-table
.. table:: AMDGPU ELF Header Enumeration Values
      ``ELFABIVERSION_AMDGPU_MESA3D`` 0
      =============================== =====
 
 ``e_ident[EI_CLASS]``
   The ELF class is:
 
   * ``ELFCLASS32`` for ``r600`` architecture.
 
-  * ``ELFCLASS64`` for ``amdgcn`` architecture which only supports 64
-    bit applications.
+  * ``ELFCLASS64`` for ``amdgcn`` architecture which only supports 64-bit
+    process address space applications.
 
 ``e_ident[EI_DATA]``
   All AMDGPU targets use ``ELFDATA2LSB`` for little-endian byte ordering.
 
 ``e_ident[EI_OSABI]``
-  One of the following AMD GPU architecture specific OS ABIs
+  One of the following AMDGPU architecture specific OS ABIs
   (see :ref:`amdgpu-os-table`):
 
   * ``ELFOSABI_NONE`` for *unknown* OS.
 
   * ``ELFOSABI_AMDGPU_HSA`` for ``amdhsa`` OS.
 
   * ``ELFOSABI_AMDGPU_PAL`` for ``amdpal`` OS.
 
   * ``ELFOSABI_AMDGPU_MESA3D`` for ``mesa3D`` OS.
 
 ``e_ident[EI_ABIVERSION]``
-  The ABI version of the AMD GPU architecture specific OS ABI to which the code
+  The ABI version of the AMDGPU architecture specific OS ABI to which the code
   object conforms:
 
   * ``ELFABIVERSION_AMDGPU_HSA`` is used to specify the version of AMD HSA
     runtime ABI.
 
   * ``ELFABIVERSION_AMDGPU_PAL`` is used to specify the version of AMD PAL
     runtime ABI.
 
   * ``ELFABIVERSION_AMDGPU_MESA3D`` is used to specify the version of AMD MESA
     3D runtime ABI.
 
 ``e_type``
   Can be one of the following values:
 
 
   ``ET_REL``
-    The type produced by the AMD GPU backend compiler as it is relocatable code
+    The type produced by the AMDGPU backend compiler as it is relocatable code
     object.
 
   ``ET_DYN``
     The type produced by the linker as it is a shared code object.
 
   The AMD HSA runtime loader requires a ``ET_DYN`` code object.
 
 ``e_machine``
Symbols include the following:
   .. table:: AMDGPU ELF Symbols
      :name: amdgpu-elf-symbols-table
 
      ===================== ================== ================ ==================
      Name                  Type               Section          Description
      ===================== ================== ================ ==================
      *link-name*           ``STT_OBJECT``     - ``.data``      Global variable
                                               - ``.rodata``
-                - ``.bss``
+                                              - ``.bss``
      *link-name*\ ``.kd``  ``STT_OBJECT``     - ``.rodata``    Kernel descriptor
      *link-name*           ``STT_FUNC``       - ``.text``      Kernel entry point
      *link-name*           ``STT_OBJECT``     - SHN_AMDGPU_LDS Global variable in LDS
      ===================== ================== ================ ==================
 
 Global variable
   Global variables both used and defined by the compilation unit.
 
   If the symbol is defined in the compilation unit then it is allocated in the
   appropriate section according to if it has initialized data or is readonly.
 
   If the symbol is external then its section is ``STN_UNDEF`` and the loader
   will resolve relocations using the definition provided by another code object
   or explicitly defined by the runtime.
 
   If the symbol resides in local/group memory (LDS) then its section is the
   special processor-specific section name ``SHN_AMDGPU_LDS``, and the
   ``st_value`` field describes alignment requirements as it does for common
   symbols.
 
-  .. TODO
+  .. TODO::
+
      Add description of linked shared object symbols. Seems undefined symbols
      are marked as STT_NOTYPE.
 
 Kernel descriptor
   Every HSA kernel has an associated kernel descriptor. It is the address of the
   kernel descriptor that is used in the AQL dispatch packet used to invoke the
   kernel, not the kernel entry point. The layout of the HSA kernel descriptor is
   defined in :ref:`amdgpu-amdhsa-kernel-descriptor`.
 
 Kernel entry point
   Every HSA kernel also has a symbol for its machine code entry point.
 
 .. _amdgpu-relocation-records:
 
 Relocation Records
 ------------------
 
 AMDGPU backend generates ``Elf64_Rela`` relocation records. Supported
 relocatable fields are:
 
 ``word32``
   This specifies a 32-bit field occupying 4 bytes with arbitrary byte
   alignment. These values use the same byte order as other word values in the
-  AMD GPU architecture.
+  AMDGPU architecture.
 
 ``word64``
   This specifies a 64-bit field occupying 8 bytes with arbitrary byte
   alignment. These values use the same byte order as other word values in the
-  AMD GPU architecture.
+  AMDGPU architecture.
 
 Following notations are used for specifying relocation calculations:
 
 **A**
   Represents the addend used to compute the value of the relocatable field.
 
 **G**
   Represents the offset into the global offset table at which the relocation
   entry's symbol will reside during execution.
 
 **GOT**
   Represents the address of the global offset table.
 
 **P**
   Represents the place (section offset for ``et_rel`` or address for ``et_dyn``)
   of the storage unit being relocated (computed using ``r_offset``).
 
 **S**
   Represents the value of the symbol whose index resides in the relocation
-  entry. Relocations not using this must specify a symbol index of ``STN_UNDEF``.
+  entry. Relocations not using this must specify a symbol index of
+  ``STN_UNDEF``.
 
 **B**
   Represents the base address of a loaded executable or shared object which is
-  the difference between the ELF address and the actual load address. Relocations
-  using this are only valid in executable or shared objects.
+  the difference between the ELF address and the actual load address.
+  Relocations using this are only valid in executable or shared objects.
 
 The following relocation types are supported:
 
   .. table:: AMDGPU ELF Relocation Records
      :name: amdgpu-elf-relocation-records-table
 
      ========================== ======= =====  ==========  ==============================
      Relocation Type            Kind    Value  Field       Calculation
.. table:: AMDGPU ELF Relocation Records
      ``R_AMDGPU_REL32_HI``      Static  11     ``word32``  (S + A - P) >> 32
      *reserved*                         12
      ``R_AMDGPU_RELATIVE64``    Dynamic 13     ``word64``  B + A
      ========================== ======= =====  ==========  ==============================
 
 ``R_AMDGPU_ABS32_LO`` and ``R_AMDGPU_ABS32_HI`` are only supported by
 the ``mesa3d`` OS, which does not support ``R_AMDGPU_ABS64``.
 
-There is no current OS loader support for 32 bit programs and so
+There is no current OS loader support for 32-bit programs and so
 ``R_AMDGPU_ABS32`` is not used.
 
 .. _amdgpu-dwarf:
 
 DWARF
 -----
 
 Standard DWARF [DWARF]_ Version 5 sections can be generated. These contain
.. table:: AMDGPU DWARF Address Space Mapping
      1                   Private (Scratch)
      2                   Local (group/LDS)
      *omitted*           Global
      *omitted*           Constant
      *omitted*           Generic (Flat)
      *not supported*     Region (GDS)
      =================== =================
 
-See :ref:`amdgpu-address-spaces` for information on the memory space terminology
-used in the table.
+See :ref:`amdgpu-address-spaces` for information on the address space
+terminology used in the table.
 
 An ``address_class`` attribute is generated on pointer type DIEs to specify the
 DWARF address space of the value of the pointer when it is in the *private* or
 *local* address space. Otherwise the attribute is omitted.
 
-An ``XDEREF`` operation is generated in location list expressions for variables
-that are allocated in the *private* and *local* address space. Otherwise no
-``XDREF`` is omitted.
+An ``DW_OP_xderef`` operation is generated in location list expressions for
+variables that are allocated in the *private* and *local* address space.
+Otherwise, ``DW_OP_xderef`` is omitted.
 
 Register Mapping
 ~~~~~~~~~~~~~~~~
 
 *This section is WIP.*
 
-.. TODO
+.. TODO::
    Define DWARF register enumeration.
 
    If want to present a wavefront state then should expose vector registers as
-   64 wide (rather than per work-item view that LLVM uses). Either as separate
-   registers, or a 64x4 byte single register. In either case use a new LANE op
-   (akin to XDREF) to select the current lane usage in a location
-   expression. This would also allow scalar register spilling to vector register
-   lanes to be expressed (currently no debug information is being generated for
-   spilling). If choose a wide single register approach then use LANE in
-   conjunction with PIECE operation to select the dword part of the register for
-   the current lane. If the separate register approach then use LANE to select
-   the register.
+   64 dword wide (rather than per work-item view that LLVM uses). Either as
+   separate registers, or a 64x4 byte single register. In either case use a new
+   ``DW_OP_lane`` op (akin to ``DW_OP_xderef``) to select the current lane usage
+   in a location expression. This would also allow scalar register spilling to
+   vector register lanes to be expressed (currently no debug information is
+   being generated for spilling). If choose a wide single register approach then
+   use ``DW_OP_lane`` in conjunction with ``DW_OP_piece`` operation to select
+   the dword part of the register for the current lane. If the separate register
+   approach then use ``DW_OP_lane`` to select the register.
 
 Source Text
 ~~~~~~~~~~~
 
 Source text for online-compiled programs (e.g. those compiled by the OpenCL
 runtime) may be embedded into the DWARF v5 line table using the ``clang
 -gembed-source`` option, described in table :ref:`amdgpu-debug-options`.
 
.. warning:: Code Object V2 is not the default code object version emitted by
   :ref:`amdgpu-amdhsa-code-object-metadata-v3`.
 
 Code object V2 metadata is specified by the ``NT_AMD_AMDGPU_METADATA`` note
 record (see :ref:`amdgpu-note-records-v2`).
 
 The metadata is specified as a YAML formatted string (see [YAML]_ and
 :doc:`YamlIO`).
 
-.. TODO
-   Is the string null terminated? It probably should not if YAML allows it to
-   contain null characters, otherwise it should be.
+.. TODO::
+
 
 The metadata is represented as a single YAML document comprised of the mapping
 defined in table :ref:`amdgpu-amdhsa-code-object-metadata-map-table-v2` and
 referenced tables.
 
 For boolean values, the string values of ``false`` and ``true`` are used for
 false and true respectively.
 
"Printf"   sequence of              Each string is encoded information
                                          encoded information is organized as
                                          fields separated by colon (':'):
 
                                          ``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``
 
                                          where:
 
                                          ``ID``
-                                           A 32 bit integer as a unique id for
+                                           A 32-bit integer as a unique id for
                                            each printf function call
 
                                          ``N``
-                                           A 32 bit integer equal to the number
+                                           A 32-bit integer equal to the number
                                            of arguments of printf function call
                                            minus 1
 
                                          ``S[i]`` (where i = 0, 1, ... , N-1)
-                                           32 bit integers for the size in bytes
+                                           32-bit integers for the size in bytes
                                            of the i-th FormatString argument of
                                            the printf function call
 
                                          FormatString
                                            The format string passed to the
                                            printf function call.
      "Kernels"  sequence of    Required  Sequence of the mappings for each
                 mapping                  kernel in the code object. See
"ValueType"       string         Required  Kernel argument value type. Only
                                                 - "F16"
                                                 - "I32"
                                                 - "U32"
                                                 - "F32"
                                                 - "I64"
                                                 - "U64"
                                                 - "F64"
 
-                                                .. TODO
+                                                .. TODO::
                                                    How can it be determined if a
                                                    vector type, and what size
                                                    vector?
      "PointeeAlign"    integer                  Alignment in bytes of pointee
                                                 type for pointer type kernel
                                                 argument. Must be a power
                                                 of 2. Only present if
                                                 "ValueKind" is
                                                 "DynamicSharedPointer".
      "AddrSpaceQual"   string                   Kernel argument address space
                                                 qualifier. Only present if
                                                 "ValueKind" is "GlobalBuffer" or
                                                 "DynamicSharedPointer". Values
                                                 are:
 
                                                 - "Private"
                                                 - "Global"
                                                 - "Constant"
                                                 - "Local"
                                                 - "Generic"
                                                 - "Region"
 
-                                                .. TODO
+                                                .. TODO::
                                                    Is GlobalBuffer only Global
                                                    or Constant? Is
                                                    DynamicSharedPointer always
                                                    Local? Can HCC allow Generic?
                                                    How can Private or Region
                                                    ever happen?
      "AccQual"         string                   Kernel argument access
                                                 qualifier. Only present if
                                                 "ValueKind" is "Image" or
                                                 "Pipe". Values
                                                 are:
 
                                                 - "ReadOnly"
                                                 - "WriteOnly"
                                                 - "ReadWrite"
 
-                                                .. TODO
+                                                .. TODO::
                                                    Does this apply to
                                                    GlobalBuffer?
      "ActualAccQual"   string                   The actual memory accesses
                                                 performed by the kernel on the
                                                 kernel argument. Only present if
                                                 "ValueKind" is "GlobalBuffer",
                                                 "Image", or "Pipe". This may be
                                                 more restrictive than indicated
"IsVolatile"      boolean                  Indicates if the kernel argument
                                                 is volatile qualified. Only
                                                 present if "ValueKind" is
                                                 "GlobalBuffer".
 
      "IsPipe"          boolean                  Indicates if the kernel argument
                                                 is pipe qualified. Only present
                                                 if "ValueKind" is "Pipe".
 
-                                                .. TODO
+                                                .. TODO::
                                                    Can GlobalBuffer be pipe
                                                    qualified?
      ================= ============== ========= ================================
 
 ..
 
   .. table:: AMDHSA Code Object V2 Kernel Code Properties Metadata Map
      :name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-map-table-v2
"amdhsa.printf"   sequence of              Each string is encoded information
                                                 encoded information is organized as
                                                 fields separated by colon (':'):
 
                                                 ``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``
 
                                                 where:
 
                                                 ``ID``
-                                                  A 32 bit integer as a unique id for
+                                                  A 32-bit integer as a unique id for
                                                   each printf function call
 
                                                 ``N``
-                                                  A 32 bit integer equal to the number
+                                                  A 32-bit integer equal to the number
                                                   of arguments of printf function call
                                                   minus 1
 
                                                 ``S[i]`` (where i = 0, 1, ... , N-1)
-                                                  32 bit integers for the size in bytes
+                                                  32-bit integers for the size in bytes
                                                   of the i-th FormatString argument of
                                                   the printf function call
 
                                                 FormatString
                                                   The format string passed to the
                                                   printf function call.
      "amdhsa.kernels"  sequence of    Required  Sequence of the maps for each
                        map                      kernel in the code object. See
".value_type"          string         Required  Kernel argument value type. Only
                                                      - "f16"
                                                      - "i32"
                                                      - "u32"
                                                      - "f32"
                                                      - "i64"
                                                      - "u64"
                                                      - "f64"
 
-                                                     .. TODO
+                                                     .. TODO::
                                                         How can it be determined if a
                                                         vector type, and what size
                                                         vector?
      ".pointee_align"       integer                  Alignment in bytes of pointee
                                                      type for pointer type kernel
                                                      argument. Must be a power
                                                      of 2. Only present if
                                                      ".value_kind" is
                                                      "dynamic_shared_pointer".
      ".address_space"       string                   Kernel argument address space
                                                      qualifier. Only present if
                                                      ".value_kind" is "global_buffer" or
                                                      "dynamic_shared_pointer". Values
                                                      are:
 
                                                      - "private"
                                                      - "global"
                                                      - "constant"
                                                      - "local"
                                                      - "generic"
                                                      - "region"
 
-                                                     .. TODO
+                                                     .. TODO::
                                                         Is "global_buffer" only "global"
                                                         or "constant"? Is
                                                         "dynamic_shared_pointer" always
                                                         "local"? Can HCC allow "generic"?
                                                         How can "private" or "region"
                                                         ever happen?
      ".access"              string                   Kernel argument access
                                                      qualifier. Only present if
                                                      ".value_kind" is "image" or
                                                      "pipe". Values
                                                      are:
 
                                                      - "read_only"
                                                      - "write_only"
                                                      - "read_write"
 
-                                                     .. TODO
+                                                     .. TODO::
                                                         Does this apply to
                                                         "global_buffer"?
      ".actual_access"       string                   The actual memory accesses
                                                      performed by the kernel on the
                                                      kernel argument. Only present if
                                                      ".value_kind" is "global_buffer",
                                                      "image", or "pipe". This may be
                                                      more restrictive than indicated
".is_volatile"         boolean                  Indicates if the kernel argument
                                                      is volatile qualified. Only
                                                      present if ".value_kind" is
                                                      "global_buffer".
 
      ".is_pipe"             boolean                  Indicates if the kernel argument
                                                      is pipe qualified. Only present
                                                      if ".value_kind" is "pipe".
 
-                                                     .. TODO
+                                                     .. TODO::
                                                         Can "global_buffer" be pipe
                                                         qualified?
      ====================== ============== ========= ================================
 
 ..
 
 Kernel Dispatch
 ~~~~~~~~~~~~~~~
 
-The HSA architected queuing language (AQL) defines a user space memory interface
-that can be used to control the dispatch of kernels, in an agent independent
-way. An agent can have zero or more AQL queues created for it using the ROCm
-runtime, in which AQL packets (all of which are 64 bytes) can be placed. See the
-*HSA Platform System Architecture Specification* [HSA]_ for the AQL queue
-mechanics and packet layouts.
+The HSA architected queuing language (AQL) defines a user space memory
+interface that can be used to control the dispatch of kernels, in an agent
+independent way. An agent can have zero or more AQL queues created for it using
+the ROCm runtime, in which AQL packets (all of which are 64 bytes) can be
+placed. See the *HSA Platform System Architecture Specification* [HSA]_ for the
+AQL queue mechanics and packet layouts.
 
 The packet processor of a kernel agent is responsible for detecting and
 dispatching HSA kernels from the AQL queues associated with it. For AMD GPUs the
 packet processor is implemented by the hardware command processor (CP),
 asynchronous dispatch controller (ADC) and shader processor input controller
 (SPI).
 
 The ROCm runtime can be used to allocate an AQL queue object. It uses the kernel
 mode driver to initialize and register the AQL queue with CP.
 
 To dispatch a kernel the following actions are performed. This can occur in the
 CPU host program, or from an HSA kernel executing on a GPU.
 
 1. A pointer to an AQL queue for the kernel agent on which the kernel is to be
    executed is obtained.
 2. A pointer to the kernel descriptor (see
-   :ref:`amdgpu-amdhsa-kernel-descriptor`) of the kernel to execute is
-   obtained. It must be for a kernel that is contained in a code object that that
-   was loaded by the ROCm runtime on the kernel agent with which the AQL queue is
+   :ref:`amdgpu-amdhsa-kernel-descriptor`) of the kernel to execute is obtained.
+   It must be for a kernel that is contained in a code object that that was
+   loaded by the ROCm runtime on the kernel agent with which the AQL queue is
    associated.
 3. Space is allocated for the kernel arguments using the ROCm runtime allocator
    for a memory region with the kernarg property for the kernel agent that will
    execute the kernel. It must be at least 16 byte aligned.
 4. Kernel argument values are assigned to the kernel argument memory
-   allocation. The layout is defined in the *HSA Programmer's Language Reference*
-   [HSA]_. For AMDGPU the kernel execution directly accesses the kernel argument
-   memory in the same way constant memory is accessed. (Note that the HSA
-   specification allows an implementation to copy the kernel argument contents to
-   another location that is accessed by the kernel.)
+   allocation. The layout is defined in the *HSA Programmer's Language
+   Reference* [HSA]_. For AMDGPU the kernel execution directly accesses the
+   kernel argument memory in the same way constant memory is accessed. (Note
+   that the HSA specification allows an implementation to copy the kernel
+   argument contents to another location that is accessed by the kernel.)
 5. An AQL kernel dispatch packet is created on the AQL queue. The ROCm runtime
-   api uses 64 bit atomic operations to reserve space in the AQL queue for the
+   api uses 64-bit atomic operations to reserve space in the AQL queue for the
    packet. The packet must be set up, and the final write must use an atomic
    store release to set the packet kind to ensure the packet contents are
    visible to the kernel agent. AQL defines a doorbell signal mechanism to
    notify the kernel agent that the AQL queue has been updated. These rules, and
    the layout of the AQL queue and kernel dispatch packet is defined in the *HSA
    System Architecture Specification* [HSA]_.
 6. A kernel dispatch packet includes information about the actual dispatch,
    such as grid and work-group size, together with information from the code
8. CP ensures that when the a wavefront starts executing the kernel machine
    register state is defined in
    :ref:`amdgpu-amdhsa-initial-kernel-execution-state`.
 9. The prolog of the kernel machine code (see
    :ref:`amdgpu-amdhsa-kernel-prolog`) sets up the machine state as necessary
    before continuing executing the machine code that corresponds to the kernel.
 10. When the kernel dispatch has completed execution, CP signals the completion
     signal specified in the kernel dispatch packet if not 0.
 
-.. _amdgpu-amdhsa-memory-spaces:
-
-Memory Spaces
-~~~~~~~~~~~~~
-
-The memory space properties are:
-
-  .. table:: AMDHSA Memory Spaces
-     :name: amdgpu-amdhsa-memory-spaces-table
-
-     ================= =========== ======== ======= ==================
-     Memory Space Name HSA Segment Hardware Address NULL Value
-                       Name        Name     Size
-     ================= =========== ======== ======= ==================
-     Private           private     scratch  32      0x00000000
-     Local             group       LDS      32      0xFFFFFFFF
-     Global            global      global   64      0x0000000000000000
-     Constant          constant    *same as 64      0x0000000000000000
-                                   global*
-     Generic           flat        flat     64      0x0000000000000000
-     Region            N/A         GDS      32      *not implemented
-                                                    for AMDHSA*
-     ================= =========== ======== ======= ==================
-
-The global and constant memory spaces both use global virtual addresses, which
-are the same virtual address space used by the CPU. However, some virtual
-addresses may only be accessible to the CPU, some only accessible by the GPU,
-and some by both.
-
-Using the constant memory space indicates that the data will not change during
-the execution of the kernel. This allows scalar read instructions to be
-used. The vector and scalar L1 caches are invalidated of volatile data before
-each kernel dispatch execution to allow constant memory to change values between
-kernel dispatches.
-
-The local memory space uses the hardware Local Data Store (LDS) which is
-automatically allocated when the hardware creates work-groups of wavefronts, and
-freed when all the wavefronts of a work-group have terminated. The data store
-(DS) instructions can be used to access it.
-
-The private memory space uses the hardware scratch memory support. If the kernel
-uses scratch, then the hardware allocates memory that is accessed using
-wavefront lane dword (4 byte) interleaving. The mapping used from private
-address to physical address is:
-
-  ``wavefront-scratch-base +
-  (private-address * wavefront-size * 4) +
-  (wavefront-lane-id * 4)``
-
-There are different ways that the wavefront scratch base address is determined
-by a wavefront (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). This
-memory can be accessed in an interleaved manner using buffer instruction with
-the scratch buffer descriptor and per wavefront scratch offset, by the scratch
-instructions, or by flat instructions. If each lane of a wavefront accesses the
-same private address, the interleaving results in adjacent dwords being accessed
-and hence requires fewer cache lines to be fetched. Multi-dword access is not
-supported except by flat and scratch instructions in GFX9-GFX10.
-
-The generic address space uses the hardware flat address support available in
-GFX7-GFX10. This uses two fixed ranges of virtual addresses (the private and
-local apertures), that are outside the range of addressible global memory, to
-map from a flat address to a private or local address.
-
-FLAT instructions can take a flat address and access global, private (scratch)
-and group (LDS) memory depending in if the address is within one of the
-aperture ranges. Flat access to scratch requires hardware aperture setup and
-setup in the kernel prologue (see :ref:`amdgpu-amdhsa-flat-scratch`). Flat
-access to LDS requires hardware aperture setup and M0 (GFX7-GFX8) register setup
-(see :ref:`amdgpu-amdhsa-m0`).
-
-To convert between a segment address and a flat address the base address of the
-apertures address can be used. For GFX7-GFX8 these are available in the
-:ref:`amdgpu-amdhsa-hsa-aql-queue` the address of which can be obtained with
-Queue Ptr SGPR (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). For
-GFX9-GFX10 the aperture base addresses are directly available as inline constant
-registers ``SRC_SHARED_BASE/LIMIT`` and ``SRC_PRIVATE_BASE/LIMIT``. In 64 bit
-address mode the aperture sizes are 2^32 bytes and the base is aligned to 2^32
-which makes it easier to convert from flat to segment or segment to flat.
-
 Image and Samplers
 ~~~~~~~~~~~~~~~~~~
 
-Image and sample handles created by the ROCm runtime are 64 bit addresses of a
+Image and sample handles created by the ROCm runtime are 64-bit addresses of a
 hardware 32 byte V# and 48 byte S# object respectively. In order to support the
 HSA ``query_sampler`` operations two extra dwords are used to store the HSA BRIG
 enumeration values for the queries that are not trivially deducible from the S#
 representation.
 
 HSA Signals
 ~~~~~~~~~~~
 
-HSA signal handles created by the ROCm runtime are 64 bit addresses of a
+HSA signal handles created by the ROCm runtime are 64-bit addresses of a
 structure allocated in memory accessible from both the CPU and GPU. The
 structure is defined by the ROCm runtime and subject to change between releases
 (see [AMD-ROCm-github]_).
 
 .. _amdgpu-amdhsa-hsa-aql-queue:
 
 HSA AQL Queue
 ~~~~~~~~~~~~~
13:12   2 bits  FLOAT_ROUND_MODE_32             Wavefront starts execution
                                                      mode values are defined in
                                                      :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
 
                                                      Used by CP to set up
                                                      ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
      15:14   2 bits  FLOAT_ROUND_MODE_16_64          Wavefront starts execution
                                                      with specified rounding
                                                      denorm mode for half/double (16
-                                                     and 64 bit) floating point
+                                                     and 64-bit) floating point
                                                      precision floating point
                                                      operations.
 
                                                      Floating point rounding
                                                      mode values are defined in
                                                      :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
 
                                                      Used by CP to set up
17:16   2 bits  FLOAT_DENORM_MODE_32            Wavefront starts execution
                                                      values are defined in
                                                      :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
 
                                                      Used by CP to set up
                                                      ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
      19:18   2 bits  FLOAT_DENORM_MODE_16_64         Wavefront starts execution
                                                      with specified denorm mode
                                                      for half/double (16
-                                                     and 64 bit) floating point
+                                                     and 64-bit) floating point
                                                      precision floating point
                                                      operations.
 
                                                      Floating point denorm mode
                                                      values are defined in
                                                      :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
 
                                                      Used by CP to set up

 The order of the SGPR registers is defined, but the compiler can specify which
 ones are actually setup in the kernel descriptor using the ``enable_sgpr_*`` bit
 fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used
 for enabled registers are dense starting at SGPR0: the first enabled register is
 SGPR0, the next enabled register is SGPR1 etc.; disabled registers do not have
 an SGPR number.
 
 The initial SGPRs comprise up to 16 User SRGPs that are set by CP and apply to
-all wavefronts of the grid. It is possible to specify more than 16 User SGPRs using
-the ``enable_sgpr_*`` bit fields, in which case only the first 16 are actually
-initialized. These are then immediately followed by the System SGPRs that are
-set up by ADC/SPI and can have different values for each wavefront of the grid
-dispatch.
+all wavefronts of the grid. It is possible to specify more than 16 User SGPRs
+using the ``enable_sgpr_*`` bit fields, in which case only the first 16 are
+actually initialized. These are then immediately followed by the System SGPRs
+that are set up by ADC/SPI and can have different values for each wavefront of
+the grid dispatch.
 
 SGPR register initial state is defined in
 :ref:`amdgpu-amdhsa-sgpr-register-set-up-order-table`.
 
   .. table:: SGPR Register Set Up Order
      :name: amdgpu-amdhsa-sgpr-register-set-up-order-table
 
      ========== ========================== ====== ==============================
      SGPR Order Name                       Number Description
                 (kernel descriptor enable  of
                 field)                     SGPRs
      ========== ========================== ====== ==============================
      First      Private Segment Buffer     4      V# that can be used, together
                 (enable_sgpr_private              with Scratch Wavefront Offset
                 _segment_buffer)                  as an offset, to access the
-                                                  private memory space using a
+                                                  private address space using a
                                                   segment address.
 
                                                   CP uses the value provided by
                                                   the runtime.
-     then       Dispatch Ptr               2      64 bit address of AQL dispatch
+     then       Dispatch Ptr               2      64-bit address of AQL dispatch
                 (enable_sgpr_dispatch_ptr)        packet for kernel dispatch
                                                   actually executing.
-     then       Queue Ptr                  2      64 bit address of amd_queue_t
+     then       Queue Ptr                  2      64-bit address of amd_queue_t
                 (enable_sgpr_queue_ptr)           object for AQL queue on which
                                                   the dispatch packet was
                                                   queued.
-     then       Kernarg Segment Ptr        2      64 bit address of Kernarg
+     then       Kernarg Segment Ptr        2      64-bit address of Kernarg
                 (enable_sgpr_kernarg              segment. This is directly
                 _segment_ptr)                     copied from the
                                                   kernarg_address in the kernel
                                                   dispatch packet.
 
                                                   Having CP load it once avoids
                                                   loading it at the beginning of
                                                   every wavefront.
-     then       Dispatch Id                2      64 bit Dispatch ID of the
+     then       Dispatch Id                2      64-bit Dispatch ID of the
                 (enable_sgpr_dispatch_id)         dispatch packet being
                                                   executed.
      then       Flat Scratch Init          2      This is 2 SGPRs:
                 (enable_sgpr_flat_scratch
                 _init)                            GFX6
                                                     Not supported.
                                                   GFX7-GFX8
-                                                    The first SGPR is a 32 bit
+                                                    The first SGPR is a 32-bit
                                                     byte offset from
                                                     ``SH_HIDDEN_PRIVATE_BASE_VIMID``
                                                     to per SPI base of memory
                                                     for scratch for the queue
                                                     executing the kernel
                                                     dispatch. CP obtains this
                                                     from the runtime. (The
                                                     Scratch Segment Buffer base
then       Flat Scratch Init          2      This is 2 SGPRs:
                                                     added to
                                                     ``SH_HIDDEN_PRIVATE_BASE_VIMID``
                                                     to calculate the per wavefront
                                                     FLAT SCRATCH BASE in flat
                                                     memory instructions that
                                                     access the scratch
                                                     aperture.
 
-                                                    The second SGPR is 32 bit
+                                                    The second SGPR is 32-bit
                                                     byte size of a single
                                                     work-item's scratch memory
                                                     usage. CP obtains this from
                                                     the runtime, and it is
                                                     always a multiple of DWORD.
                                                     CP checks that the value in
                                                     the kernel dispatch packet
                                                     Private Segment Byte Size is
                                                     not larger, and requests the
                                                     runtime to increase the
                                                     queue's scratch size if
                                                     necessary. The kernel code
                                                     must move it to
                                                     FLAT_SCRATCH_LO which is
                                                     SGPRn-3 on GFX7 and SGPRn-5
                                                     on GFX8. FLAT_SCRATCH_LO is
                                                     used as the FLAT SCRATCH
                                                     SIZE in flat memory
                                                     instructions. Having CP load
                                                     it once avoids loading it at
                                                     the beginning of every
                                                     wavefront.
                                                   GFX9-GFX10
                                                     This is the
-                                                    64 bit base address of the
+                                                    64-bit base address of the
                                                     per SPI scratch backing
                                                     memory managed by SPI for
                                                     the queue executing the
                                                     kernel dispatch. CP obtains
                                                     this from the runtime (and
                                                     divides it if there are
                                                     multiple Shader Arrays each
                                                     with its own SPI). The value
                                                     of Scratch Wavefront Offset must
                                                     be added by the kernel
                                                     machine code and the result
                                                     moved to the FLAT_SCRATCH
                                                     SGPR which is SGPRn-6 and
                                                     SGPRn-5. It is used as the
                                                     FLAT SCRATCH BASE in flat
                                                     memory instructions.
-     then       Private Segment Size       1      The 32 bit byte size of a
+     then       Private Segment Size       1      The 32-bit byte size of a
                                                   (enable_sgpr_private single
                                                   work-item's
                                                   scratch_segment_size) memory
                                                   allocation. This is the
                                                   value from the kernel
                                                   dispatch packet Private
                                                   Segment Byte Size rounded up
                                                   by CP to a multiple of
                                                   DWORD.
 
                                                   Having CP load it once avoids
                                                   loading it at the beginning of
                                                   every wavefront.
 
                                                   This is not used for
                                                   GFX7-GFX8 since it is the same
                                                   value as the second SGPR of
                                                   Flat Scratch Init. However, it
                                                   may be needed for GFX9-GFX10 which
                                                   changes the meaning of the
                                                   Flat Scratch Init value.
-     then       Grid Work-Group Count X    1      32 bit count of the number of
+     then       Grid Work-Group Count X    1      32-bit count of the number of
                 (enable_sgpr_grid                 work-groups in the X dimension
                 _workgroup_count_X)               for the grid being
                                                   executed. Computed from the
                                                   fields in the kernel dispatch
                                                   packet as ((grid_size.x +
                                                   workgroup_size.x - 1) /
                                                   workgroup_size.x).
-     then       Grid Work-Group Count Y    1      32 bit count of the number of
+     then       Grid Work-Group Count Y    1      32-bit count of the number of
                 (enable_sgpr_grid                 work-groups in the Y dimension
                 _workgroup_count_Y &&             for the grid being
                 less than 16 previous             executed. Computed from the
                 SGPRs)                            fields in the kernel dispatch
                                                   packet as ((grid_size.y +
                                                   workgroup_size.y - 1) /
                                                   workgroupSize.y).
 
                                                   Only initialized if <16
                                                   previous SGPRs initialized.
-     then       Grid Work-Group Count Z    1      32 bit count of the number of
+     then       Grid Work-Group Count Z    1      32-bit count of the number of
                 (enable_sgpr_grid                 work-groups in the Z dimension
                 _workgroup_count_Z &&             for the grid being
                 less than 16 previous             executed. Computed from the
                 SGPRs)                            fields in the kernel dispatch
                                                   packet as ((grid_size.z +
                                                   workgroup_size.z - 1) /
                                                   workgroupSize.z).
 
                                                   Only initialized if <16
                                                   previous SGPRs initialized.
-     then       Work-Group Id X            1      32 bit work-group id in X
+     then       Work-Group Id X            1      32-bit work-group id in X
                 (enable_sgpr_workgroup_id         dimension of grid for
                 _X)                               wavefront.
-     then       Work-Group Id Y            1      32 bit work-group id in Y
+     then       Work-Group Id Y            1      32-bit work-group id in Y
                 (enable_sgpr_workgroup_id         dimension of grid for
                 _Y)                               wavefront.
-     then       Work-Group Id Z            1      32 bit work-group id in Z
+     then       Work-Group Id Z            1      32-bit work-group id in Z
                 (enable_sgpr_workgroup_id         dimension of grid for
                 _Z)                               wavefront.
      then       Work-Group Info            1      {first_wavefront, 14'b0000,
                 (enable_sgpr_workgroup            ordered_append_term[10:0],
                 _info)                            threadgroup_size_in_wavefronts[5:0]}
-     then       Scratch Wavefront Offset   1      32 bit byte offset from base
+     then       Scratch Wavefront Offset   1      32-bit byte offset from base
                 (enable_sgpr_private              of scratch base of queue
                 _segment_wavefront_offset)        executing the kernel
                                                   dispatch. Must be used as an
                                                   offset with Private
                                                   segment address when using
                                                   Scratch Segment Buffer. It
                                                   must be used to set up FLAT
                                                   SCRATCH for flat addressing
:ref:`amdgpu-amdhsa-vgpr-register-set-up-order-table`.
   .. table:: VGPR Register Set Up Order
      :name: amdgpu-amdhsa-vgpr-register-set-up-order-table
 
      ========== ========================== ====== ==============================
      VGPR Order Name                       Number Description
                 (kernel descriptor enable  of
                 field)                     VGPRs
      ========== ========================== ====== ==============================
-     First      Work-Item Id X             1      32 bit work item id in X
+     First      Work-Item Id X             1      32-bit work item id in X
                 (Always initialized)              dimension of work-group for
                                                   wavefront lane.
-     then       Work-Item Id Y             1      32 bit work item id in Y
+     then       Work-Item Id Y             1      32-bit work item id in Y
                 (enable_vgpr_workitem_id          dimension of work-group for
                 > 0)                              wavefront lane.
-     then       Work-Item Id Z             1      32 bit work item id in Z
+     then       Work-Item Id Z             1      32-bit work item id in Z
                 (enable_vgpr_workitem_id          dimension of work-group for
                 > 1)                              wavefront lane.
      ========== ========================== ====== ==============================
 
 The setting of registers is done by GPU CP/ADC/SPI hardware as follows:
 
 1. SGPRs before the Work-Group Ids are set by CP using the 16 User Data
    registers.
 2. Work-group Id registers X, Y, Z are set by ADC which supports any
    combination including none.
 3. Scratch Wavefront Offset is set by SPI in a per wavefront basis which is why
-   its value cannot included with the flat scratch init value which is per queue.
+   its value cannot included with the flat scratch init value which is per
+   queue.
 4. The VGPRs are set by SPI which only supports specifying either (X), (X, Y)
    or (X, Y, Z).
 
-Flat Scratch register pair are adjacent SGRRs so they can be moved as a 64 bit
+Flat Scratch register pair are adjacent SGRRs so they can be moved as a 64-bit
 value to the hardware required SGPRn-3 and SGPRn-4 respectively.
 
 The global segment can be accessed either using buffer instructions (GFX6 which
-has V# 64 bit address support), flat instructions (GFX7-GFX10), or global
+has V# 64-bit address support), flat instructions (GFX7-GFX10), or global
 instructions (GFX9-GFX10).
 
 If buffer operations are used then the compiler can generate a V# with the
 following properties:
 
 * base address of 0
 * no swizzle
 * ATC: 1 if IOMMU present (such as APU)

 
 .. _amdgpu-amdhsa-flat-scratch:
 
 Flat Scratch
 ++++++++++++
 
 If the kernel may use flat operations to access scratch memory, the prolog code
 must set up FLAT_SCRATCH register pair (FLAT_SCRATCH_LO/FLAT_SCRATCH_HI which
-are in SGPRn-4/SGPRn-3). Initialization uses Flat Scratch Init and Scratch Wavefront
-Offset SGPR registers (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`):
+are in SGPRn-4/SGPRn-3). Initialization uses Flat Scratch Init and Scratch
+Wavefront Offset SGPR registers (see
+:ref:`amdgpu-amdhsa-initial-kernel-execution-state`):
 
 GFX6
   Flat scratch is not supported.
 
 GFX7-GFX8
-  1. The low word of Flat Scratch Init is 32 bit byte offset from
+
+  1. The low word of Flat Scratch Init is 32-bit byte offset from
      ``SH_HIDDEN_PRIVATE_BASE_VIMID`` to the base of scratch backing memory
      being managed by SPI for the queue executing the kernel dispatch. This is
      the same value used in the Scratch Segment Buffer V# base address. The
-     prolog must add the value of Scratch Wavefront Offset to get the wavefront's byte
-     scratch backing memory offset from ``SH_HIDDEN_PRIVATE_BASE_VIMID``. Since
-     FLAT_SCRATCH_LO is in units of 256 bytes, the offset must be right shifted
-     by 8 before moving into FLAT_SCRATCH_LO.
-  2. The second word of Flat Scratch Init is 32 bit byte size of a single
+     prolog must add the value of Scratch Wavefront Offset to get the
+     wavefront's byte scratch backing memory offset from
+     ``SH_HIDDEN_PRIVATE_BASE_VIMID``. Since FLAT_SCRATCH_LO is in units of 256
+     bytes, the offset must be right shifted by 8 before moving into
+     FLAT_SCRATCH_LO.
+  2. The second word of Flat Scratch Init is 32-bit byte size of a single
      work-items scratch memory usage. This is directly loaded from the kernel
      dispatch packet Private Segment Byte Size and rounded up to a multiple of
      DWORD. Having CP load it once avoids loading it at the beginning of every
-     wavefront. The prolog must move it to FLAT_SCRATCH_LO for use as FLAT SCRATCH
-     SIZE.
+     wavefront. The prolog must move it to FLAT_SCRATCH_LO for use as FLAT
+     SCRATCH SIZE.
 
 GFX9-GFX10
-  The Flat Scratch Init is the 64 bit address of the base of scratch backing
+  The Flat Scratch Init is the 64-bit address of the base of scratch backing
   memory being managed by SPI for the queue executing the kernel dispatch. The
-  prolog must add the value of Scratch Wavefront Offset and moved to the FLAT_SCRATCH
-  pair for use as the flat scratch base in flat memory instructions.
+  prolog must add the value of Scratch Wavefront Offset and moved to the
+  FLAT_SCRATCH pair for use as the flat scratch base in flat memory
+  instructions.
 
 .. _amdgpu-amdhsa-memory-model:
 
 Memory Model
 ~~~~~~~~~~~~
 
 This section describes the mapping of LLVM memory model onto AMDGPU machine code
-(see :ref:`memmodel`). *The implementation is WIP.*
-
-.. TODO
-   Update when implementation complete.
+(see :ref:`memmodel`).
 
 The AMDGPU backend supports the memory synchronization scopes specified in
 :ref:`amdgpu-memory-scopes`.
 
 The code sequences used to implement the memory model are defined in table
 :ref:`amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx10-table`.
 
 The sequences specify the order of instructions that a single thread must
The AMDGPU backend supports the following memory models:
   HSA Memory Model [HSA]_
     The HSA memory model uses a single happens-before relation for all address
     spaces (see :ref:`amdgpu-address-spaces`).
   OpenCL Memory Model [OpenCL]_
     The OpenCL memory model which has separate happens-before relations for the
     global and local address spaces. Only a fence specifying both global and
     local address space, and seq_cst instructions join the relationships. Since
     the LLVM ``memfence`` instruction does not allow an address space to be
-    specified the OpenCL fence has to convervatively assume both local and
+    specified the OpenCL fence has to conservatively assume both local and
     global address space was specified. However, optimizations can often be
     done to eliminate the additional ``s_waitcnt`` instructions when there are
     no intervening memory instructions which access the corresponding address
     space. The code sequences in the table indicate what can be omitted for the
     OpenCL memory. The target triple environment is used to determine if the
     source language is OpenCL (see :ref:`amdgpu-opencl`).
 
 ``ds/flat_load/store/atomic`` instructions to local memory are termed LDS

 * The wavefronts for a single work-group are executed in the same CU but may be
   executed by different SIMDs.
 * Each CU has a single LDS memory shared by the wavefronts of the work-groups
   executing on it.
 * All LDS operations of a CU are performed as wavefront wide operations in a
   global order and involve no caching. Completion is reported to a wavefront in
   execution order.
 * The LDS memory has multiple request queues shared by the SIMDs of a
-  CU. Therefore, the LDS operations performed by different wavefronts of a work-group
-  can be reordered relative to each other, which can result in reordering the
-  visibility of vector memory operations with respect to LDS operations of other
-  wavefronts in the same work-group. A ``s_waitcnt lgkmcnt(0)`` is required to
-  ensure synchronization between LDS operations and vector memory operations
-  between wavefronts of a work-group, but not between operations performed by the
-  same wavefront.
+  CU. Therefore, the LDS operations performed by different wavefronts of a
+  work-group can be reordered relative to each other, which can result in
+  reordering the visibility of vector memory operations with respect to LDS
+  operations of other wavefronts in the same work-group. A ``s_waitcnt
+  lgkmcnt(0)`` is required to ensure synchronization between LDS operations and
+  vector memory operations between wavefronts of a work-group, but not between
+  operations performed by the same wavefront.
 * The vector memory operations are performed as wavefront wide operations and
   completion is reported to a wavefront in execution order. The exception is
   that for GFX7-GFX9 ``flat_load/store/atomic`` instructions can report out of
   vector memory order if they access LDS memory, and out of LDS operation order
   if they access global memory.
 * The vector memory operations access a single vector L1 cache shared by all
   SIMDs a CU. Therefore, no special action is required for coherence between the
   lanes of a single wavefront, or for coherence between wavefronts in the same
-  work-group. A ``buffer_wbinvl1_vol`` is required for coherence between wavefronts
-  executing in different work-groups as they may be executing on different CUs.
+  work-group. A ``buffer_wbinvl1_vol`` is required for coherence between
+  wavefronts executing in different work-groups as they may be executing on
+  different CUs.
 * The scalar memory operations access a scalar L1 cache shared by all wavefronts
   on a group of CUs. The scalar and vector L1 caches are not coherent. However,
   scalar operations are used in a restricted way so do not impact the memory
-  model. See :ref:`amdgpu-amdhsa-memory-spaces`.
+  model. See :ref:`amdgpu-address-spaces`.
 * The vector and scalar memory operations use an L2 cache shared by all CUs on
   the same agent.
 * The L2 cache has independent channels to service disjoint ranges of virtual
   addresses.
 * Each CU has a separate request queue per channel. Therefore, the vector and
-  scalar memory operations performed by wavefronts executing in different work-groups
-  (which may be executing on different CUs) of an agent can be reordered
-  relative to each other. A ``s_waitcnt vmcnt(0)`` is required to ensure
-  synchronization between vector memory operations of different CUs. It ensures a
-  previous vector memory operation has completed before executing a subsequent
-  vector memory or LDS operation and so can be used to meet the requirements of
-  acquire and release.
+  scalar memory operations performed by wavefronts executing in different
+  work-groups (which may be executing on different CUs) of an agent can be
+  reordered relative to each other. A ``s_waitcnt vmcnt(0)`` is required to
+  ensure synchronization between vector memory operations of different CUs. It
+  ensures a previous vector memory operation has completed before executing a
+  subsequent vector memory or LDS operation and so can be used to meet the
+  requirements of acquire and release.
 * The L2 cache can be kept coherent with other agents on some targets, or ranges
   of virtual addresses can be set up to bypass it to ensure system coherence.
 
 For GFX10:
 
 * Each agent has multiple shader arrays (SA).
 * Each SA has multiple work-group processors (WGP).
 * Each WGP has multiple compute units (CU).
 * Each CU has multiple SIMDs that execute wavefronts.
 * The wavefronts for a single work-group are executed in the same
   WGP. In CU wavefront execution mode the wavefronts may be executed by
   different SIMDs in the same CU. In WGP wavefront execution mode the
   wavefronts may be executed by different SIMDs in different CUs in the same
   WGP.
 * Each WGP has a single LDS memory shared by the wavefronts of the work-groups
   executing on it.
 * All LDS operations of a WGP are performed as wavefront wide operations in a
   global order and involve no caching. Completion is reported to a wavefront in
   execution order.
 * The LDS memory has multiple request queues shared by the SIMDs of a
-  WGP. Therefore, the LDS operations performed by different wavefronts of a work-group
-  can be reordered relative to each other, which can result in reordering the
-  visibility of vector memory operations with respect to LDS operations of other
-  wavefronts in the same work-group. A ``s_waitcnt lgkmcnt(0)`` is required to
-  ensure synchronization between LDS operations and vector memory operations
-  between wavefronts of a work-group, but not between operations performed by the
-  same wavefront.
+  WGP. Therefore, the LDS operations performed by different wavefronts of a
+  work-group can be reordered relative to each other, which can result in
+  reordering the visibility of vector memory operations with respect to LDS
+  operations of other wavefronts in the same work-group. A ``s_waitcnt
+  lgkmcnt(0)`` is required to ensure synchronization between LDS operations and
+  vector memory operations between wavefronts of a work-group, but not between
+  operations performed by the same wavefront.
 * The vector memory operations are performed as wavefront wide operations.
   Completion of load/store/sample operations are reported to a wavefront in
   execution order of other load/store/sample operations performed by that
   wavefront.
 * The vector memory operations access a vector L0 cache. There is a single L0
-  cache per CU. Each SIMD of a CU accesses the same L0 cache.
-  Therefore, no special action is required for coherence between the lanes of a
-  single wavefront. However, a ``BUFFER_GL0_INV`` is required for coherence
-  between wavefronts executing in the same work-group as they may be executing on
-  SIMDs of different CUs that access different L0s. A ``BUFFER_GL0_INV`` is also
-  required for coherence between wavefronts executing in different work-groups as
-  they may be executing on different WGPs.
+  cache per CU. Each SIMD of a CU accesses the same L0 cache. Therefore, no
+  special action is required for coherence between the lanes of a single
+  wavefront. However, a ``BUFFER_GL0_INV`` is required for coherence between
+  wavefronts executing in the same work-group as they may be executing on SIMDs
+  of different CUs that access different L0s. A ``BUFFER_GL0_INV`` is also
+  required for coherence between wavefronts executing in different work-groups
+  as they may be executing on different WGPs.
 * The scalar memory operations access a scalar L0 cache shared by all wavefronts
   on a WGP. The scalar and vector L0 caches are not coherent. However, scalar
   operations are used in a restricted way so do not impact the memory model. See
-  :ref:`amdgpu-amdhsa-memory-spaces`.
+  :ref:`amdgpu-address-spaces`.
 * The vector and scalar memory L0 caches use an L1 cache shared by all WGPs on
   the same SA. Therefore, no special action is required for coherence between
   the wavefronts of a single work-group. However, a ``BUFFER_GL1_INV`` is
-  required for coherence between wavefronts executing in different work-groups as
-  they may be executing on different SAs that access different L1s.
+  required for coherence between wavefronts executing in different work-groups
+  as they may be executing on different SAs that access different L1s.
 * The L1 caches have independent quadrants to service disjoint ranges of virtual
   addresses.
 * Each L0 cache has a separate request queue per L1 quadrant. Therefore, the
   vector and scalar memory operations performed by different wavefronts, whether
   executing in the same or different work-groups (which may be executing on
   different CUs accessing different L0s), can be reordered relative to each
-  other. A ``s_waitcnt vmcnt(0) & vscnt(0)`` is required to ensure synchronization
-  between vector memory operations of different wavefronts. It ensures a previous
-  vector memory operation has completed before executing a subsequent vector
-  memory or LDS operation and so can be used to meet the requirements of acquire,
-  release and sequential consistency.
+  other. A ``s_waitcnt vmcnt(0) & vscnt(0)`` is required to ensure
+  synchronization between vector memory operations of different wavefronts. It
+  ensures a previous vector memory operation has completed before executing a
+  subsequent vector memory or LDS operation and so can be used to meet the
+  requirements of acquire, release and sequential consistency.
 * The L1 caches use an L2 cache shared by all SAs on the same agent.
 * The L2 cache has independent channels to service disjoint ranges of virtual
   addresses.
 * Each L1 quadrant of a single SA accesses a different L2 channel. Each L1
   quadrant has a separate request queue per L2 channel. Therefore, the vector
   and scalar memory operations performed by wavefronts executing in different
   work-groups (which may be executing on different SAs) of an agent can be
   reordered relative to each other. A ``s_waitcnt vmcnt(0) & vscnt(0)`` is
   required to ensure synchronization between vector memory operations of
   different SAs. It ensures a previous vector memory operation has completed
   before executing a subsequent vector memory and so can be used to meet the
   requirements of acquire, release and sequential consistency.
 * The L2 cache can be kept coherent with other agents on some targets, or ranges
   of virtual addresses can be set up to bypass it to ensure system coherence.
 
-Private address space uses ``buffer_load/store`` using the scratch V# (GFX6-GFX8),
-or ``scratch_load/store`` (GFX9-GFX10). Since only a single thread is accessing the
-memory, atomic memory orderings are not meaningful and all accesses are treated
-as non-atomic.
+Private address space uses ``buffer_load/store`` using the scratch V#
+(GFX6-GFX8), or ``scratch_load/store`` (GFX9-GFX10). Since only a single thread
+is accessing the memory, atomic memory orderings are not meaningful and all
+accesses are treated as non-atomic.
 
 Constant address space uses ``buffer/global_load`` instructions (or equivalent
 scalar memory instructions). Since the constant address space contents do not
 change during the execution of a kernel dispatch it is not legal to perform
 stores, and atomic memory orderings are not meaningful and all access are
 treated as non-atomic.
 
 A memory synchronization scope wider than work-group is not meaningful for the

 
 AMDGPU backend only uses scalar memory operations to access memory that is
 proven to not change during the execution of the kernel dispatch. This includes
 constant address space and global address space for program scope const
 variables. Therefore the kernel machine code does not have to maintain the
 scalar L1 cache to ensure it is coherent with the vector L1 cache. The scalar
 and vector L1 caches are invalidated between kernel dispatches by CP since
 constant address space data may change between kernel dispatch executions. See
-:ref:`amdgpu-amdhsa-memory-spaces`.
+:ref:`amdgpu-address-spaces`.
 
-The one execption is if scalar writes are used to spill SGPR registers. In this
+The one exception is if scalar writes are used to spill SGPR registers. In this
 case the AMDGPU backend ensures the memory location used to spill is never
 accessed by vector memory operations at the same time. If scalar writes are used
 then a ``s_dcache_wb`` is inserted before the ``s_endpgm`` and before a function
 return since the locations may be used for vector memory instructions by a
-future wavefront that uses the same scratch area, or a function call that creates a
-frame at the same address, respectively. There is no need for a ``s_dcache_inv``
-as all scalar writes are write-before-read in the same thread.
-
-For GFX6-GFX9, scratch backing memory (which is used for the private address space)
-is accessed with MTYPE NC_NV (non-coherenent non-volatile). Since the private
-address space is only accessed by a single thread, and is always
+future wavefront that uses the same scratch area, or a function call that
+creates a frame at the same address, respectively. There is no need for a
+``s_dcache_inv`` as all scalar writes are write-before-read in the same thread.
+
+For GFX6-GFX9, scratch backing memory (which is used for the private address
+space) is accessed with MTYPE NC_NV (non-coherent non-volatile). Since the
+private address space is only accessed by a single thread, and is always
 write-before-read, there is never a need to invalidate these entries from the L1
 cache. Hence all cache invalidates are done as ``*_vol`` to only invalidate the
 volatile cache lines.
 
 For GFX10, scratch backing memory (which is used for the private address space)
-is accessed with MTYPE NC (non-coherenent). Since the private address space is
+is accessed with MTYPE NC (non-coherent). Since the private address space is
 only accessed by a single thread, and is always write-before-read, there is
 never a need to invalidate these entries from the L0 or L1 caches.
 
-For GFX10, wavefronts are executed in native mode with in-order reporting of loads
-and sample instructions. In this mode vmcnt reports completion of load, atomic
-with return and sample instructions in order, and the vscnt reports the
-completion of store and atomic without return in order. See ``MEM_ORDERED`` field
-in :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
+For GFX10, wavefronts are executed in native mode with in-order reporting of
+loads and sample instructions. In this mode vmcnt reports completion of load,
+atomic with return and sample instructions in order, and the vscnt reports the
+completion of store and atomic without return in order. See ``MEM_ORDERED``
+field in :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
 
 In GFX10, wavefronts can be executed in WGP or CU wavefront execution mode:
 
 * In WGP wavefront execution mode the wavefronts of a work-group are executed
   on the SIMDs of both CUs of the WGP. Therefore, explicit management of the per
-  CU L0 caches is required for work-group synchronization. Also accesses to L1 at
-  work-group scope need to be expicitly ordered as the accesses from different
-  CUs are not ordered.
+  CU L0 caches is required for work-group synchronization. Also accesses to L1
+  at work-group scope need to be explicitly ordered as the accesses from
+  different CUs are not ordered.
 * In CU wavefront execution mode the wavefronts of a work-group are executed on
   the SIMDs of a single CU of the WGP. Therefore, all global memory access by
   the work-group access the same L0 which in turn ensures L1 accesses are
   ordered and so do not require explicit management of the caches for
   work-group synchronization.
 
-See ``WGP_MODE`` field in :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`
-and :ref:`amdgpu-target-features`.
+See ``WGP_MODE`` field in
+:ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table` and
+:ref:`amdgpu-target-features`.
 
 On dGPU the kernarg backing memory is accessed as UC (uncached) to avoid needing
 to invalidate the L2 cache. For GFX6-GFX9, this also causes it to be treated as
 non-volatile and so is not invalidated by ``*_vol``. On APU it is accessed as CC
 (cache coherent) and so the L2 cache will be coherent with the CPU and other
 agents.
 
   .. table:: AMDHSA Memory Model Code Sequences GFX6-GFX10
.. table:: AMDHSA Memory Model Code Sequences GFX6-GFX10
 
      load         *none*       *none*         - local    1. ds_load                      1. ds_load
      store        *none*       *none*         - global   - !nontemporal                  - !nontemporal
                                               - generic
                                               - private    1. buffer/global/flat_store     1. buffer/global/flat_store
                                               - constant
                                                          - nontemporal                   - nontemporal
 
-                                                           1. buffer/global/flat_stote      1. buffer/global/flat_store
+                                                           1. buffer/global/flat_store      1. buffer/global/flat_store
                                                               glc=1 slc=1                      slc=1
 
      store        *none*       *none*         - local    1. ds_store                     1. ds_store
      **Unordered Atomic**
      ----------------------------------------------------------------------------------------------------------------------
      load atomic  unordered    *any*          *any*      *Same as non-atomic*.           *Same as non-atomic*.
      store atomic unordered    *any*          *any*      *Same as non-atomic*.           *Same as non-atomic*.
      atomicrmw    unordered    *any*          *any*      *Same as monotonic              *Same as monotonic
load atomic  acquire      - workgroup    - global   1. buffer/global/flat_load      1. buffer/global_load glc=1
                                                                                          2. s_waitcnt vmcnt(0)
 
                                                                                            - If CU wavefront execution mode, omit.
                                                                                            - Must happen before
                                                                                              the following buffer_gl0_inv
                                                                                              and before any following
                                                                                              global/generic
                                                                                              load/load
-                                                                                             atomic/stote/store
+                                                                                             atomic/store/store
                                                                                              atomic/atomicrmw.
 
                                                                                          3. buffer_gl0_inv
 
                                                                                            - If CU wavefront execution mode, omit.
                                                                                            - Ensures that
                                                                                              following
                                                                                              loads will not see
atomicrmw    acquire      - workgroup    - global   1. buffer/global/flat_atomic    1. buffer/global_atomic
                                                                                            - Use vmcnt if atomic with
                                                                                              return and vscnt if atomic
                                                                                              with no-return.
                                                                                            - Must happen before
                                                                                              the following buffer_gl0_inv
                                                                                              and before any following
                                                                                              global/generic
                                                                                              load/load
-                                                                                             atomic/stote/store
+                                                                                             atomic/store/store
                                                                                              atomic/atomicrmw.
 
                                                                                          3. buffer_gl0_inv
 
                                                                                            - If CU wavefront execution mode, omit.
                                                                                            - Ensures that
                                                                                              following
                                                                                              loads will not see
.. table:: PAL Graphics Shader User Data Registers
      +             Spill Table (32-bit pointer)
      +             Draw Index (First Stage Only)
      +             Vertex Offset (First Stage Only)
      +             Instance Offset (First Stage Only)
      ============= ================================
 
   The placement of the global internal table remains fixed in the first *user
   data SGPR register*. Otherwise all parameters are optional, and can be mapped
-  to any desired *user data SGPR register*, with the following regstrictions:
+  to any desired *user data SGPR register*, with the following restrictions:
 
   * Draw Index, Vertex Offset, and Instance Offset can only be used by the first
-    activehardware stage in a graphics pipeline (i.e. where the API vertex
+    active hardware stage in a graphics pipeline (i.e. where the API vertex
     shader runs).
 
   * Application-controlled user data must be mapped into a contiguous range of
     user data registers.
 
   * The application-controlled user data range supports compaction remapping, so
     only *entries* that are actually consumed by the shader must be assigned to
     corresponding *registers*. Note that in order to support an efficient runtime
     implementation, the remapping must pack *registers* in the same order as
     *entries*, with unused *entries* removed.
 
 .. _pal_global_internal_table:
 
 Global Internal Table
 ~~~~~~~~~~~~~~~~~~~~~
 
-The global internal table is a table of *shader resource descriptors* (SRDs) that
-define how certain engine-wide, runtime-managed resources should be accessed
-from a shader. The majority of these resources have HW-defined formats, and it
-is up to the compiler to write/read data as required by the target hardware.
+The global internal table is a table of *shader resource descriptors* (SRDs)
+that define how certain engine-wide, runtime-managed resources should be
+accessed from a shader. The majority of these resources have HW-defined formats,
+and it is up to the compiler to write/read data as required by the target
+hardware.
 
 The following table illustrates the required format:
 
   .. table:: PAL Global Internal Table
      :name: pal-git-table
 
      ============= ================================
      Offset        Description
.. toctree::
    AMDGPU/AMDGPUAsmGFX10
    AMDGPUModifierSyntax
    AMDGPUOperandSyntax
    AMDGPUInstructionSyntax
    AMDGPUInstructionNotation
 
 An instruction has the following :doc:`syntax<AMDGPUInstructionSyntax>`:
 
-    ``<``\ *opcode*\ ``>    <``\ *operand0*\ ``>, <``\ *operand1*\ ``>,...    <``\ *modifier0*\ ``> <``\ *modifier1*\ ``>...``
+  | ``<``\ *opcode*\ ``> <``\ *operand0*\ ``>, <``\ *operand1*\ ``>,...
+    <``\ *modifier0*\ ``> <``\ *modifier1*\ ``>...``
 
 :doc:`Operands<AMDGPUOperandSyntax>` are normally comma-separated while
 :doc:`modifiers<AMDGPUModifierSyntax>` are space-separated.
 
 The order of *operands* and *modifiers* is fixed.
 Most *modifiers* are optional and may be omitted.
 
-See detailed instruction syntax description for :doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`,
-:doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`, :doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`
-and :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>`.
+See detailed instruction syntax description for
+:doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`, :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`,
+:doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`, and :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>`.
 
 Note that features under development are not included in this description.
 
-For more information about instructions, their semantics and supported combinations of
-operands, refer to one of instruction set architecture manuals
+For more information about instructions, their semantics and supported
+combinations of operands, refer to one of instruction set architecture manuals
 [AMD-GCN-GFX6]_, [AMD-GCN-GFX7]_, [AMD-GCN-GFX8]_, [AMD-GCN-GFX9]_ and
 [AMD-GCN-GFX10]_.
 
 Operands
 ~~~~~~~~
 
 Detailed description of operands may be found :doc:`here<AMDGPUOperandSyntax>`.
 
 Modifiers
 ~~~~~~~~~
 
-Detailed description of modifiers may be found :doc:`here<AMDGPUModifierSyntax>`.
+Detailed description of modifiers may be found
+:doc:`here<AMDGPUModifierSyntax>`.
 
 Instruction Examples
 ~~~~~~~~~~~~~~~~~~~~
 
 DS
 ++
 
 .. code-block:: nasm
 
   ds_add_u32 v2, v4 offset:16
   ds_write_src2_b64 v2 offset0:4 offset1:8
   ds_cmpst_f32 v2, v4, v6
   ds_min_rtn_f64 v[8:9], v2, v[4:5]
 
-
-For full list of supported instructions, refer to "LDS/GDS instructions" in ISA Manual.
+For full list of supported instructions, refer to "LDS/GDS instructions" in ISA
+Manual.
 
 FLAT
 ++++
 
 .. code-block:: nasm
 
   flat_load_dword v1, v[3:4]
   flat_store_dwordx3 v[3:4], v[5:7]
   flat_atomic_swap v1, v[3:4], v5 glc
   flat_atomic_cmpswap v1, v[3:4], v[5:6] glc slc
   flat_atomic_fmax_x2 v[1:2], v[3:4], v[5:6] glc
 
-For full list of supported instructions, refer to "FLAT instructions" in ISA Manual.
+For full list of supported instructions, refer to "FLAT instructions" in ISA
+Manual.
 
 MUBUF
 +++++
 
 .. code-block:: nasm
 
   buffer_load_dword v1, off, s[4:7], s1
   buffer_store_dwordx4 v[1:4], v2, ttmp[4:7], s1 offen offset:4 glc tfe
   buffer_store_format_xy v[1:2], off, s[4:7], s1
   buffer_wbinvl1
   buffer_atomic_inc v1, v2, s[8:11], s4 idxen offset:4 slc
 
-For full list of supported instructions, refer to "MUBUF Instructions" in ISA Manual.
+For full list of supported instructions, refer to "MUBUF Instructions" in ISA
+Manual.
 
 SMRD/SMEM
 +++++++++
 
 .. code-block:: nasm
 
   s_load_dword s1, s[2:3], 0xfc
   s_load_dwordx8 s[8:15], s[2:3], s4
   s_load_dwordx16 s[88:103], s[2:3], s4
   s_dcache_inv_vol
   s_memtime s[4:5]
 
-For full list of supported instructions, refer to "Scalar Memory Operations" in ISA Manual.
+For full list of supported instructions, refer to "Scalar Memory Operations" in
+ISA Manual.
 
 SOP1
 ++++
 
 .. code-block:: nasm
 
   s_mov_b32 s1, s2
   s_mov_b64 s[0:1], 0x80000000
   s_cmov_b32 s1, 200
   s_wqm_b64 s[2:3], s[4:5]
   s_bcnt0_i32_b64 s1, s[2:3]
   s_swappc_b64 s[2:3], s[4:5]
   s_cbranch_join s[4:5]
 
-For full list of supported instructions, refer to "SOP1 Instructions" in ISA Manual.
+For full list of supported instructions, refer to "SOP1 Instructions" in ISA
+Manual.
 
 SOP2
 ++++
 
 .. code-block:: nasm
 
   s_add_u32 s1, s2, s3
   s_and_b64 s[2:3], s[4:5], s[6:7]
   s_cselect_b32 s1, s2, s3
   s_andn2_b32 s2, s4, s6
   s_lshr_b64 s[2:3], s[4:5], s6
   s_ashr_i32 s2, s4, s6
   s_bfm_b64 s[2:3], s4, s6
   s_bfe_i64 s[2:3], s[4:5], s6
   s_cbranch_g_fork s[4:5], s[6:7]
 
-For full list of supported instructions, refer to "SOP2 Instructions" in ISA Manual.
+For full list of supported instructions, refer to "SOP2 Instructions" in ISA
+Manual.
 
 SOPC
 ++++
 
 .. code-block:: nasm
 
   s_cmp_eq_i32 s1, s2
   s_bitcmp1_b32 s1, s2
   s_bitcmp0_b64 s[2:3], s4
   s_setvskip s3, s5
 
-For full list of supported instructions, refer to "SOPC Instructions" in ISA Manual.
+For full list of supported instructions, refer to "SOPC Instructions" in ISA
+Manual.
 
 SOPP
 ++++
 
 .. code-block:: nasm
 
   s_barrier
   s_nop 2
   s_endpgm
   s_waitcnt 0 ; Wait for all counters to be 0
   s_waitcnt vmcnt(0) & expcnt(0) & lgkmcnt(0) ; Equivalent to above
   s_waitcnt vmcnt(1) ; Wait for vmcnt counter to be 1.
   s_sethalt 9
   s_sleep 10
   s_sendmsg 0x1
   s_sendmsg sendmsg(MSG_INTERRUPT)
   s_trap 1
 
-For full list of supported instructions, refer to "SOPP Instructions" in ISA Manual.
+For full list of supported instructions, refer to "SOPP Instructions" in ISA
+Manual.
 
 Unless otherwise mentioned, little verification is performed on the operands
 of SOPP Instructions, so it is up to the programmer to be familiar with the
 range or acceptable values.
 
 VALU
 ++++
 
 For vector ALU instruction opcodes (VOP1, VOP2, VOP3, VOPC, VOP_DPP, VOP_SDWA),
-the assembler will automatically use optimal encoding based on its operands.
-To force specific encoding, one can add a suffix to the opcode of the instruction:
+the assembler will automatically use optimal encoding based on its operands. To
+force specific encoding, one can add a suffix to the opcode of the instruction:
 
 * _e32 for 32-bit VOP1/VOP2/VOPC
 * _e64 for 64-bit VOP3
 * _dpp for VOP_DPP
 * _sdwa for VOP_SDWA
 
 VOP1/VOP2/VOP3/VOPC examples:
 
.. code-block:: nasm
   v_mov_b32 v1, v2 dst_sel:BYTE_0 dst_unused:UNUSED_PRESERVE src0_sel:DWORD
   v_min_u32 v200, v200, v1 dst_sel:WORD_1 dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD
   v_sin_f32 v0, v0 dst_unused:UNUSED_PAD src0_sel:WORD_1
   v_fract_f32 v0, |v0| dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
   v_cmpx_le_u32 vcc, v1, v2 src0_sel:BYTE_2 src1_sel:WORD_0
 
 For full list of supported instructions, refer to "Vector ALU instructions".
 
-.. TODO
-   Remove once we switch to code object v3 by default.
+.. TODO::
+
 
 .. _amdgpu-amdhsa-assembler-predefined-symbols-v2:
 
 Code Object V2 Predefined Symbols (-mattr=-code-object-v3)
 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
 .. warning:: Code Object V2 is not the default code object version emitted by
   this version of LLVM. For a description of the predefined symbols available

 
 .hsa_code_object_isa [major, minor, stepping, vendor, arch]
 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
 
 
 *major*, *minor*, and *stepping* are all integers that describe the instruction
 set architecture (ISA) version of the assembly program.
 
-*vendor* and *arch* are quoted strings.  *vendor* should always be equal to
+*vendor* and *arch* are quoted strings. *vendor* should always be equal to
 "AMD" and *arch* should always be equal to "AMDGPU".
 
 By default, the assembler will derive the ISA version, *vendor*, and *arch*
 from the value of the -mcpu option that is passed to the assembler.
 
 .. _amdgpu-amdhsa-assembler-directive-amdgpu_hsa_kernel:
 
 .amdgpu_hsa_kernel (name)
 +++++++++++++++++++++++++
 
-This directives specifies that the symbol with given name is a kernel entry point
-(label) and the object should contain corresponding symbol of type STT_AMDGPU_HSA_KERNEL.
+This directives specifies that the symbol with given name is a kernel entry
+point (label) and the object should contain corresponding symbol of type
+STT_AMDGPU_HSA_KERNEL.
 
 .amd_kernel_code_t
 ++++++++++++++++++
 
 This directive marks the beginning of a list of key / value pairs that are used
 to specify the amd_kernel_code_t object that will be emitted by the assembler.
-The list must be terminated by the *.end_amd_kernel_code_t* directive.  For
-any amd_kernel_code_t values that are unspecified a default value will be
-used.  The default value for all keys is 0, with the following exceptions:
+The list must be terminated by the *.end_amd_kernel_code_t* directive. For any
+amd_kernel_code_t values that are unspecified a default value will be used. The
+default value for all keys is 0, with the following exceptions:
 
 - *amd_code_version_major* defaults to 1.
 - *amd_kernel_code_version_minor* defaults to 2.
 - *amd_machine_kind* defaults to 1.
 - *amd_machine_version_major*, *machine_version_minor*, and
   *amd_machine_version_stepping* are derived from the value of the -mcpu option
   that is passed to the assembler.
 - *kernel_code_entry_byte_offset* defaults to 256.

 
 .. warning:: Code Object V2 is not the default code object version emitted by
   this version of LLVM. For a description of the directives supported with
   the default configuration (Code Object V3) see
   :ref:`amdgpu-amdhsa-assembler-example-v3`.
 
 Here is an example of a minimal assembly source file, defining one HSA kernel:
 
-.. code-block:: none
+.. code::
+   :number-lines:
 
    .hsa_code_object_version 1,0
    .hsa_code_object_isa
 
    .hsatext
    .globl  hello_world
    .p2align 8
    .amdgpu_hsa_kernel hello_world

 
 .. _amdgpu-amdhsa-assembler-example-v3:
 
 Code Object V3 Example Source Code (-mattr=+code-object-v3)
 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 
 Here is an example of a minimal assembly source file, defining one HSA kernel:
 
-.. code-block:: none
+.. code::
+   :number-lines:
 
    .amdgcn_target "amdgcn-amd-amdhsa--gfx900+xnack" // optional
 
    .text
    .globl hello_world
    .p2align 8
    .type hello_world,@function
    hello_world:
      s_load_dwordx2 s[0:1], s[0:1] 0x0
      v_mov_b32 v0, 3.14159
      s_waitcnt lgkmcnt(0)
      v_mov_b32 v1, s0
      v_mov_b32 v2, s1
      flat_store_dword v[1:2], v0
      s_endpgm
    .Lfunc_end0:
      .size   hello_world, .Lfunc_end0-hello_world
 
    .rodata
    .p2align 6
    .amdhsa_kernel hello_world
      .amdhsa_user_sgpr_kernarg_segment_ptr 1
      .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
      .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
    .end_amdhsa_kernel
 
    .amdgpu_metadata
    ---
    amdhsa.version:
      - 1
      - 0
    amdhsa.kernels:
      - .name: hello_world
        .symbol: hello_world.kd
        .kernarg_segment_size: 48
        .group_segment_fixed_size: 0
        .private_segment_fixed_size: 0
        .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
+.. code::
+   :number-lines:
 
    .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>`__
 .. [AMD-GCN-GFX7] `AMD Sea Islands Series ISA <http://developer.amd.com/wordpress/media/2013/07/AMD_Sea_Islands_Instruction_Set_Architecture.pdf>`_
 .. [AMD-GCN-GFX8] `AMD GCN3 Instruction Set Architecture <http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/12/AMD_GCN3_Instruction_Set_Architecture_rev1.1.pdf>`__
 .. [AMD-GCN-GFX9] `AMD "Vega" Instruction Set Architecture <http://developer.amd.com/wordpress/media/2013/12/Vega_Shader_ISA_28July2017.pdf>`__
-.. [AMD-GCN-GFX10] AMD "Navi" Instruction Set Architecture *TBA*
-.. TODO
-   ttye Add link when made public.
+.. [AMD-GCN-GFX10] `AMD "RDNA 1.0" Instruction Set Architecture <https://gpuopen.com/wp-content/uploads/2019/08/RDNA_Shader_ISA_5August2019.pdf>`__
 .. [AMD-ROCm] `ROCm: Open Platform for Development, Discovery and Education Around GPU Computing <http://gpuopen.com/compute-product/rocm/>`__
 .. [AMD-ROCm-github] `ROCm github <http://github.com/RadeonOpenCompute>`__
 .. [HSA] `Heterogeneous System Architecture (HSA) Foundation <http://www.hsafoundation.com/>`__
 .. [ELF] `Executable and Linkable Format (ELF) <http://www.sco.com/developers/gabi/>`__
 .. [DWARF] `DWARF Debugging Information Format <http://dwarfstd.org/>`__
 .. [YAML] `YAML Ain't Markup Language (YAML™) Version 1.2 <http://www.yaml.org/spec/1.2/spec.html>`__
 .. [MsgPack] `Message Pack <http://www.msgpack.org/>`__
 .. [OpenCL] `The OpenCL Specification Version 2.0 <http://www.khronos.org/registry/cl/specs/opencl-2.0.pdf>`__
 .. [HRF] `Heterogeneous-race-free Memory Models <http://benedictgaster.org/wp-content/uploads/2014/01/asplos269-FINAL.pdf>`__
 .. [CLANG-ATTR] `Attributes in Clang <http://clang.llvm.org/docs/AttributeReference.html>`__