[PATCH] D99223: [NFC][AMDGPU] Corrections to AMD GPU initial kernel launch documentation

Tony Tye via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Thu Mar 25 19:06:08 PDT 2021


This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG850fcedb272f: [NFC][AMDGPU] Corrections to AMD GPU initial kernel launch documentation (authored by t-tye).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D99223/new/

https://reviews.llvm.org/D99223

Files:
  llvm/docs/AMDGPUUsage.rst


Index: llvm/docs/AMDGPUUsage.rst
===================================================================
--- llvm/docs/AMDGPUUsage.rst
+++ llvm/docs/AMDGPUUsage.rst
@@ -4280,12 +4280,11 @@
                 (enable_sgpr_dispatch_id)         dispatch packet being
                                                   executed.
      then       Flat Scratch Init          2      See
-                                                  :ref:`amdgpu-amdhsa-kernel-prolog-flat-scratch`.
+                (enable_sgpr_flat_scratch         :ref:`amdgpu-amdhsa-kernel-prolog-flat-scratch`.
+                _init)
      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
+                (enable_sgpr_private              single work-item's memory
+                _segment_size)                    allocation. This is the
                                                   value from the kernel
                                                   dispatch packet Private
                                                   Segment Byte Size rounded up
@@ -4303,36 +4302,6 @@
                                                   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
-                (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
-                (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
-                (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
                 (enable_sgpr_workgroup_id         dimension of grid for
                 _X)                               wavefront.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D99223.333485.patch
Type: text/x-patch
Size: 3982 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20210326/a3c56e68/attachment.bin>


More information about the llvm-commits mailing list