This is an archive of the discontinued LLVM Phabricator instance.

Differential D38902

AMDGPU: Add AMDGPU HSA Kernel Descriptor
ClosedPublic

Authored by kzhuravl on Oct 13 2017, 1:08 PM.

Details

Reviewers
t-tye
b-sumner
arsenm
Commits
rG13376a4bdfb2: AMDGPU: Add AMDGPU HSA Kernel Descriptor
rL315822: AMDGPU: Add AMDGPU HSA Kernel Descriptor
Summary
  • Update docs to match llvm coding style
  • Add missing FP16_OVFL bit for gfx9
  • Fix the size of the kernel descriptor in the docs
  • Generation in separate patch

Diff Detail

Event Timeline

kzhuravl created this revision.Oct 13 2017, 1:08 PM
Herald added subscribers: tpr, dstuttard, yaxunl and 2 others. · View Herald TranscriptOct 13 2017, 1:08 PM
kzhuravl updated this revision to Diff 118965.Oct 13 2017, 1:14 PM

Update FP16_OVFL wording.

t-tye accepted this revision.Oct 13 2017, 8:13 PM

LGTM with minor layout suggestion.

docs/AMDGPUUsage.rst
1489

Maybe Size will fit on the line above now? Same comment for following rows.

This revision is now accepted and ready to land.Oct 13 2017, 8:13 PM
Closed by commit rL315822: AMDGPU: Add AMDGPU HSA Kernel Descriptor (authored by kzhuravl). · Explain WhyOct 14 2017, 12:17 PM
This revision was automatically updated to reflect the committed changes.
kzhuravl marked an inline comment as done.

Revision Contents

PathSize
docs/
205 lines
include/
llvm/
Support/
139 lines

Diff 118965

docs/AMDGPUUsage.rst

