From 382c4642c9a5376ef33ae96d92d681abe7e19d46 Mon Sep 17 00:00:00 2001 From: Tony Date: Tue, 23 Mar 2021 22:38:10 +0000 Subject: [PATCH] [NFC][AMDGPU] Corrections to AMD GPU initial kernel launch documentation Reviewed By: rampitec Differential Revision: https://reviews.llvm.org/D99223 --- docs/AMDGPUUsage.rst | 39 ++++----------------------------------- 1 file changed, 4 insertions(+), 35 deletions(-) diff --git a/docs/AMDGPUUsage.rst b/docs/AMDGPUUsage.rst index f397d7542d2..51fd90e058a 100644 --- a/docs/AMDGPUUsage.rst +++ b/docs/AMDGPUUsage.rst @@ -4280,12 +4280,11 @@ SGPR register initial state is defined in (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 @@ SGPR register initial state is defined in 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.