This is an archive of the discontinued LLVM Phabricator instance.

[AMDGPU] Update code object metadata format documentation
AcceptedPublic

Authored by t-tye on May 30 2018, 11:50 AM.

Download Raw Diff

Details

Reviewers

scott.linder
kzhuravl
b-sumner

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.

Diff Detail

Event Timeline

t-tye created this revision.May 30 2018, 11:50 AM

Herald added subscribers: tpr, dstuttard, yaxunl and 2 others. · View Herald TranscriptMay 30 2018, 11:50 AM

Rename note record enmerator from NT_AMD_AMDGPU_* to NT_AMDGPU_* to match the vendor name change from "AMD" to "AMDGPU".

looks good to me

This revision is now accepted and ready to land.Jun 13 2018, 6:34 PM

kzhuravl mentioned this in D48103: [AMDGPU] Update code object metadata format documentation.Jun 13 2018, 6:35 PM

scott.linder added a child revision: D48179: [AMDGPU] Emit MessagePack HSA Metadata for v3 code object.Jun 14 2018, 10:02 AM

scott.linder removed a child revision: D48179: [AMDGPU] Emit MessagePack HSA Metadata for v3 code object.Nov 14 2018, 9:32 AM

Revision Contents

Path

Size

docs/

AMDGPUUsage.rst

843 lines

Diff 151207

docs/AMDGPUUsage.rst

Show First 20 Lines • Show All 626 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		The AMDGPU backend code object uses the following ELF note record in the
``.note`` section. The Description column specifies the layout of the note		``.note`` section.
record's ``desc`` field. All fields are consecutive bytes. Note records with
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.

Additional note records can be present.		Additional note records can be present.

.. table:: AMDGPU ELF Note Records		.. table:: AMDGPU ELF Note Records
:name: amdgpu-elf-note-records-table		:name: amdgpu-elf-note-records-table

===== ============================== ======================================		======== ============================== ======================================
Name Type Description		Name Type Description
===== ============================== ======================================		======== ============================== ======================================
"AMD" ``NT_AMD_AMDGPU_HSA_METADATA`` <metadata null terminated string>		"AMDGPU" ``NT_AMDGPU_METADATA`` Metadata in Message Pack [MsgPack]_
===== ============================== ======================================		binary format.
		======== ============================== ======================================

..		..

.. table:: AMDGPU ELF Note Record Enumeration Values		.. table:: AMDGPU ELF Note Record Enumeration Values
:name: amdgpu-elf-note-record-enumeration-values-table		:name: amdgpu-elf-note-record-enumeration-values-table

============================== =====		============================== =====
Name Value		Name Value
============================== =====		============================== =====
reserved 0-9		reserved 0
``NT_AMD_AMDGPU_HSA_METADATA`` 10		``NT_AMDGPU_METADATA`` 1
reserved 11
============================== =====		============================== =====