Show First 20 Lines • Show All 1,411 Lines • ▼ Show 20 Lines
CP microcode requires the Kernel descritor to be allocated on 64 byte alignment. CP microcode requires the Kernel descritor to be allocated on 64 byte alignment.
.. table:: Kernel Descriptor for GFX6-GFX9 .. table:: Kernel Descriptor for GFX6-GFX9
:name: amdgpu-amdhsa-kernel-descriptor-gfx6-gfx9-table :name: amdgpu-amdhsa-kernel-descriptor-gfx6-gfx9-table
======= ======= =============================== =========================== ======= ======= =============================== ===========================
Bits Size Field Name Description Bits Size Field Name Description
======= ======= =============================== =========================== ======= ======= =============================== ===========================
31:0 4 bytes group_segment_fixed_size The amount of fixed local 31:0 4 bytes GroupSegmentFixedSize The amount of fixed local
address space memory address space memory
required for a work-group required for a work-group
in bytes. This does not in bytes. This does not
include any dynamically include any dynamically
allocated local address allocated local address
space memory that may be space memory that may be
added when the kernel is added when the kernel is
dispatched. dispatched.
63:32 4 bytes private_segment_fixed_size The amount of fixed 63:32 4 bytes PrivateSegmentFixedSize The amount of fixed
private address space private address space
memory required for a memory required for a
work-item in bytes. If work-item in bytes. If
is_dynamic_callstack is 1 is_dynamic_callstack is 1
then additional space must then additional space must
be added to this value for be added to this value for
the call stack. the call stack.
95:64 4 bytes max_flat_workgroup_size Maximum flat work-group 95:64 4 bytes MaxFlatWorkgroupSize Maximum flat work-group
size supported by the size supported by the
kernel in work-items. kernel in work-items.
96 1 bit is_dynamic_call_stack Indicates if the generated 96 1 bit IsDynamicCallStack Indicates if the generated
machine code is using a machine code is using a
dynamically sized call dynamically sized call
stack. stack.
97 1 bit is_xnack_enabled Indicates if the generated 97 1 bit IsXNACKEnabled Indicates if the generated
machine code is capable of machine code is capable of
suppoting XNACK. suppoting XNACK.
127:98 30 bits Reserved. Must be 0. 127:98 30 bits Reserved. Must be 0.
191:128 8 bytes kernel_code_entry_byte_offset Byte offset (possibly 191:128 8 bytes KernelCodeEntryByteOffset Byte offset (possibly
negative) from base negative) from base
address of kernel address of kernel
descriptor to kernel's descriptor to kernel's
entry point instruction entry point instruction
which must be 256 byte which must be 256 byte
aligned. aligned.
383:192 24 Reserved. Must be 0. 383:192 24 Reserved. Must be 0.
bytes bytes
415:384 4 bytes compute_pgm_rsrc1 Compute Shader (CS) 415:384 4 bytes ComputePgmRsrc1 Compute Shader (CS)
program settings used by program settings used by
CP to set up CP to set up
``COMPUTE_PGM_RSRC1`` ``COMPUTE_PGM_RSRC1``
configuration configuration
register. See register. See
:ref:`amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table`. :ref:`amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table`.
447:416 4 bytes compute_pgm_rsrc2 Compute Shader (CS) 447:416 4 bytes ComputePgmRsrc2 Compute Shader (CS)
program settings used by program settings used by
CP to set up CP to set up
``COMPUTE_PGM_RSRC2`` ``COMPUTE_PGM_RSRC2``
configuration configuration
register. See register. See
:ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`. :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`.
448 1 bit enable_sgpr_private_segment Enable the setup of the 448 1 bit EnableSGPRPrivateSegment Enable the setup of the
_buffer SGPR user data registers Buffer SGPR user data registers
(see (see
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`). :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
The total number of SGPR The total number of SGPR
user data registers user data registers
requested must not exceed requested must not exceed
16 and match value in 16 and match value in
``compute_pgm_rsrc2.user_sgpr.user_sgpr_count``. ``compute_pgm_rsrc2.user_sgpr.user_sgpr_count``.
Any requests beyond 16 Any requests beyond 16
will be ignored. will be ignored.
449 1 bit enable_sgpr_dispatch_ptr *see above* 449 1 bit EnableSGPRDispatchPtr *see above*
450 1 bit enable_sgpr_queue_ptr *see above* 450 1 bit EnableSGPRQueuePtr *see above*
451 1 bit enable_sgpr_kernarg_segment_ptr *see above* 451 1 bit EnableSGPRKernargSegmentPtr *see above*
452 1 bit enable_sgpr_dispatch_id *see above* 452 1 bit EnableSGPRDispatchID *see above*
453 1 bit enable_sgpr_flat_scratch_init *see above* 453 1 bit EnableSGPRFlatScratchInit *see above*
454 1 bit enable_sgpr_private_segment *see above* 454 1 bit EnableSGPRPrivateSegment *see above*
_size Size
t-tyeUnsubmitted
Done
ReplyInline Actions

Maybe Size will fit on the line above now? Same comment for following rows.

t-tye: Maybe Size will fit on the line above now? Same comment for following rows.
455 1 bit enable_sgpr_grid_workgroup Not implemented in CP and 455 1 bit EnableSGPRGridWorkgroup Not implemented in CP and
_count_X should always be 0. CountX should always be 0.
456 1 bit enable_sgpr_grid_workgroup Not implemented in CP and 456 1 bit EnableSGPRGridWorkgroup Not implemented in CP and
_count_Y should always be 0. CountY should always be 0.
457 1 bit enable_sgpr_grid_workgroup Not implemented in CP and 457 1 bit EnableSGPRGridWorkgroup Not implemented in CP and
_count_Z should always be 0. CountZ should always be 0.
463:458 6 bits Reserved. Must be 0. 463:458 6 bits Reserved. Must be 0.
511:464 4 Reserved. Must be 0. 511:464 6 Reserved. Must be 0.
bytes bytes
512 **Total size 64 bytes.** 512 **Total size 64 bytes.**
======= =================================================================== ======= ===================================================================
.. ..
.. table:: compute_pgm_rsrc1 for GFX6-GFX9 .. table:: compute_pgm_rsrc1 for GFX6-GFX9
:name: amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table :name: amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table
======= ======= =============================== =========================================================================== ======= ======= =============================== ===========================================================================
Bits Size Field Name Description Bits Size Field Name Description
======= ======= =============================== =========================================================================== ======= ======= =============================== ===========================================================================
5:0 6 bits granulated_workitem_vgpr_count Number of vector registers 5:0 6 bits GRANULATED_WORKITEM_VGPR_COUNT Number of vector registers
used by each work-item, used by each work-item,
granularity is device granularity is device
specific: specific:
GFX6-9 GFX6-9
roundup((max-vgpg + 1) roundup((max-vgpg + 1)
/ 4) - 1 / 4) - 1
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.VGPRS``. ``COMPUTE_PGM_RSRC1.VGPRS``.
9:6 4 bits granulated_wavefront_sgpr_count Number of scalar registers 9:6 4 bits GRANULATED_WAVEFRONT_SGPR_COUNT Number of scalar registers
used by a wavefront, used by a wavefront,
granularity is device granularity is device
specific: specific:
GFX6-8 GFX6-8
roundup((max-sgpg + 1) roundup((max-sgpg + 1)
/ 8) - 1 / 8) - 1
GFX9 GFX9
roundup((max-sgpg + 1) roundup((max-sgpg + 1)
/ 16) - 1 / 16) - 1
Includes the special SGPRs Includes the special SGPRs
for VCC, Flat Scratch (for for VCC, Flat Scratch (for
GFX7 onwards) and XNACK GFX7 onwards) and XNACK
(for GFX8 onwards). It does (for GFX8 onwards). It does
not include the 16 SGPR not include the 16 SGPR
added if a trap handler is added if a trap handler is
enabled. enabled.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.SGPRS``. ``COMPUTE_PGM_RSRC1.SGPRS``.
11:10 2 bits priority Must be 0. 11:10 2 bits PRIORITY Must be 0.
Start executing wavefront Start executing wavefront
at the specified priority. at the specified priority.
CP is responsible for CP is responsible for
filling in filling in
``COMPUTE_PGM_RSRC1.PRIORITY``. ``COMPUTE_PGM_RSRC1.PRIORITY``.
13:12 2 bits float_mode_round_32 Wavefront starts execution 13:12 2 bits FLOAT_ROUND_MODE_32 Wavefront starts execution
with specified rounding with specified rounding
mode for single (32 mode for single (32
bit) floating point bit) floating point
precision floating point precision floating point
operations. operations.
Floating point rounding Floating point rounding
mode values are defined in mode values are defined in
:ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`. :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.FLOAT_MODE``. ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
15:14 2 bits float_mode_round_16_64 Wavefront starts execution 15:14 2 bits FLOAT_ROUND_MODE_16_64 Wavefront starts execution
with specified rounding with specified rounding
denorm mode for half/double (16 denorm mode for half/double (16
and 64 bit) floating point and 64 bit) floating point
precision floating point precision floating point
operations. operations.
Floating point rounding Floating point rounding
mode values are defined in mode values are defined in
:ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`. :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.FLOAT_MODE``. ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
17:16 2 bits float_mode_denorm_32 Wavefront starts execution 17:16 2 bits FLOAT_DENORM_MODE_32 Wavefront starts execution
with specified denorm mode with specified denorm mode
for single (32 for single (32
bit) floating point bit) floating point
precision floating point precision floating point
operations. operations.
Floating point denorm mode Floating point denorm mode
values are defined in values are defined in
:ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`. :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.FLOAT_MODE``. ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
19:18 2 bits float_mode_denorm_16_64 Wavefront starts execution 19:18 2 bits FLOAT_DENORM_MODE_16_64 Wavefront starts execution
with specified denorm mode with specified denorm mode
for half/double (16 for half/double (16
and 64 bit) floating point and 64 bit) floating point
precision floating point precision floating point
operations. operations.
Floating point denorm mode Floating point denorm mode
values are defined in values are defined in
:ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`. :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.FLOAT_MODE``. ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
20 1 bit priv Must be 0. 20 1 bit PRIV Must be 0.
Start executing wavefront Start executing wavefront
in privilege trap handler in privilege trap handler
mode. mode.
CP is responsible for CP is responsible for
filling in filling in
``COMPUTE_PGM_RSRC1.PRIV``. ``COMPUTE_PGM_RSRC1.PRIV``.
21 1 bit enable_dx10_clamp Wavefront starts execution 21 1 bit ENABLE_DX10_CLAMP Wavefront starts execution
with DX10 clamp mode with DX10 clamp mode
enabled. Used by the vector enabled. Used by the vector
ALU to force DX-10 style ALU to force DX-10 style
treatment of NaN's (when treatment of NaN's (when
set, clamp NaN to zero, set, clamp NaN to zero,
otherwise pass NaN otherwise pass NaN
through). through).
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.DX10_CLAMP``. ``COMPUTE_PGM_RSRC1.DX10_CLAMP``.
22 1 bit debug_mode Must be 0. 22 1 bit DEBUG_MODE Must be 0.
Start executing wavefront Start executing wavefront
in single step mode. in single step mode.
CP is responsible for CP is responsible for
filling in filling in
``COMPUTE_PGM_RSRC1.DEBUG_MODE``. ``COMPUTE_PGM_RSRC1.DEBUG_MODE``.
23 1 bit enable_ieee_mode Wavefront starts execution 23 1 bit ENABLE_IEEE_MODE Wavefront starts execution
with IEEE mode with IEEE mode
enabled. Floating point enabled. Floating point
opcodes that support opcodes that support
exception flag gathering exception flag gathering
will quiet and propagate will quiet and propagate
signaling-NaN inputs per signaling-NaN inputs per
IEEE 754-2008. Min_dx10 and IEEE 754-2008. Min_dx10 and
max_dx10 become IEEE max_dx10 become IEEE
754-2008 compliant due to 754-2008 compliant due to
signaling-NaN propagation signaling-NaN propagation
and quieting. and quieting.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC1.IEEE_MODE``. ``COMPUTE_PGM_RSRC1.IEEE_MODE``.
24 1 bit bulky Must be 0. 24 1 bit BULKY Must be 0.
Only one work-group allowed Only one work-group allowed
to execute on a compute to execute on a compute
unit. unit.
CP is responsible for CP is responsible for
filling in filling in
``COMPUTE_PGM_RSRC1.BULKY``. ``COMPUTE_PGM_RSRC1.BULKY``.
25 1 bit cdbg_user Must be 0. 25 1 bit CDBG_USER Must be 0.
Flag that can be used to Flag that can be used to
control debugging code. control debugging code.
CP is responsible for CP is responsible for
filling in filling in
``COMPUTE_PGM_RSRC1.CDBG_USER``. ``COMPUTE_PGM_RSRC1.CDBG_USER``.
31:26 6 bits Reserved. Must be 0. 26 1 bit FP16_OVFL GFX6-8:
Reserved. Must be 0.
GFX9:
Wavefront starts
execution with specified
fp16 overflow mode.
- If 0, then fp16
overflow generates
+/-INF values.
- If 1, then fp16
overflow that is the
result of an +/-INF
input value or divide
by 0 generates a
+/-INF, otherwise
clamps computed
overflow to +/-MAX_FP16
as appropriate.
Used by CP to set up
``COMPUTE_PGM_RSRC1.FP16_OVFL``.
31:27 5 bits Reserved. Must be 0.
32 **Total size 4 bytes** 32 **Total size 4 bytes**
======= =================================================================================================================== ======= ===================================================================================================================
.. ..
.. table:: compute_pgm_rsrc2 for GFX6-GFX9 .. table:: compute_pgm_rsrc2 for GFX6-GFX9
:name: amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table :name: amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table
======= ======= =============================== =========================================================================== ======= ======= =============================== ===========================================================================
Bits Size Field Name Description Bits Size Field Name Description
======= ======= =============================== =========================================================================== ======= ======= =============================== ===========================================================================
0 1 bit enable_sgpr_private_segment Enable the setup of the 0 1 bit ENABLE_SGPR_PRIVATE_SEGMENT Enable the setup of the
_wave_offset SGPR wave scratch offset _WAVE_OFFSET SGPR wave scratch offset
system register (see system register (see
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`). :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.SCRATCH_EN``. ``COMPUTE_PGM_RSRC2.SCRATCH_EN``.
5:1 5 bits user_sgpr_count The total number of SGPR 5:1 5 bits USER_SGPR_COUNT The total number of SGPR
user data registers user data registers
requested. This number must requested. This number must
match the number of user match the number of user
data registers enabled. data registers enabled.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.USER_SGPR``. ``COMPUTE_PGM_RSRC2.USER_SGPR``.
6 1 bit enable_trap_handler Set to 1 if code contains a 6 1 bit ENABLE_TRAP_HANDLER Set to 1 if code contains a
TRAP instruction which TRAP instruction which
requires a trap handler to requires a trap handler to
be enabled. be enabled.
CP sets CP sets
``COMPUTE_PGM_RSRC2.TRAP_PRESENT`` ``COMPUTE_PGM_RSRC2.TRAP_PRESENT``
if the runtime has if the runtime has
installed a trap handler installed a trap handler
regardless of the setting regardless of the setting
of this field. of this field.
7 1 bit enable_sgpr_workgroup_id_x Enable the setup of the 7 1 bit ENABLE_SGPR_WORKGROUP_ID_X Enable the setup of the
system SGPR register for system SGPR register for
the work-group id in the X the work-group id in the X
dimension (see dimension (see
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`). :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.TGID_X_EN``. ``COMPUTE_PGM_RSRC2.TGID_X_EN``.
8 1 bit enable_sgpr_workgroup_id_y Enable the setup of the 8 1 bit ENABLE_SGPR_WORKGROUP_ID_Y Enable the setup of the
system SGPR register for system SGPR register for
the work-group id in the Y the work-group id in the Y
dimension (see dimension (see
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`). :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.TGID_Y_EN``. ``COMPUTE_PGM_RSRC2.TGID_Y_EN``.
9 1 bit enable_sgpr_workgroup_id_z Enable the setup of the 9 1 bit ENABLE_SGPR_WORKGROUP_ID_Z Enable the setup of the
system SGPR register for system SGPR register for
the work-group id in the Z the work-group id in the Z
dimension (see dimension (see
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`). :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.TGID_Z_EN``. ``COMPUTE_PGM_RSRC2.TGID_Z_EN``.
10 1 bit enable_sgpr_workgroup_info Enable the setup of the 10 1 bit ENABLE_SGPR_WORKGROUP_INFO Enable the setup of the
system SGPR register for system SGPR register for
work-group information (see work-group information (see
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`). :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.TGID_SIZE_EN``. ``COMPUTE_PGM_RSRC2.TGID_SIZE_EN``.
12:11 2 bits enable_vgpr_workitem_id Enable the setup of the 12:11 2 bits ENABLE_VGPR_WORKITEM_ID Enable the setup of the
VGPR system registers used VGPR system registers used
for the work-item ID. for the work-item ID.
:ref:`amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table` :ref:`amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table`
defines the values. defines the values.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.TIDIG_CMP_CNT``. ``COMPUTE_PGM_RSRC2.TIDIG_CMP_CNT``.
13 1 bit enable_exception_address_watch Must be 0. 13 1 bit ENABLE_EXCEPTION_ADDRESS_WATCH Must be 0.
Wavefront starts execution Wavefront starts execution
with address watch with address watch
exceptions enabled which exceptions enabled which
are generated when L1 has are generated when L1 has
witnessed a thread access witnessed a thread access
an *address of an *address of
interest*. interest*.
CP is responsible for CP is responsible for
filling in the address filling in the address
watch bit in watch bit in
``COMPUTE_PGM_RSRC2.EXCP_EN_MSB`` ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB``
according to what the according to what the
runtime requests. runtime requests.
14 1 bit enable_exception_memory Must be 0. 14 1 bit ENABLE_EXCEPTION_MEMORY Must be 0.
Wavefront starts execution Wavefront starts execution
with memory violation with memory violation
exceptions exceptions exceptions exceptions
enabled which are generated enabled which are generated
when a memory violation has when a memory violation has
occurred for this wave from occurred for this wave from
L1 or LDS L1 or LDS
(write-to-read-only-memory, (write-to-read-only-memory,
mis-aligned atomic, LDS mis-aligned atomic, LDS
address out of range, address out of range,
illegal address, etc.). illegal address, etc.).
CP sets the memory CP sets the memory
violation bit in violation bit in
``COMPUTE_PGM_RSRC2.EXCP_EN_MSB`` ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB``
according to what the according to what the
runtime requests. runtime requests.
23:15 9 bits granulated_lds_size Must be 0. 23:15 9 bits GRANULATED_LDS_SIZE Must be 0.
CP uses the rounded value CP uses the rounded value
from the dispatch packet, from the dispatch packet,
not this value, as the not this value, as the
dispatch may contain dispatch may contain
dynamically allocated group dynamically allocated group
segment memory. CP writes segment memory. CP writes
directly to directly to
``COMPUTE_PGM_RSRC2.LDS_SIZE``. ``COMPUTE_PGM_RSRC2.LDS_SIZE``.
Amount of group segment Amount of group segment
(LDS) to allocate for each (LDS) to allocate for each
work-group. Granularity is work-group. Granularity is
device specific: device specific:
GFX6: GFX6:
roundup(lds-size / (64 * 4)) roundup(lds-size / (64 * 4))
GFX7-GFX9: GFX7-GFX9:
roundup(lds-size / (128 * 4)) roundup(lds-size / (128 * 4))
24 1 bit enable_exception_ieee_754_fp Wavefront starts execution 24 1 bit ENABLE_EXCEPTION_IEEE_754_FP Wavefront starts execution
_invalid_operation with specified exceptions _INVALID_OPERATION with specified exceptions
enabled. enabled.
Used by CP to set up Used by CP to set up
``COMPUTE_PGM_RSRC2.EXCP_EN`` ``COMPUTE_PGM_RSRC2.EXCP_EN``
(set from bits 0..6). (set from bits 0..6).
IEEE 754 FP Invalid IEEE 754 FP Invalid
Operation Operation
25 1 bit enable_exception_fp_denormal FP Denormal one or more 25 1 bit ENABLE_EXCEPTION_FP_DENORMAL FP Denormal one or more
_source input operands is a _SOURCE input operands is a
denormal number denormal number
26 1 bit enable_exception_ieee_754_fp IEEE 754 FP Division by 26 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP Division by
_division_by_zero Zero _DIVISION_BY_ZERO Zero
27 1 bit enable_exception_ieee_754_fp IEEE 754 FP FP Overflow 27 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP FP Overflow
_overflow _OVERFLOW
28 1 bit enable_exception_ieee_754_fp IEEE 754 FP Underflow 28 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP Underflow
_underflow _UNDERFLOW
29 1 bit enable_exception_ieee_754_fp IEEE 754 FP Inexact 29 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP Inexact
_inexact _INEXACT
30 1 bit enable_exception_int_divide_by Integer Division by Zero 30 1 bit ENABLE_EXCEPTION_INT_DIVIDE_BY Integer Division by Zero
_zero (rcp_iflag_f32 instruction _ZERO (rcp_iflag_f32 instruction
only) only)
31 1 bit Reserved. Must be 0. 31 1 bit Reserved. Must be 0.
32 **Total size 4 bytes.** 32 **Total size 4 bytes.**
======= =================================================================================================================== ======= ===================================================================================================================
.. ..
.. table:: Floating Point Rounding Mode Enumeration Values .. table:: Floating Point Rounding Mode Enumeration Values
:name: amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table :name: amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table
===================================== ===== =============================== ====================================== ===== ==============================
Enumeration Name Value Description Enumeration Name Value Description
===================================== ===== =============================== ====================================== ===== ==============================
AMD_FLOAT_ROUND_MODE_NEAR_EVEN 0 Round Ties To Even AMDGPU_FLOAT_ROUND_MODE_NEAR_EVEN 0 Round Ties To Even
AMD_FLOAT_ROUND_MODE_PLUS_INFINITY 1 Round Toward +infinity AMDGPU_FLOAT_ROUND_MODE_PLUS_INFINITY 1 Round Toward +infinity
AMD_FLOAT_ROUND_MODE_MINUS_INFINITY 2 Round Toward -infinity AMDGPU_FLOAT_ROUND_MODE_MINUS_INFINITY 2 Round Toward -infinity
AMD_FLOAT_ROUND_MODE_ZERO 3 Round Toward 0 AMDGPU_FLOAT_ROUND_MODE_ZERO 3 Round Toward 0
===================================== ===== =============================== ====================================== ===== ==============================
.. ..
.. table:: Floating Point Denorm Mode Enumeration Values .. table:: Floating Point Denorm Mode Enumeration Values
:name: amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table :name: amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table
===================================== ===== =============================== ====================================== ===== ==============================
Enumeration Name Value Description Enumeration Name Value Description
===================================== ===== =============================== ====================================== ===== ==============================
AMD_FLOAT_DENORM_MODE_FLUSH_SRC_DST 0 Flush Source and Destination AMDGPU_FLOAT_DENORM_MODE_FLUSH_SRC_DST 0 Flush Source and Destination
Denorms Denorms
AMD_FLOAT_DENORM_MODE_FLUSH_DST 1 Flush Output Denorms AMDGPU_FLOAT_DENORM_MODE_FLUSH_DST 1 Flush Output Denorms
AMD_FLOAT_DENORM_MODE_FLUSH_SRC 2 Flush Source Denorms AMDGPU_FLOAT_DENORM_MODE_FLUSH_SRC 2 Flush Source Denorms
AMD_FLOAT_DENORM_MODE_FLUSH_NONE 3 No Flush AMDGPU_FLOAT_DENORM_MODE_FLUSH_NONE 3 No Flush
===================================== ===== =============================== ====================================== ===== ==============================
.. ..
.. table:: System VGPR Work-Item ID Enumeration Values .. table:: System VGPR Work-Item ID Enumeration Values
:name: amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table :name: amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table
===================================== ===== =============================== ======================================== ===== ============================
Enumeration Name Value Description Enumeration Name Value Description
===================================== ===== =============================== ======================================== ===== ============================
AMD_SYSTEM_VGPR_WORKITEM_ID_X 0 Set work-item X dimension ID. AMDGPU_SYSTEM_VGPR_WORKITEM_ID_X 0 Set work-item X dimension
AMD_SYSTEM_VGPR_WORKITEM_ID_X_Y 1 Set work-item X and Y ID.
AMDGPU_SYSTEM_VGPR_WORKITEM_ID_X_Y 1 Set work-item X and Y
dimensions ID. dimensions ID.
AMD_SYSTEM_VGPR_WORKITEM_ID_X_Y_Z 2 Set work-item X, Y and Z AMDGPU_SYSTEM_VGPR_WORKITEM_ID_X_Y_Z 2 Set work-item X, Y and Z
dimensions ID. dimensions ID.
AMD_SYSTEM_VGPR_WORKITEM_ID_UNDEFINED 3 Undefined. AMDGPU_SYSTEM_VGPR_WORKITEM_ID_UNDEFINED 3 Undefined.
===================================== ===== =============================== ======================================== ===== ============================
.. _amdgpu-amdhsa-initial-kernel-execution-state: .. _amdgpu-amdhsa-initial-kernel-execution-state:
Initial Kernel Execution State Initial Kernel Execution State
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
This section defines the register state that will be set up by the packet This section defines the register state that will be set up by the packet
processor prior to the start of execution of every wavefront. This is limited by processor prior to the start of execution of every wavefront. This is limited by
▲ Show 20 Lines • Show All 1,954 Lines • Show Last 20 Lines

include/llvm/Support/AMDGPUKernelDescriptor.h

  • This file was added.
//===--- AMDGPUKernelDescriptor.h -------------------------------*- C++ -*-===//
//
// The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
//
/// \file
/// \brief AMDGPU kernel descriptor definitions. For more information, visit
/// https://llvm.org/docs/AMDGPUUsage.html#kernel-descriptor-for-gfx6-gfx9
//
//===----------------------------------------------------------------------===//
#ifndef LLVM_SUPPORT_AMDGPUKERNELDESCRIPTOR_H
#define LLVM_SUPPORT_AMDGPUKERNELDESCRIPTOR_H
#include <cstdint>
// Creates enumeration entries used for packing bits into integers. Enumeration
// entries include bit shift amount, bit width, and bit mask.
#define AMDGPU_BITS_ENUM_ENTRY(name, shift, width) \
name ## _SHIFT = (shift), \
name ## _WIDTH = (width), \
name = (((1 << (width)) - 1) << (shift)) \
// Gets bits for specified bit mask from specified source.
#define AMDGPU_BITS_GET(src, mask) \
((src & mask) >> mask ## _SHIFT) \
// Sets bits for specified bit mask in specified destination.
#define AMDGPU_BITS_SET(dst, mask, val) \
dst &= (~(1 << mask ## _SHIFT) & ~mask); \
dst |= (((val) << mask ## _SHIFT) & mask) \
namespace llvm {
namespace AMDGPU {
namespace HSAKD {
/// \brief Floating point rounding modes.
enum : uint8_t {
AMDGPU_FLOAT_ROUND_MODE_NEAR_EVEN = 0,
AMDGPU_FLOAT_ROUND_MODE_PLUS_INFINITY = 1,
AMDGPU_FLOAT_ROUND_MODE_MINUS_INFINITY = 2,
AMDGPU_FLOAT_ROUND_MODE_ZERO = 3,
};
/// \brief Floating point denorm modes.
enum : uint8_t {
AMDGPU_FLOAT_DENORM_MODE_FLUSH_SRC_DST = 0,
AMDGPU_FLOAT_DENORM_MODE_FLUSH_DST = 1,
AMDGPU_FLOAT_DENORM_MODE_FLUSH_SRC = 2,
AMDGPU_FLOAT_DENORM_MODE_FLUSH_NONE = 3,
};
/// \brief System VGPR workitem IDs.
enum : uint8_t {
AMDGPU_SYSTEM_VGPR_WORKITEM_ID_X = 0,
AMDGPU_SYSTEM_VGPR_WORKITEM_ID_X_Y = 1,
AMDGPU_SYSTEM_VGPR_WORKITEM_ID_X_Y_Z = 2,
AMDGPU_SYSTEM_VGPR_WORKITEM_ID_UNDEFINED = 3,
};
/// \brief Compute program resource register one layout.
enum ComputePgmRsrc1 {
AMDGPU_BITS_ENUM_ENTRY(GRANULATED_WORKITEM_VGPR_COUNT, 0, 6),
AMDGPU_BITS_ENUM_ENTRY(GRANULATED_WAVEFRONT_SGPR_COUNT, 6, 4),
AMDGPU_BITS_ENUM_ENTRY(PRIORITY, 10, 2),
AMDGPU_BITS_ENUM_ENTRY(FLOAT_ROUND_MODE_32, 12, 2),
AMDGPU_BITS_ENUM_ENTRY(FLOAT_ROUND_MODE_16_64, 14, 2),
AMDGPU_BITS_ENUM_ENTRY(FLOAT_DENORM_MODE_32, 16, 2),
AMDGPU_BITS_ENUM_ENTRY(FLOAT_DENORM_MODE_16_64, 18, 2),
AMDGPU_BITS_ENUM_ENTRY(PRIV, 20, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_DX10_CLAMP, 21, 1),
AMDGPU_BITS_ENUM_ENTRY(DEBUG_MODE, 22, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_IEEE_MODE, 23, 1),
AMDGPU_BITS_ENUM_ENTRY(BULKY, 24, 1),
AMDGPU_BITS_ENUM_ENTRY(CDBG_USER, 25, 1),
AMDGPU_BITS_ENUM_ENTRY(FP16_OVFL, 26, 1),
AMDGPU_BITS_ENUM_ENTRY(RESERVED0, 27, 5),
};
/// \brief Compute program resource register two layout.
enum ComputePgmRsrc2 {
AMDGPU_BITS_ENUM_ENTRY(ENABLE_SGPR_PRIVATE_SEGMENT_WAVE_OFFSET, 0, 1),
AMDGPU_BITS_ENUM_ENTRY(USER_SGPR_COUNT, 1, 5),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_TRAP_HANDLER, 6, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_SGPR_WORKGROUP_ID_X, 7, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_SGPR_WORKGROUP_ID_Y, 8, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_SGPR_WORKGROUP_ID_Z, 9, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_SGPR_WORKGROUP_INFO, 10, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_VGPR_WORKITEM_ID, 11, 2),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_ADDRESS_WATCH, 13, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_MEMORY, 14, 1),
AMDGPU_BITS_ENUM_ENTRY(GRANULATED_LDS_SIZE, 15, 9),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_IEEE_754_FP_INVALID_OPERATION, 24, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_FP_DENORMAL_SOURCE, 25, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_IEEE_754_FP_DIVISION_BY_ZERO, 26, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_IEEE_754_FP_OVERFLOW, 27, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_IEEE_754_FP_UNDERFLOW, 28, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_IEEE_754_FP_INEXACT, 29, 1),
AMDGPU_BITS_ENUM_ENTRY(ENABLE_EXCEPTION_INT_DIVIDE_BY_ZERO, 30, 1),
AMDGPU_BITS_ENUM_ENTRY(RESERVED1, 31, 1),
};
/// \brief Kernel descriptor layout. This layout should be kept backwards
/// compatible as it is consumed by the command processor.
struct KernelDescriptor final {
uint32_t GroupSegmentFixedSize;
uint32_t PrivateSegmentFixedSize;
uint32_t MaxFlatWorkgroupSize;
uint64_t IsDynamicCallStack : 1;
uint64_t IsXNACKEnabled : 1;
uint64_t Reserved0 : 30;
int64_t KernelCodeEntryByteOffset;
uint64_t Reserved1[3];
uint32_t ComputePgmRsrc1;
uint32_t ComputePgmRsrc2;
uint64_t EnableSGPRPrivateSegmentBuffer : 1;
uint64_t EnableSGPRDispatchPtr : 1;
uint64_t EnableSGPRQueuePtr : 1;
uint64_t EnableSGPRKernargSegmentPtr : 1;
uint64_t EnableSGPRDispatchID : 1;
uint64_t EnableSGPRFlatScratchInit : 1;
uint64_t EnableSGPRPrivateSegmentSize : 1;
uint64_t EnableSGPRGridWorkgroupCountX : 1;
uint64_t EnableSGPRGridWorkgroupCountY : 1;
uint64_t EnableSGPRGridWorkgroupCountZ : 1;
uint64_t Reserved2 : 54;
KernelDescriptor() = default;
};
} // end namespace HSAKD
} // end namespace AMDGPU
} // end namespace llvm
#endif // LLVM_SUPPORT_AMDGPUKERNELDESCRIPTOR_H