Skip to content

Commit 850fced

Browse files
committed
[NFC][AMDGPU] Corrections to AMD GPU initial kernel launch documentation
Reviewed By: rampitec Differential Revision: https://reviews.llvm.org/D99223
1 parent 19e402d commit 850fced

File tree

1 file changed

+4
-35
lines changed

1 file changed

+4
-35
lines changed

llvm/docs/AMDGPUUsage.rst

Lines changed: 4 additions & 35 deletions
Original file line numberDiff line numberDiff line change
@@ -4280,12 +4280,11 @@ SGPR register initial state is defined in
42804280
(enable_sgpr_dispatch_id) dispatch packet being
42814281
executed.
42824282
then Flat Scratch Init 2 See
4283-
:ref:`amdgpu-amdhsa-kernel-prolog-flat-scratch`.
4283+
(enable_sgpr_flat_scratch :ref:`amdgpu-amdhsa-kernel-prolog-flat-scratch`.
4284+
_init)
42844285
then Private Segment Size 1 The 32-bit byte size of a
4285-
(enable_sgpr_private single
4286-
work-item's
4287-
scratch_segment_size) memory
4288-
allocation. This is the
4286+
(enable_sgpr_private single work-item's memory
4287+
_segment_size) allocation. This is the
42894288
value from the kernel
42904289
dispatch packet Private
42914290
Segment Byte Size rounded up
@@ -4303,36 +4302,6 @@ SGPR register initial state is defined in
43034302
may be needed for GFX9-GFX10 which
43044303
changes the meaning of the
43054304
Flat Scratch Init value.
4306-
then Grid Work-Group Count X 1 32-bit count of the number of
4307-
(enable_sgpr_grid work-groups in the X dimension
4308-
_workgroup_count_X) for the grid being
4309-
executed. Computed from the
4310-
fields in the kernel dispatch
4311-
packet as ((grid_size.x +
4312-
workgroup_size.x - 1) /
4313-
workgroup_size.x).
4314-
then Grid Work-Group Count Y 1 32-bit count of the number of
4315-
(enable_sgpr_grid work-groups in the Y dimension
4316-
_workgroup_count_Y && for the grid being
4317-
less than 16 previous executed. Computed from the
4318-
SGPRs) fields in the kernel dispatch
4319-
packet as ((grid_size.y +
4320-
workgroup_size.y - 1) /
4321-
workgroupSize.y).
4322-
4323-
Only initialized if <16
4324-
previous SGPRs initialized.
4325-
then Grid Work-Group Count Z 1 32-bit count of the number of
4326-
(enable_sgpr_grid work-groups in the Z dimension
4327-
_workgroup_count_Z && for the grid being
4328-
less than 16 previous executed. Computed from the
4329-
SGPRs) fields in the kernel dispatch
4330-
packet as ((grid_size.z +
4331-
workgroup_size.z - 1) /
4332-
workgroupSize.z).
4333-
4334-
Only initialized if <16
4335-
previous SGPRs initialized.
43364305
then Work-Group Id X 1 32-bit work-group id in X
43374306
(enable_sgpr_workgroup_id dimension of grid for
43384307
_X) wavefront.

0 commit comments

Comments
 (0)