``NT_AMD_AMDGPU_HSA_METADATA``		``NT_AMDGPU_METADATA``
Specifies extensible metadata associated with the code objects executed on HSA		Specifies extensible metadata associated with an AMDGPU code
[HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. It is required when		object. It is encoded as a map in the Message Pack [MsgPack]_ binary
the target triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`). See		data format. See :ref:`amdgpu-amdhsa-code-object-metadata` for the
:ref:`amdgpu-amdhsa-hsa-code-object-metadata` for the syntax of the code		map keys defined for the ``amdhsa`` OS.
object metadata string.

.. _amdgpu-symbols:		.. _amdgpu-symbols:

Symbols		Symbols
-------		-------

Symbols include the following:		Symbols include the following:

▲ Show 20 Lines • Show All 240 Lines • ▼ Show 20 Lines
(see :ref:`amdgpu-target-triples`).		(see :ref:`amdgpu-target-triples`).

AMDHSA		AMDHSA
------		------

This section provides code conventions used when the target triple OS is		This section provides code conventions used when the target triple OS is
``amdhsa`` (see :ref:`amdgpu-target-triples`).		``amdhsa`` (see :ref:`amdgpu-target-triples`).

.. _amdgpu-amdhsa-hsa-code-object-metadata:		.. _amdgpu-amdhsa-code-object-target-identification:

Code Object Target Identification		Code Object Target Identification
~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~		~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~

The AMDHSA OS uses the following syntax to specify the code object		The AMDHSA OS uses the following syntax to specify the code object
target as a single string:		target as a single string:

``<Architecture>-<Vendor>-<OS>-<Environment>-<Processor><Target Features>``		``<Architecture>-<Vendor>-<OS>-<Environment>-<Processor><Target Features>``
Show All 13 Lines	- ``<Target Features>`` is a list of the enabled Target Features
in the table :ref:`amdgpu-target-feature-table`. Note that *Target		in the table :ref:`amdgpu-target-feature-table`. Note that *Target
Features* must be included in the list if they are enabled even if		Features* must be included in the list if they are enabled even if
that is the default for Processor.		that is the default for Processor.

For example:		For example:

``"amdgcn-amd-amdhsa--gfx902+xnack"``		``"amdgcn-amd-amdhsa--gfx902+xnack"``

		.. _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 by the ``NT_AMD_AMDGPU_METADATA`` note record
(see :ref:`amdgpu-note-records`) and is required when the target triple OS is		(see :ref:`amdgpu-note-records`) and is required when the target triple OS is
``amdhsa`` (see :ref:`amdgpu-target-triples`). It must contain the minimum		``amdhsa`` (see :ref:`amdgpu-target-triples`). It must contain the minimum
information necessary to support the ROCM kernel queries. For example, the		information necessary to support the ROCM kernel queries. For example, the
segment sizes needed in a dispatch packet. In addition, a high level language		segment sizes needed in a dispatch packet. In addition, a high level language
runtime may require other information to be included. For example, the AMD		runtime may require other information to be included. For example, the AMD
OpenCL runtime records kernel argument information.		OpenCL runtime records kernel argument information.

The metadata is specified as a YAML formatted string (see [YAML]_ and		The metadata is represented as Message Pack formatted binary data (see
:doc:`YamlIO`).		[MsgPack]_). The top level is a Message Pack map that includes the
		keys defined in table
.. TODO		:ref:`amdgpu-amdhsa-code-object-metadata-map-table` and referenced
Is the string null terminated? It probably should not if YAML allows it to		tables.
contain null characters, otherwise it should be.
		Additional information can be added to the maps. To avoid conflicts,
The metadata is represented as a single YAML document comprised of the mapping		any key names should be prefixed by "vendor-name." where
defined in table :ref:`amdgpu-amdhsa-code-object-metadata-mapping-table` and		``vendor-name`` can be the the name of the vendor and specific vendor
referenced tables.		tool that generates the information. The prefix is abbreviated to
		simply "." when it appears within a map that has been added by the
For boolean values, the string values of ``false`` and ``true`` are used for		same vendor-name.
false and true respectively.

Additional information can be added to the mappings. To avoid conflicts, any		.. table:: AMDHSA Code Object Metadata Map
non-AMD key names should be prefixed by "vendor-name.".		:name: amdgpu-amdhsa-code-object-metadata-map-table

.. table:: AMDHSA Code Object Metadata Mapping		================= ============== ========= =======================================
:name: amdgpu-amdhsa-code-object-metadata-mapping-table

========== ============== ========= =======================================
String Key Value Type Required? Description		String Key Value Type Required? Description
========== ============== ========= =======================================		================= ============== ========= =======================================
"Version" sequence of Required - The first integer is the major		"amdhsa.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.
"Printf" sequence of Each string is encoded information		"amdhsa.printf" sequence of Each string is encoded information
strings about a printf function call. The		strings about a printf function call. The
encoded information is organized as		encoded information is organized as
fields separated by colon (':'):		fields separated by colon (':'):

``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``		``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``

where:		where:

``ID``		``ID``
A 32 bit integer as a unique id for		A 32 bit integer as a unique id for
each printf function call		each printf function call

``N``		``N``
A 32 bit integer equal to the number		A 32 bit integer equal to the number
of arguments of printf function call		of arguments of printf function call
minus 1		minus 1

``S[i]`` (where i = 0, 1, ... , N-1)		``S[i]`` (where i = 0, 1, ... , N-1)
32 bit integers for the size in bytes		32 bit integers for the size in bytes
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		"amdhsa.kernels" sequence of Required Sequence of the maps for each
mapping kernel in the code object. See		map kernel in the code object. See
:ref:`amdgpu-amdhsa-code-object-kernel-metadata-mapping-table`		:ref:`amdgpu-amdhsa-code-object-kernel-metadata-map-table`
for the definition of the mapping.		for the definition of the keys included
========== ============== ========= =======================================		in that map.
		================= ============== ========= =======================================

..		..

.. table:: AMDHSA Code Object Kernel Metadata Mapping		.. table:: AMDHSA Code Object Kernel Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-metadata-mapping-table		:name: amdgpu-amdhsa-code-object-kernel-metadata-map-table

================= ============== ========= ================================		=================================== ============== ========= ================================
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		".symbol" 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"
		- "HIP"
- "OpenMP"		- "OpenMP"
		- "Assembler"

"LanguageVersion" sequence of - The first integer is the major		".language_version" 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.		".args" sequence of Sequence of maps of the
See		map kernel arguments. See
:ref:`amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table`		:ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table`
for the mapping definition.		for the definition of the keys
"Args" sequence of Sequence of mappings of the		included in that map.
mapping kernel arguments. See		".reqd_workgroup_size" sequence of If not 0, 0, 0 then all values
:ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table`
for the definition of the mapping.
"CodeProps" mapping Mapping of properties related to
the kernel code. See
:ref:`amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table`
for the mapping definition.
================= ============== ========= ================================

..

.. table:: AMDHSA Code Object Kernel Attribute Metadata Mapping
:name: amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table

=================== ============== ========= ==============================
String Key Value Type Required? Description
=================== ============== ========= ==============================
"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
values. Defaults to 0, 0, 0.		values. Defaults to 0, 0, 0.

Corresponds to the OpenCL		Corresponds to the OpenCL
``reqd_work_group_size``		``reqd_work_group_size``
attribute.		attribute.
"WorkGroupSizeHint" sequence of The dispatch work-group size		".workgroup_size_hint" sequence of The dispatch work-group size
3 integers X, Y, Z is likely to be the		3 integers X, Y, Z is likely to be the
specified values.		specified values.

Corresponds to the OpenCL		Corresponds to the OpenCL
``work_group_size_hint``		``work_group_size_hint``
attribute.		attribute.
"VecTypeHint" string The name of a scalar or vector		".vec_type_hint" string The name of a scalar or vector
type.		type.

Corresponds to the OpenCL		Corresponds to the OpenCL
``vec_type_hint`` attribute.		``vec_type_hint`` attribute.

"RuntimeHandle" string The external symbol name		".device_enqueue_symbol" string The external symbol name
associated with a kernel.		associated with a kernel.
OpenCL runtime allocates a		OpenCL runtime allocates a
global buffer for the symbol		global buffer for the symbol
and saves the kernel's address		and saves the kernel's address
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.
=================== ============== ========= ==============================		".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 Kernel Argument Metadata Mapping		.. table:: AMDHSA Code Object Kernel Argument Metadata Map
:name: amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table		:name: amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table

================= ============== ========= ================================		====================== ============== ========= ================================
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.		".type_name" 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		".offset" integer Required Kernel argument offset in
bytes. Must be a power of two.		bytes. The offset must be a
"ValueKind" string Required Kernel argument kind that		multiple of the alignment
		required by the argument.
		".value_kind" string Required Kernel argument kind that
specifies how to set up the		specifies how to set up the
corresponding argument.		corresponding argument.
Values include:		Values include:

"ByValue"		"by_value"
The argument is copied		The argument is copied
directly into the kernarg.		directly into the kernarg.

"GlobalBuffer"		"global_buffer"
A global address space pointer		A global address space pointer
to the buffer data is passed		to the buffer data is passed
in the kernarg.		in the kernarg.

"DynamicSharedPointer"		"dynamic_shared_pointer"
A group address space pointer		A group address space pointer
to dynamically allocated LDS		to dynamically allocated LDS
is passed in the kernarg.		is passed in the kernarg.

"Sampler"		"sampler"
A global address space		A global address space
pointer to a S# is passed in		pointer to a S# is passed in
the kernarg.		the kernarg.

"Image"		"image"
A global address space		A global address space
pointer to a T# is passed in		pointer to a T# is passed in
the kernarg.		the kernarg.

"Pipe"		"pipe"
A global address space pointer		A global address space pointer
to an OpenCL pipe is passed in		to an OpenCL pipe is passed in
the kernarg.		the kernarg.

"Queue"		"queue"
A global address space pointer		A global address space pointer
to an OpenCL device enqueue		to an OpenCL device enqueue
queue is passed in the		queue is passed in the
kernarg.		kernarg.

"HiddenGlobalOffsetX"		"hidden_global_offset_x"
The OpenCL grid dispatch		The OpenCL grid dispatch
global offset for the X		global offset for the X
dimension is passed in the		dimension is passed in the
kernarg.		kernarg.

"HiddenGlobalOffsetY"		"hidden_global_offset_y"
The OpenCL grid dispatch		The OpenCL grid dispatch
global offset for the Y		global offset for the Y
dimension is passed in the		dimension is passed in the
kernarg.		kernarg.

"HiddenGlobalOffsetZ"		"hidden_global_offset_z"
The OpenCL grid dispatch		The OpenCL grid dispatch
global offset for the Z		global offset for the Z
dimension is passed in the		dimension is passed in the
kernarg.		kernarg.

"HiddenNone"		"hidden_none"
An argument that is not used		An argument that is not used
by the kernel. Space needs to		by the kernel. Space needs to
be left for it, but it does		be left for it, but it does
not need to be set up.		not need to be set up.

"HiddenPrintfBuffer"		"hidden_printf_buffer"
A global address space pointer		A global address space pointer
to the runtime printf buffer		to the runtime printf buffer
is passed in kernarg.		is passed in kernarg.

"HiddenDefaultQueue"		"hidden_default_queue"
A global address space pointer		A global address space pointer
to the OpenCL device enqueue		to the OpenCL device enqueue
queue that should be used by		queue that should be used by
the kernel by default is		the kernel by default is
passed in the kernarg.		passed in the kernarg.

"HiddenCompletionAction"		"hidden_completion_action"
A global address space pointer		A global address space pointer
to help link enqueued kernels into		to help link enqueued kernels into
the ancestor tree for determining		the ancestor tree for determining
when the parent kernel has finished.		when the parent kernel has finished.

"ValueType" string Required Kernel argument value type. Only		".value_type" string Required Kernel argument value type. Only
present if "ValueKind" is		present if ".value_kind" is
"ByValue". For vector data		"by_value". For vector data
types, the value is for the		types, the value is for the
element type. Values include:		element type. Values include:

- "Struct"		- "struct"
- "I8"		- "i8"
- "U8"		- "u8"
- "I16"		- "i16"
- "U16"		- "u16"
- "F16"		- "f16"
- "I32"		- "i32"
- "U32"		- "u32"
- "F32"		- "f32"
- "I64"		- "i64"
- "U64"		- "u64"
- "F64"		- "f64"

.. TODO		.. TODO
How can it be determined if a		How can it be determined if a
vector type, and what size		vector type, and what size
vector?		vector?
"PointeeAlign" integer Alignment in bytes of pointee		".pointee_align" integer Alignment in bytes of pointee
type for pointer type kernel		type for pointer type kernel
argument. Must be a power		argument. Must be a power
of 2. Only present if		of 2. Only present if
"ValueKind" is		".value_kind" is
"DynamicSharedPointer".		"dynamic_shared_pointer".
"AddrSpaceQual" string Kernel argument address space		".address_space" string Kernel argument address space
qualifier. Only present if		qualifier. Only present if
"ValueKind" is "GlobalBuffer" or		".value_kind" is "global_buffer" or
"DynamicSharedPointer". Values		"dynamic_shared_pointer". Values
are:		are:

- "Private"		- "private"
- "Global"		- "global"
- "Constant"		- "constant"
- "Local"		- "local"
- "Generic"		- "generic"
- "Region"		- "region"

.. TODO		.. TODO
Is GlobalBuffer only Global		Is "global_buffer" only "global"
or Constant? Is		or "constant"? Is
DynamicSharedPointer always		"dynamic_shared_pointer" always
Local? Can HCC allow Generic?		"local"? Can HCC allow "generic"?
How can Private or Region		How can "private" or "region"
ever happen?		ever happen?
"AccQual" string Kernel argument access		".access" string Kernel argument access
qualifier. Only present if		qualifier. Only present if
"ValueKind" is "Image" or		".value_kind" is "image" or
"Pipe". Values		"pipe". Values
are:		are:

- "ReadOnly"		- "read_only"
- "WriteOnly"		- "write_only"
- "ReadWrite"		- "read_write"

.. TODO		.. TODO
Does this apply to		Does this apply to
GlobalBuffer?		"global_buffer"?
"ActualAccQual" string The actual memory accesses		".actual_access" string The actual memory accesses
performed by the kernel on the		performed by the kernel on the
kernel argument. Only present if		kernel argument. Only present if
"ValueKind" is "GlobalBuffer",		".value_kind" is "global_buffer",
"Image", or "Pipe". This may be		"image", or "pipe". This may be
more restrictive than indicated		more restrictive than indicated
by "AccQual" to reflect what the		by ".access" to reflect what the
kernel actual does. If not		kernel actual does. If not
present then the runtime must		present then the runtime must
assume what is implied by		assume what is implied by
"AccQual" and "IsConst". Values		".access" and ".is_const" . Values
are:		are:

- "ReadOnly"		- "read_only"
- "WriteOnly"		- "write_only"
- "ReadWrite"		- "read_write"

"IsConst" boolean Indicates if the kernel argument		".is_const" boolean Indicates if the kernel argument
is const qualified. Only present		is const qualified. Only present
if "ValueKind" is		if ".value_kind" is
"GlobalBuffer".		"global_buffer".

"IsRestrict" boolean Indicates if the kernel argument		".is_restrict" boolean Indicates if the kernel argument
is restrict qualified. Only		is restrict qualified. Only
present if "ValueKind" is		present if ".value_kind" is
"GlobalBuffer".		"global_buffer".

"IsVolatile" boolean Indicates if the kernel argument		".is_volatile" boolean Indicates if the kernel argument
is volatile qualified. Only		is volatile qualified. Only
present if "ValueKind" is		present if ".value_kind" is
"GlobalBuffer".		"global_buffer".

"IsPipe" boolean Indicates if the kernel argument		".is_pipe" boolean Indicates if the kernel argument
is pipe qualified. Only present		is pipe qualified. Only present
if "ValueKind" is "Pipe".		if ".value_kind" is "pipe".

.. TODO		.. TODO
Can GlobalBuffer be pipe		Can "global_buffer" be pipe
qualified?		qualified?
================= ============== ========= ================================		====================== ============== ========= ================================

..

.. table:: AMDHSA Code Object Kernel Code Properties Metadata Mapping
:name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table

============================ ============== ========= =====================
String Key Value Type Required? Description
============================ ============== ========= =====================
"KernargSegmentSize" integer Required The size in bytes of
the kernarg segment
that holds the values
of the arguments to
the kernel.
"GroupSegmentFixedSize" 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.
"PrivateSegmentFixedSize" 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.
"KernargSegmentAlign" integer Required The maximum byte
alignment of
arguments in the
kernarg segment. Must
be a power of 2.
"WavefrontSize" integer Required Wavefront size. Must
be a power of 2.
"NumSGPRs" integer Required Number of scalar
registers used by a
wavefront for
GFX6-GFX9. 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.
"NumVGPRs" integer Required Number of vector
registers used by
each work-item for
GFX6-GFX9
"MaxFlatWorkGroupSize" 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.
"NumSpilledSGPRs" integer Number of stores from
a scalar register to
a register allocator
created spill
location.
"NumSpilledVGPRs" integer Number of stores from
a vector register to
a register allocator
created spill
location.
============================ ============== ========= =====================

..		..

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
Show All 37 Lines	5. An AQL kernel dispatch packet is created on the AQL queue. The ROCm runtime
visible to the kernel agent. AQL defines a doorbell signal mechanism to		visible to the kernel agent. AQL defines a doorbell signal mechanism to
notify the kernel agent that the AQL queue has been updated. These rules, and		notify the kernel agent that the AQL queue has been updated. These rules, and
the layout of the AQL queue and kernel dispatch packet is defined in the *HSA		the layout of the AQL queue and kernel dispatch packet is defined in the *HSA
System Architecture Specification* [HSA]_.		System Architecture Specification* [HSA]_.
6. A kernel dispatch packet includes information about the actual dispatch,		6. A kernel dispatch packet includes information about the actual dispatch,
such as grid and work-group size, together with information from the code		such as grid and work-group size, together with information from the code
object about the kernel, such as segment sizes. The ROCm runtime queries on		object about the kernel, such as segment sizes. The ROCm runtime queries on
the kernel symbol can be used to obtain the code object values which are		the kernel symbol can be used to obtain the code object values which are
recorded in the :ref:`amdgpu-amdhsa-hsa-code-object-metadata`.		recorded in the :ref:`amdgpu-amdhsa-code-object-metadata`.
7. CP executes micro-code and is responsible for detecting and setting up the		7. CP executes micro-code and is responsible for detecting and setting up the
GPU to execute the wavefronts of a kernel dispatch.		GPU to execute the wavefronts of a kernel dispatch.
8. CP ensures that when the a wavefront starts executing the kernel machine		8. CP ensures that when the a wavefront starts executing the kernel machine
code, the scalar general purpose registers (SGPR) and vector general purpose		code, the scalar general purpose registers (SGPR) and vector general purpose
registers (VGPR) are set up as required by the machine code. The required		registers (VGPR) are set up as required by the machine code. The required
setup is defined in the :ref:`amdgpu-amdhsa-kernel-descriptor`. The initial		setup is defined in the :ref:`amdgpu-amdhsa-kernel-descriptor`. The initial
register state is defined in		register state is defined in
:ref:`amdgpu-amdhsa-initial-kernel-execution-state`.		:ref:`amdgpu-amdhsa-initial-kernel-execution-state`.
▲ Show 20 Lines • Show All 2,518 Lines • ▼ Show 20 Lines

When the language is OpenCL the following differences occur:		When the language is OpenCL the following differences occur:

1. The OpenCL memory model is used (see :ref:`amdgpu-amdhsa-memory-model`).		1. The OpenCL memory model is used (see :ref:`amdgpu-amdhsa-memory-model`).
2. The AMDGPU backend appends additional arguments to the kernel's explicit		2. The AMDGPU backend appends additional arguments to the kernel's explicit
arguments for the AMDHSA OS (see		arguments for the AMDHSA OS (see
:ref:`opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table`).		:ref:`opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table`).
3. Additional metadata is generated		3. Additional metadata is generated
(see :ref:`amdgpu-amdhsa-hsa-code-object-metadata`).		(see :ref:`amdgpu-amdhsa-code-object-metadata`).

.. table:: OpenCL kernel implicit arguments appended for AMDHSA OS		.. table:: OpenCL kernel implicit arguments appended for AMDHSA OS
:name: opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table		:name: opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table

======== ==== ========= ===========================================		======== ==== ========= ===========================================
Position Byte Byte Description		Position Byte Byte Description
Size Alignment		Size Alignment
======== ==== ========= ===========================================		======== ==== ========= ===========================================
▲ Show 20 Lines • Show All 292 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 42 Lines • ▼ Show 20 Lines
.. [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>`__		.. [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>`__