This is an archive of the discontinued LLVM Phabricator instance.

Differential D53445

[AMDGPU] Update code object metadata format documentation
ClosedPublic

Authored by scott.linder on Oct 19 2018, 11:56 AM.

Details

Reviewers
t-tye
kzhuravl
b-sumner
Commits
rG8d5a36a8394c: [AMDGPU] Update code object metadata format documentation
rL346992: [AMDGPU] Update code object metadata format documentation
Summary
  • Add amdhsa prefix to names to allow other tools to use the metadata without collision.
  • Make names consistent.
  • Simplify structure.
  • Change note record ID.
  • Switch from YAML to MsgPack format.
  • Document metadata assembler directive.

Patch by Tony Tye (https://reviews.llvm.org/D47549 rebased with added description of assembler directive)

Diff Detail

Repository
rL LLVM

Event Timeline

scott.linder created this revision.Oct 19 2018, 11:56 AM
Herald added subscribers: llvm-commits, tpr, dstuttard and 4 others. · View Herald TranscriptOct 19 2018, 11:56 AM
scott.linder updated this revision to Diff 170255.Oct 19 2018, 2:26 PM

Update note type value.

kzhuravl added a comment.EditedOct 22 2018, 12:44 PM

In v2 asm, we had the following:

.option.machine_version_major
.option.machine_version_minor
.option.machine_version_stepping

What are the equivalents for v3 asm? The updates you made do not seem to describe those... Unless I missed it.

kzhuravl added inline comments.Oct 22 2018, 12:47 PM
docs/AMDGPUUsage.rst
688 ↗(On Diff #170255)

6->32

scott.linder marked an inline comment as done.Oct 23 2018, 7:07 AM
In D53445#1271220, @kzhuravl wrote:

In v2 asm, we had the following:

.option.machine_version_major
.option.machine_version_minor
.option.machine_version_stepping

What are the equivalents for v3 asm? The updates you made do not seem to describe those... Unless I missed it.

This was replaced with .amdgcn.gfx_generation_number which is documented under "Predefined Symbols". I believe we did not want to expose minor/stepping in the assembler, and we wanted to give a more meaningful name to the major version number.

scott.linder updated this revision to Diff 170635.Oct 23 2018, 7:08 AM

Update note type value

scott.linder added a comment.Oct 29 2018, 7:03 AM

Ping

scott.linder added a comment.Nov 5 2018, 12:55 PM

Ping

t-tye added a comment.Nov 5 2018, 1:05 PM

Should the v2 version of the metadata be moved to a separate section rather than be deleted since it is still generated when the v2 code object format is requested?

scott.linder added a child revision: D48179: [AMDGPU] Emit MessagePack HSA Metadata for v3 code object.Nov 14 2018, 9:32 AM
scott.linder updated this revision to Diff 174262.Nov 15 2018, 11:50 AM

Preserve docs for code object v2 metadata

t-tye accepted this revision.Nov 15 2018, 12:27 PM

LGTM

This revision is now accepted and ready to land.Nov 15 2018, 12:27 PM
Closed by commit rL346992: [AMDGPU] Update code object metadata format documentation (authored by scott.linder). · Explain WhyNov 15 2018, 12:49 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
llvm/
trunk/
docs/
553 lines

Diff 174270

llvm/trunk/docs/AMDGPUUsage.rst

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 673 Lines • ▼ Show 20 Lines
------------ ------------
As required by ``ELFCLASS32`` and ``ELFCLASS64``, minimal zero byte padding must As required by ``ELFCLASS32`` and ``ELFCLASS64``, minimal zero byte padding must
be generated after the ``name`` field to ensure the ``desc`` field is 4 byte be generated after the ``name`` field to ensure the ``desc`` field is 4 byte
aligned. In addition, minimal zero byte padding must be generated to ensure the aligned. In addition, minimal zero byte padding must be generated to ensure the
``desc`` field size is a multiple of 4 bytes. The ``sh_addralign`` field of the ``desc`` field size is a multiple of 4 bytes. The ``sh_addralign`` field of the
``.note`` section must be at least 4 to indicate at least 8 byte alignment. ``.note`` section must be at least 4 to indicate at least 8 byte alignment.
The AMDGPU backend code object uses the following ELF note records in the .. _amdgpu-note-records-v2:
``.note`` section. The *Description* column specifies the layout of the note
record's ``desc`` field. All fields are consecutive bytes. Note records with Code Object V2 Note Records (-mattr=-code-object-v3)
variable size strings have a corresponding ``*_size`` field that specifies the ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
number of bytes, including the terminating null character, in the string. The
string(s) come immediately after the preceding fields. The AMDGPU backend code object uses the following ELF note record in the
``.note`` section.
Additional note records can be present. Additional note records can be present.
.. table:: AMDGPU ELF Note Records .. table:: AMDGPU Code Object V2 ELF Note Records
:name: amdgpu-elf-note-records-table :name: amdgpu-elf-note-records-table-v2
===== ============================== ====================================== ===== ============================== ======================================
Name Type Description Name Type Description
===== ============================== ====================================== ===== ============================== ======================================
"AMD" ``NT_AMD_AMDGPU_HSA_METADATA`` <metadata null terminated string> "AMD" ``NT_AMD_AMDGPU_HSA_METADATA`` <metadata null terminated string>
===== ============================== ====================================== ===== ============================== ======================================
.. ..
.. table:: AMDGPU ELF Note Record Enumeration Values .. table:: AMDGPU Code Object V2 ELF Note Record Enumeration Values
:name: amdgpu-elf-note-record-enumeration-values-table :name: amdgpu-elf-note-record-enumeration-values-table-v2
============================== ===== ============================== =====
Name Value Name Value
============================== ===== ============================== =====
*reserved* 0-9 *reserved* 0-9
``NT_AMD_AMDGPU_HSA_METADATA`` 10 ``NT_AMD_AMDGPU_HSA_METADATA`` 10
*reserved* 11 *reserved* 11
============================== ===== ============================== =====
``NT_AMD_AMDGPU_HSA_METADATA`` ``NT_AMD_AMDGPU_HSA_METADATA``
Specifies extensible metadata associated with the code objects executed on HSA Specifies extensible metadata associated with the code objects executed on HSA
[HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. It is required when [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. It is required when
the target triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`). See the target triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`). See
:ref:`amdgpu-amdhsa-code-object-metadata` for the syntax of the code :ref:`amdgpu-amdhsa-code-object-metadata-v2` for the syntax of the code
object metadata string. object metadata string.
.. _amdgpu-note-records-v3:
Code Object V3 Note Records (-mattr=+code-object-v3)
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
The AMDGPU backend code object uses the following ELF note record in the
``.note`` section.
Additional note records can be present.
.. table:: AMDGPU Code Object V3 ELF Note Records
:name: amdgpu-elf-note-records-table-v3
======== ============================== ======================================
Name Type Description
======== ============================== ======================================
"AMDGPU" ``NT_AMDGPU_METADATA`` Metadata in Message Pack [MsgPack]_
binary format.
======== ============================== ======================================
..
.. table:: AMDGPU Code Object V3 ELF Note Record Enumeration Values
:name: amdgpu-elf-note-record-enumeration-values-table-v3
============================== =====
Name Value
============================== =====
*reserved* 0-31
``NT_AMDGPU_METADATA`` 32
============================== =====
``NT_AMDGPU_METADATA``
Specifies extensible metadata associated with an AMDGPU code
object. It is encoded as a map in the Message Pack [MsgPack]_ binary
data format. See :ref:`amdgpu-amdhsa-code-object-metadata-v3` for the
map keys defined for the ``amdhsa`` OS.
.. _amdgpu-symbols: .. _amdgpu-symbols:
Symbols Symbols
------- -------
Symbols include the following: Symbols include the following:
.. table:: AMDGPU ELF Symbols .. table:: AMDGPU ELF Symbols
▲ Show 20 Lines • Show All 276 Lines • ▼ Show 20 Lines
.. _amdgpu-amdhsa-code-object-metadata: .. _amdgpu-amdhsa-code-object-metadata:
Code Object Metadata Code Object Metadata
~~~~~~~~~~~~~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~
The code object metadata specifies extensible metadata associated with the code The code object metadata specifies extensible metadata associated with the code
objects executed on HSA [HSA]_ compatible runtimes such as AMD's ROCm objects executed on HSA [HSA]_ compatible runtimes such as AMD's ROCm
[AMD-ROCm]_. It is specified by the ``NT_AMD_AMDGPU_HSA_METADATA`` note record [AMD-ROCm]_. It is specified in a note record (see :ref:`amdgpu-note-records`)
(see :ref:`amdgpu-note-records`) and is required when the target triple OS is and is required when the target triple OS is ``amdhsa`` (see
``amdhsa`` (see :ref:`amdgpu-target-triples`). It must contain the minimum :ref:`amdgpu-target-triples`). It must contain the minimum information
information necessary to support the ROCM kernel queries. For example, the necessary to support the ROCM kernel queries. For example, the segment sizes
segment sizes needed in a dispatch packet. In addition, a high level language needed in a dispatch packet. In addition, a high level language runtime may
runtime may require other information to be included. For example, the AMD require other information to be included. For example, the AMD OpenCL runtime
OpenCL runtime records kernel argument information. records kernel argument information.
.. _amdgpu-amdhsa-code-object-metadata-v2:
Code Object V2 Metadata (-mattr=-code-object-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 The metadata is specified as a YAML formatted string (see [YAML]_ and
:doc:`YamlIO`). :doc:`YamlIO`).
.. TODO .. TODO
Is the string null terminated? It probably should not if YAML allows it to Is the string null terminated? It probably should not if YAML allows it to
contain null characters, otherwise it should be. contain null characters, otherwise it should be.
The metadata is represented as a single YAML document comprised of the mapping The metadata is represented as a single YAML document comprised of the mapping
defined in table :ref:`amdgpu-amdhsa-code-object-metadata-mapping-table` and defined in table :ref:`amdgpu-amdhsa-code-object-metadata-map-table-v2` and
referenced tables. referenced tables.
For boolean values, the string values of ``false`` and ``true`` are used for For boolean values, the string values of ``false`` and ``true`` are used for
false and true respectively. false and true respectively.
Additional information can be added to the mappings. To avoid conflicts, any Additional information can be added to the mappings. To avoid conflicts, any
non-AMD key names should be prefixed by "*vendor-name*.". non-AMD key names should be prefixed by "*vendor-name*.".
.. table:: AMDHSA Code Object Metadata Mapping .. table:: AMDHSA Code Object V2 Metadata Map
:name: amdgpu-amdhsa-code-object-metadata-mapping-table :name: amdgpu-amdhsa-code-object-metadata-map-table-v2
========== ============== ========= ======================================= ========== ============== ========= =======================================
String Key Value Type Required? Description String Key Value Type Required? Description
========== ============== ========= ======================================= ========== ============== ========= =======================================
"Version" sequence of Required - The first integer is the major "Version" sequence of Required - The first integer is the major
2 integers version. Currently 1. 2 integers version. Currently 1.
- The second integer is the minor - The second integer is the minor
version. Currently 0. version. Currently 0.
Show All 20 Lines "Printf" sequence of Each string is encoded information
of the i-th FormatString argument of of the i-th FormatString argument of
the printf function call the printf function call
FormatString FormatString
The format string passed to the The format string passed to the
printf function call. printf function call.
"Kernels" sequence of Required Sequence of the mappings for each "Kernels" sequence of Required Sequence of the mappings for each
mapping kernel in the code object. See mapping kernel in the code object. See
:ref:`amdgpu-amdhsa-code-object-kernel-metadata-mapping-table` :ref:`amdgpu-amdhsa-code-object-kernel-metadata-map-table-v2`
for the definition of the mapping. for the definition of the mapping.
========== ============== ========= ======================================= ========== ============== ========= =======================================
.. ..
.. table:: AMDHSA Code Object Kernel Metadata Mapping .. table:: AMDHSA Code Object V2 Kernel Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-metadata-mapping-table :name: amdgpu-amdhsa-code-object-kernel-metadata-map-table-v2
================= ============== ========= ================================ ================= ============== ========= ================================
String Key Value Type Required? Description String Key Value Type Required? Description
================= ============== ========= ================================ ================= ============== ========= ================================
"Name" string Required Source name of the kernel. "Name" string Required Source name of the kernel.
"SymbolName" string Required Name of the kernel "SymbolName" string Required Name of the kernel
descriptor ELF symbol. descriptor ELF symbol.
"Language" string Source language of the kernel. "Language" string Source language of the kernel.
Values include: Values include:
- "OpenCL C" - "OpenCL C"
- "OpenCL C++" - "OpenCL C++"
- "HCC" - "HCC"
- "OpenMP" - "OpenMP"
"LanguageVersion" sequence of - The first integer is the major "LanguageVersion" sequence of - The first integer is the major
2 integers version. 2 integers version.
- The second integer is the - The second integer is the
minor version. minor version.
"Attrs" mapping Mapping of kernel attributes. "Attrs" mapping Mapping of kernel attributes.
See See
:ref:`amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table` :ref:`amdgpu-amdhsa-code-object-kernel-attribute-metadata-map-table-v2`
for the mapping definition. for the mapping definition.
"Args" sequence of Sequence of mappings of the "Args" sequence of Sequence of mappings of the
mapping kernel arguments. See mapping kernel arguments. See
:ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table` :ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v2`
for the definition of the mapping. for the definition of the mapping.
"CodeProps" mapping Mapping of properties related to "CodeProps" mapping Mapping of properties related to
the kernel code. See the kernel code. See
:ref:`amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table` :ref:`amdgpu-amdhsa-code-object-kernel-code-properties-metadata-map-table-v2`
for the mapping definition. for the mapping definition.
================= ============== ========= ================================ ================= ============== ========= ================================
.. ..
.. table:: AMDHSA Code Object Kernel Attribute Metadata Mapping .. table:: AMDHSA Code Object V2 Kernel Attribute Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table :name: amdgpu-amdhsa-code-object-kernel-attribute-metadata-map-table-v2
=================== ============== ========= ============================== =================== ============== ========= ==============================
String Key Value Type Required? Description String Key Value Type Required? Description
=================== ============== ========= ============================== =================== ============== ========= ==============================
"ReqdWorkGroupSize" sequence of If not 0, 0, 0 then all values "ReqdWorkGroupSize" sequence of If not 0, 0, 0 then all values
3 integers must be >=1 and the dispatch 3 integers must be >=1 and the dispatch
work-group size X, Y, Z must work-group size X, Y, Z must
correspond to the specified correspond to the specified
Show All 23 Lines "RuntimeHandle" string The external symbol name
to it, which is used for to it, which is used for
device side enqueueing. Only device side enqueueing. Only
available for device side available for device side
enqueued kernels. enqueued kernels.
=================== ============== ========= ============================== =================== ============== ========= ==============================
.. ..
.. table:: AMDHSA Code Object Kernel Argument Metadata Mapping .. table:: AMDHSA Code Object V2 Kernel Argument Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table :name: amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v2
================= ============== ========= ================================ ================= ============== ========= ================================
String Key Value Type Required? Description String Key Value Type Required? Description
================= ============== ========= ================================ ================= ============== ========= ================================
"Name" string Kernel argument name. "Name" string Kernel argument name.
"TypeName" string Kernel argument type name. "TypeName" string Kernel argument type name.
"Size" integer Required Kernel argument size in bytes. "Size" integer Required Kernel argument size in bytes.
"Align" integer Required Kernel argument alignment in "Align" integer Required Kernel argument alignment in
▲ Show 20 Lines • Show All 180 Lines • ▼ Show 20 Lines "IsPipe" boolean Indicates if the kernel argument
.. TODO .. TODO
Can GlobalBuffer be pipe Can GlobalBuffer be pipe
qualified? qualified?
================= ============== ========= ================================ ================= ============== ========= ================================
.. ..
.. table:: AMDHSA Code Object Kernel Code Properties Metadata Mapping .. table:: AMDHSA Code Object V2 Kernel Code Properties Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table :name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-map-table-v2
============================ ============== ========= ===================== ============================ ============== ========= =====================
String Key Value Type Required? Description String Key Value Type Required? Description
============================ ============== ========= ===================== ============================ ============== ========= =====================
"KernargSegmentSize" integer Required The size in bytes of "KernargSegmentSize" integer Required The size in bytes of
the kernarg segment the kernarg segment
that holds the values that holds the values
of the arguments to of the arguments to
▲ Show 20 Lines • Show All 61 Lines • ▼ Show 20 Lines "NumSpilledSGPRs" integer Number of stores from
location. location.
"NumSpilledVGPRs" integer Number of stores from "NumSpilledVGPRs" integer Number of stores from
a vector register to a vector register to
a register allocator a register allocator
created spill created spill
location. location.
============================ ============== ========= ===================== ============================ ============== ========= =====================
.. _amdgpu-amdhsa-code-object-metadata-v3:
Code Object V3 Metadata (-mattr=+code-object-v3)
++++++++++++++++++++++++++++++++++++++++++++++++
Code object V3 metadata is specified by the ``NT_AMDGPU_METADATA`` note record
(see :ref:`amdgpu-note-records-v3`).
The metadata is represented as Message Pack formatted binary data (see
[MsgPack]_). The top level is a Message Pack map that includes the
keys defined in table
:ref:`amdgpu-amdhsa-code-object-metadata-map-table-v3` and referenced
tables.
Additional information can be added to the maps. To avoid conflicts,
any key names should be prefixed by "*vendor-name*." where
``vendor-name`` can be the the name of the vendor and specific vendor
tool that generates the information. The prefix is abbreviated to
simply "." when it appears within a map that has been added by the
same *vendor-name*.
.. table:: AMDHSA Code Object V3 Metadata Map
:name: amdgpu-amdhsa-code-object-metadata-map-table-v3
================= ============== ========= =======================================
String Key Value Type Required? Description
================= ============== ========= =======================================
"amdhsa.version" sequence of Required - The first integer is the major
2 integers version. Currently 1.
- The second integer is the minor
version. Currently 0.
"amdhsa.printf" sequence of Each string is encoded information
strings about a printf function call. The
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
each printf function call
``N``
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
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
:ref:`amdgpu-amdhsa-code-object-kernel-metadata-map-table-v3`
for the definition of the keys included
in that map.
================= ============== ========= =======================================
..
.. table:: AMDHSA Code Object V3 Kernel Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-metadata-map-table-v3
=================================== ============== ========= ================================
String Key Value Type Required? Description
=================================== ============== ========= ================================
".name" string Required Source name of the kernel.
".symbol" string Required Name of the kernel
descriptor ELF symbol.
".language" string Source language of the kernel.
Values include:
- "OpenCL C"
- "OpenCL C++"
- "HCC"
- "HIP"
- "OpenMP"
- "Assembler"
".language_version" sequence of - The first integer is the major
2 integers version.
- The second integer is the
minor version.
".args" sequence of Sequence of maps of the
map kernel arguments. See
:ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3`
for the definition of the keys
included in that map.
".reqd_workgroup_size" sequence of If not 0, 0, 0 then all values
3 integers must be >=1 and the dispatch
work-group size X, Y, Z must
correspond to the specified
values. Defaults to 0, 0, 0.
Corresponds to the OpenCL
``reqd_work_group_size``
attribute.
".workgroup_size_hint" sequence of The dispatch work-group size
3 integers X, Y, Z is likely to be the
specified values.
Corresponds to the OpenCL
``work_group_size_hint``
attribute.
".vec_type_hint" string The name of a scalar or vector
type.
Corresponds to the OpenCL
``vec_type_hint`` attribute.
".device_enqueue_symbol" string The external symbol name
associated with a kernel.
OpenCL runtime allocates a
global buffer for the symbol
and saves the kernel's address
to it, which is used for
device side enqueueing. Only
available for device side
enqueued kernels.
".kernarg_segment_size" integer Required The size in bytes of
the kernarg segment
that holds the values
of the arguments to
the kernel.
".group_segment_fixed_size" integer Required The amount of group
segment memory
required by a
work-group in
bytes. This does not
include any
dynamically allocated
group segment memory
that may be added
when the kernel is
dispatched.
".private_segment_fixed_size" integer Required The amount of fixed
private address space
memory required for a
work-item in
bytes. If the kernel
uses a dynamic call
stack then additional
space must be added
to this value for the
call stack.
".kernarg_segment_align" integer Required The maximum byte
alignment of
arguments in the
kernarg segment. Must
be a power of 2.
".wavefront_size" integer Required Wavefront size. Must
be a power of 2.
".sgpr_count" integer Required Number of scalar
registers required by a
wavefront for
GFX6-GFX9. A register
is required if it is
used explicitly, or
if a higher numbered
register is used
explicitly. This
includes the special
SGPRs for VCC, Flat
Scratch (GFX7-GFX9)
and XNACK (for
GFX8-GFX9). It does
not include the 16
SGPR added if a trap
handler is
enabled. It is not
rounded up to the
allocation
granularity.
".vgpr_count" integer Required Number of vector
registers required by
each work-item for
GFX6-GFX9. A register
is required if it is
used explicitly, or
if a higher numbered
register is used
explicitly.
".max_flat_workgroup_size" integer Required Maximum flat
work-group size
supported by the
kernel in work-items.
Must be >=1 and
consistent with
ReqdWorkGroupSize if
not 0, 0, 0.
".sgpr_spill_count" integer Number of stores from
a scalar register to
a register allocator
created spill
location.
".vgpr_spill_count" integer Number of stores from
a vector register to
a register allocator
created spill
location.
=================================== ============== ========= ================================
..
.. table:: AMDHSA Code Object V3 Kernel Argument Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3
====================== ============== ========= ================================
String Key Value Type Required? Description
====================== ============== ========= ================================
".name" string Kernel argument name.
".type_name" string Kernel argument type name.
".size" integer Required Kernel argument size in bytes.
".offset" integer Required Kernel argument offset in
bytes. The offset must be a
multiple of the alignment
required by the argument.
".value_kind" string Required Kernel argument kind that
specifies how to set up the
corresponding argument.
Values include:
"by_value"
The argument is copied
directly into the kernarg.
"global_buffer"
A global address space pointer
to the buffer data is passed
in the kernarg.
"dynamic_shared_pointer"
A group address space pointer
to dynamically allocated LDS
is passed in the kernarg.
"sampler"
A global address space
pointer to a S# is passed in
the kernarg.
"image"
A global address space
pointer to a T# is passed in
the kernarg.
"pipe"
A global address space pointer
to an OpenCL pipe is passed in
the kernarg.
"queue"
A global address space pointer
to an OpenCL device enqueue
queue is passed in the
kernarg.
"hidden_global_offset_x"
The OpenCL grid dispatch
global offset for the X
dimension is passed in the
kernarg.
"hidden_global_offset_y"
The OpenCL grid dispatch
global offset for the Y
dimension is passed in the
kernarg.
"hidden_global_offset_z"
The OpenCL grid dispatch
global offset for the Z
dimension is passed in the
kernarg.
"hidden_none"
An argument that is not used
by the kernel. Space needs to
be left for it, but it does
not need to be set up.
"hidden_printf_buffer"
A global address space pointer
to the runtime printf buffer
is passed in kernarg.
"hidden_default_queue"
A global address space pointer
to the OpenCL device enqueue
queue that should be used by
the kernel by default is
passed in the kernarg.
"hidden_completion_action"
A global address space pointer
to help link enqueued kernels into
the ancestor tree for determining
when the parent kernel has finished.
".value_type" string Required Kernel argument value type. Only
present if ".value_kind" is
"by_value". For vector data
types, the value is for the
element type. Values include:
- "struct"
- "i8"
- "u8"
- "i16"
- "u16"
- "f16"
- "i32"
- "u32"
- "f32"
- "i64"
- "u64"
- "f64"
.. 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
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
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
by ".access" to reflect what the
kernel actual does. If not
present then the runtime must
assume what is implied by
".access" and ".is_const" . Values
are:
- "read_only"
- "write_only"
- "read_write"
".is_const" boolean Indicates if the kernel argument
is const qualified. Only present
if ".value_kind" is
"global_buffer".
".is_restrict" boolean Indicates if the kernel argument
is restrict qualified. Only
present if ".value_kind" is
"global_buffer".
".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
Can "global_buffer" be pipe
qualified?
====================== ============== ========= ================================
.. ..
Kernel Dispatch Kernel Dispatch
~~~~~~~~~~~~~~~ ~~~~~~~~~~~~~~~
The HSA architected queuing language (AQL) defines a user space memory interface 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 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 way. An agent can have zero or more AQL queues created for it using the ROCm
▲ Show 20 Lines • Show All 2,924 Lines • ▼ Show 20 Lines
- *machine_kind* defaults to 1. - *machine_kind* defaults to 1.
- *machine_version_major*, *machine_version_minor*, and - *machine_version_major*, *machine_version_minor*, and
*machine_version_stepping* are derived from the value of the -mcpu option *machine_version_stepping* are derived from the value of the -mcpu option
that is passed to the assembler. that is passed to the assembler.
- *kernel_code_entry_byte_offset* defaults to 256. - *kernel_code_entry_byte_offset* defaults to 256.
- *wavefront_size* defaults to 6. - *wavefront_size* defaults to 6.
- *kernarg_segment_alignment*, *group_segment_alignment*, and - *kernarg_segment_alignment*, *group_segment_alignment*, and
*private_segment_alignment* default to 4. Note that alignments are specified *private_segment_alignment* default to 4. Note that alignments are specified
as a power of two, so a value of **n** means an alignment of 2^ **n**. as a power of 2, so a value of **n** means an alignment of 2^ **n**.
The *.amd_kernel_code_t* directive must be placed immediately after the The *.amd_kernel_code_t* directive must be placed immediately after the
function label and before any instructions. function label and before any instructions.
For a full list of amd_kernel_code_t keys, refer to AMDGPU ABI document, For a full list of amd_kernel_code_t keys, refer to AMDGPU ABI document,
comments in lib/Target/AMDGPU/AmdKernelCodeT.h and test/CodeGen/AMDGPU/hsa.s. comments in lib/Target/AMDGPU/AmdKernelCodeT.h and test/CodeGen/AMDGPU/hsa.s.
Here is an example of a minimal amd_kernel_code_t specification: Here is an example of a minimal amd_kernel_code_t specification:
▲ Show 20 Lines • Show All 196 Lines • ▼ Show 20 Lines .. table:: AMDHSA Kernel Assembler Directives
``.amdhsa_exception_fp_ieee_underflow`` 0 GFX6-GFX9 Controls ENABLE_EXCEPTION_IEEE_754_FP_UNDERFLOW in ``.amdhsa_exception_fp_ieee_underflow`` 0 GFX6-GFX9 Controls ENABLE_EXCEPTION_IEEE_754_FP_UNDERFLOW in
:ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`. :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`.
``.amdhsa_exception_fp_ieee_inexact`` 0 GFX6-GFX9 Controls ENABLE_EXCEPTION_IEEE_754_FP_INEXACT in ``.amdhsa_exception_fp_ieee_inexact`` 0 GFX6-GFX9 Controls ENABLE_EXCEPTION_IEEE_754_FP_INEXACT in
:ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`. :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`.
``.amdhsa_exception_int_div_zero`` 0 GFX6-GFX9 Controls ENABLE_EXCEPTION_INT_DIVIDE_BY_ZERO in ``.amdhsa_exception_int_div_zero`` 0 GFX6-GFX9 Controls ENABLE_EXCEPTION_INT_DIVIDE_BY_ZERO in
:ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`. :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`.
======================================================== ================ ============ =================== ======================================================== ================ ============ ===================
.amdgpu_metadata
++++++++++++++++
Optional directive which declares the contents of the ``NT_AMDGPU_METADATA``
note record (see :ref:`amdgpu-elf-note-records-table-v3`).
The contents must be in the [YAML]_ markup format, with the same structure and
semantics described in :ref:`amdgpu-amdhsa-code-object-metadata-v3`.
This directive is terminated by an ``.end_amdgpu_metadata`` directive.
Example HSA Source Code (-mattr=+code-object-v3) Example HSA Source Code (-mattr=+code-object-v3)
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
Here is an example of a minimal assembly source file, defining one HSA kernel: Here is an example of a minimal assembly source file, defining one HSA kernel:
.. code-block:: none .. code-block:: none
.amdgcn_target "amdgcn-amd-amdhsa--gfx900+xnack" // optional .amdgcn_target "amdgcn-amd-amdhsa--gfx900+xnack" // optional
Show All 16 Lines.. code-block:: none
.rodata .rodata
.p2align 6 .p2align 6
.amdhsa_kernel hello_world .amdhsa_kernel hello_world
.amdhsa_user_sgpr_kernarg_segment_ptr 1 .amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_next_free_vgpr .amdgcn.next_free_vgpr .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
.amdhsa_next_free_sgpr .amdgcn.next_free_sgpr .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
.end_amdhsa_kernel .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
Additional Documentation 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-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-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-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-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-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-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-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-GFX9] `AMD "Vega" Instruction Set Architecture <http://developer.amd.com/wordpress/media/2013/12/Vega_Shader_ISA_28July2017.pdf>`__
.. [AMD-ROCm] `ROCm: Open Platform for Development, Discovery and Education Around GPU Computing <http://gpuopen.com/compute-product/rocm/>`__ .. [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>`__ .. [AMD-ROCm-github] `ROCm github <http://github.com/RadeonOpenCompute>`__
.. [HSA] `Heterogeneous System Architecture (HSA) Foundation <http://www.hsafoundation.com/>`__ .. [HSA] `Heterogeneous System Architecture (HSA) Foundation <http://www.hsafoundation.com/>`__
.. [ELF] `Executable and Linkable Format (ELF) <http://www.sco.com/developers/gabi/>`__ .. [ELF] `Executable and Linkable Format (ELF) <http://www.sco.com/developers/gabi/>`__
.. [DWARF] `DWARF Debugging Information Format <http://dwarfstd.org/>`__ .. [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>`__ .. [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>`__ .. [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>`__ .. [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>`__ .. [CLANG-ATTR] `Attributes in Clang <http://clang.llvm.org/docs/AttributeReference.html>`__