| blob | fc624a724a861fce522d0ab917c55474e690bf0f |
1 =============================
2 User Guide for AMDGPU Backend
3 =============================
5 .. contents::
6 :local:
8 Introduction
9 ============
11 The AMDGPU backend provides ISA code generation for AMD GPUs, starting with the
12 R600 family up until the current GCN families. It lives in the
13 ``lib/Target/AMDGPU`` directory.
15 LLVM
16 ====
18 .. _amdgpu-target-triples:
20 Target Triples
21 --------------
23 Use the ``clang -target <Architecture>-<Vendor>-<OS>-<Environment>`` option to
24 specify the target triple:
26 .. table:: AMDGPU Architectures
27 :name: amdgpu-architecture-table
29 ============ ==============================================================
30 Architecture Description
31 ============ ==============================================================
32 ``r600`` AMD GPUs HD2XXX-HD6XXX for graphics and compute shaders.
33 ``amdgcn`` AMD GPUs GCN GFX6 onwards for graphics and compute shaders.
34 ============ ==============================================================
36 .. table:: AMDGPU Vendors
37 :name: amdgpu-vendor-table
39 ============ ==============================================================
40 Vendor Description
41 ============ ==============================================================
42 ``amd`` Can be used for all AMD GPU usage.
43 ``mesa3d`` Can be used if the OS is ``mesa3d``.
44 ============ ==============================================================
46 .. table:: AMDGPU Operating Systems
47 :name: amdgpu-os-table
49 ============== ============================================================
50 OS Description
51 ============== ============================================================
52 *<empty>* Defaults to the *unknown* OS.
53 ``amdhsa`` Compute kernels executed on HSA [HSA]_ compatible runtimes
54 such as AMD's ROCm [AMD-ROCm]_.
55 ``amdpal`` Graphic shaders and compute kernels executed on AMD PAL
56 runtime.
57 ``mesa3d`` Graphic shaders and compute kernels executed on Mesa 3D
58 runtime.
59 ============== ============================================================
61 .. table:: AMDGPU Environments
62 :name: amdgpu-environment-table
64 ============ ==============================================================
65 Environment Description
66 ============ ==============================================================
67 *<empty>* Default.
68 ============ ==============================================================
70 .. _amdgpu-processors:
72 Processors
73 ----------
75 Use the ``clang -mcpu <Processor>`` option to specify the AMD GPU processor. The
76 names from both the *Processor* and *Alternative Processor* can be used.
78 .. table:: AMDGPU Processors
79 :name: amdgpu-processor-table
81 =========== =============== ============ ===== ================= ======= ======================
82 Processor Alternative Target dGPU/ Target ROCm Example
83 Processor Triple APU Features Support Products
84 Architecture Supported
85 [Default]
86 =========== =============== ============ ===== ================= ======= ======================
87 **Radeon HD 2000/3000 Series (R600)** [AMD-RADEON-HD-2000-3000]_
88 -----------------------------------------------------------------------------------------------
89 ``r600`` ``r600`` dGPU
90 ``r630`` ``r600`` dGPU
91 ``rs880`` ``r600`` dGPU
92 ``rv670`` ``r600`` dGPU
93 **Radeon HD 4000 Series (R700)** [AMD-RADEON-HD-4000]_
94 -----------------------------------------------------------------------------------------------
95 ``rv710`` ``r600`` dGPU
96 ``rv730`` ``r600`` dGPU
97 ``rv770`` ``r600`` dGPU
98 **Radeon HD 5000 Series (Evergreen)** [AMD-RADEON-HD-5000]_
99 -----------------------------------------------------------------------------------------------
100 ``cedar`` ``r600`` dGPU
101 ``cypress`` ``r600`` dGPU
102 ``juniper`` ``r600`` dGPU
103 ``redwood`` ``r600`` dGPU
104 ``sumo`` ``r600`` dGPU
105 **Radeon HD 6000 Series (Northern Islands)** [AMD-RADEON-HD-6000]_
106 -----------------------------------------------------------------------------------------------
107 ``barts`` ``r600`` dGPU
108 ``caicos`` ``r600`` dGPU
109 ``cayman`` ``r600`` dGPU
110 ``turks`` ``r600`` dGPU
111 **GCN GFX6 (Southern Islands (SI))** [AMD-GCN-GFX6]_
112 -----------------------------------------------------------------------------------------------
113 ``gfx600`` - ``tahiti`` ``amdgcn`` dGPU
114 ``gfx601`` - ``hainan`` ``amdgcn`` dGPU
115 - ``oland``
116 - ``pitcairn``
117 - ``verde``
118 **GCN GFX7 (Sea Islands (CI))** [AMD-GCN-GFX7]_
119 -----------------------------------------------------------------------------------------------
120 ``gfx700`` - ``kaveri`` ``amdgcn`` APU - A6-7000
121 - A6 Pro-7050B
122 - A8-7100
123 - A8 Pro-7150B
124 - A10-7300
125 - A10 Pro-7350B
126 - FX-7500
127 - A8-7200P
128 - A10-7400P
129 - FX-7600P
130 ``gfx701`` - ``hawaii`` ``amdgcn`` dGPU ROCm - FirePro W8100
131 - FirePro W9100
132 - FirePro S9150
133 - FirePro S9170
134 ``gfx702`` ``amdgcn`` dGPU ROCm - Radeon R9 290
135 - Radeon R9 290x
136 - Radeon R390
137 - Radeon R390x
138 ``gfx703`` - ``kabini`` ``amdgcn`` APU - E1-2100
139 - ``mullins`` - E1-2200
140 - E1-2500
141 - E2-3000
142 - E2-3800
143 - A4-5000
144 - A4-5100
145 - A6-5200
146 - A4 Pro-3340B
147 ``gfx704`` - ``bonaire`` ``amdgcn`` dGPU - Radeon HD 7790
148 - Radeon HD 8770
149 - R7 260
150 - R7 260X
151 **GCN GFX8 (Volcanic Islands (VI))** [AMD-GCN-GFX8]_
152 -----------------------------------------------------------------------------------------------
153 ``gfx801`` - ``carrizo`` ``amdgcn`` APU - xnack - A6-8500P
154 [on] - Pro A6-8500B
155 - A8-8600P
156 - Pro A8-8600B
157 - FX-8800P
158 - Pro A12-8800B
159 \ ``amdgcn`` APU - xnack ROCm - A10-8700P
160 [on] - Pro A10-8700B
161 - A10-8780P
162 \ ``amdgcn`` APU - xnack - A10-9600P
163 [on] - A10-9630P
164 - A12-9700P
165 - A12-9730P
166 - FX-9800P
167 - FX-9830P
168 \ ``amdgcn`` APU - xnack - E2-9010
169 [on] - A6-9210
170 - A9-9410
171 ``gfx802`` - ``iceland`` ``amdgcn`` dGPU - xnack ROCm - FirePro S7150
172 - ``tonga`` [off] - FirePro S7100
173 - FirePro W7100
174 - Radeon R285
175 - Radeon R9 380
176 - Radeon R9 385
177 - Mobile FirePro
178 M7170
179 ``gfx803`` - ``fiji`` ``amdgcn`` dGPU - xnack ROCm - Radeon R9 Nano
180 [off] - Radeon R9 Fury
181 - Radeon R9 FuryX
182 - Radeon Pro Duo
183 - FirePro S9300x2
184 - Radeon Instinct MI8
185 \ - ``polaris10`` ``amdgcn`` dGPU - xnack ROCm - Radeon RX 470
186 [off] - Radeon RX 480
187 - Radeon Instinct MI6
188 \ - ``polaris11`` ``amdgcn`` dGPU - xnack ROCm - Radeon RX 460
189 [off]
190 ``gfx810`` - ``stoney`` ``amdgcn`` APU - xnack
191 [on]
192 **GCN GFX9** [AMD-GCN-GFX9]_
193 -----------------------------------------------------------------------------------------------
194 ``gfx900`` ``amdgcn`` dGPU - xnack ROCm - Radeon Vega
195 [off] Frontier Edition
196 - Radeon RX Vega 56
197 - Radeon RX Vega 64
198 - Radeon RX Vega 64
199 Liquid
200 - Radeon Instinct MI25
201 ``gfx902`` ``amdgcn`` APU - xnack - Ryzen 3 2200G
202 [on] - Ryzen 5 2400G
203 ``gfx904`` ``amdgcn`` dGPU - xnack *TBA*
204 [off]
205 .. TODO
206 Add product
207 names.
208 ``gfx906`` ``amdgcn`` dGPU - xnack - Radeon Instinct MI50
209 [off] - Radeon Instinct MI60
210 ``gfx908`` ``amdgcn`` dGPU - xnack *TBA*
211 [off]
212 sram-ecc
213 [on]
214 ``gfx909`` ``amdgcn`` APU - xnack *TBA* (Raven Ridge 2)
215 [on]
216 .. TODO
217 Add product
218 names.
219 **GCN GFX10** [AMD-GCN-GFX10]_
220 -----------------------------------------------------------------------------------------------
221 ``gfx1010`` ``amdgcn`` dGPU - xnack *TBA*
222 [off]
223 - wavefrontsize64
224 [off]
225 - cumode
226 [off]
227 .. TODO
228 Add product
229 names.
230 ``gfx1011`` ``amdgcn`` dGPU - xnack *TBA*
231 [off]
232 - wavefrontsize64
233 [off]
234 - cumode
235 [off]
236 .. TODO
237 Add product
238 names.
239 ``gfx1012`` ``amdgcn`` dGPU - xnack *TBA*
240 [off]
241 - wavefrontsize64
242 [off]
243 - cumode
244 [off]
245 .. TODO
246 Add product
247 names.
248 =========== =============== ============ ===== ================= ======= ======================
250 .. _amdgpu-target-features:
252 Target Features
253 ---------------
255 Target features control how code is generated to support certain
256 processor specific features. Not all target features are supported by
257 all processors. The runtime must ensure that the features supported by
258 the device used to execute the code match the features enabled when
259 generating the code. A mismatch of features may result in incorrect
260 execution, or a reduction in performance.
262 The target features supported by each processor, and the default value
263 used if not specified explicitly, is listed in
264 :ref:`amdgpu-processor-table`.
266 Use the ``clang -m[no-]<TargetFeature>`` option to specify the AMD GPU
267 target features.
269 For example:
271 ``-mxnack``
272 Enable the ``xnack`` feature.
273 ``-mno-xnack``
274 Disable the ``xnack`` feature.
276 .. table:: AMDGPU Target Features
277 :name: amdgpu-target-feature-table
279 ====================== ==================================================
280 Target Feature Description
281 ====================== ==================================================
282 -m[no-]xnack Enable/disable generating code that has
283 memory clauses that are compatible with
284 having XNACK replay enabled.
286 This is used for demand paging and page
287 migration. If XNACK replay is enabled in
288 the device, then if a page fault occurs
289 the code may execute incorrectly if the
290 ``xnack`` feature is not enabled. Executing
291 code that has the feature enabled on a
292 device that does not have XNACK replay
293 enabled will execute correctly, but may
294 be less performant than code with the
295 feature disabled.
297 -m[no-]sram-ecc Enable/disable generating code that assumes SRAM
298 ECC is enabled/disabled.
300 -m[no-]wavefrontsize64 Control the default wavefront size used when
301 generating code for kernels. When disabled
302 native wavefront size 32 is used, when enabled
303 wavefront size 64 is used.
305 -m[no-]cumode Control the default wavefront execution mode used
306 when generating code for kernels. When disabled
307 native WGP wavefront execution mode is used,
308 when enabled CU wavefront execution mode is used
309 (see :ref:`amdgpu-amdhsa-memory-model`).
310 ====================== ==================================================
312 .. _amdgpu-address-spaces:
314 Address Spaces
315 --------------
317 The AMDGPU backend uses the following address space mappings.
319 The memory space names used in the table, aside from the region memory space, is
320 from the OpenCL standard.
322 LLVM Address Space number is used throughout LLVM (for example, in LLVM IR).
324 .. table:: Address Space Mapping
325 :name: amdgpu-address-space-mapping-table
327 ================== =================================
328 LLVM Address Space Memory Space
329 ================== =================================
330 0 Generic (Flat)
331 1 Global
332 2 Region (GDS)
333 3 Local (group/LDS)
334 4 Constant
335 5 Private (Scratch)
336 6 Constant 32-bit
337 7 Buffer Fat Pointer (experimental)
338 ================== =================================
340 The buffer fat pointer is an experimental address space that is currently
341 unsupported in the backend. It exposes a non-integral pointer that is in future
342 intended to support the modelling of 128-bit buffer descriptors + a 32-bit
343 offset into the buffer descriptor (in total encapsulating a 160-bit 'pointer'),
344 allowing us to use normal LLVM load/store/atomic operations to model the buffer
345 descriptors used heavily in graphics workloads targeting the backend.
347 .. _amdgpu-memory-scopes:
349 Memory Scopes
350 -------------
352 This section provides LLVM memory synchronization scopes supported by the AMDGPU
353 backend memory model when the target triple OS is ``amdhsa`` (see
354 :ref:`amdgpu-amdhsa-memory-model` and :ref:`amdgpu-target-triples`).
356 The memory model supported is based on the HSA memory model [HSA]_ which is
357 based in turn on HRF-indirect with scope inclusion [HRF]_. The happens-before
358 relation is transitive over the synchonizes-with relation independent of scope,
359 and synchonizes-with allows the memory scope instances to be inclusive (see
360 table :ref:`amdgpu-amdhsa-llvm-sync-scopes-table`).
362 This is different to the OpenCL [OpenCL]_ memory model which does not have scope
363 inclusion and requires the memory scopes to exactly match. However, this
364 is conservatively correct for OpenCL.
366 .. table:: AMDHSA LLVM Sync Scopes
367 :name: amdgpu-amdhsa-llvm-sync-scopes-table
369 ======================= ===================================================
370 LLVM Sync Scope Description
371 ======================= ===================================================
372 *none* The default: ``system``.
374 Synchronizes with, and participates in modification
375 and seq_cst total orderings with, other operations
376 (except image operations) for all address spaces
377 (except private, or generic that accesses private)
378 provided the other operation's sync scope is:
380 - ``system``.
381 - ``agent`` and executed by a thread on the same
382 agent.
383 - ``workgroup`` and executed by a thread in the
384 same workgroup.
385 - ``wavefront`` and executed by a thread in the
386 same wavefront.
388 ``agent`` Synchronizes with, and participates in modification
389 and seq_cst total orderings with, other operations
390 (except image operations) for all address spaces
391 (except private, or generic that accesses private)
392 provided the other operation's sync scope is:
394 - ``system`` or ``agent`` and executed by a thread
395 on the same agent.
396 - ``workgroup`` and executed by a thread in the
397 same workgroup.
398 - ``wavefront`` and executed by a thread in the
399 same wavefront.
401 ``workgroup`` Synchronizes with, and participates in modification
402 and seq_cst total orderings with, other operations
403 (except image operations) for all address spaces
404 (except private, or generic that accesses private)
405 provided the other operation's sync scope is:
407 - ``system``, ``agent`` or ``workgroup`` and
408 executed by a thread in the same workgroup.
409 - ``wavefront`` and executed by a thread in the
410 same wavefront.
412 ``wavefront`` Synchronizes with, and participates in modification
413 and seq_cst total orderings with, other operations
414 (except image operations) for all address spaces
415 (except private, or generic that accesses private)
416 provided the other operation's sync scope is:
418 - ``system``, ``agent``, ``workgroup`` or
419 ``wavefront`` and executed by a thread in the
420 same wavefront.
422 ``singlethread`` Only synchronizes with, and participates in
423 modification and seq_cst total orderings with,
424 other operations (except image operations) running
425 in the same thread for all address spaces (for
426 example, in signal handlers).
428 ``one-as`` Same as ``system`` but only synchronizes with other
429 operations within the same address space.
431 ``agent-one-as`` Same as ``agent`` but only synchronizes with other
432 operations within the same address space.
434 ``workgroup-one-as`` Same as ``workgroup`` but only synchronizes with
435 other operations within the same address space.
437 ``wavefront-one-as`` Same as ``wavefront`` but only synchronizes with
438 other operations within the same address space.
440 ``singlethread-one-as`` Same as ``singlethread`` but only synchronizes with
441 other operations within the same address space.
442 ======================= ===================================================
444 AMDGPU Intrinsics
445 -----------------
447 The AMDGPU backend implements the following LLVM IR intrinsics.
449 *This section is WIP.*
451 .. TODO
452 List AMDGPU intrinsics
454 AMDGPU Attributes
455 -----------------
457 The AMDGPU backend supports the following LLVM IR attributes.
459 .. table:: AMDGPU LLVM IR Attributes
460 :name: amdgpu-llvm-ir-attributes-table
462 ======================================= ==========================================================
463 LLVM Attribute Description
464 ======================================= ==========================================================
465 "amdgpu-flat-work-group-size"="min,max" Specify the minimum and maximum flat work group sizes that
466 will be specified when the kernel is dispatched. Generated
467 by the ``amdgpu_flat_work_group_size`` CLANG attribute [CLANG-ATTR]_.
468 "amdgpu-implicitarg-num-bytes"="n" Number of kernel argument bytes to add to the kernel
469 argument block size for the implicit arguments. This
470 varies by OS and language (for OpenCL see
471 :ref:`opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table`).
472 "amdgpu-num-sgpr"="n" Specifies the number of SGPRs to use. Generated by
473 the ``amdgpu_num_sgpr`` CLANG attribute [CLANG-ATTR]_.
474 "amdgpu-num-vgpr"="n" Specifies the number of VGPRs to use. Generated by the
475 ``amdgpu_num_vgpr`` CLANG attribute [CLANG-ATTR]_.
476 "amdgpu-waves-per-eu"="m,n" Specify the minimum and maximum number of waves per
477 execution unit. Generated by the ``amdgpu_waves_per_eu``
478 CLANG attribute [CLANG-ATTR]_.
479 "amdgpu-ieee" true/false. Specify whether the function expects the IEEE field of the
480 mode register to be set on entry. Overrides the default for
481 the calling convention.
482 "amdgpu-dx10-clamp" true/false. Specify whether the function expects the DX10_CLAMP field of
483 the mode register to be set on entry. Overrides the default
484 for the calling convention.
485 ======================================= ==========================================================
487 Code Object
488 ===========
490 The AMDGPU backend generates a standard ELF [ELF]_ relocatable code object that
491 can be linked by ``lld`` to produce a standard ELF shared code object which can
492 be loaded and executed on an AMDGPU target.
494 Header
495 ------
497 The AMDGPU backend uses the following ELF header:
499 .. table:: AMDGPU ELF Header
500 :name: amdgpu-elf-header-table
502 ========================== ===============================
503 Field Value
504 ========================== ===============================
505 ``e_ident[EI_CLASS]`` ``ELFCLASS64``
506 ``e_ident[EI_DATA]`` ``ELFDATA2LSB``
507 ``e_ident[EI_OSABI]`` - ``ELFOSABI_NONE``
508 - ``ELFOSABI_AMDGPU_HSA``
509 - ``ELFOSABI_AMDGPU_PAL``
510 - ``ELFOSABI_AMDGPU_MESA3D``
511 ``e_ident[EI_ABIVERSION]`` - ``ELFABIVERSION_AMDGPU_HSA``
512 - ``ELFABIVERSION_AMDGPU_PAL``
513 - ``ELFABIVERSION_AMDGPU_MESA3D``
514 ``e_type`` - ``ET_REL``
515 - ``ET_DYN``
516 ``e_machine`` ``EM_AMDGPU``
517 ``e_entry`` 0
518 ``e_flags`` See :ref:`amdgpu-elf-header-e_flags-table`
519 ========================== ===============================
521 ..
523 .. table:: AMDGPU ELF Header Enumeration Values
524 :name: amdgpu-elf-header-enumeration-values-table
526 =============================== =====
527 Name Value
528 =============================== =====
529 ``EM_AMDGPU`` 224
530 ``ELFOSABI_NONE`` 0
531 ``ELFOSABI_AMDGPU_HSA`` 64
532 ``ELFOSABI_AMDGPU_PAL`` 65
533 ``ELFOSABI_AMDGPU_MESA3D`` 66
534 ``ELFABIVERSION_AMDGPU_HSA`` 1
535 ``ELFABIVERSION_AMDGPU_PAL`` 0
536 ``ELFABIVERSION_AMDGPU_MESA3D`` 0
537 =============================== =====
539 ``e_ident[EI_CLASS]``
540 The ELF class is:
542 * ``ELFCLASS32`` for ``r600`` architecture.
544 * ``ELFCLASS64`` for ``amdgcn`` architecture which only supports 64
545 bit applications.
547 ``e_ident[EI_DATA]``
548 All AMDGPU targets use ``ELFDATA2LSB`` for little-endian byte ordering.
550 ``e_ident[EI_OSABI]``
551 One of the following AMD GPU architecture specific OS ABIs
552 (see :ref:`amdgpu-os-table`):
554 * ``ELFOSABI_NONE`` for *unknown* OS.
556 * ``ELFOSABI_AMDGPU_HSA`` for ``amdhsa`` OS.
558 * ``ELFOSABI_AMDGPU_PAL`` for ``amdpal`` OS.
560 * ``ELFOSABI_AMDGPU_MESA3D`` for ``mesa3D`` OS.
562 ``e_ident[EI_ABIVERSION]``
563 The ABI version of the AMD GPU architecture specific OS ABI to which the code
564 object conforms:
566 * ``ELFABIVERSION_AMDGPU_HSA`` is used to specify the version of AMD HSA
567 runtime ABI.
569 * ``ELFABIVERSION_AMDGPU_PAL`` is used to specify the version of AMD PAL
570 runtime ABI.
572 * ``ELFABIVERSION_AMDGPU_MESA3D`` is used to specify the version of AMD MESA
573 3D runtime ABI.
575 ``e_type``
576 Can be one of the following values:
579 ``ET_REL``
580 The type produced by the AMD GPU backend compiler as it is relocatable code
581 object.
583 ``ET_DYN``
584 The type produced by the linker as it is a shared code object.
586 The AMD HSA runtime loader requires a ``ET_DYN`` code object.
588 ``e_machine``
589 The value ``EM_AMDGPU`` is used for the machine for all processors supported
590 by the ``r600`` and ``amdgcn`` architectures (see
591 :ref:`amdgpu-processor-table`). The specific processor is specified in the
592 ``EF_AMDGPU_MACH`` bit field of the ``e_flags`` (see
593 :ref:`amdgpu-elf-header-e_flags-table`).
595 ``e_entry``
596 The entry point is 0 as the entry points for individual kernels must be
597 selected in order to invoke them through AQL packets.
599 ``e_flags``
600 The AMDGPU backend uses the following ELF header flags:
602 .. table:: AMDGPU ELF Header ``e_flags``
603 :name: amdgpu-elf-header-e_flags-table
605 ================================= ========== =============================
606 Name Value Description
607 ================================= ========== =============================
608 **AMDGPU Processor Flag** See :ref:`amdgpu-processor-table`.
609 -------------------------------------------- -----------------------------
610 ``EF_AMDGPU_MACH`` 0x000000ff AMDGPU processor selection
611 mask for
612 ``EF_AMDGPU_MACH_xxx`` values
613 defined in
614 :ref:`amdgpu-ef-amdgpu-mach-table`.
615 ``EF_AMDGPU_XNACK`` 0x00000100 Indicates if the ``xnack``
616 target feature is
617 enabled for all code
618 contained in the code object.
619 If the processor
620 does not support the
621 ``xnack`` target
622 feature then must
623 be 0.
624 See
625 :ref:`amdgpu-target-features`.
626 ``EF_AMDGPU_SRAM_ECC`` 0x00000200 Indicates if the ``sram-ecc``
627 target feature is
628 enabled for all code
629 contained in the code object.
630 If the processor
631 does not support the
632 ``sram-ecc`` target
633 feature then must
634 be 0.
635 See
636 :ref:`amdgpu-target-features`.
637 ================================= ========== =============================
639 .. table:: AMDGPU ``EF_AMDGPU_MACH`` Values
640 :name: amdgpu-ef-amdgpu-mach-table
642 ================================= ========== =============================
643 Name Value Description (see
644 :ref:`amdgpu-processor-table`)
645 ================================= ========== =============================
646 ``EF_AMDGPU_MACH_NONE`` 0x000 *not specified*
647 ``EF_AMDGPU_MACH_R600_R600`` 0x001 ``r600``
648 ``EF_AMDGPU_MACH_R600_R630`` 0x002 ``r630``
649 ``EF_AMDGPU_MACH_R600_RS880`` 0x003 ``rs880``
650 ``EF_AMDGPU_MACH_R600_RV670`` 0x004 ``rv670``
651 ``EF_AMDGPU_MACH_R600_RV710`` 0x005 ``rv710``
652 ``EF_AMDGPU_MACH_R600_RV730`` 0x006 ``rv730``
653 ``EF_AMDGPU_MACH_R600_RV770`` 0x007 ``rv770``
654 ``EF_AMDGPU_MACH_R600_CEDAR`` 0x008 ``cedar``
655 ``EF_AMDGPU_MACH_R600_CYPRESS`` 0x009 ``cypress``
656 ``EF_AMDGPU_MACH_R600_JUNIPER`` 0x00a ``juniper``
657 ``EF_AMDGPU_MACH_R600_REDWOOD`` 0x00b ``redwood``
658 ``EF_AMDGPU_MACH_R600_SUMO`` 0x00c ``sumo``
659 ``EF_AMDGPU_MACH_R600_BARTS`` 0x00d ``barts``
660 ``EF_AMDGPU_MACH_R600_CAICOS`` 0x00e ``caicos``
661 ``EF_AMDGPU_MACH_R600_CAYMAN`` 0x00f ``cayman``
662 ``EF_AMDGPU_MACH_R600_TURKS`` 0x010 ``turks``
663 *reserved* 0x011 - Reserved for ``r600``
664 0x01f architecture processors.
665 ``EF_AMDGPU_MACH_AMDGCN_GFX600`` 0x020 ``gfx600``
666 ``EF_AMDGPU_MACH_AMDGCN_GFX601`` 0x021 ``gfx601``
667 ``EF_AMDGPU_MACH_AMDGCN_GFX700`` 0x022 ``gfx700``
668 ``EF_AMDGPU_MACH_AMDGCN_GFX701`` 0x023 ``gfx701``
669 ``EF_AMDGPU_MACH_AMDGCN_GFX702`` 0x024 ``gfx702``
670 ``EF_AMDGPU_MACH_AMDGCN_GFX703`` 0x025 ``gfx703``
671 ``EF_AMDGPU_MACH_AMDGCN_GFX704`` 0x026 ``gfx704``
672 *reserved* 0x027 Reserved.
673 ``EF_AMDGPU_MACH_AMDGCN_GFX801`` 0x028 ``gfx801``
674 ``EF_AMDGPU_MACH_AMDGCN_GFX802`` 0x029 ``gfx802``
675 ``EF_AMDGPU_MACH_AMDGCN_GFX803`` 0x02a ``gfx803``
676 ``EF_AMDGPU_MACH_AMDGCN_GFX810`` 0x02b ``gfx810``
677 ``EF_AMDGPU_MACH_AMDGCN_GFX900`` 0x02c ``gfx900``
678 ``EF_AMDGPU_MACH_AMDGCN_GFX902`` 0x02d ``gfx902``
679 ``EF_AMDGPU_MACH_AMDGCN_GFX904`` 0x02e ``gfx904``
680 ``EF_AMDGPU_MACH_AMDGCN_GFX906`` 0x02f ``gfx906``
681 ``EF_AMDGPU_MACH_AMDGCN_GFX908`` 0x030 ``gfx908``
682 ``EF_AMDGPU_MACH_AMDGCN_GFX909`` 0x031 ``gfx909``
683 *reserved* 0x032 Reserved.
684 ``EF_AMDGPU_MACH_AMDGCN_GFX1010`` 0x033 ``gfx1010``
685 ``EF_AMDGPU_MACH_AMDGCN_GFX1011`` 0x034 ``gfx1011``
686 ``EF_AMDGPU_MACH_AMDGCN_GFX1012`` 0x035 ``gfx1012``
687 ================================= ========== =============================
689 Sections
690 --------
692 An AMDGPU target ELF code object has the standard ELF sections which include:
694 .. table:: AMDGPU ELF Sections
695 :name: amdgpu-elf-sections-table
697 ================== ================ =================================
698 Name Type Attributes
699 ================== ================ =================================
700 ``.bss`` ``SHT_NOBITS`` ``SHF_ALLOC`` + ``SHF_WRITE``
701 ``.data`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_WRITE``
702 ``.debug_``\ *\** ``SHT_PROGBITS`` *none*
703 ``.dynamic`` ``SHT_DYNAMIC`` ``SHF_ALLOC``
704 ``.dynstr`` ``SHT_PROGBITS`` ``SHF_ALLOC``
705 ``.dynsym`` ``SHT_PROGBITS`` ``SHF_ALLOC``
706 ``.got`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_WRITE``
707 ``.hash`` ``SHT_HASH`` ``SHF_ALLOC``
708 ``.note`` ``SHT_NOTE`` *none*
709 ``.rela``\ *name* ``SHT_RELA`` *none*
710 ``.rela.dyn`` ``SHT_RELA`` *none*
711 ``.rodata`` ``SHT_PROGBITS`` ``SHF_ALLOC``
712 ``.shstrtab`` ``SHT_STRTAB`` *none*
713 ``.strtab`` ``SHT_STRTAB`` *none*
714 ``.symtab`` ``SHT_SYMTAB`` *none*
715 ``.text`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_EXECINSTR``
716 ================== ================ =================================
718 These sections have their standard meanings (see [ELF]_) and are only generated
719 if needed.
721 ``.debug``\ *\**
722 The standard DWARF sections. See :ref:`amdgpu-dwarf` for information on the
723 DWARF produced by the AMDGPU backend.
725 ``.dynamic``, ``.dynstr``, ``.dynsym``, ``.hash``
726 The standard sections used by a dynamic loader.
728 ``.note``
729 See :ref:`amdgpu-note-records` for the note records supported by the AMDGPU
730 backend.
732 ``.rela``\ *name*, ``.rela.dyn``
733 For relocatable code objects, *name* is the name of the section that the
734 relocation records apply. For example, ``.rela.text`` is the section name for
735 relocation records associated with the ``.text`` section.
737 For linked shared code objects, ``.rela.dyn`` contains all the relocation
738 records from each of the relocatable code object's ``.rela``\ *name* sections.
740 See :ref:`amdgpu-relocation-records` for the relocation records supported by
741 the AMDGPU backend.
743 ``.text``
744 The executable machine code for the kernels and functions they call. Generated
745 as position independent code. See :ref:`amdgpu-code-conventions` for
746 information on conventions used in the isa generation.
748 .. _amdgpu-note-records:
750 Note Records
751 ------------
753 The AMDGPU backend code object contains ELF note records in the ``.note``
754 section. The set of generated notes and their semantics depend on the code
755 object version; see :ref:`amdgpu-note-records-v2` and
756 :ref:`amdgpu-note-records-v3`.
758 As required by ``ELFCLASS32`` and ``ELFCLASS64``, minimal zero byte padding
759 must be generated after the ``name`` field to ensure the ``desc`` field is 4
760 byte aligned. In addition, minimal zero byte padding must be generated to
761 ensure the ``desc`` field size is a multiple of 4 bytes. The ``sh_addralign``
762 field of the ``.note`` section must be at least 4 to indicate at least 8 byte
763 alignment.
765 .. _amdgpu-note-records-v2:
767 Code Object V2 Note Records (-mattr=-code-object-v3)
768 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
770 .. warning:: Code Object V2 is not the default code object version emitted by
771 this version of LLVM. For a description of the notes generated with the
772 default configuration (Code Object V3) see :ref:`amdgpu-note-records-v3`.
774 The AMDGPU backend code object uses the following ELF note record in the
775 ``.note`` section when compiling for Code Object V2 (-mattr=-code-object-v3).
777 Additional note records may be present, but any which are not documented here
778 are deprecated and should not be used.
780 .. table:: AMDGPU Code Object V2 ELF Note Records
781 :name: amdgpu-elf-note-records-table-v2
783 ===== ============================== ======================================
784 Name Type Description
785 ===== ============================== ======================================
786 "AMD" ``NT_AMD_AMDGPU_HSA_METADATA`` <metadata null terminated string>
787 ===== ============================== ======================================
789 ..
791 .. table:: AMDGPU Code Object V2 ELF Note Record Enumeration Values
792 :name: amdgpu-elf-note-record-enumeration-values-table-v2
794 ============================== =====
795 Name Value
796 ============================== =====
797 *reserved* 0-9
798 ``NT_AMD_AMDGPU_HSA_METADATA`` 10
799 *reserved* 11
800 ============================== =====
802 ``NT_AMD_AMDGPU_HSA_METADATA``
803 Specifies extensible metadata associated with the code objects executed on HSA
804 [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. It is required when
805 the target triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`). See
806 :ref:`amdgpu-amdhsa-code-object-metadata-v2` for the syntax of the code
807 object metadata string.
809 .. _amdgpu-note-records-v3:
811 Code Object V3 Note Records (-mattr=+code-object-v3)
812 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
814 The AMDGPU backend code object uses the following ELF note record in the
815 ``.note`` section when compiling for Code Object V3 (-mattr=+code-object-v3).
817 Additional note records may be present, but any which are not documented here
818 are deprecated and should not be used.
820 .. table:: AMDGPU Code Object V3 ELF Note Records
821 :name: amdgpu-elf-note-records-table-v3
823 ======== ============================== ======================================
824 Name Type Description
825 ======== ============================== ======================================
826 "AMDGPU" ``NT_AMDGPU_METADATA`` Metadata in Message Pack [MsgPack]_
827 binary format.
828 ======== ============================== ======================================
830 ..
832 .. table:: AMDGPU Code Object V3 ELF Note Record Enumeration Values
833 :name: amdgpu-elf-note-record-enumeration-values-table-v3
835 ============================== =====
836 Name Value
837 ============================== =====
838 *reserved* 0-31
839 ``NT_AMDGPU_METADATA`` 32
840 ============================== =====
842 ``NT_AMDGPU_METADATA``
843 Specifies extensible metadata associated with an AMDGPU code
844 object. It is encoded as a map in the Message Pack [MsgPack]_ binary
845 data format. See :ref:`amdgpu-amdhsa-code-object-metadata-v3` for the
846 map keys defined for the ``amdhsa`` OS.
848 .. _amdgpu-symbols:
850 Symbols
851 -------
853 Symbols include the following:
855 .. table:: AMDGPU ELF Symbols
856 :name: amdgpu-elf-symbols-table
858 ===================== ================== ================ ==================
859 Name Type Section Description
860 ===================== ================== ================ ==================
861 *link-name* ``STT_OBJECT`` - ``.data`` Global variable
862 - ``.rodata``
863 - ``.bss``
864 *link-name*\ ``.kd`` ``STT_OBJECT`` - ``.rodata`` Kernel descriptor
865 *link-name* ``STT_FUNC`` - ``.text`` Kernel entry point
866 *link-name* ``STT_OBJECT`` - SHN_AMDGPU_LDS Global variable in LDS
867 ===================== ================== ================ ==================
869 Global variable
870 Global variables both used and defined by the compilation unit.
872 If the symbol is defined in the compilation unit then it is allocated in the
873 appropriate section according to if it has initialized data or is readonly.
875 If the symbol is external then its section is ``STN_UNDEF`` and the loader
876 will resolve relocations using the definition provided by another code object
877 or explicitly defined by the runtime.
879 If the symbol resides in local/group memory (LDS) then its section is the
880 special processor-specific section name ``SHN_AMDGPU_LDS``, and the
881 ``st_value`` field describes alignment requirements as it does for common
882 symbols.
884 .. TODO
885 Add description of linked shared object symbols. Seems undefined symbols
886 are marked as STT_NOTYPE.
888 Kernel descriptor
889 Every HSA kernel has an associated kernel descriptor. It is the address of the
890 kernel descriptor that is used in the AQL dispatch packet used to invoke the
891 kernel, not the kernel entry point. The layout of the HSA kernel descriptor is
892 defined in :ref:`amdgpu-amdhsa-kernel-descriptor`.
894 Kernel entry point
895 Every HSA kernel also has a symbol for its machine code entry point.
897 .. _amdgpu-relocation-records:
899 Relocation Records
900 ------------------
902 AMDGPU backend generates ``Elf64_Rela`` relocation records. Supported
903 relocatable fields are:
905 ``word32``
906 This specifies a 32-bit field occupying 4 bytes with arbitrary byte
907 alignment. These values use the same byte order as other word values in the
908 AMD GPU architecture.
910 ``word64``
911 This specifies a 64-bit field occupying 8 bytes with arbitrary byte
912 alignment. These values use the same byte order as other word values in the
913 AMD GPU architecture.
915 Following notations are used for specifying relocation calculations:
917 **A**
918 Represents the addend used to compute the value of the relocatable field.
920 **G**
921 Represents the offset into the global offset table at which the relocation
922 entry's symbol will reside during execution.
924 **GOT**
925 Represents the address of the global offset table.
927 **P**
928 Represents the place (section offset for ``et_rel`` or address for ``et_dyn``)
929 of the storage unit being relocated (computed using ``r_offset``).
931 **S**
932 Represents the value of the symbol whose index resides in the relocation
933 entry. Relocations not using this must specify a symbol index of ``STN_UNDEF``.
935 **B**
936 Represents the base address of a loaded executable or shared object which is
937 the difference between the ELF address and the actual load address. Relocations
938 using this are only valid in executable or shared objects.
940 The following relocation types are supported:
942 .. table:: AMDGPU ELF Relocation Records
943 :name: amdgpu-elf-relocation-records-table
945 ========================== ======= ===== ========== ==============================
946 Relocation Type Kind Value Field Calculation
947 ========================== ======= ===== ========== ==============================
948 ``R_AMDGPU_NONE`` 0 *none* *none*
949 ``R_AMDGPU_ABS32_LO`` Static, 1 ``word32`` (S + A) & 0xFFFFFFFF
950 Dynamic
951 ``R_AMDGPU_ABS32_HI`` Static, 2 ``word32`` (S + A) >> 32
952 Dynamic
953 ``R_AMDGPU_ABS64`` Static, 3 ``word64`` S + A
954 Dynamic
955 ``R_AMDGPU_REL32`` Static 4 ``word32`` S + A - P
956 ``R_AMDGPU_REL64`` Static 5 ``word64`` S + A - P
957 ``R_AMDGPU_ABS32`` Static, 6 ``word32`` S + A
958 Dynamic
959 ``R_AMDGPU_GOTPCREL`` Static 7 ``word32`` G + GOT + A - P
960 ``R_AMDGPU_GOTPCREL32_LO`` Static 8 ``word32`` (G + GOT + A - P) & 0xFFFFFFFF
961 ``R_AMDGPU_GOTPCREL32_HI`` Static 9 ``word32`` (G + GOT + A - P) >> 32
962 ``R_AMDGPU_REL32_LO`` Static 10 ``word32`` (S + A - P) & 0xFFFFFFFF
963 ``R_AMDGPU_REL32_HI`` Static 11 ``word32`` (S + A - P) >> 32
964 *reserved* 12
965 ``R_AMDGPU_RELATIVE64`` Dynamic 13 ``word64`` B + A
966 ========================== ======= ===== ========== ==============================
968 ``R_AMDGPU_ABS32_LO`` and ``R_AMDGPU_ABS32_HI`` are only supported by
969 the ``mesa3d`` OS, which does not support ``R_AMDGPU_ABS64``.
971 There is no current OS loader support for 32 bit programs and so
972 ``R_AMDGPU_ABS32`` is not used.
974 .. _amdgpu-dwarf:
976 DWARF
977 -----
979 Standard DWARF [DWARF]_ Version 5 sections can be generated. These contain
980 information that maps the code object executable code and data to the source
981 language constructs. It can be used by tools such as debuggers and profilers.
983 Address Space Mapping
984 ~~~~~~~~~~~~~~~~~~~~~
986 The following address space mapping is used:
988 .. table:: AMDGPU DWARF Address Space Mapping
989 :name: amdgpu-dwarf-address-space-mapping-table
991 =================== =================
992 DWARF Address Space Memory Space
993 =================== =================
994 1 Private (Scratch)
995 2 Local (group/LDS)
996 *omitted* Global
997 *omitted* Constant
998 *omitted* Generic (Flat)
999 *not supported* Region (GDS)
1000 =================== =================
1002 See :ref:`amdgpu-address-spaces` for information on the memory space terminology
1003 used in the table.
1005 An ``address_class`` attribute is generated on pointer type DIEs to specify the
1006 DWARF address space of the value of the pointer when it is in the *private* or
1007 *local* address space. Otherwise the attribute is omitted.
1009 An ``XDEREF`` operation is generated in location list expressions for variables
1010 that are allocated in the *private* and *local* address space. Otherwise no
1011 ``XDREF`` is omitted.
1013 Register Mapping
1014 ~~~~~~~~~~~~~~~~
1016 *This section is WIP.*
1018 .. TODO
1019 Define DWARF register enumeration.
1021 If want to present a wavefront state then should expose vector registers as
1022 64 wide (rather than per work-item view that LLVM uses). Either as separate
1023 registers, or a 64x4 byte single register. In either case use a new LANE op
1024 (akin to XDREF) to select the current lane usage in a location
1025 expression. This would also allow scalar register spilling to vector register
1026 lanes to be expressed (currently no debug information is being generated for
1027 spilling). If choose a wide single register approach then use LANE in
1028 conjunction with PIECE operation to select the dword part of the register for
1029 the current lane. If the separate register approach then use LANE to select
1030 the register.
1032 Source Text
1033 ~~~~~~~~~~~
1035 Source text for online-compiled programs (e.g. those compiled by the OpenCL
1036 runtime) may be embedded into the DWARF v5 line table using the ``clang
1037 -gembed-source`` option, described in table :ref:`amdgpu-debug-options`.
1039 For example:
1041 ``-gembed-source``
1042 Enable the embedded source DWARF v5 extension.
1043 ``-gno-embed-source``
1044 Disable the embedded source DWARF v5 extension.
1046 .. table:: AMDGPU Debug Options
1047 :name: amdgpu-debug-options
1049 ==================== ==================================================
1050 Debug Flag Description
1051 ==================== ==================================================
1052 -g[no-]embed-source Enable/disable embedding source text in DWARF
1053 debug sections. Useful for environments where
1054 source cannot be written to disk, such as
1055 when performing online compilation.
1056 ==================== ==================================================
1058 This option enables one extended content types in the DWARF v5 Line Number
1059 Program Header, which is used to encode embedded source.
1061 .. table:: AMDGPU DWARF Line Number Program Header Extended Content Types
1062 :name: amdgpu-dwarf-extended-content-types
1064 ============================ ======================
1065 Content Type Form
1066 ============================ ======================
1067 ``DW_LNCT_LLVM_source`` ``DW_FORM_line_strp``
1068 ============================ ======================
1070 The source field will contain the UTF-8 encoded, null-terminated source text
1071 with ``'\n'`` line endings. When the source field is present, consumers can use
1072 the embedded source instead of attempting to discover the source on disk. When
1073 the source field is absent, consumers can access the file to get the source
1074 text.
1076 The above content type appears in the ``file_name_entry_format`` field of the
1077 line table prologue, and its corresponding value appear in the ``file_names``
1078 field. The current encoding of the content type is documented in table
1079 :ref:`amdgpu-dwarf-extended-content-types-encoding`
1081 .. table:: AMDGPU DWARF Line Number Program Header Extended Content Types Encoding
1082 :name: amdgpu-dwarf-extended-content-types-encoding
1084 ============================ ====================
1085 Content Type Value
1086 ============================ ====================
1087 ``DW_LNCT_LLVM_source`` 0x2001
1088 ============================ ====================
1090 .. _amdgpu-code-conventions:
1092 Code Conventions
1093 ================
1095 This section provides code conventions used for each supported target triple OS
1096 (see :ref:`amdgpu-target-triples`).
1098 AMDHSA
1099 ------
1101 This section provides code conventions used when the target triple OS is
1102 ``amdhsa`` (see :ref:`amdgpu-target-triples`).
1104 .. _amdgpu-amdhsa-code-object-target-identification:
1106 Code Object Target Identification
1107 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
1109 The AMDHSA OS uses the following syntax to specify the code object
1110 target as a single string:
1112 ``<Architecture>-<Vendor>-<OS>-<Environment>-<Processor><Target Features>``
1114 Where:
1116 - ``<Architecture>``, ``<Vendor>``, ``<OS>`` and ``<Environment>``
1117 are the same as the *Target Triple* (see
1118 :ref:`amdgpu-target-triples`).
1120 - ``<Processor>`` is the same as the *Processor* (see
1121 :ref:`amdgpu-processors`).
1123 - ``<Target Features>`` is a list of the enabled *Target Features*
1124 (see :ref:`amdgpu-target-features`), each prefixed by a plus, that
1125 apply to *Processor*. The list must be in the same order as listed
1126 in the table :ref:`amdgpu-target-feature-table`. Note that *Target
1127 Features* must be included in the list if they are enabled even if
1128 that is the default for *Processor*.
1130 For example:
1132 ``"amdgcn-amd-amdhsa--gfx902+xnack"``
1134 .. _amdgpu-amdhsa-code-object-metadata:
1136 Code Object Metadata
1137 ~~~~~~~~~~~~~~~~~~~~
1139 The code object metadata specifies extensible metadata associated with the code
1140 objects executed on HSA [HSA]_ compatible runtimes such as AMD's ROCm
1141 [AMD-ROCm]_. The encoding and semantics of this metadata depends on the code
1142 object version; see :ref:`amdgpu-amdhsa-code-object-metadata-v2` and
1143 :ref:`amdgpu-amdhsa-code-object-metadata-v3`.
1145 Code object metadata is specified in a note record (see
1146 :ref:`amdgpu-note-records`) and is required when the target triple OS is
1147 ``amdhsa`` (see :ref:`amdgpu-target-triples`). It must contain the minimum
1148 information necessary to support the ROCM kernel queries. For example, the
1149 segment sizes needed in a dispatch packet. In addition, a high level language
1150 runtime may require other information to be included. For example, the AMD
1151 OpenCL runtime records kernel argument information.
1153 .. _amdgpu-amdhsa-code-object-metadata-v2:
1155 Code Object V2 Metadata (-mattr=-code-object-v3)
1156 ++++++++++++++++++++++++++++++++++++++++++++++++
1158 .. warning:: Code Object V2 is not the default code object version emitted by
1159 this version of LLVM. For a description of the metadata generated with the
1160 default configuration (Code Object V3) see
1161 :ref:`amdgpu-amdhsa-code-object-metadata-v3`.
1163 Code object V2 metadata is specified by the ``NT_AMD_AMDGPU_METADATA`` note
1164 record (see :ref:`amdgpu-note-records-v2`).
1166 The metadata is specified as a YAML formatted string (see [YAML]_ and
1167 :doc:`YamlIO`).
1169 .. TODO
1170 Is the string null terminated? It probably should not if YAML allows it to
1171 contain null characters, otherwise it should be.
1173 The metadata is represented as a single YAML document comprised of the mapping
1174 defined in table :ref:`amdgpu-amdhsa-code-object-metadata-map-table-v2` and
1175 referenced tables.
1177 For boolean values, the string values of ``false`` and ``true`` are used for
1178 false and true respectively.
1180 Additional information can be added to the mappings. To avoid conflicts, any
1181 non-AMD key names should be prefixed by "*vendor-name*.".
1183 .. table:: AMDHSA Code Object V2 Metadata Map
1184 :name: amdgpu-amdhsa-code-object-metadata-map-table-v2
1186 ========== ============== ========= =======================================
1187 String Key Value Type Required? Description
1188 ========== ============== ========= =======================================
1189 "Version" sequence of Required - The first integer is the major
1190 2 integers version. Currently 1.
1191 - The second integer is the minor
1192 version. Currently 0.
1193 "Printf" sequence of Each string is encoded information
1194 strings about a printf function call. The
1195 encoded information is organized as
1196 fields separated by colon (':'):
1198 ``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``
1200 where:
1202 ``ID``
1203 A 32 bit integer as a unique id for
1204 each printf function call
1206 ``N``
1207 A 32 bit integer equal to the number
1208 of arguments of printf function call
1209 minus 1
1211 ``S[i]`` (where i = 0, 1, ... , N-1)
1212 32 bit integers for the size in bytes
1213 of the i-th FormatString argument of
1214 the printf function call
1216 FormatString
1217 The format string passed to the
1218 printf function call.
1219 "Kernels" sequence of Required Sequence of the mappings for each
1220 mapping kernel in the code object. See
1221 :ref:`amdgpu-amdhsa-code-object-kernel-metadata-map-table-v2`
1222 for the definition of the mapping.
1223 ========== ============== ========= =======================================
1225 ..
1227 .. table:: AMDHSA Code Object V2 Kernel Metadata Map
1228 :name: amdgpu-amdhsa-code-object-kernel-metadata-map-table-v2
1230 ================= ============== ========= ================================
1231 String Key Value Type Required? Description
1232 ================= ============== ========= ================================
1233 "Name" string Required Source name of the kernel.
1234 "SymbolName" string Required Name of the kernel
1235 descriptor ELF symbol.
1236 "Language" string Source language of the kernel.
1237 Values include:
1239 - "OpenCL C"
1240 - "OpenCL C++"
1241 - "HCC"
1242 - "OpenMP"
1244 "LanguageVersion" sequence of - The first integer is the major
1245 2 integers version.
1246 - The second integer is the
1247 minor version.
1248 "Attrs" mapping Mapping of kernel attributes.
1249 See
1250 :ref:`amdgpu-amdhsa-code-object-kernel-attribute-metadata-map-table-v2`
1251 for the mapping definition.
1252 "Args" sequence of Sequence of mappings of the
1253 mapping kernel arguments. See
1254 :ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v2`
1255 for the definition of the mapping.
1256 "CodeProps" mapping Mapping of properties related to
1257 the kernel code. See
1258 :ref:`amdgpu-amdhsa-code-object-kernel-code-properties-metadata-map-table-v2`
1259 for the mapping definition.
1260 ================= ============== ========= ================================
1262 ..
1264 .. table:: AMDHSA Code Object V2 Kernel Attribute Metadata Map
1265 :name: amdgpu-amdhsa-code-object-kernel-attribute-metadata-map-table-v2
1267 =================== ============== ========= ==============================
1268 String Key Value Type Required? Description
1269 =================== ============== ========= ==============================
1270 "ReqdWorkGroupSize" sequence of If not 0, 0, 0 then all values
1271 3 integers must be >=1 and the dispatch
1272 work-group size X, Y, Z must
1273 correspond to the specified
1274 values. Defaults to 0, 0, 0.
1276 Corresponds to the OpenCL
1277 ``reqd_work_group_size``
1278 attribute.
1279 "WorkGroupSizeHint" sequence of The dispatch work-group size
1280 3 integers X, Y, Z is likely to be the
1281 specified values.
1283 Corresponds to the OpenCL
1284 ``work_group_size_hint``
1285 attribute.
1286 "VecTypeHint" string The name of a scalar or vector
1287 type.
1289 Corresponds to the OpenCL
1290 ``vec_type_hint`` attribute.
1292 "RuntimeHandle" string The external symbol name
1293 associated with a kernel.
1294 OpenCL runtime allocates a
1295 global buffer for the symbol
1296 and saves the kernel's address
1297 to it, which is used for
1298 device side enqueueing. Only
1299 available for device side
1300 enqueued kernels.
1301 =================== ============== ========= ==============================
1303 ..
1305 .. table:: AMDHSA Code Object V2 Kernel Argument Metadata Map
1306 :name: amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v2
1308 ================= ============== ========= ================================
1309 String Key Value Type Required? Description
1310 ================= ============== ========= ================================
1311 "Name" string Kernel argument name.
1312 "TypeName" string Kernel argument type name.
1313 "Size" integer Required Kernel argument size in bytes.
1314 "Align" integer Required Kernel argument alignment in
1315 bytes. Must be a power of two.
1316 "ValueKind" string Required Kernel argument kind that
1317 specifies how to set up the
1318 corresponding argument.
1319 Values include:
1321 "ByValue"
1322 The argument is copied
1323 directly into the kernarg.
1325 "GlobalBuffer"
1326 A global address space pointer
1327 to the buffer data is passed
1328 in the kernarg.
1330 "DynamicSharedPointer"
1331 A group address space pointer
1332 to dynamically allocated LDS
1333 is passed in the kernarg.
1335 "Sampler"
1336 A global address space
1337 pointer to a S# is passed in
1338 the kernarg.
1340 "Image"
1341 A global address space
1342 pointer to a T# is passed in
1343 the kernarg.
1345 "Pipe"
1346 A global address space pointer
1347 to an OpenCL pipe is passed in
1348 the kernarg.
1350 "Queue"
1351 A global address space pointer
1352 to an OpenCL device enqueue
1353 queue is passed in the
1354 kernarg.
1356 "HiddenGlobalOffsetX"
1357 The OpenCL grid dispatch
1358 global offset for the X
1359 dimension is passed in the
1360 kernarg.
1362 "HiddenGlobalOffsetY"
1363 The OpenCL grid dispatch
1364 global offset for the Y
1365 dimension is passed in the
1366 kernarg.
1368 "HiddenGlobalOffsetZ"
1369 The OpenCL grid dispatch
1370 global offset for the Z
1371 dimension is passed in the
1372 kernarg.
1374 "HiddenNone"
1375 An argument that is not used
1376 by the kernel. Space needs to
1377 be left for it, but it does
1378 not need to be set up.
1380 "HiddenPrintfBuffer"
1381 A global address space pointer
1382 to the runtime printf buffer
1383 is passed in kernarg.
1385 "HiddenDefaultQueue"
1386 A global address space pointer
1387 to the OpenCL device enqueue
1388 queue that should be used by
1389 the kernel by default is
1390 passed in the kernarg.
1392 "HiddenCompletionAction"
1393 A global address space pointer
1394 to help link enqueued kernels into
1395 the ancestor tree for determining
1396 when the parent kernel has finished.
1398 "HiddenMultiGridSyncArg"
1399 A global address space pointer for
1400 multi-grid synchronization is
1401 passed in the kernarg.
1403 "ValueType" string Required Kernel argument value type. Only
1404 present if "ValueKind" is
1405 "ByValue". For vector data
1406 types, the value is for the
1407 element type. Values include:
1409 - "Struct"
1410 - "I8"
1411 - "U8"
1412 - "I16"
1413 - "U16"
1414 - "F16"
1415 - "I32"
1416 - "U32"
1417 - "F32"
1418 - "I64"
1419 - "U64"
1420 - "F64"
1422 .. TODO
1423 How can it be determined if a
1424 vector type, and what size
1425 vector?
1426 "PointeeAlign" integer Alignment in bytes of pointee
1427 type for pointer type kernel
1428 argument. Must be a power
1429 of 2. Only present if
1430 "ValueKind" is
1431 "DynamicSharedPointer".
1432 "AddrSpaceQual" string Kernel argument address space
1433 qualifier. Only present if
1434 "ValueKind" is "GlobalBuffer" or
1435 "DynamicSharedPointer". Values
1436 are:
1438 - "Private"
1439 - "Global"
1440 - "Constant"
1441 - "Local"
1442 - "Generic"
1443 - "Region"
1445 .. TODO
1446 Is GlobalBuffer only Global
1447 or Constant? Is
1448 DynamicSharedPointer always
1449 Local? Can HCC allow Generic?
1450 How can Private or Region
1451 ever happen?
1452 "AccQual" string Kernel argument access
1453 qualifier. Only present if
1454 "ValueKind" is "Image" or
1455 "Pipe". Values
1456 are:
1458 - "ReadOnly"
1459 - "WriteOnly"
1460 - "ReadWrite"
1462 .. TODO
1463 Does this apply to
1464 GlobalBuffer?
1465 "ActualAccQual" string The actual memory accesses
1466 performed by the kernel on the
1467 kernel argument. Only present if
1468 "ValueKind" is "GlobalBuffer",
1469 "Image", or "Pipe". This may be
1470 more restrictive than indicated
1471 by "AccQual" to reflect what the
1472 kernel actual does. If not
1473 present then the runtime must
1474 assume what is implied by
1475 "AccQual" and "IsConst". Values
1476 are:
1478 - "ReadOnly"
1479 - "WriteOnly"
1480 - "ReadWrite"
1482 "IsConst" boolean Indicates if the kernel argument
1483 is const qualified. Only present
1484 if "ValueKind" is
1485 "GlobalBuffer".
1487 "IsRestrict" boolean Indicates if the kernel argument
1488 is restrict qualified. Only
1489 present if "ValueKind" is
1490 "GlobalBuffer".
1492 "IsVolatile" boolean Indicates if the kernel argument
1493 is volatile qualified. Only
1494 present if "ValueKind" is
1495 "GlobalBuffer".
1497 "IsPipe" boolean Indicates if the kernel argument
1498 is pipe qualified. Only present
1499 if "ValueKind" is "Pipe".
1501 .. TODO
1502 Can GlobalBuffer be pipe
1503 qualified?
1504 ================= ============== ========= ================================
1506 ..
1508 .. table:: AMDHSA Code Object V2 Kernel Code Properties Metadata Map
1509 :name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-map-table-v2
1511 ============================ ============== ========= =====================
1512 String Key Value Type Required? Description
1513 ============================ ============== ========= =====================
1514 "KernargSegmentSize" integer Required The size in bytes of
1515 the kernarg segment
1516 that holds the values
1517 of the arguments to
1518 the kernel.
1519 "GroupSegmentFixedSize" integer Required The amount of group
1520 segment memory
1521 required by a
1522 work-group in
1523 bytes. This does not
1524 include any
1525 dynamically allocated
1526 group segment memory
1527 that may be added
1528 when the kernel is
1529 dispatched.
1530 "PrivateSegmentFixedSize" integer Required The amount of fixed
1531 private address space
1532 memory required for a
1533 work-item in
1534 bytes. If the kernel
1535 uses a dynamic call
1536 stack then additional
1537 space must be added
1538 to this value for the
1539 call stack.
1540 "KernargSegmentAlign" integer Required The maximum byte
1541 alignment of
1542 arguments in the
1543 kernarg segment. Must
1544 be a power of 2.
1545 "WavefrontSize" integer Required Wavefront size. Must
1546 be a power of 2.
1547 "NumSGPRs" integer Required Number of scalar
1548 registers used by a
1549 wavefront for
1550 GFX6-GFX10. This
1551 includes the special
1552 SGPRs for VCC, Flat
1553 Scratch (GFX7-GFX10)
1554 and XNACK (for
1555 GFX8-GFX10). It does
1556 not include the 16
1557 SGPR added if a trap
1558 handler is
1559 enabled. It is not
1560 rounded up to the
1561 allocation
1562 granularity.
1563 "NumVGPRs" integer Required Number of vector
1564 registers used by
1565 each work-item for
1566 GFX6-GFX10
1567 "MaxFlatWorkGroupSize" integer Required Maximum flat
1568 work-group size
1569 supported by the
1570 kernel in work-items.
1571 Must be >=1 and
1572 consistent with
1573 ReqdWorkGroupSize if
1574 not 0, 0, 0.
1575 "NumSpilledSGPRs" integer Number of stores from
1576 a scalar register to
1577 a register allocator
1578 created spill
1579 location.
1580 "NumSpilledVGPRs" integer Number of stores from
1581 a vector register to
1582 a register allocator
1583 created spill
1584 location.
1585 ============================ ============== ========= =====================
1587 .. _amdgpu-amdhsa-code-object-metadata-v3:
1589 Code Object V3 Metadata (-mattr=+code-object-v3)
1590 ++++++++++++++++++++++++++++++++++++++++++++++++
1592 Code object V3 metadata is specified by the ``NT_AMDGPU_METADATA`` note record
1593 (see :ref:`amdgpu-note-records-v3`).
1595 The metadata is represented as Message Pack formatted binary data (see
1596 [MsgPack]_). The top level is a Message Pack map that includes the
1597 keys defined in table
1598 :ref:`amdgpu-amdhsa-code-object-metadata-map-table-v3` and referenced
1599 tables.
1601 Additional information can be added to the maps. To avoid conflicts,
1602 any key names should be prefixed by "*vendor-name*." where
1603 ``vendor-name`` can be the the name of the vendor and specific vendor
1604 tool that generates the information. The prefix is abbreviated to
1605 simply "." when it appears within a map that has been added by the
1606 same *vendor-name*.
1608 .. table:: AMDHSA Code Object V3 Metadata Map
1609 :name: amdgpu-amdhsa-code-object-metadata-map-table-v3
1611 ================= ============== ========= =======================================
1612 String Key Value Type Required? Description
1613 ================= ============== ========= =======================================
1614 "amdhsa.version" sequence of Required - The first integer is the major
1615 2 integers version. Currently 1.
1616 - The second integer is the minor
1617 version. Currently 0.
1618 "amdhsa.printf" sequence of Each string is encoded information
1619 strings about a printf function call. The
1620 encoded information is organized as
1621 fields separated by colon (':'):
1623 ``ID:N:S[0]:S[1]:...:S[N-1]:FormatString``
1625 where:
1627 ``ID``
1628 A 32 bit integer as a unique id for
1629 each printf function call
1631 ``N``
1632 A 32 bit integer equal to the number
1633 of arguments of printf function call
1634 minus 1
1636 ``S[i]`` (where i = 0, 1, ... , N-1)
1637 32 bit integers for the size in bytes
1638 of the i-th FormatString argument of
1639 the printf function call
1641 FormatString
1642 The format string passed to the
1643 printf function call.
1644 "amdhsa.kernels" sequence of Required Sequence of the maps for each
1645 map kernel in the code object. See
1646 :ref:`amdgpu-amdhsa-code-object-kernel-metadata-map-table-v3`
1647 for the definition of the keys included
1648 in that map.
1649 ================= ============== ========= =======================================
1651 ..
1653 .. table:: AMDHSA Code Object V3 Kernel Metadata Map
1654 :name: amdgpu-amdhsa-code-object-kernel-metadata-map-table-v3
1656 =================================== ============== ========= ================================
1657 String Key Value Type Required? Description
1658 =================================== ============== ========= ================================
1659 ".name" string Required Source name of the kernel.
1660 ".symbol" string Required Name of the kernel
1661 descriptor ELF symbol.
1662 ".language" string Source language of the kernel.
1663 Values include:
1665 - "OpenCL C"
1666 - "OpenCL C++"
1667 - "HCC"
1668 - "HIP"
1669 - "OpenMP"
1670 - "Assembler"
1672 ".language_version" sequence of - The first integer is the major
1673 2 integers version.
1674 - The second integer is the
1675 minor version.
1676 ".args" sequence of Sequence of maps of the
1677 map kernel arguments. See
1678 :ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3`
1679 for the definition of the keys
1680 included in that map.
1681 ".reqd_workgroup_size" sequence of If not 0, 0, 0 then all values
1682 3 integers must be >=1 and the dispatch
1683 work-group size X, Y, Z must
1684 correspond to the specified
1685 values. Defaults to 0, 0, 0.
1687 Corresponds to the OpenCL
1688 ``reqd_work_group_size``
1689 attribute.
1690 ".workgroup_size_hint" sequence of The dispatch work-group size
1691 3 integers X, Y, Z is likely to be the
1692 specified values.
1694 Corresponds to the OpenCL
1695 ``work_group_size_hint``
1696 attribute.
1697 ".vec_type_hint" string The name of a scalar or vector
1698 type.
1700 Corresponds to the OpenCL
1701 ``vec_type_hint`` attribute.
1703 ".device_enqueue_symbol" string The external symbol name
1704 associated with a kernel.
1705 OpenCL runtime allocates a
1706 global buffer for the symbol
1707 and saves the kernel's address
1708 to it, which is used for
1709 device side enqueueing. Only
1710 available for device side
1711 enqueued kernels.
1712 ".kernarg_segment_size" integer Required The size in bytes of
1713 the kernarg segment
1714 that holds the values
1715 of the arguments to
1716 the kernel.
1717 ".group_segment_fixed_size" integer Required The amount of group
1718 segment memory
1719 required by a
1720 work-group in
1721 bytes. This does not
1722 include any
1723 dynamically allocated
1724 group segment memory
1725 that may be added
1726 when the kernel is
1727 dispatched.
1728 ".private_segment_fixed_size" integer Required The amount of fixed
1729 private address space
1730 memory required for a
1731 work-item in
1732 bytes. If the kernel
1733 uses a dynamic call
1734 stack then additional
1735 space must be added
1736 to this value for the
1737 call stack.
1738 ".kernarg_segment_align" integer Required The maximum byte
1739 alignment of
1740 arguments in the
1741 kernarg segment. Must
1742 be a power of 2.
1743 ".wavefront_size" integer Required Wavefront size. Must
1744 be a power of 2.
1745 ".sgpr_count" integer Required Number of scalar
1746 registers required by a
1747 wavefront for
1748 GFX6-GFX9. A register
1749 is required if it is
1750 used explicitly, or
1751 if a higher numbered
1752 register is used
1753 explicitly. This
1754 includes the special
1755 SGPRs for VCC, Flat
1756 Scratch (GFX7-GFX9)
1757 and XNACK (for
1758 GFX8-GFX9). It does
1759 not include the 16
1760 SGPR added if a trap
1761 handler is
1762 enabled. It is not
1763 rounded up to the
1764 allocation
1765 granularity.
1766 ".vgpr_count" integer Required Number of vector
1767 registers required by
1768 each work-item for
1769 GFX6-GFX9. A register
1770 is required if it is
1771 used explicitly, or
1772 if a higher numbered
1773 register is used
1774 explicitly.
1775 ".max_flat_workgroup_size" integer Required Maximum flat
1776 work-group size
1777 supported by the
1778 kernel in work-items.
1779 Must be >=1 and
1780 consistent with
1781 ReqdWorkGroupSize if
1782 not 0, 0, 0.
1783 ".sgpr_spill_count" integer Number of stores from
1784 a scalar register to
1785 a register allocator
1786 created spill
1787 location.
1788 ".vgpr_spill_count" integer Number of stores from
1789 a vector register to
1790 a register allocator
1791 created spill
1792 location.
1793 =================================== ============== ========= ================================
1795 ..
1797 .. table:: AMDHSA Code Object V3 Kernel Argument Metadata Map
1798 :name: amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3
1800 ====================== ============== ========= ================================
1801 String Key Value Type Required? Description
1802 ====================== ============== ========= ================================
1803 ".name" string Kernel argument name.
1804 ".type_name" string Kernel argument type name.
1805 ".size" integer Required Kernel argument size in bytes.
1806 ".offset" integer Required Kernel argument offset in
1807 bytes. The offset must be a
1808 multiple of the alignment
1809 required by the argument.
1810 ".value_kind" string Required Kernel argument kind that
1811 specifies how to set up the
1812 corresponding argument.
1813 Values include:
1815 "by_value"
1816 The argument is copied
1817 directly into the kernarg.
1819 "global_buffer"
1820 A global address space pointer
1821 to the buffer data is passed
1822 in the kernarg.
1824 "dynamic_shared_pointer"
1825 A group address space pointer
1826 to dynamically allocated LDS
1827 is passed in the kernarg.
1829 "sampler"
1830 A global address space
1831 pointer to a S# is passed in
1832 the kernarg.
1834 "image"
1835 A global address space
1836 pointer to a T# is passed in
1837 the kernarg.
1839 "pipe"
1840 A global address space pointer
1841 to an OpenCL pipe is passed in
1842 the kernarg.
1844 "queue"
1845 A global address space pointer
1846 to an OpenCL device enqueue
1847 queue is passed in the
1848 kernarg.
1850 "hidden_global_offset_x"
1851 The OpenCL grid dispatch
1852 global offset for the X
1853 dimension is passed in the
1854 kernarg.
1856 "hidden_global_offset_y"
1857 The OpenCL grid dispatch
1858 global offset for the Y
1859 dimension is passed in the
1860 kernarg.
1862 "hidden_global_offset_z"
1863 The OpenCL grid dispatch
1864 global offset for the Z
1865 dimension is passed in the
1866 kernarg.
1868 "hidden_none"
1869 An argument that is not used
1870 by the kernel. Space needs to
1871 be left for it, but it does
1872 not need to be set up.
1874 "hidden_printf_buffer"
1875 A global address space pointer
1876 to the runtime printf buffer
1877 is passed in kernarg.
1879 "hidden_default_queue"
1880 A global address space pointer
1881 to the OpenCL device enqueue
1882 queue that should be used by
1883 the kernel by default is
1884 passed in the kernarg.
1886 "hidden_completion_action"
1887 A global address space pointer
1888 to help link enqueued kernels into
1889 the ancestor tree for determining
1890 when the parent kernel has finished.
1892 "hidden_multigrid_sync_arg"
1893 A global address space pointer for
1894 multi-grid synchronization is
1895 passed in the kernarg.
1897 ".value_type" string Required Kernel argument value type. Only
1898 present if ".value_kind" is
1899 "by_value". For vector data
1900 types, the value is for the
1901 element type. Values include:
1903 - "struct"
1904 - "i8"
1905 - "u8"
1906 - "i16"
1907 - "u16"
1908 - "f16"
1909 - "i32"
1910 - "u32"
1911 - "f32"
1912 - "i64"
1913 - "u64"
1914 - "f64"
1916 .. TODO
1917 How can it be determined if a
1918 vector type, and what size
1919 vector?
1920 ".pointee_align" integer Alignment in bytes of pointee
1921 type for pointer type kernel
1922 argument. Must be a power
1923 of 2. Only present if
1924 ".value_kind" is
1925 "dynamic_shared_pointer".
1926 ".address_space" string Kernel argument address space
1927 qualifier. Only present if
1928 ".value_kind" is "global_buffer" or
1929 "dynamic_shared_pointer". Values
1930 are:
1932 - "private"
1933 - "global"
1934 - "constant"
1935 - "local"
1936 - "generic"
1937 - "region"
1939 .. TODO
1940 Is "global_buffer" only "global"
1941 or "constant"? Is
1942 "dynamic_shared_pointer" always
1943 "local"? Can HCC allow "generic"?
1944 How can "private" or "region"
1945 ever happen?
1946 ".access" string Kernel argument access
1947 qualifier. Only present if
1948 ".value_kind" is "image" or
1949 "pipe". Values
1950 are:
1952 - "read_only"
1953 - "write_only"
1954 - "read_write"
1956 .. TODO
1957 Does this apply to
1958 "global_buffer"?
1959 ".actual_access" string The actual memory accesses
1960 performed by the kernel on the
1961 kernel argument. Only present if
1962 ".value_kind" is "global_buffer",
1963 "image", or "pipe". This may be
1964 more restrictive than indicated
1965 by ".access" to reflect what the
1966 kernel actual does. If not
1967 present then the runtime must
1968 assume what is implied by
1969 ".access" and ".is_const" . Values
1970 are:
1972 - "read_only"
1973 - "write_only"
1974 - "read_write"
1976 ".is_const" boolean Indicates if the kernel argument
1977 is const qualified. Only present
1978 if ".value_kind" is
1979 "global_buffer".
1981 ".is_restrict" boolean Indicates if the kernel argument
1982 is restrict qualified. Only
1983 present if ".value_kind" is
1984 "global_buffer".
1986 ".is_volatile" boolean Indicates if the kernel argument
1987 is volatile qualified. Only
1988 present if ".value_kind" is
1989 "global_buffer".
1991 ".is_pipe" boolean Indicates if the kernel argument
1992 is pipe qualified. Only present
1993 if ".value_kind" is "pipe".
1995 .. TODO
1996 Can "global_buffer" be pipe
1997 qualified?
1998 ====================== ============== ========= ================================
2000 ..
2002 Kernel Dispatch
2003 ~~~~~~~~~~~~~~~
2005 The HSA architected queuing language (AQL) defines a user space memory interface
2006 that can be used to control the dispatch of kernels, in an agent independent
2007 way. An agent can have zero or more AQL queues created for it using the ROCm
2008 runtime, in which AQL packets (all of which are 64 bytes) can be placed. See the
2009 *HSA Platform System Architecture Specification* [HSA]_ for the AQL queue
2010 mechanics and packet layouts.
2012 The packet processor of a kernel agent is responsible for detecting and
2013 dispatching HSA kernels from the AQL queues associated with it. For AMD GPUs the
2014 packet processor is implemented by the hardware command processor (CP),
2015 asynchronous dispatch controller (ADC) and shader processor input controller
2016 (SPI).
2018 The ROCm runtime can be used to allocate an AQL queue object. It uses the kernel
2019 mode driver to initialize and register the AQL queue with CP.
2021 To dispatch a kernel the following actions are performed. This can occur in the
2022 CPU host program, or from an HSA kernel executing on a GPU.
2024 1. A pointer to an AQL queue for the kernel agent on which the kernel is to be
2025 executed is obtained.
2026 2. A pointer to the kernel descriptor (see
2027 :ref:`amdgpu-amdhsa-kernel-descriptor`) of the kernel to execute is
2028 obtained. It must be for a kernel that is contained in a code object that that
2029 was loaded by the ROCm runtime on the kernel agent with which the AQL queue is
2030 associated.
2031 3. Space is allocated for the kernel arguments using the ROCm runtime allocator
2032 for a memory region with the kernarg property for the kernel agent that will
2033 execute the kernel. It must be at least 16 byte aligned.
2034 4. Kernel argument values are assigned to the kernel argument memory
2035 allocation. The layout is defined in the *HSA Programmer's Language Reference*
2036 [HSA]_. For AMDGPU the kernel execution directly accesses the kernel argument
2037 memory in the same way constant memory is accessed. (Note that the HSA
2038 specification allows an implementation to copy the kernel argument contents to
2039 another location that is accessed by the kernel.)
2040 5. An AQL kernel dispatch packet is created on the AQL queue. The ROCm runtime
2041 api uses 64 bit atomic operations to reserve space in the AQL queue for the
2042 packet. The packet must be set up, and the final write must use an atomic
2043 store release to set the packet kind to ensure the packet contents are
2044 visible to the kernel agent. AQL defines a doorbell signal mechanism to
2045 notify the kernel agent that the AQL queue has been updated. These rules, and
2046 the layout of the AQL queue and kernel dispatch packet is defined in the *HSA
2047 System Architecture Specification* [HSA]_.
2048 6. A kernel dispatch packet includes information about the actual dispatch,
2049 such as grid and work-group size, together with information from the code
2050 object about the kernel, such as segment sizes. The ROCm runtime queries on
2051 the kernel symbol can be used to obtain the code object values which are
2052 recorded in the :ref:`amdgpu-amdhsa-code-object-metadata`.
2053 7. CP executes micro-code and is responsible for detecting and setting up the
2054 GPU to execute the wavefronts of a kernel dispatch.
2055 8. CP ensures that when the a wavefront starts executing the kernel machine
2056 code, the scalar general purpose registers (SGPR) and vector general purpose
2057 registers (VGPR) are set up as required by the machine code. The required
2058 setup is defined in the :ref:`amdgpu-amdhsa-kernel-descriptor`. The initial
2059 register state is defined in
2060 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`.
2061 9. The prolog of the kernel machine code (see
2062 :ref:`amdgpu-amdhsa-kernel-prolog`) sets up the machine state as necessary
2063 before continuing executing the machine code that corresponds to the kernel.
2064 10. When the kernel dispatch has completed execution, CP signals the completion
2065 signal specified in the kernel dispatch packet if not 0.
2067 .. _amdgpu-amdhsa-memory-spaces:
2069 Memory Spaces
2070 ~~~~~~~~~~~~~
2072 The memory space properties are:
2074 .. table:: AMDHSA Memory Spaces
2075 :name: amdgpu-amdhsa-memory-spaces-table
2077 ================= =========== ======== ======= ==================
2078 Memory Space Name HSA Segment Hardware Address NULL Value
2079 Name Name Size
2080 ================= =========== ======== ======= ==================
2081 Private private scratch 32 0x00000000
2082 Local group LDS 32 0xFFFFFFFF
2083 Global global global 64 0x0000000000000000
2084 Constant constant *same as 64 0x0000000000000000
2085 global*
2086 Generic flat flat 64 0x0000000000000000
2087 Region N/A GDS 32 *not implemented
2088 for AMDHSA*
2089 ================= =========== ======== ======= ==================
2091 The global and constant memory spaces both use global virtual addresses, which
2092 are the same virtual address space used by the CPU. However, some virtual
2093 addresses may only be accessible to the CPU, some only accessible by the GPU,
2094 and some by both.
2096 Using the constant memory space indicates that the data will not change during
2097 the execution of the kernel. This allows scalar read instructions to be
2098 used. The vector and scalar L1 caches are invalidated of volatile data before
2099 each kernel dispatch execution to allow constant memory to change values between
2100 kernel dispatches.
2102 The local memory space uses the hardware Local Data Store (LDS) which is
2103 automatically allocated when the hardware creates work-groups of wavefronts, and
2104 freed when all the wavefronts of a work-group have terminated. The data store
2105 (DS) instructions can be used to access it.
2107 The private memory space uses the hardware scratch memory support. If the kernel
2108 uses scratch, then the hardware allocates memory that is accessed using
2109 wavefront lane dword (4 byte) interleaving. The mapping used from private
2110 address to physical address is:
2112 ``wavefront-scratch-base +
2113 (private-address * wavefront-size * 4) +
2114 (wavefront-lane-id * 4)``
2116 There are different ways that the wavefront scratch base address is determined
2117 by a wavefront (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). This
2118 memory can be accessed in an interleaved manner using buffer instruction with
2119 the scratch buffer descriptor and per wavefront scratch offset, by the scratch
2120 instructions, or by flat instructions. If each lane of a wavefront accesses the
2121 same private address, the interleaving results in adjacent dwords being accessed
2122 and hence requires fewer cache lines to be fetched. Multi-dword access is not
2123 supported except by flat and scratch instructions in GFX9-GFX10.
2125 The generic address space uses the hardware flat address support available in
2126 GFX7-GFX10. This uses two fixed ranges of virtual addresses (the private and
2127 local appertures), that are outside the range of addressible global memory, to
2128 map from a flat address to a private or local address.
2130 FLAT instructions can take a flat address and access global, private (scratch)
2131 and group (LDS) memory depending in if the address is within one of the
2132 apperture ranges. Flat access to scratch requires hardware aperture setup and
2133 setup in the kernel prologue (see :ref:`amdgpu-amdhsa-flat-scratch`). Flat
2134 access to LDS requires hardware aperture setup and M0 (GFX7-GFX8) register setup
2135 (see :ref:`amdgpu-amdhsa-m0`).
2137 To convert between a segment address and a flat address the base address of the
2138 appertures address can be used. For GFX7-GFX8 these are available in the
2139 :ref:`amdgpu-amdhsa-hsa-aql-queue` the address of which can be obtained with
2140 Queue Ptr SGPR (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). For
2141 GFX9-GFX10 the appature base addresses are directly available as inline constant
2142 registers ``SRC_SHARED_BASE/LIMIT`` and ``SRC_PRIVATE_BASE/LIMIT``. In 64 bit
2143 address mode the apperture sizes are 2^32 bytes and the base is aligned to 2^32
2144 which makes it easier to convert from flat to segment or segment to flat.
2146 Image and Samplers
2147 ~~~~~~~~~~~~~~~~~~
2149 Image and sample handles created by the ROCm runtime are 64 bit addresses of a
2150 hardware 32 byte V# and 48 byte S# object respectively. In order to support the
2151 HSA ``query_sampler`` operations two extra dwords are used to store the HSA BRIG
2152 enumeration values for the queries that are not trivially deducible from the S#
2153 representation.
2155 HSA Signals
2156 ~~~~~~~~~~~
2158 HSA signal handles created by the ROCm runtime are 64 bit addresses of a
2159 structure allocated in memory accessible from both the CPU and GPU. The
2160 structure is defined by the ROCm runtime and subject to change between releases
2161 (see [AMD-ROCm-github]_).
2163 .. _amdgpu-amdhsa-hsa-aql-queue:
2165 HSA AQL Queue
2166 ~~~~~~~~~~~~~
2168 The HSA AQL queue structure is defined by the ROCm runtime and subject to change
2169 between releases (see [AMD-ROCm-github]_). For some processors it contains
2170 fields needed to implement certain language features such as the flat address
2171 aperture bases. It also contains fields used by CP such as managing the
2172 allocation of scratch memory.
2174 .. _amdgpu-amdhsa-kernel-descriptor:
2176 Kernel Descriptor
2177 ~~~~~~~~~~~~~~~~~
2179 A kernel descriptor consists of the information needed by CP to initiate the
2180 execution of a kernel, including the entry point address of the machine code
2181 that implements the kernel.
2183 Kernel Descriptor for GFX6-GFX10
2184 ++++++++++++++++++++++++++++++++
2186 CP microcode requires the Kernel descriptor to be allocated on 64 byte
2187 alignment.
2189 .. table:: Kernel Descriptor for GFX6-GFX10
2190 :name: amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table
2192 ======= ======= =============================== ============================
2193 Bits Size Field Name Description
2194 ======= ======= =============================== ============================
2195 31:0 4 bytes GROUP_SEGMENT_FIXED_SIZE The amount of fixed local
2196 address space memory
2197 required for a work-group
2198 in bytes. This does not
2199 include any dynamically
2200 allocated local address
2201 space memory that may be
2202 added when the kernel is
2203 dispatched.
2204 63:32 4 bytes PRIVATE_SEGMENT_FIXED_SIZE The amount of fixed
2205 private address space
2206 memory required for a
2207 work-item in bytes. If
2208 is_dynamic_callstack is 1
2209 then additional space must
2210 be added to this value for
2211 the call stack.
2212 127:64 8 bytes Reserved, must be 0.
2213 191:128 8 bytes KERNEL_CODE_ENTRY_BYTE_OFFSET Byte offset (possibly
2214 negative) from base
2215 address of kernel
2216 descriptor to kernel's
2217 entry point instruction
2218 which must be 256 byte
2219 aligned.
2220 351:272 20 Reserved, must be 0.
2221 bytes
2222 383:352 4 bytes COMPUTE_PGM_RSRC3 GFX6-9
2223 Reserved, must be 0.
2224 GFX10
2225 Compute Shader (CS)
2226 program settings used by
2227 CP to set up
2228 ``COMPUTE_PGM_RSRC3``
2229 configuration
2230 register. See
2231 :ref:`amdgpu-amdhsa-compute_pgm_rsrc3-gfx10-table`.
2232 415:384 4 bytes COMPUTE_PGM_RSRC1 Compute Shader (CS)
2233 program settings used by
2234 CP to set up
2235 ``COMPUTE_PGM_RSRC1``
2236 configuration
2237 register. See
2238 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
2239 447:416 4 bytes COMPUTE_PGM_RSRC2 Compute Shader (CS)
2240 program settings used by
2241 CP to set up
2242 ``COMPUTE_PGM_RSRC2``
2243 configuration
2244 register. See
2245 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
2246 448 1 bit ENABLE_SGPR_PRIVATE_SEGMENT Enable the setup of the
2247 _BUFFER SGPR user data registers
2248 (see
2249 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
2251 The total number of SGPR
2252 user data registers
2253 requested must not exceed
2254 16 and match value in
2255 ``compute_pgm_rsrc2.user_sgpr.user_sgpr_count``.
2256 Any requests beyond 16
2257 will be ignored.
2258 449 1 bit ENABLE_SGPR_DISPATCH_PTR *see above*
2259 450 1 bit ENABLE_SGPR_QUEUE_PTR *see above*
2260 451 1 bit ENABLE_SGPR_KERNARG_SEGMENT_PTR *see above*
2261 452 1 bit ENABLE_SGPR_DISPATCH_ID *see above*
2262 453 1 bit ENABLE_SGPR_FLAT_SCRATCH_INIT *see above*
2263 454 1 bit ENABLE_SGPR_PRIVATE_SEGMENT *see above*
2264 _SIZE
2265 457:455 3 bits Reserved, must be 0.
2266 458 1 bit ENABLE_WAVEFRONT_SIZE32 GFX6-9
2267 Reserved, must be 0.
2268 GFX10
2269 - If 0 execute in
2270 wavefront size 64 mode.
2271 - If 1 execute in
2272 native wavefront size
2273 32 mode.
2274 463:459 5 bits Reserved, must be 0.
2275 511:464 6 bytes Reserved, must be 0.
2276 512 **Total size 64 bytes.**
2277 ======= ====================================================================
2279 ..
2281 .. table:: compute_pgm_rsrc1 for GFX6-GFX10
2282 :name: amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table
2284 ======= ======= =============================== ===========================================================================
2285 Bits Size Field Name Description
2286 ======= ======= =============================== ===========================================================================
2287 5:0 6 bits GRANULATED_WORKITEM_VGPR_COUNT Number of vector register
2288 blocks used by each work-item;
2289 granularity is device
2290 specific:
2292 GFX6-GFX9
2293 - vgprs_used 0..256
2294 - max(0, ceil(vgprs_used / 4) - 1)
2295 GFX10 (wavefront size 64)
2296 - max_vgpr 1..256
2297 - max(0, ceil(vgprs_used / 4) - 1)
2298 GFX10 (wavefront size 32)
2299 - max_vgpr 1..256
2300 - max(0, ceil(vgprs_used / 8) - 1)
2302 Where vgprs_used is defined
2303 as the highest VGPR number
2304 explicitly referenced plus
2305 one.
2307 Used by CP to set up
2308 ``COMPUTE_PGM_RSRC1.VGPRS``.
2310 The
2311 :ref:`amdgpu-assembler`
2312 calculates this
2313 automatically for the
2314 selected processor from
2315 values provided to the
2316 `.amdhsa_kernel` directive
2317 by the
2318 `.amdhsa_next_free_vgpr`
2319 nested directive (see
2320 :ref:`amdhsa-kernel-directives-table`).
2321 9:6 4 bits GRANULATED_WAVEFRONT_SGPR_COUNT Number of scalar register
2322 blocks used by a wavefront;
2323 granularity is device
2324 specific:
2326 GFX6-GFX8
2327 - sgprs_used 0..112
2328 - max(0, ceil(sgprs_used / 8) - 1)
2329 GFX9
2330 - sgprs_used 0..112
2331 - 2 * max(0, ceil(sgprs_used / 16) - 1)
2332 GFX10
2333 Reserved, must be 0.
2334 (128 SGPRs always
2335 allocated.)
2337 Where sgprs_used is
2338 defined as the highest
2339 SGPR number explicitly
2340 referenced plus one, plus
2341 a target-specific number
2342 of additional special
2343 SGPRs for VCC,
2344 FLAT_SCRATCH (GFX7+) and
2345 XNACK_MASK (GFX8+), and
2346 any additional
2347 target-specific
2348 limitations. It does not
2349 include the 16 SGPRs added
2350 if a trap handler is
2351 enabled.
2353 The target-specific
2354 limitations and special
2355 SGPR layout are defined in
2356 the hardware
2357 documentation, which can
2358 be found in the
2359 :ref:`amdgpu-processors`
2360 table.
2362 Used by CP to set up
2363 ``COMPUTE_PGM_RSRC1.SGPRS``.
2365 The
2366 :ref:`amdgpu-assembler`
2367 calculates this
2368 automatically for the
2369 selected processor from
2370 values provided to the
2371 `.amdhsa_kernel` directive
2372 by the
2373 `.amdhsa_next_free_sgpr`
2374 and `.amdhsa_reserve_*`
2375 nested directives (see
2376 :ref:`amdhsa-kernel-directives-table`).
2377 11:10 2 bits PRIORITY Must be 0.
2379 Start executing wavefront
2380 at the specified priority.
2382 CP is responsible for
2383 filling in
2384 ``COMPUTE_PGM_RSRC1.PRIORITY``.
2385 13:12 2 bits FLOAT_ROUND_MODE_32 Wavefront starts execution
2386 with specified rounding
2387 mode for single (32
2388 bit) floating point
2389 precision floating point
2390 operations.
2392 Floating point rounding
2393 mode values are defined in
2394 :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
2396 Used by CP to set up
2397 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
2398 15:14 2 bits FLOAT_ROUND_MODE_16_64 Wavefront starts execution
2399 with specified rounding
2400 denorm mode for half/double (16
2401 and 64 bit) floating point
2402 precision floating point
2403 operations.
2405 Floating point rounding
2406 mode values are defined in
2407 :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
2409 Used by CP to set up
2410 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
2411 17:16 2 bits FLOAT_DENORM_MODE_32 Wavefront starts execution
2412 with specified denorm mode
2413 for single (32
2414 bit) floating point
2415 precision floating point
2416 operations.
2418 Floating point denorm mode
2419 values are defined in
2420 :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
2422 Used by CP to set up
2423 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
2424 19:18 2 bits FLOAT_DENORM_MODE_16_64 Wavefront starts execution
2425 with specified denorm mode
2426 for half/double (16
2427 and 64 bit) floating point
2428 precision floating point
2429 operations.
2431 Floating point denorm mode
2432 values are defined in
2433 :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
2435 Used by CP to set up
2436 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``.
2437 20 1 bit PRIV Must be 0.
2439 Start executing wavefront
2440 in privilege trap handler
2441 mode.
2443 CP is responsible for
2444 filling in
2445 ``COMPUTE_PGM_RSRC1.PRIV``.
2446 21 1 bit ENABLE_DX10_CLAMP Wavefront starts execution
2447 with DX10 clamp mode
2448 enabled. Used by the vector
2449 ALU to force DX10 style
2450 treatment of NaN's (when
2451 set, clamp NaN to zero,
2452 otherwise pass NaN
2453 through).
2455 Used by CP to set up
2456 ``COMPUTE_PGM_RSRC1.DX10_CLAMP``.
2457 22 1 bit DEBUG_MODE Must be 0.
2459 Start executing wavefront
2460 in single step mode.
2462 CP is responsible for
2463 filling in
2464 ``COMPUTE_PGM_RSRC1.DEBUG_MODE``.
2465 23 1 bit ENABLE_IEEE_MODE Wavefront starts execution
2466 with IEEE mode
2467 enabled. Floating point
2468 opcodes that support
2469 exception flag gathering
2470 will quiet and propagate
2471 signaling-NaN inputs per
2472 IEEE 754-2008. Min_dx10 and
2473 max_dx10 become IEEE
2474 754-2008 compliant due to
2475 signaling-NaN propagation
2476 and quieting.
2478 Used by CP to set up
2479 ``COMPUTE_PGM_RSRC1.IEEE_MODE``.
2480 24 1 bit BULKY Must be 0.
2482 Only one work-group allowed
2483 to execute on a compute
2484 unit.
2486 CP is responsible for
2487 filling in
2488 ``COMPUTE_PGM_RSRC1.BULKY``.
2489 25 1 bit CDBG_USER Must be 0.
2491 Flag that can be used to
2492 control debugging code.
2494 CP is responsible for
2495 filling in
2496 ``COMPUTE_PGM_RSRC1.CDBG_USER``.
2497 26 1 bit FP16_OVFL GFX6-GFX8
2498 Reserved, must be 0.
2499 GFX9-GFX10
2500 Wavefront starts execution
2501 with specified fp16 overflow
2502 mode.
2504 - If 0, fp16 overflow generates
2505 +/-INF values.
2506 - If 1, fp16 overflow that is the
2507 result of an +/-INF input value
2508 or divide by 0 produces a +/-INF,
2509 otherwise clamps computed
2510 overflow to +/-MAX_FP16 as
2511 appropriate.
2513 Used by CP to set up
2514 ``COMPUTE_PGM_RSRC1.FP16_OVFL``.
2515 28:27 2 bits Reserved, must be 0.
2516 29 1 bit WGP_MODE GFX6-GFX9
2517 Reserved, must be 0.
2518 GFX10
2519 - If 0 execute work-groups in
2520 CU wavefront execution mode.
2521 - If 1 execute work-groups on
2522 in WGP wavefront execution mode.
2524 See :ref:`amdgpu-amdhsa-memory-model`.
2526 Used by CP to set up
2527 ``COMPUTE_PGM_RSRC1.WGP_MODE``.
2528 30 1 bit MEM_ORDERED GFX6-9
2529 Reserved, must be 0.
2530 GFX10
2531 Controls the behavior of the
2532 waitcnt's vmcnt and vscnt
2533 counters.
2535 - If 0 vmcnt reports completion
2536 of load and atomic with return
2537 out of order with sample
2538 instructions, and the vscnt
2539 reports the completion of
2540 store and atomic without
2541 return in order.
2542 - If 1 vmcnt reports completion
2543 of load, atomic with return
2544 and sample instructions in
2545 order, and the vscnt reports
2546 the completion of store and
2547 atomic without return in order.
2549 Used by CP to set up
2550 ``COMPUTE_PGM_RSRC1.MEM_ORDERED``.
2551 31 1 bit FWD_PROGRESS GFX6-9
2552 Reserved, must be 0.
2553 GFX10
2554 - If 0 execute SIMD wavefronts
2555 using oldest first policy.
2556 - If 1 execute SIMD wavefronts to
2557 ensure wavefronts will make some
2558 forward progress.
2560 Used by CP to set up
2561 ``COMPUTE_PGM_RSRC1.FWD_PROGRESS``.
2562 32 **Total size 4 bytes**
2563 ======= ===================================================================================================================
2565 ..
2567 .. table:: compute_pgm_rsrc2 for GFX6-GFX10
2568 :name: amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table
2570 ======= ======= =============================== ===========================================================================
2571 Bits Size Field Name Description
2572 ======= ======= =============================== ===========================================================================
2573 0 1 bit ENABLE_SGPR_PRIVATE_SEGMENT Enable the setup of the
2574 _WAVEFRONT_OFFSET SGPR wavefront scratch offset
2575 system register (see
2576 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
2578 Used by CP to set up
2579 ``COMPUTE_PGM_RSRC2.SCRATCH_EN``.
2580 5:1 5 bits USER_SGPR_COUNT The total number of SGPR
2581 user data registers
2582 requested. This number must
2583 match the number of user
2584 data registers enabled.
2586 Used by CP to set up
2587 ``COMPUTE_PGM_RSRC2.USER_SGPR``.
2588 6 1 bit ENABLE_TRAP_HANDLER Must be 0.
2590 This bit represents
2591 ``COMPUTE_PGM_RSRC2.TRAP_PRESENT``,
2592 which is set by the CP if
2593 the runtime has installed a
2594 trap handler.
2595 7 1 bit ENABLE_SGPR_WORKGROUP_ID_X Enable the setup of the
2596 system SGPR register for
2597 the work-group id in the X
2598 dimension (see
2599 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
2601 Used by CP to set up
2602 ``COMPUTE_PGM_RSRC2.TGID_X_EN``.
2603 8 1 bit ENABLE_SGPR_WORKGROUP_ID_Y Enable the setup of the
2604 system SGPR register for
2605 the work-group id in the Y
2606 dimension (see
2607 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
2609 Used by CP to set up
2610 ``COMPUTE_PGM_RSRC2.TGID_Y_EN``.
2611 9 1 bit ENABLE_SGPR_WORKGROUP_ID_Z Enable the setup of the
2612 system SGPR register for
2613 the work-group id in the Z
2614 dimension (see
2615 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
2617 Used by CP to set up
2618 ``COMPUTE_PGM_RSRC2.TGID_Z_EN``.
2619 10 1 bit ENABLE_SGPR_WORKGROUP_INFO Enable the setup of the
2620 system SGPR register for
2621 work-group information (see
2622 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`).
2624 Used by CP to set up
2625 ``COMPUTE_PGM_RSRC2.TGID_SIZE_EN``.
2626 12:11 2 bits ENABLE_VGPR_WORKITEM_ID Enable the setup of the
2627 VGPR system registers used
2628 for the work-item ID.
2629 :ref:`amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table`
2630 defines the values.
2632 Used by CP to set up
2633 ``COMPUTE_PGM_RSRC2.TIDIG_CMP_CNT``.
2634 13 1 bit ENABLE_EXCEPTION_ADDRESS_WATCH Must be 0.
2636 Wavefront starts execution
2637 with address watch
2638 exceptions enabled which
2639 are generated when L1 has
2640 witnessed a thread access
2641 an *address of
2642 interest*.
2644 CP is responsible for
2645 filling in the address
2646 watch bit in
2647 ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB``
2648 according to what the
2649 runtime requests.
2650 14 1 bit ENABLE_EXCEPTION_MEMORY Must be 0.
2652 Wavefront starts execution
2653 with memory violation
2654 exceptions exceptions
2655 enabled which are generated
2656 when a memory violation has
2657 occurred for this wavefront from
2658 L1 or LDS
2659 (write-to-read-only-memory,
2660 mis-aligned atomic, LDS
2661 address out of range,
2662 illegal address, etc.).
2664 CP sets the memory
2665 violation bit in
2666 ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB``
2667 according to what the
2668 runtime requests.
2669 23:15 9 bits GRANULATED_LDS_SIZE Must be 0.
2671 CP uses the rounded value
2672 from the dispatch packet,
2673 not this value, as the
2674 dispatch may contain
2675 dynamically allocated group
2676 segment memory. CP writes
2677 directly to
2678 ``COMPUTE_PGM_RSRC2.LDS_SIZE``.
2680 Amount of group segment
2681 (LDS) to allocate for each
2682 work-group. Granularity is
2683 device specific:
2685 GFX6:
2686 roundup(lds-size / (64 * 4))
2687 GFX7-GFX10:
2688 roundup(lds-size / (128 * 4))
2690 24 1 bit ENABLE_EXCEPTION_IEEE_754_FP Wavefront starts execution
2691 _INVALID_OPERATION with specified exceptions
2692 enabled.
2694 Used by CP to set up
2695 ``COMPUTE_PGM_RSRC2.EXCP_EN``
2696 (set from bits 0..6).
2698 IEEE 754 FP Invalid
2699 Operation
2700 25 1 bit ENABLE_EXCEPTION_FP_DENORMAL FP Denormal one or more
2701 _SOURCE input operands is a
2702 denormal number
2703 26 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP Division by
2704 _DIVISION_BY_ZERO Zero
2705 27 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP FP Overflow
2706 _OVERFLOW
2707 28 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP Underflow
2708 _UNDERFLOW
2709 29 1 bit ENABLE_EXCEPTION_IEEE_754_FP IEEE 754 FP Inexact
2710 _INEXACT
2711 30 1 bit ENABLE_EXCEPTION_INT_DIVIDE_BY Integer Division by Zero
2712 _ZERO (rcp_iflag_f32 instruction
2713 only)
2714 31 1 bit Reserved, must be 0.
2715 32 **Total size 4 bytes.**
2716 ======= ===================================================================================================================
2718 ..
2720 .. table:: compute_pgm_rsrc3 for GFX10
2721 :name: amdgpu-amdhsa-compute_pgm_rsrc3-gfx10-table
2723 ======= ======= =============================== ===========================================================================
2724 Bits Size Field Name Description
2725 ======= ======= =============================== ===========================================================================
2726 3:0 4 bits SHARED_VGPR_COUNT Number of shared VGPRs for wavefront size 64. Granularity 8. Value 0-120.
2727 compute_pgm_rsrc1.vgprs + shared_vgpr_cnt cannot exceed 64.
2728 31:4 28 Reserved, must be 0.
2729 bits
2730 32 **Total size 4 bytes.**
2731 ======= ===================================================================================================================
2733 ..
2735 .. table:: Floating Point Rounding Mode Enumeration Values
2736 :name: amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table
2738 ====================================== ===== ==============================
2739 Enumeration Name Value Description
2740 ====================================== ===== ==============================
2741 FLOAT_ROUND_MODE_NEAR_EVEN 0 Round Ties To Even
2742 FLOAT_ROUND_MODE_PLUS_INFINITY 1 Round Toward +infinity
2743 FLOAT_ROUND_MODE_MINUS_INFINITY 2 Round Toward -infinity
2744 FLOAT_ROUND_MODE_ZERO 3 Round Toward 0
2745 ====================================== ===== ==============================
2747 ..
2749 .. table:: Floating Point Denorm Mode Enumeration Values
2750 :name: amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table
2752 ====================================== ===== ==============================
2753 Enumeration Name Value Description
2754 ====================================== ===== ==============================
2755 FLOAT_DENORM_MODE_FLUSH_SRC_DST 0 Flush Source and Destination
2756 Denorms
2757 FLOAT_DENORM_MODE_FLUSH_DST 1 Flush Output Denorms
2758 FLOAT_DENORM_MODE_FLUSH_SRC 2 Flush Source Denorms
2759 FLOAT_DENORM_MODE_FLUSH_NONE 3 No Flush
2760 ====================================== ===== ==============================
2762 ..
2764 .. table:: System VGPR Work-Item ID Enumeration Values
2765 :name: amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table
2767 ======================================== ===== ============================
2768 Enumeration Name Value Description
2769 ======================================== ===== ============================
2770 SYSTEM_VGPR_WORKITEM_ID_X 0 Set work-item X dimension
2771 ID.
2772 SYSTEM_VGPR_WORKITEM_ID_X_Y 1 Set work-item X and Y
2773 dimensions ID.
2774 SYSTEM_VGPR_WORKITEM_ID_X_Y_Z 2 Set work-item X, Y and Z
2775 dimensions ID.
2776 SYSTEM_VGPR_WORKITEM_ID_UNDEFINED 3 Undefined.
2777 ======================================== ===== ============================
2779 .. _amdgpu-amdhsa-initial-kernel-execution-state:
2781 Initial Kernel Execution State
2782 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
2784 This section defines the register state that will be set up by the packet
2785 processor prior to the start of execution of every wavefront. This is limited by
2786 the constraints of the hardware controllers of CP/ADC/SPI.
2788 The order of the SGPR registers is defined, but the compiler can specify which
2789 ones are actually setup in the kernel descriptor using the ``enable_sgpr_*`` bit
2790 fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used
2791 for enabled registers are dense starting at SGPR0: the first enabled register is
2792 SGPR0, the next enabled register is SGPR1 etc.; disabled registers do not have
2793 an SGPR number.
2795 The initial SGPRs comprise up to 16 User SRGPs that are set by CP and apply to
2796 all wavefronts of the grid. It is possible to specify more than 16 User SGPRs using
2797 the ``enable_sgpr_*`` bit fields, in which case only the first 16 are actually
2798 initialized. These are then immediately followed by the System SGPRs that are
2799 set up by ADC/SPI and can have different values for each wavefront of the grid
2800 dispatch.
2802 SGPR register initial state is defined in
2803 :ref:`amdgpu-amdhsa-sgpr-register-set-up-order-table`.
2805 .. table:: SGPR Register Set Up Order
2806 :name: amdgpu-amdhsa-sgpr-register-set-up-order-table
2808 ========== ========================== ====== ==============================
2809 SGPR Order Name Number Description
2810 (kernel descriptor enable of
2811 field) SGPRs
2812 ========== ========================== ====== ==============================
2813 First Private Segment Buffer 4 V# that can be used, together
2814 (enable_sgpr_private with Scratch Wavefront Offset
2815 _segment_buffer) as an offset, to access the
2816 private memory space using a
2817 segment address.
2819 CP uses the value provided by
2820 the runtime.
2821 then Dispatch Ptr 2 64 bit address of AQL dispatch
2822 (enable_sgpr_dispatch_ptr) packet for kernel dispatch
2823 actually executing.
2824 then Queue Ptr 2 64 bit address of amd_queue_t
2825 (enable_sgpr_queue_ptr) object for AQL queue on which
2826 the dispatch packet was
2827 queued.
2828 then Kernarg Segment Ptr 2 64 bit address of Kernarg
2829 (enable_sgpr_kernarg segment. This is directly
2830 _segment_ptr) copied from the
2831 kernarg_address in the kernel
2832 dispatch packet.
2834 Having CP load it once avoids
2835 loading it at the beginning of
2836 every wavefront.
2837 then Dispatch Id 2 64 bit Dispatch ID of the
2838 (enable_sgpr_dispatch_id) dispatch packet being
2839 executed.
2840 then Flat Scratch Init 2 This is 2 SGPRs:
2841 (enable_sgpr_flat_scratch
2842 _init) GFX6
2843 Not supported.
2844 GFX7-GFX8
2845 The first SGPR is a 32 bit
2846 byte offset from
2847 ``SH_HIDDEN_PRIVATE_BASE_VIMID``
2848 to per SPI base of memory
2849 for scratch for the queue
2850 executing the kernel
2851 dispatch. CP obtains this
2852 from the runtime. (The
2853 Scratch Segment Buffer base
2854 address is
2855 ``SH_HIDDEN_PRIVATE_BASE_VIMID``
2856 plus this offset.) The value
2857 of Scratch Wavefront Offset must
2858 be added to this offset by
2859 the kernel machine code,
2860 right shifted by 8, and
2861 moved to the FLAT_SCRATCH_HI
2862 SGPR register.
2863 FLAT_SCRATCH_HI corresponds
2864 to SGPRn-4 on GFX7, and
2865 SGPRn-6 on GFX8 (where SGPRn
2866 is the highest numbered SGPR
2867 allocated to the wavefront).
2868 FLAT_SCRATCH_HI is
2869 multiplied by 256 (as it is
2870 in units of 256 bytes) and
2871 added to
2872 ``SH_HIDDEN_PRIVATE_BASE_VIMID``
2873 to calculate the per wavefront
2874 FLAT SCRATCH BASE in flat
2875 memory instructions that
2876 access the scratch
2877 apperture.
2879 The second SGPR is 32 bit
2880 byte size of a single
2881 work-item's scratch memory
2882 usage. CP obtains this from
2883 the runtime, and it is
2884 always a multiple of DWORD.
2885 CP checks that the value in
2886 the kernel dispatch packet
2887 Private Segment Byte Size is
2888 not larger, and requests the
2889 runtime to increase the
2890 queue's scratch size if
2891 necessary. The kernel code
2892 must move it to
2893 FLAT_SCRATCH_LO which is
2894 SGPRn-3 on GFX7 and SGPRn-5
2895 on GFX8. FLAT_SCRATCH_LO is
2896 used as the FLAT SCRATCH
2897 SIZE in flat memory
2898 instructions. Having CP load
2899 it once avoids loading it at
2900 the beginning of every
2901 wavefront.
2902 GFX9-GFX10
2903 This is the
2904 64 bit base address of the
2905 per SPI scratch backing
2906 memory managed by SPI for
2907 the queue executing the
2908 kernel dispatch. CP obtains
2909 this from the runtime (and
2910 divides it if there are
2911 multiple Shader Arrays each
2912 with its own SPI). The value
2913 of Scratch Wavefront Offset must
2914 be added by the kernel
2915 machine code and the result
2916 moved to the FLAT_SCRATCH
2917 SGPR which is SGPRn-6 and
2918 SGPRn-5. It is used as the
2919 FLAT SCRATCH BASE in flat
2920 memory instructions.
2921 then Private Segment Size 1 The 32 bit byte size of a
2922 (enable_sgpr_private single
2923 work-item's
2924 scratch_segment_size) memory
2925 allocation. This is the
2926 value from the kernel
2927 dispatch packet Private
2928 Segment Byte Size rounded up
2929 by CP to a multiple of
2930 DWORD.
2932 Having CP load it once avoids
2933 loading it at the beginning of
2934 every wavefront.
2936 This is not used for
2937 GFX7-GFX8 since it is the same
2938 value as the second SGPR of
2939 Flat Scratch Init. However, it
2940 may be needed for GFX9-GFX10 which
2941 changes the meaning of the
2942 Flat Scratch Init value.
2943 then Grid Work-Group Count X 1 32 bit count of the number of
2944 (enable_sgpr_grid work-groups in the X dimension
2945 _workgroup_count_X) for the grid being
2946 executed. Computed from the
2947 fields in the kernel dispatch
2948 packet as ((grid_size.x +
2949 workgroup_size.x - 1) /
2950 workgroup_size.x).
2951 then Grid Work-Group Count Y 1 32 bit count of the number of
2952 (enable_sgpr_grid work-groups in the Y dimension
2953 _workgroup_count_Y && for the grid being
2954 less than 16 previous executed. Computed from the
2955 SGPRs) fields in the kernel dispatch
2956 packet as ((grid_size.y +
2957 workgroup_size.y - 1) /
2958 workgroupSize.y).
2960 Only initialized if <16
2961 previous SGPRs initialized.
2962 then Grid Work-Group Count Z 1 32 bit count of the number of
2963 (enable_sgpr_grid work-groups in the Z dimension
2964 _workgroup_count_Z && for the grid being
2965 less than 16 previous executed. Computed from the
2966 SGPRs) fields in the kernel dispatch
2967 packet as ((grid_size.z +
2968 workgroup_size.z - 1) /
2969 workgroupSize.z).
2971 Only initialized if <16
2972 previous SGPRs initialized.
2973 then Work-Group Id X 1 32 bit work-group id in X
2974 (enable_sgpr_workgroup_id dimension of grid for
2975 _X) wavefront.
2976 then Work-Group Id Y 1 32 bit work-group id in Y
2977 (enable_sgpr_workgroup_id dimension of grid for
2978 _Y) wavefront.
2979 then Work-Group Id Z 1 32 bit work-group id in Z
2980 (enable_sgpr_workgroup_id dimension of grid for
2981 _Z) wavefront.
2982 then Work-Group Info 1 {first_wavefront, 14'b0000,
2983 (enable_sgpr_workgroup ordered_append_term[10:0],
2984 _info) threadgroup_size_in_wavefronts[5:0]}
2985 then Scratch Wavefront Offset 1 32 bit byte offset from base
2986 (enable_sgpr_private of scratch base of queue
2987 _segment_wavefront_offset) executing the kernel
2988 dispatch. Must be used as an
2989 offset with Private
2990 segment address when using
2991 Scratch Segment Buffer. It
2992 must be used to set up FLAT
2993 SCRATCH for flat addressing
2994 (see
2995 :ref:`amdgpu-amdhsa-flat-scratch`).
2996 ========== ========================== ====== ==============================
2998 The order of the VGPR registers is defined, but the compiler can specify which
2999 ones are actually setup in the kernel descriptor using the ``enable_vgpr*`` bit
3000 fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used
3001 for enabled registers are dense starting at VGPR0: the first enabled register is
3002 VGPR0, the next enabled register is VGPR1 etc.; disabled registers do not have a
3003 VGPR number.
3005 VGPR register initial state is defined in
3006 :ref:`amdgpu-amdhsa-vgpr-register-set-up-order-table`.
3008 .. table:: VGPR Register Set Up Order
3009 :name: amdgpu-amdhsa-vgpr-register-set-up-order-table
3011 ========== ========================== ====== ==============================
3012 VGPR Order Name Number Description
3013 (kernel descriptor enable of
3014 field) VGPRs
3015 ========== ========================== ====== ==============================
3016 First Work-Item Id X 1 32 bit work item id in X
3017 (Always initialized) dimension of work-group for
3018 wavefront lane.
3019 then Work-Item Id Y 1 32 bit work item id in Y
3020 (enable_vgpr_workitem_id dimension of work-group for
3021 > 0) wavefront lane.
3022 then Work-Item Id Z 1 32 bit work item id in Z
3023 (enable_vgpr_workitem_id dimension of work-group for
3024 > 1) wavefront lane.
3025 ========== ========================== ====== ==============================
3027 The setting of registers is done by GPU CP/ADC/SPI hardware as follows:
3029 1. SGPRs before the Work-Group Ids are set by CP using the 16 User Data
3030 registers.
3031 2. Work-group Id registers X, Y, Z are set by ADC which supports any
3032 combination including none.
3033 3. Scratch Wavefront Offset is set by SPI in a per wavefront basis which is why
3034 its value cannot included with the flat scratch init value which is per queue.
3035 4. The VGPRs are set by SPI which only supports specifying either (X), (X, Y)
3036 or (X, Y, Z).
3038 Flat Scratch register pair are adjacent SGRRs so they can be moved as a 64 bit
3039 value to the hardware required SGPRn-3 and SGPRn-4 respectively.
3041 The global segment can be accessed either using buffer instructions (GFX6 which
3042 has V# 64 bit address support), flat instructions (GFX7-GFX10), or global
3043 instructions (GFX9-GFX10).
3045 If buffer operations are used then the compiler can generate a V# with the
3046 following properties:
3048 * base address of 0
3049 * no swizzle
3050 * ATC: 1 if IOMMU present (such as APU)
3051 * ptr64: 1
3052 * MTYPE set to support memory coherence that matches the runtime (such as CC for
3053 APU and NC for dGPU).
3055 .. _amdgpu-amdhsa-kernel-prolog:
3057 Kernel Prolog
3058 ~~~~~~~~~~~~~
3060 .. _amdgpu-amdhsa-m0:
3062 M0
3063 ++
3065 GFX6-GFX8
3066 The M0 register must be initialized with a value at least the total LDS size
3067 if the kernel may access LDS via DS or flat operations. Total LDS size is
3068 available in dispatch packet. For M0, it is also possible to use maximum
3069 possible value of LDS for given target (0x7FFF for GFX6 and 0xFFFF for
3070 GFX7-GFX8).
3071 GFX9-GFX10
3072 The M0 register is not used for range checking LDS accesses and so does not
3073 need to be initialized in the prolog.
3075 .. _amdgpu-amdhsa-flat-scratch:
3077 Flat Scratch
3078 ++++++++++++
3080 If the kernel may use flat operations to access scratch memory, the prolog code
3081 must set up FLAT_SCRATCH register pair (FLAT_SCRATCH_LO/FLAT_SCRATCH_HI which
3082 are in SGPRn-4/SGPRn-3). Initialization uses Flat Scratch Init and Scratch Wavefront
3083 Offset SGPR registers (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`):
3085 GFX6
3086 Flat scratch is not supported.
3088 GFX7-GFX8
3089 1. The low word of Flat Scratch Init is 32 bit byte offset from
3090 ``SH_HIDDEN_PRIVATE_BASE_VIMID`` to the base of scratch backing memory
3091 being managed by SPI for the queue executing the kernel dispatch. This is
3092 the same value used in the Scratch Segment Buffer V# base address. The
3093 prolog must add the value of Scratch Wavefront Offset to get the wavefront's byte
3094 scratch backing memory offset from ``SH_HIDDEN_PRIVATE_BASE_VIMID``. Since
3095 FLAT_SCRATCH_LO is in units of 256 bytes, the offset must be right shifted
3096 by 8 before moving into FLAT_SCRATCH_LO.
3097 2. The second word of Flat Scratch Init is 32 bit byte size of a single
3098 work-items scratch memory usage. This is directly loaded from the kernel
3099 dispatch packet Private Segment Byte Size and rounded up to a multiple of
3100 DWORD. Having CP load it once avoids loading it at the beginning of every
3101 wavefront. The prolog must move it to FLAT_SCRATCH_LO for use as FLAT SCRATCH
3102 SIZE.
3104 GFX9-GFX10
3105 The Flat Scratch Init is the 64 bit address of the base of scratch backing
3106 memory being managed by SPI for the queue executing the kernel dispatch. The
3107 prolog must add the value of Scratch Wavefront Offset and moved to the FLAT_SCRATCH
3108 pair for use as the flat scratch base in flat memory instructions.
3110 .. _amdgpu-amdhsa-memory-model:
3112 Memory Model
3113 ~~~~~~~~~~~~
3115 This section describes the mapping of LLVM memory model onto AMDGPU machine code
3116 (see :ref:`memmodel`). *The implementation is WIP.*
3118 .. TODO
3119 Update when implementation complete.
3121 The AMDGPU backend supports the memory synchronization scopes specified in
3122 :ref:`amdgpu-memory-scopes`.
3124 The code sequences used to implement the memory model are defined in table
3125 :ref:`amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx10-table`.
3127 The sequences specify the order of instructions that a single thread must
3128 execute. The ``s_waitcnt`` and ``buffer_wbinvl1_vol`` are defined with respect
3129 to other memory instructions executed by the same thread. This allows them to be
3130 moved earlier or later which can allow them to be combined with other instances
3131 of the same instruction, or hoisted/sunk out of loops to improve
3132 performance. Only the instructions related to the memory model are given;
3133 additional ``s_waitcnt`` instructions are required to ensure registers are
3134 defined before being used. These may be able to be combined with the memory
3135 model ``s_waitcnt`` instructions as described above.
3137 The AMDGPU backend supports the following memory models:
3139 HSA Memory Model [HSA]_
3140 The HSA memory model uses a single happens-before relation for all address
3141 spaces (see :ref:`amdgpu-address-spaces`).
3142 OpenCL Memory Model [OpenCL]_
3143 The OpenCL memory model which has separate happens-before relations for the
3144 global and local address spaces. Only a fence specifying both global and
3145 local address space, and seq_cst instructions join the relationships. Since
3146 the LLVM ``memfence`` instruction does not allow an address space to be
3147 specified the OpenCL fence has to convervatively assume both local and
3148 global address space was specified. However, optimizations can often be
3149 done to eliminate the additional ``s_waitcnt`` instructions when there are
3150 no intervening memory instructions which access the corresponding address
3151 space. The code sequences in the table indicate what can be omitted for the
3152 OpenCL memory. The target triple environment is used to determine if the
3153 source language is OpenCL (see :ref:`amdgpu-opencl`).
3155 ``ds/flat_load/store/atomic`` instructions to local memory are termed LDS
3156 operations.
3158 ``buffer/global/flat_load/store/atomic`` instructions to global memory are
3159 termed vector memory operations.
3161 For GFX6-GFX9:
3163 * Each agent has multiple shader arrays (SA).
3164 * Each SA has multiple compute units (CU).
3165 * Each CU has multiple SIMDs that execute wavefronts.
3166 * The wavefronts for a single work-group are executed in the same CU but may be
3167 executed by different SIMDs.
3168 * Each CU has a single LDS memory shared by the wavefronts of the work-groups
3169 executing on it.
3170 * All LDS operations of a CU are performed as wavefront wide operations in a
3171 global order and involve no caching. Completion is reported to a wavefront in
3172 execution order.
3173 * The LDS memory has multiple request queues shared by the SIMDs of a
3174 CU. Therefore, the LDS operations performed by different wavefronts of a work-group
3175 can be reordered relative to each other, which can result in reordering the
3176 visibility of vector memory operations with respect to LDS operations of other
3177 wavefronts in the same work-group. A ``s_waitcnt lgkmcnt(0)`` is required to
3178 ensure synchronization between LDS operations and vector memory operations
3179 between wavefronts of a work-group, but not between operations performed by the
3180 same wavefront.
3181 * The vector memory operations are performed as wavefront wide operations and
3182 completion is reported to a wavefront in execution order. The exception is
3183 that for GFX7-GFX9 ``flat_load/store/atomic`` instructions can report out of
3184 vector memory order if they access LDS memory, and out of LDS operation order
3185 if they access global memory.
3186 * The vector memory operations access a single vector L1 cache shared by all
3187 SIMDs a CU. Therefore, no special action is required for coherence between the
3188 lanes of a single wavefront, or for coherence between wavefronts in the same
3189 work-group. A ``buffer_wbinvl1_vol`` is required for coherence between wavefronts
3190 executing in different work-groups as they may be executing on different CUs.
3191 * The scalar memory operations access a scalar L1 cache shared by all wavefronts
3192 on a group of CUs. The scalar and vector L1 caches are not coherent. However,
3193 scalar operations are used in a restricted way so do not impact the memory
3194 model. See :ref:`amdgpu-amdhsa-memory-spaces`.
3195 * The vector and scalar memory operations use an L2 cache shared by all CUs on
3196 the same agent.
3197 * The L2 cache has independent channels to service disjoint ranges of virtual
3198 addresses.
3199 * Each CU has a separate request queue per channel. Therefore, the vector and
3200 scalar memory operations performed by wavefronts executing in different work-groups
3201 (which may be executing on different CUs) of an agent can be reordered
3202 relative to each other. A ``s_waitcnt vmcnt(0)`` is required to ensure
3203 synchronization between vector memory operations of different CUs. It ensures a
3204 previous vector memory operation has completed before executing a subsequent
3205 vector memory or LDS operation and so can be used to meet the requirements of
3206 acquire and release.
3207 * The L2 cache can be kept coherent with other agents on some targets, or ranges
3208 of virtual addresses can be set up to bypass it to ensure system coherence.
3210 For GFX10:
3212 * Each agent has multiple shader arrays (SA).
3213 * Each SA has multiple work-group processors (WGP).
3214 * Each WGP has multiple compute units (CU).
3215 * Each CU has multiple SIMDs that execute wavefronts.
3216 * The wavefronts for a single work-group are executed in the same
3217 WGP. In CU wavefront execution mode the wavefronts may be executed by
3218 different SIMDs in the same CU. In WGP wavefront execution mode the
3219 wavefronts may be executed by different SIMDs in different CUs in the same
3220 WGP.
3221 * Each WGP has a single LDS memory shared by the wavefronts of the work-groups
3222 executing on it.
3223 * All LDS operations of a WGP are performed as wavefront wide operations in a
3224 global order and involve no caching. Completion is reported to a wavefront in
3225 execution order.
3226 * The LDS memory has multiple request queues shared by the SIMDs of a
3227 WGP. Therefore, the LDS operations performed by different wavefronts of a work-group
3228 can be reordered relative to each other, which can result in reordering the
3229 visibility of vector memory operations with respect to LDS operations of other
3230 wavefronts in the same work-group. A ``s_waitcnt lgkmcnt(0)`` is required to
3231 ensure synchronization between LDS operations and vector memory operations
3232 between wavefronts of a work-group, but not between operations performed by the
3233 same wavefront.
3234 * The vector memory operations are performed as wavefront wide operations.
3235 Completion of load/store/sample operations are reported to a wavefront in
3236 execution order of other load/store/sample operations performed by that
3237 wavefront.
3238 * The vector memory operations access a vector L0 cache. There is a single L0
3239 cache per CU. Each SIMD of a CU accesses the same L0 cache.
3240 Therefore, no special action is required for coherence between the lanes of a
3241 single wavefront. However, a ``BUFFER_GL0_INV`` is required for coherence
3242 between wavefronts executing in the same work-group as they may be executing on
3243 SIMDs of different CUs that access different L0s. A ``BUFFER_GL0_INV`` is also
3244 required for coherence between wavefronts executing in different work-groups as
3245 they may be executing on different WGPs.
3246 * The scalar memory operations access a scalar L0 cache shared by all wavefronts
3247 on a WGP. The scalar and vector L0 caches are not coherent. However, scalar
3248 operations are used in a restricted way so do not impact the memory model. See
3249 :ref:`amdgpu-amdhsa-memory-spaces`.
3250 * The vector and scalar memory L0 caches use an L1 cache shared by all WGPs on
3251 the same SA. Therefore, no special action is required for coherence between
3252 the wavefronts of a single work-group. However, a ``BUFFER_GL1_INV`` is
3253 required for coherence between wavefronts executing in different work-groups as
3254 they may be executing on different SAs that access different L1s.
3255 * The L1 caches have independent quadrants to service disjoint ranges of virtual
3256 addresses.
3257 * Each L0 cache has a separate request queue per L1 quadrant. Therefore, the
3258 vector and scalar memory operations performed by different wavefronts, whether
3259 executing in the same or different work-groups (which may be executing on
3260 different CUs accessing different L0s), can be reordered relative to each
3261 other. A ``s_waitcnt vmcnt(0) & vscnt(0)`` is required to ensure synchronization
3262 between vector memory operations of different wavefronts. It ensures a previous
3263 vector memory operation has completed before executing a subsequent vector
3264 memory or LDS operation and so can be used to meet the requirements of acquire,
3265 release and sequential consistency.
3266 * The L1 caches use an L2 cache shared by all SAs on the same agent.
3267 * The L2 cache has independent channels to service disjoint ranges of virtual
3268 addresses.
3269 * Each L1 quadrant of a single SA accesses a different L2 channel. Each L1
3270 quadrant has a separate request queue per L2 channel. Therefore, the vector
3271 and scalar memory operations performed by wavefronts executing in different
3272 work-groups (which may be executing on different SAs) of an agent can be
3273 reordered relative to each other. A ``s_waitcnt vmcnt(0) & vscnt(0)`` is
3274 required to ensure synchronization between vector memory operations of
3275 different SAs. It ensures a previous vector memory operation has completed
3276 before executing a subsequent vector memory and so can be used to meet the
3277 requirements of acquire, release and sequential consistency.
3278 * The L2 cache can be kept coherent with other agents on some targets, or ranges
3279 of virtual addresses can be set up to bypass it to ensure system coherence.
3281 Private address space uses ``buffer_load/store`` using the scratch V# (GFX6-GFX8),
3282 or ``scratch_load/store`` (GFX9-GFX10). Since only a single thread is accessing the
3283 memory, atomic memory orderings are not meaningful and all accesses are treated
3284 as non-atomic.
3286 Constant address space uses ``buffer/global_load`` instructions (or equivalent
3287 scalar memory instructions). Since the constant address space contents do not
3288 change during the execution of a kernel dispatch it is not legal to perform
3289 stores, and atomic memory orderings are not meaningful and all access are
3290 treated as non-atomic.
3292 A memory synchronization scope wider than work-group is not meaningful for the
3293 group (LDS) address space and is treated as work-group.
3295 The memory model does not support the region address space which is treated as
3296 non-atomic.
3298 Acquire memory ordering is not meaningful on store atomic instructions and is
3299 treated as non-atomic.
3301 Release memory ordering is not meaningful on load atomic instructions and is
3302 treated a non-atomic.
3304 Acquire-release memory ordering is not meaningful on load or store atomic
3305 instructions and is treated as acquire and release respectively.
3307 AMDGPU backend only uses scalar memory operations to access memory that is
3308 proven to not change during the execution of the kernel dispatch. This includes
3309 constant address space and global address space for program scope const
3310 variables. Therefore the kernel machine code does not have to maintain the
3311 scalar L1 cache to ensure it is coherent with the vector L1 cache. The scalar
3312 and vector L1 caches are invalidated between kernel dispatches by CP since
3313 constant address space data may change between kernel dispatch executions. See
3314 :ref:`amdgpu-amdhsa-memory-spaces`.
3316 The one execption is if scalar writes are used to spill SGPR registers. In this
3317 case the AMDGPU backend ensures the memory location used to spill is never
3318 accessed by vector memory operations at the same time. If scalar writes are used
3319 then a ``s_dcache_wb`` is inserted before the ``s_endpgm`` and before a function
3320 return since the locations may be used for vector memory instructions by a
3321 future wavefront that uses the same scratch area, or a function call that creates a
3322 frame at the same address, respectively. There is no need for a ``s_dcache_inv``
3323 as all scalar writes are write-before-read in the same thread.
3325 For GFX6-GFX9, scratch backing memory (which is used for the private address space)
3326 is accessed with MTYPE NC_NV (non-coherenent non-volatile). Since the private
3327 address space is only accessed by a single thread, and is always
3328 write-before-read, there is never a need to invalidate these entries from the L1
3329 cache. Hence all cache invalidates are done as ``*_vol`` to only invalidate the
3330 volatile cache lines.
3332 For GFX10, scratch backing memory (which is used for the private address space)
3333 is accessed with MTYPE NC (non-coherenent). Since the private address space is
3334 only accessed by a single thread, and is always write-before-read, there is
3335 never a need to invalidate these entries from the L0 or L1 caches.
3337 For GFX10, wavefronts are executed in native mode with in-order reporting of loads
3338 and sample instructions. In this mode vmcnt reports completion of load, atomic
3339 with return and sample instructions in order, and the vscnt reports the
3340 completion of store and atomic without return in order. See ``MEM_ORDERED`` field
3341 in :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
3343 In GFX10, wavefronts can be executed in WGP or CU wavefront execution mode:
3345 * In WGP wavefront execution mode the wavefronts of a work-group are executed
3346 on the SIMDs of both CUs of the WGP. Therefore, explicit management of the per
3347 CU L0 caches is required for work-group synchronization. Also accesses to L1 at
3348 work-group scope need to be expicitly ordered as the accesses from different
3349 CUs are not ordered.
3350 * In CU wavefront execution mode the wavefronts of a work-group are executed on
3351 the SIMDs of a single CU of the WGP. Therefore, all global memory access by
3352 the work-group access the same L0 which in turn ensures L1 accesses are
3353 ordered and so do not require explicit management of the caches for
3354 work-group synchronization.
3356 See ``WGP_MODE`` field in :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`
3357 and :ref:`amdgpu-target-features`.
3359 On dGPU the kernarg backing memory is accessed as UC (uncached) to avoid needing
3360 to invalidate the L2 cache. For GFX6-GFX9, this also causes it to be treated as
3361 non-volatile and so is not invalidated by ``*_vol``. On APU it is accessed as CC
3362 (cache coherent) and so the L2 cache will be coherent with the CPU and other
3363 agents.
3365 .. table:: AMDHSA Memory Model Code Sequences GFX6-GFX10
3366 :name: amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx10-table
3368 ============ ============ ============== ========== =============================== ==================================
3369 LLVM Instr LLVM Memory LLVM Memory AMDGPU AMDGPU Machine Code AMDGPU Machine Code
3370 Ordering Sync Scope Address GFX6-9 GFX10
3371 Space
3372 ============ ============ ============== ========== =============================== ==================================
3373 **Non-Atomic**
3374 ----------------------------------------------------------------------------------------------------------------------
3375 load *none* *none* - global - !volatile & !nontemporal - !volatile & !nontemporal
3376 - generic
3377 - private 1. buffer/global/flat_load 1. buffer/global/flat_load
3378 - constant
3379 - volatile & !nontemporal - volatile & !nontemporal
3381 1. buffer/global/flat_load 1. buffer/global/flat_load
3382 glc=1 glc=1 dlc=1
3384 - nontemporal - nontemporal
3386 1. buffer/global/flat_load 1. buffer/global/flat_load
3387 glc=1 slc=1 slc=1
3389 load *none* *none* - local 1. ds_load 1. ds_load
3390 store *none* *none* - global - !nontemporal - !nontemporal
3391 - generic
3392 - private 1. buffer/global/flat_store 1. buffer/global/flat_store
3393 - constant
3394 - nontemporal - nontemporal
3396 1. buffer/global/flat_stote 1. buffer/global/flat_store
3397 glc=1 slc=1 slc=1
3399 store *none* *none* - local 1. ds_store 1. ds_store
3400 **Unordered Atomic**
3401 ----------------------------------------------------------------------------------------------------------------------
3402 load atomic unordered *any* *any* *Same as non-atomic*. *Same as non-atomic*.
3403 store atomic unordered *any* *any* *Same as non-atomic*. *Same as non-atomic*.
3404 atomicrmw unordered *any* *any* *Same as monotonic *Same as monotonic
3405 atomic*. atomic*.
3406 **Monotonic Atomic**
3407 ----------------------------------------------------------------------------------------------------------------------
3408 load atomic monotonic - singlethread - global 1. buffer/global/flat_load 1. buffer/global/flat_load
3409 - wavefront - generic
3410 load atomic monotonic - workgroup - global 1. buffer/global/flat_load 1. buffer/global/flat_load
3411 - generic glc=1
3413 - If CU wavefront execution mode, omit glc=1.
3415 load atomic monotonic - singlethread - local 1. ds_load 1. ds_load
3416 - wavefront
3417 - workgroup
3418 load atomic monotonic - agent - global 1. buffer/global/flat_load 1. buffer/global/flat_load
3419 - system - generic glc=1 glc=1 dlc=1
3420 store atomic monotonic - singlethread - global 1. buffer/global/flat_store 1. buffer/global/flat_store
3421 - wavefront - generic
3422 - workgroup
3423 - agent
3424 - system
3425 store atomic monotonic - singlethread - local 1. ds_store 1. ds_store
3426 - wavefront
3427 - workgroup
3428 atomicrmw monotonic - singlethread - global 1. buffer/global/flat_atomic 1. buffer/global/flat_atomic
3429 - wavefront - generic
3430 - workgroup
3431 - agent
3432 - system
3433 atomicrmw monotonic - singlethread - local 1. ds_atomic 1. ds_atomic
3434 - wavefront
3435 - workgroup
3436 **Acquire Atomic**
3437 ----------------------------------------------------------------------------------------------------------------------
3438 load atomic acquire - singlethread - global 1. buffer/global/ds/flat_load 1. buffer/global/ds/flat_load
3439 - wavefront - local
3440 - generic
3441 load atomic acquire - workgroup - global 1. buffer/global/flat_load 1. buffer/global_load glc=1
3443 - If CU wavefront execution mode, omit glc=1.
3445 2. s_waitcnt vmcnt(0)
3447 - If CU wavefront execution mode, omit.
3448 - Must happen before
3449 the following buffer_gl0_inv
3450 and before any following
3451 global/generic
3452 load/load
3453 atomic/stote/store
3454 atomic/atomicrmw.
3456 3. buffer_gl0_inv
3458 - If CU wavefront execution mode, omit.
3459 - Ensures that
3460 following
3461 loads will not see
3462 stale data.
3464 load atomic acquire - workgroup - local 1. ds_load 1. ds_load
3465 2. s_waitcnt lgkmcnt(0) 2. s_waitcnt lgkmcnt(0)
3467 - If OpenCL, omit. - If OpenCL, omit.
3468 - Must happen before - Must happen before
3469 any following the following buffer_gl0_inv
3470 global/generic and before any following
3471 load/load global/generic load/load
3472 atomic/store/store atomic/store/store
3473 atomic/atomicrmw. atomic/atomicrmw.
3474 - Ensures any - Ensures any
3475 following global following global
3476 data read is no data read is no
3477 older than the load older than the load
3478 atomic value being atomic value being
3479 acquired. acquired.
3481 3. buffer_gl0_inv
3483 - If CU wavefront execution mode, omit.
3484 - If OpenCL, omit.
3485 - Ensures that
3486 following
3487 loads will not see
3488 stale data.
3490 load atomic acquire - workgroup - generic 1. flat_load 1. flat_load glc=1
3492 - If CU wavefront execution mode, omit glc=1.
3494 2. s_waitcnt lgkmcnt(0) 2. s_waitcnt lgkmcnt(0) &
3495 vmcnt(0)
3497 - If CU wavefront execution mode, omit vmcnt.
3498 - If OpenCL, omit. - If OpenCL, omit
3499 lgkmcnt(0).
3500 - Must happen before - Must happen before
3501 any following the following
3502 global/generic buffer_gl0_inv and any
3503 load/load following global/generic
3504 atomic/store/store load/load
3505 atomic/atomicrmw. atomic/store/store
3506 atomic/atomicrmw.
3507 - Ensures any - Ensures any
3508 following global following global
3509 data read is no data read is no
3510 older than the load older than the load
3511 atomic value being atomic value being
3512 acquired. acquired.
3514 3. buffer_gl0_inv
3516 - If CU wavefront execution mode, omit.
3517 - Ensures that
3518 following
3519 loads will not see
3520 stale data.
3522 load atomic acquire - agent - global 1. buffer/global/flat_load 1. buffer/global_load
3523 - system glc=1 glc=1 dlc=1
3524 2. s_waitcnt vmcnt(0) 2. s_waitcnt vmcnt(0)
3526 - Must happen before - Must happen before
3527 following following
3528 buffer_wbinvl1_vol. buffer_gl*_inv.
3529 - Ensures the load - Ensures the load
3530 has completed has completed
3531 before invalidating before invalidating
3532 the cache. the caches.
3534 3. buffer_wbinvl1_vol 3. buffer_gl0_inv;
3535 buffer_gl1_inv
3537 - Must happen before - Must happen before
3538 any following any following
3539 global/generic global/generic
3540 load/load load/load
3541 atomic/atomicrmw. atomic/atomicrmw.
3542 - Ensures that - Ensures that
3543 following following
3544 loads will not see loads will not see
3545 stale global data. stale global data.
3547 load atomic acquire - agent - generic 1. flat_load glc=1 1. flat_load glc=1 dlc=1
3548 - system 2. s_waitcnt vmcnt(0) & 2. s_waitcnt vmcnt(0) &
3549 lgkmcnt(0) lgkmcnt(0)
3551 - If OpenCL omit - If OpenCL omit
3552 lgkmcnt(0). lgkmcnt(0).
3553 - Must happen before - Must happen before
3554 following following
3555 buffer_wbinvl1_vol. buffer_gl*_invl.
3556 - Ensures the flat_load - Ensures the flat_load
3557 has completed has completed
3558 before invalidating before invalidating
3559 the cache. the caches.
3561 3. buffer_wbinvl1_vol 3. buffer_gl0_inv;
3562 buffer_gl1_inv
3564 - Must happen before - Must happen before
3565 any following any following
3566 global/generic global/generic
3567 load/load load/load
3568 atomic/atomicrmw. atomic/atomicrmw.
3569 - Ensures that - Ensures that
3570 following loads following loads
3571 will not see stale will not see stale
3572 global data. global data.
3574 atomicrmw acquire - singlethread - global 1. buffer/global/ds/flat_atomic 1. buffer/global/ds/flat_atomic
3575 - wavefront - local
3576 - generic
3577 atomicrmw acquire - workgroup - global 1. buffer/global/flat_atomic 1. buffer/global_atomic
3578 2. s_waitcnt vm/vscnt(0)
3580 - If CU wavefront execution mode, omit.
3581 - Use vmcnt if atomic with
3582 return and vscnt if atomic
3583 with no-return.
3584 - Must happen before
3585 the following buffer_gl0_inv
3586 and before any following
3587 global/generic
3588 load/load
3589 atomic/stote/store
3590 atomic/atomicrmw.
3592 3. buffer_gl0_inv
3594 - If CU wavefront execution mode, omit.
3595 - Ensures that
3596 following
3597 loads will not see
3598 stale data.
3600 atomicrmw acquire - workgroup - local 1. ds_atomic 1. ds_atomic
3601 2. waitcnt lgkmcnt(0) 2. waitcnt lgkmcnt(0)
3603 - If OpenCL, omit. - If OpenCL, omit.
3604 - Must happen before - Must happen before
3605 any following the following
3606 global/generic buffer_gl0_inv.
3607 load/load
3608 atomic/store/store
3609 atomic/atomicrmw.
3610 - Ensures any - Ensures any
3611 following global following global
3612 data read is no data read is no
3613 older than the older than the
3614 atomicrmw value atomicrmw value
3615 being acquired. being acquired.
3617 3. buffer_gl0_inv
3619 - If OpenCL omit.
3620 - Ensures that
3621 following
3622 loads will not see
3623 stale data.
3625 atomicrmw acquire - workgroup - generic 1. flat_atomic 1. flat_atomic
3626 2. waitcnt lgkmcnt(0) 2. waitcnt lgkmcnt(0) &
3627 vm/vscnt(0)
3629 - If CU wavefront execution mode, omit vm/vscnt.
3630 - If OpenCL, omit. - If OpenCL, omit
3631 waitcnt lgkmcnt(0)..
3632 - Use vmcnt if atomic with
3633 return and vscnt if atomic
3634 with no-return.
3635 waitcnt lgkmcnt(0).
3636 - Must happen before - Must happen before
3637 any following the following
3638 global/generic buffer_gl0_inv.
3639 load/load
3640 atomic/store/store
3641 atomic/atomicrmw.
3642 - Ensures any - Ensures any
3643 following global following global
3644 data read is no data read is no
3645 older than the older than the
3646 atomicrmw value atomicrmw value
3647 being acquired. being acquired.
3649 3. buffer_gl0_inv
3651 - If CU wavefront execution mode, omit.
3652 - Ensures that
3653 following
3654 loads will not see
3655 stale data.
3657 atomicrmw acquire - agent - global 1. buffer/global/flat_atomic 1. buffer/global_atomic
3658 - system 2. s_waitcnt vmcnt(0) 2. s_waitcnt vm/vscnt(0)
3660 - Use vmcnt if atomic with
3661 return and vscnt if atomic
3662 with no-return.
3663 waitcnt lgkmcnt(0).
3664 - Must happen before - Must happen before
3665 following following
3666 buffer_wbinvl1_vol. buffer_gl*_inv.
3667 - Ensures the - Ensures the
3668 atomicrmw has atomicrmw has
3669 completed before completed before
3670 invalidating the invalidating the
3671 cache. caches.
3673 3. buffer_wbinvl1_vol 3. buffer_gl0_inv;
3674 buffer_gl1_inv
3676 - Must happen before - Must happen before
3677 any following any following
3678 global/generic global/generic
3679 load/load load/load
3680 atomic/atomicrmw. atomic/atomicrmw.
3681 - Ensures that - Ensures that
3682 following loads following loads
3683 will not see stale will not see stale
3684 global data. global data.
3686 atomicrmw acquire - agent - generic 1. flat_atomic 1. flat_atomic
3687 - system 2. s_waitcnt vmcnt(0) & 2. s_waitcnt vm/vscnt(0) &
3688 lgkmcnt(0) lgkmcnt(0)
3690 - If OpenCL, omit - If OpenCL, omit
3691 lgkmcnt(0). lgkmcnt(0).
3692 - Use vmcnt if atomic with
3693 return and vscnt if atomic
3694 with no-return.
3695 - Must happen before - Must happen before
3696 following following
3697 buffer_wbinvl1_vol. buffer_gl*_inv.
3698 - Ensures the - Ensures the
3699 atomicrmw has atomicrmw has
3700 completed before completed before
3701 invalidating the invalidating the
3702 cache. caches.
3704 3. buffer_wbinvl1_vol 3. buffer_gl0_inv;
3705 buffer_gl1_inv
3707 - Must happen before - Must happen before
3708 any following any following
3709 global/generic global/generic
3710 load/load load/load
3711 atomic/atomicrmw. atomic/atomicrmw.
3712 - Ensures that - Ensures that
3713 following loads following loads
3714 will not see stale will not see stale
3715 global data. global data.
3717 fence acquire - singlethread *none* *none* *none*
3718 - wavefront
3719 fence acquire - workgroup *none* 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
3720 vmcnt(0) & vscnt(0)
3722 - If CU wavefront execution mode, omit vmcnt and
3723 vscnt.
3724 - If OpenCL and - If OpenCL and
3725 address space is address space is
3726 not generic, omit. not generic, omit
3727 lgkmcnt(0).
3728 - If OpenCL and
3729 address space is
3730 local, omit
3731 vmcnt(0) and vscnt(0).
3732 - However, since LLVM - However, since LLVM
3733 currently has no currently has no
3734 address space on address space on
3735 the fence need to the fence need to
3736 conservatively conservatively
3737 always generate. If always generate. If
3738 fence had an fence had an
3739 address space then address space then
3740 set to address set to address
3741 space of OpenCL space of OpenCL
3742 fence flag, or to fence flag, or to
3743 generic if both generic if both
3744 local and global local and global
3745 flags are flags are
3746 specified. specified.
3747 - Must happen after
3748 any preceding
3749 local/generic load
3750 atomic/atomicrmw
3751 with an equal or
3752 wider sync scope
3753 and memory ordering
3754 stronger than
3755 unordered (this is
3756 termed the
3757 fence-paired-atomic).
3758 - Must happen before
3759 any following
3760 global/generic
3761 load/load
3762 atomic/store/store
3763 atomic/atomicrmw.
3764 - Ensures any
3765 following global
3766 data read is no
3767 older than the
3768 value read by the
3769 fence-paired-atomic.
3770 - Could be split into
3771 separate s_waitcnt
3772 vmcnt(0), s_waitcnt
3773 vscnt(0) and s_waitcnt
3774 lgkmcnt(0) to allow
3775 them to be
3776 independently moved
3777 according to the
3778 following rules.
3779 - s_waitcnt vmcnt(0)
3780 must happen after
3781 any preceding
3782 global/generic load
3783 atomic/
3784 atomicrmw-with-return-value
3785 with an equal or
3786 wider sync scope
3787 and memory ordering
3788 stronger than
3789 unordered (this is
3790 termed the
3791 fence-paired-atomic).
3792 - s_waitcnt vscnt(0)
3793 must happen after
3794 any preceding
3795 global/generic
3796 atomicrmw-no-return-value
3797 with an equal or
3798 wider sync scope
3799 and memory ordering
3800 stronger than
3801 unordered (this is
3802 termed the
3803 fence-paired-atomic).
3804 - s_waitcnt lgkmcnt(0)
3805 must happen after
3806 any preceding
3807 local/generic load
3808 atomic/atomicrmw
3809 with an equal or
3810 wider sync scope
3811 and memory ordering
3812 stronger than
3813 unordered (this is
3814 termed the
3815 fence-paired-atomic).
3816 - Must happen before
3817 the following
3818 buffer_gl0_inv.
3819 - Ensures that the
3820 fence-paired atomic
3821 has completed
3822 before invalidating
3823 the
3824 cache. Therefore
3825 any following
3826 locations read must
3827 be no older than
3828 the value read by
3829 the
3830 fence-paired-atomic.
3832 3. buffer_gl0_inv
3834 - If CU wavefront execution mode, omit.
3835 - Ensures that
3836 following
3837 loads will not see
3838 stale data.
3840 fence acquire - agent *none* 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
3841 - system vmcnt(0) vmcnt(0) & vscnt(0)
3843 - If OpenCL and - If OpenCL and
3844 address space is address space is
3845 not generic, omit not generic, omit
3846 lgkmcnt(0). lgkmcnt(0).
3847 - If OpenCL and
3848 address space is
3849 local, omit
3850 vmcnt(0) and vscnt(0).
3851 - However, since LLVM - However, since LLVM
3852 currently has no currently has no
3853 address space on address space on
3854 the fence need to the fence need to
3855 conservatively conservatively
3856 always generate always generate
3857 (see comment for (see comment for
3858 previous fence). previous fence).
3859 - Could be split into
3860 separate s_waitcnt
3861 vmcnt(0) and
3862 s_waitcnt
3863 lgkmcnt(0) to allow
3864 them to be
3865 independently moved
3866 according to the
3867 following rules.
3868 - s_waitcnt vmcnt(0)
3869 must happen after
3870 any preceding
3871 global/generic load
3872 atomic/atomicrmw
3873 with an equal or
3874 wider sync scope
3875 and memory ordering
3876 stronger than
3877 unordered (this is
3878 termed the
3879 fence-paired-atomic).
3880 - s_waitcnt lgkmcnt(0)
3881 must happen after
3882 any preceding
3883 local/generic load
3884 atomic/atomicrmw
3885 with an equal or
3886 wider sync scope
3887 and memory ordering
3888 stronger than
3889 unordered (this is
3890 termed the
3891 fence-paired-atomic).
3892 - Must happen before
3893 the following
3894 buffer_wbinvl1_vol.
3895 - Ensures that the
3896 fence-paired atomic
3897 has completed
3898 before invalidating
3899 the
3900 cache. Therefore
3901 any following
3902 locations read must
3903 be no older than
3904 the value read by
3905 the
3906 fence-paired-atomic.
3907 - Could be split into
3908 separate s_waitcnt
3909 vmcnt(0), s_waitcnt
3910 vscnt(0) and s_waitcnt
3911 lgkmcnt(0) to allow
3912 them to be
3913 independently moved
3914 according to the
3915 following rules.
3916 - s_waitcnt vmcnt(0)
3917 must happen after
3918 any preceding
3919 global/generic load
3920 atomic/
3921 atomicrmw-with-return-value
3922 with an equal or
3923 wider sync scope
3924 and memory ordering
3925 stronger than
3926 unordered (this is
3927 termed the
3928 fence-paired-atomic).
3929 - s_waitcnt vscnt(0)
3930 must happen after
3931 any preceding
3932 global/generic
3933 atomicrmw-no-return-value
3934 with an equal or
3935 wider sync scope
3936 and memory ordering
3937 stronger than
3938 unordered (this is
3939 termed the
3940 fence-paired-atomic).
3941 - s_waitcnt lgkmcnt(0)
3942 must happen after
3943 any preceding
3944 local/generic load
3945 atomic/atomicrmw
3946 with an equal or
3947 wider sync scope
3948 and memory ordering
3949 stronger than
3950 unordered (this is
3951 termed the
3952 fence-paired-atomic).
3953 - Must happen before
3954 the following
3955 buffer_gl*_inv.
3956 - Ensures that the
3957 fence-paired atomic
3958 has completed
3959 before invalidating
3960 the
3961 caches. Therefore
3962 any following
3963 locations read must
3964 be no older than
3965 the value read by
3966 the
3967 fence-paired-atomic.
3969 2. buffer_wbinvl1_vol 2. buffer_gl0_inv;
3970 buffer_gl1_inv
3972 - Must happen before any - Must happen before any
3973 following global/generic following global/generic
3974 load/load load/load
3975 atomic/store/store atomic/store/store
3976 atomic/atomicrmw. atomic/atomicrmw.
3977 - Ensures that - Ensures that
3978 following loads following loads
3979 will not see stale will not see stale
3980 global data. global data.
3982 **Release Atomic**
3983 ----------------------------------------------------------------------------------------------------------------------
3984 store atomic release - singlethread - global 1. buffer/global/ds/flat_store 1. buffer/global/ds/flat_store
3985 - wavefront - local
3986 - generic
3987 store atomic release - workgroup - global 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
3988 vmcnt(0) & vscnt(0)
3990 - If CU wavefront execution mode, omit vmcnt and
3991 vscnt.
3992 - If OpenCL, omit. - If OpenCL, omit
3993 lgkmcnt(0).
3994 - Must happen after
3995 any preceding
3996 local/generic
3997 load/store/load
3998 atomic/store
3999 atomic/atomicrmw.
4000 - Could be split into
4001 separate s_waitcnt
4002 vmcnt(0), s_waitcnt
4003 vscnt(0) and s_waitcnt
4004 lgkmcnt(0) to allow
4005 them to be
4006 independently moved
4007 according to the
4008 following rules.
4009 - s_waitcnt vmcnt(0)
4010 must happen after
4011 any preceding
4012 global/generic load/load
4013 atomic/
4014 atomicrmw-with-return-value.
4015 - s_waitcnt vscnt(0)
4016 must happen after
4017 any preceding
4018 global/generic
4019 store/store
4020 atomic/
4021 atomicrmw-no-return-value.
4022 - s_waitcnt lgkmcnt(0)
4023 must happen after
4024 any preceding
4025 local/generic
4026 load/store/load
4027 atomic/store
4028 atomic/atomicrmw.
4029 - Must happen before - Must happen before
4030 the following the following
4031 store. store.
4032 - Ensures that all - Ensures that all
4033 memory operations memory operations
4034 to local have have
4035 completed before completed before
4036 performing the performing the
4037 store that is being store that is being
4038 released. released.
4040 2. buffer/global/flat_store 2. buffer/global_store
4041 store atomic release - workgroup - local 1. waitcnt vmcnt(0) & vscnt(0)
4043 - If CU wavefront execution mode, omit.
4044 - If OpenCL, omit.
4045 - Could be split into
4046 separate s_waitcnt
4047 vmcnt(0) and s_waitcnt
4048 vscnt(0) to allow
4049 them to be
4050 independently moved
4051 according to the
4052 following rules.
4053 - s_waitcnt vmcnt(0)
4054 must happen after
4055 any preceding
4056 global/generic load/load
4057 atomic/
4058 atomicrmw-with-return-value.
4059 - s_waitcnt vscnt(0)
4060 must happen after
4061 any preceding
4062 global/generic
4063 store/store atomic/
4064 atomicrmw-no-return-value.
4065 - Must happen before
4066 the following
4067 store.
4068 - Ensures that all
4069 global memory
4070 operations have
4071 completed before
4072 performing the
4073 store that is being
4074 released.
4076 1. ds_store 2. ds_store
4077 store atomic release - workgroup - generic 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4078 vmcnt(0) & vscnt(0)
4080 - If CU wavefront execution mode, omit vmcnt and
4081 vscnt.
4082 - If OpenCL, omit. - If OpenCL, omit
4083 lgkmcnt(0).
4084 - Must happen after
4085 any preceding
4086 local/generic
4087 load/store/load
4088 atomic/store
4089 atomic/atomicrmw.
4090 - Could be split into
4091 separate s_waitcnt
4092 vmcnt(0), s_waitcnt
4093 vscnt(0) and s_waitcnt
4094 lgkmcnt(0) to allow
4095 them to be
4096 independently moved
4097 according to the
4098 following rules.
4099 - s_waitcnt vmcnt(0)
4100 must happen after
4101 any preceding
4102 global/generic load/load
4103 atomic/
4104 atomicrmw-with-return-value.
4105 - s_waitcnt vscnt(0)
4106 must happen after
4107 any preceding
4108 global/generic
4109 store/store
4110 atomic/
4111 atomicrmw-no-return-value.
4112 - s_waitcnt lgkmcnt(0)
4113 must happen after
4114 any preceding
4115 local/generic load/store/load
4116 atomic/store atomic/atomicrmw.
4117 - Must happen before - Must happen before
4118 the following the following
4119 store. store.
4120 - Ensures that all - Ensures that all
4121 memory operations memory operations
4122 to local have have
4123 completed before completed before
4124 performing the performing the
4125 store that is being store that is being
4126 released. released.
4128 2. flat_store 2. flat_store
4129 store atomic release - agent - global 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
4130 - system - generic vmcnt(0) vmcnt(0) & vscnt(0)
4132 - If OpenCL, omit - If OpenCL, omit
4133 lgkmcnt(0). lgkmcnt(0).
4134 - Could be split into - Could be split into
4135 separate s_waitcnt separate s_waitcnt
4136 vmcnt(0) and vmcnt(0), s_waitcnt vscnt(0)
4137 s_waitcnt and s_waitcnt
4138 lgkmcnt(0) to allow lgkmcnt(0) to allow
4139 them to be them to be
4140 independently moved independently moved
4141 according to the according to the
4142 following rules. following rules.
4143 - s_waitcnt vmcnt(0) - s_waitcnt vmcnt(0)
4144 must happen after must happen after
4145 any preceding any preceding
4146 global/generic global/generic
4147 load/store/load load/load
4148 atomic/store atomic/
4149 atomic/atomicrmw. atomicrmw-with-return-value.
4150 - s_waitcnt vscnt(0)
4151 must happen after
4152 any preceding
4153 global/generic
4154 store/store atomic/
4155 atomicrmw-no-return-value.
4156 - s_waitcnt lgkmcnt(0) - s_waitcnt lgkmcnt(0)
4157 must happen after must happen after
4158 any preceding any preceding
4159 local/generic local/generic
4160 load/store/load load/store/load
4161 atomic/store atomic/store
4162 atomic/atomicrmw. atomic/atomicrmw.
4163 - Must happen before - Must happen before
4164 the following the following
4165 store. store.
4166 - Ensures that all - Ensures that all
4167 memory operations memory operations
4168 to memory have to memory have
4169 completed before completed before
4170 performing the performing the
4171 store that is being store that is being
4172 released. released.
4174 2. buffer/global/ds/flat_store 2. buffer/global/ds/flat_store
4175 atomicrmw release - singlethread - global 1. buffer/global/ds/flat_atomic 1. buffer/global/ds/flat_atomic
4176 - wavefront - local
4177 - generic
4178 atomicrmw release - workgroup - global 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4179 vmcnt(0) & vscnt(0)
4181 - If CU wavefront execution mode, omit vmcnt and
4182 vscnt.
4183 - If OpenCL, omit.
4185 - Must happen after
4186 any preceding
4187 local/generic
4188 load/store/load
4189 atomic/store
4190 atomic/atomicrmw.
4191 - Could be split into
4192 separate s_waitcnt
4193 vmcnt(0), s_waitcnt
4194 vscnt(0) and s_waitcnt
4195 lgkmcnt(0) to allow
4196 them to be
4197 independently moved
4198 according to the
4199 following rules.
4200 - s_waitcnt vmcnt(0)
4201 must happen after
4202 any preceding
4203 global/generic load/load
4204 atomic/
4205 atomicrmw-with-return-value.
4206 - s_waitcnt vscnt(0)
4207 must happen after
4208 any preceding
4209 global/generic
4210 store/store
4211 atomic/
4212 atomicrmw-no-return-value.
4213 - s_waitcnt lgkmcnt(0)
4214 must happen after
4215 any preceding
4216 local/generic
4217 load/store/load
4218 atomic/store
4219 atomic/atomicrmw.
4220 - Must happen before - Must happen before
4221 the following the following
4222 atomicrmw. atomicrmw.
4223 - Ensures that all - Ensures that all
4224 memory operations memory operations
4225 to local have have
4226 completed before completed before
4227 performing the performing the
4228 atomicrmw that is atomicrmw that is
4229 being released. being released.
4231 2. buffer/global/flat_atomic 2. buffer/global_atomic
4232 atomicrmw release - workgroup - local 1. waitcnt vmcnt(0) & vscnt(0)
4234 - If CU wavefront execution mode, omit.
4235 - If OpenCL, omit.
4236 - Could be split into
4237 separate s_waitcnt
4238 vmcnt(0) and s_waitcnt
4239 vscnt(0) to allow
4240 them to be
4241 independently moved
4242 according to the
4243 following rules.
4244 - s_waitcnt vmcnt(0)
4245 must happen after
4246 any preceding
4247 global/generic load/load
4248 atomic/
4249 atomicrmw-with-return-value.
4250 - s_waitcnt vscnt(0)
4251 must happen after
4252 any preceding
4253 global/generic
4254 store/store atomic/
4255 atomicrmw-no-return-value.
4256 - Must happen before
4257 the following
4258 store.
4259 - Ensures that all
4260 global memory
4261 operations have
4262 completed before
4263 performing the
4264 store that is being
4265 released.
4267 1. ds_atomic 2. ds_atomic
4268 atomicrmw release - workgroup - generic 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4269 vmcnt(0) & vscnt(0)
4271 - If CU wavefront execution mode, omit vmcnt and
4272 vscnt.
4273 - If OpenCL, omit. - If OpenCL, omit
4274 waitcnt lgkmcnt(0).
4275 - Must happen after
4276 any preceding
4277 local/generic
4278 load/store/load
4279 atomic/store
4280 atomic/atomicrmw.
4281 - Could be split into
4282 separate s_waitcnt
4283 vmcnt(0), s_waitcnt
4284 vscnt(0) and s_waitcnt
4285 lgkmcnt(0) to allow
4286 them to be
4287 independently moved
4288 according to the
4289 following rules.
4290 - s_waitcnt vmcnt(0)
4291 must happen after
4292 any preceding
4293 global/generic load/load
4294 atomic/
4295 atomicrmw-with-return-value.
4296 - s_waitcnt vscnt(0)
4297 must happen after
4298 any preceding
4299 global/generic
4300 store/store
4301 atomic/
4302 atomicrmw-no-return-value.
4303 - s_waitcnt lgkmcnt(0)
4304 must happen after
4305 any preceding
4306 local/generic load/store/load
4307 atomic/store atomic/atomicrmw.
4308 - Must happen before - Must happen before
4309 the following the following
4310 atomicrmw. atomicrmw.
4311 - Ensures that all - Ensures that all
4312 memory operations memory operations
4313 to local have have
4314 completed before completed before
4315 performing the performing the
4316 atomicrmw that is atomicrmw that is
4317 being released. being released.
4319 2. flat_atomic 2. flat_atomic
4320 atomicrmw release - agent - global 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lkkmcnt(0) &
4321 - system - generic vmcnt(0) vmcnt(0) & vscnt(0)
4323 - If OpenCL, omit - If OpenCL, omit
4324 lgkmcnt(0). lgkmcnt(0).
4325 - Could be split into - Could be split into
4326 separate s_waitcnt separate s_waitcnt
4327 vmcnt(0) and vmcnt(0), s_waitcnt
4328 s_waitcnt vscnt(0) and s_waitcnt
4329 lgkmcnt(0) to allow lgkmcnt(0) to allow
4330 them to be them to be
4331 independently moved independently moved
4332 according to the according to the
4333 following rules. following rules.
4334 - s_waitcnt vmcnt(0) - s_waitcnt vmcnt(0)
4335 must happen after must happen after
4336 any preceding any preceding
4337 global/generic global/generic
4338 load/store/load load/load atomic/
4339 atomic/store atomicrmw-with-return-value.
4340 atomic/atomicrmw.
4341 - s_waitcnt vscnt(0)
4342 must happen after
4343 any preceding
4344 global/generic
4345 store/store atomic/
4346 atomicrmw-no-return-value.
4347 - s_waitcnt lgkmcnt(0) - s_waitcnt lgkmcnt(0)
4348 must happen after must happen after
4349 any preceding any preceding
4350 local/generic local/generic
4351 load/store/load load/store/load
4352 atomic/store atomic/store
4353 atomic/atomicrmw. atomic/atomicrmw.
4354 - Must happen before - Must happen before
4355 the following the following
4356 atomicrmw. atomicrmw.
4357 - Ensures that all - Ensures that all
4358 memory operations memory operations
4359 to global and local to global and local
4360 have completed have completed
4361 before performing before performing
4362 the atomicrmw that the atomicrmw that
4363 is being released. is being released.
4365 2. buffer/global/ds/flat_atomic 2. buffer/global/ds/flat_atomic
4366 fence release - singlethread *none* *none* *none*
4367 - wavefront
4368 fence release - workgroup *none* 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4369 vmcnt(0) & vscnt(0)
4371 - If CU wavefront execution mode, omit vmcnt and
4372 vscnt.
4373 - If OpenCL and - If OpenCL and
4374 address space is address space is
4375 not generic, omit. not generic, omit
4376 lgkmcnt(0).
4377 - If OpenCL and
4378 address space is
4379 local, omit
4380 vmcnt(0) and vscnt(0).
4381 - However, since LLVM - However, since LLVM
4382 currently has no currently has no
4383 address space on address space on
4384 the fence need to the fence need to
4385 conservatively conservatively
4386 always generate. If always generate. If
4387 fence had an fence had an
4388 address space then address space then
4389 set to address set to address
4390 space of OpenCL space of OpenCL
4391 fence flag, or to fence flag, or to
4392 generic if both generic if both
4393 local and global local and global
4394 flags are flags are
4395 specified. specified.
4396 - Must happen after
4397 any preceding
4398 local/generic
4399 load/load
4400 atomic/store/store
4401 atomic/atomicrmw.
4402 - Could be split into
4403 separate s_waitcnt
4404 vmcnt(0), s_waitcnt
4405 vscnt(0) and s_waitcnt
4406 lgkmcnt(0) to allow
4407 them to be
4408 independently moved
4409 according to the
4410 following rules.
4411 - s_waitcnt vmcnt(0)
4412 must happen after
4413 any preceding
4414 global/generic
4415 load/load
4416 atomic/
4417 atomicrmw-with-return-value.
4418 - s_waitcnt vscnt(0)
4419 must happen after
4420 any preceding
4421 global/generic
4422 store/store atomic/
4423 atomicrmw-no-return-value.
4424 - s_waitcnt lgkmcnt(0)
4425 must happen after
4426 any preceding
4427 local/generic
4428 load/store/load
4429 atomic/store atomic/
4430 atomicrmw.
4431 - Must happen before - Must happen before
4432 any following store any following store
4433 atomic/atomicrmw atomic/atomicrmw
4434 with an equal or with an equal or
4435 wider sync scope wider sync scope
4436 and memory ordering and memory ordering
4437 stronger than stronger than
4438 unordered (this is unordered (this is
4439 termed the termed the
4440 fence-paired-atomic). fence-paired-atomic).
4441 - Ensures that all - Ensures that all
4442 memory operations memory operations
4443 to local have have
4444 completed before completed before
4445 performing the performing the
4446 following following
4447 fence-paired-atomic. fence-paired-atomic.
4449 fence release - agent *none* 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
4450 - system vmcnt(0) vmcnt(0) & vscnt(0)
4452 - If OpenCL and - If OpenCL and
4453 address space is address space is
4454 not generic, omit not generic, omit
4455 lgkmcnt(0). lgkmcnt(0).
4456 - If OpenCL and - If OpenCL and
4457 address space is address space is
4458 local, omit local, omit
4459 vmcnt(0). vmcnt(0) and vscnt(0).
4460 - However, since LLVM - However, since LLVM
4461 currently has no currently has no
4462 address space on address space on
4463 the fence need to the fence need to
4464 conservatively conservatively
4465 always generate. If always generate. If
4466 fence had an fence had an
4467 address space then address space then
4468 set to address set to address
4469 space of OpenCL space of OpenCL
4470 fence flag, or to fence flag, or to
4471 generic if both generic if both
4472 local and global local and global
4473 flags are flags are
4474 specified. specified.
4475 - Could be split into - Could be split into
4476 separate s_waitcnt separate s_waitcnt
4477 vmcnt(0) and vmcnt(0), s_waitcnt
4478 s_waitcnt vscnt(0) and s_waitcnt
4479 lgkmcnt(0) to allow lgkmcnt(0) to allow
4480 them to be them to be
4481 independently moved independently moved
4482 according to the according to the
4483 following rules. following rules.
4484 - s_waitcnt vmcnt(0) - s_waitcnt vmcnt(0)
4485 must happen after must happen after
4486 any preceding any preceding
4487 global/generic global/generic
4488 load/store/load load/load atomic/
4489 atomic/store atomicrmw-with-return-value.
4490 atomic/atomicrmw.
4491 - s_waitcnt vscnt(0)
4492 must happen after
4493 any preceding
4494 global/generic
4495 store/store atomic/
4496 atomicrmw-no-return-value.
4497 - s_waitcnt lgkmcnt(0) - s_waitcnt lgkmcnt(0)
4498 must happen after must happen after
4499 any preceding any preceding
4500 local/generic local/generic
4501 load/store/load load/store/load
4502 atomic/store atomic/store
4503 atomic/atomicrmw. atomic/atomicrmw.
4504 - Must happen before - Must happen before
4505 any following store any following store
4506 atomic/atomicrmw atomic/atomicrmw
4507 with an equal or with an equal or
4508 wider sync scope wider sync scope
4509 and memory ordering and memory ordering
4510 stronger than stronger than
4511 unordered (this is unordered (this is
4512 termed the termed the
4513 fence-paired-atomic). fence-paired-atomic).
4514 - Ensures that all - Ensures that all
4515 memory operations memory operations
4516 have have
4517 completed before completed before
4518 performing the performing the
4519 following following
4520 fence-paired-atomic. fence-paired-atomic.
4522 **Acquire-Release Atomic**
4523 ----------------------------------------------------------------------------------------------------------------------
4524 atomicrmw acq_rel - singlethread - global 1. buffer/global/ds/flat_atomic 1. buffer/global/ds/flat_atomic
4525 - wavefront - local
4526 - generic
4527 atomicrmw acq_rel - workgroup - global 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4528 vmcnt(0) & vscnt(0)
4530 - If CU wavefront execution mode, omit vmcnt and
4531 vscnt.
4532 - If OpenCL, omit. - If OpenCL, omit
4533 s_waitcnt lgkmcnt(0).
4534 - Must happen after - Must happen after
4535 any preceding any preceding
4536 local/generic local/generic
4537 load/store/load load/store/load
4538 atomic/store atomic/store
4539 atomic/atomicrmw. atomic/atomicrmw.
4540 - Could be split into
4541 separate s_waitcnt
4542 vmcnt(0), s_waitcnt
4543 vscnt(0) and s_waitcnt
4544 lgkmcnt(0) to allow
4545 them to be
4546 independently moved
4547 according to the
4548 following rules.
4549 - s_waitcnt vmcnt(0)
4550 must happen after
4551 any preceding
4552 global/generic load/load
4553 atomic/
4554 atomicrmw-with-return-value.
4555 - s_waitcnt vscnt(0)
4556 must happen after
4557 any preceding
4558 global/generic
4559 store/store
4560 atomic/
4561 atomicrmw-no-return-value.
4562 - s_waitcnt lgkmcnt(0)
4563 must happen after
4564 any preceding
4565 local/generic load/store/load
4566 atomic/store atomic/atomicrmw.
4567 - Must happen before - Must happen before
4568 the following the following
4569 atomicrmw. atomicrmw.
4570 - Ensures that all - Ensures that all
4571 memory operations memory operations
4572 to local have have
4573 completed before completed before
4574 performing the performing the
4575 atomicrmw that is atomicrmw that is
4576 being released. being released.
4578 2. buffer/global/flat_atomic 2. buffer/global_atomic
4579 3. s_waitcnt vm/vscnt(0)
4581 - If CU wavefront execution mode, omit vm/vscnt.
4582 - Use vmcnt if atomic with
4583 return and vscnt if atomic
4584 with no-return.
4585 waitcnt lgkmcnt(0).
4586 - Must happen before
4587 the following
4588 buffer_gl0_inv.
4589 - Ensures any
4590 following global
4591 data read is no
4592 older than the
4593 atomicrmw value
4594 being acquired.
4596 4. buffer_gl0_inv
4598 - If CU wavefront execution mode, omit.
4599 - Ensures that
4600 following
4601 loads will not see
4602 stale data.
4604 atomicrmw acq_rel - workgroup - local 1. waitcnt vmcnt(0) & vscnt(0)
4606 - If CU wavefront execution mode, omit.
4607 - If OpenCL, omit.
4608 - Could be split into
4609 separate s_waitcnt
4610 vmcnt(0) and s_waitcnt
4611 vscnt(0) to allow
4612 them to be
4613 independently moved
4614 according to the
4615 following rules.
4616 - s_waitcnt vmcnt(0)
4617 must happen after
4618 any preceding
4619 global/generic load/load
4620 atomic/
4621 atomicrmw-with-return-value.
4622 - s_waitcnt vscnt(0)
4623 must happen after
4624 any preceding
4625 global/generic
4626 store/store atomic/
4627 atomicrmw-no-return-value.
4628 - Must happen before
4629 the following
4630 store.
4631 - Ensures that all
4632 global memory
4633 operations have
4634 completed before
4635 performing the
4636 store that is being
4637 released.
4639 1. ds_atomic 2. ds_atomic
4640 2. s_waitcnt lgkmcnt(0) 3. s_waitcnt lgkmcnt(0)
4642 - If OpenCL, omit. - If OpenCL, omit.
4643 - Must happen before - Must happen before
4644 any following the following
4645 global/generic buffer_gl0_inv.
4646 load/load
4647 atomic/store/store
4648 atomic/atomicrmw.
4649 - Ensures any - Ensures any
4650 following global following global
4651 data read is no data read is no
4652 older than the load older than the load
4653 atomic value being atomic value being
4654 acquired. acquired.
4656 4. buffer_gl0_inv
4658 - If CU wavefront execution mode, omit.
4659 - If OpenCL omit.
4660 - Ensures that
4661 following
4662 loads will not see
4663 stale data.
4665 atomicrmw acq_rel - workgroup - generic 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4666 vmcnt(0) & vscnt(0)
4668 - If CU wavefront execution mode, omit vmcnt and
4669 vscnt.
4670 - If OpenCL, omit. - If OpenCL, omit
4671 waitcnt lgkmcnt(0).
4672 - Must happen after
4673 any preceding
4674 local/generic
4675 load/store/load
4676 atomic/store
4677 atomic/atomicrmw.
4678 - Could be split into
4679 separate s_waitcnt
4680 vmcnt(0), s_waitcnt
4681 vscnt(0) and s_waitcnt
4682 lgkmcnt(0) to allow
4683 them to be
4684 independently moved
4685 according to the
4686 following rules.
4687 - s_waitcnt vmcnt(0)
4688 must happen after
4689 any preceding
4690 global/generic load/load
4691 atomic/
4692 atomicrmw-with-return-value.
4693 - s_waitcnt vscnt(0)
4694 must happen after
4695 any preceding
4696 global/generic
4697 store/store
4698 atomic/
4699 atomicrmw-no-return-value.
4700 - s_waitcnt lgkmcnt(0)
4701 must happen after
4702 any preceding
4703 local/generic load/store/load
4704 atomic/store atomic/atomicrmw.
4705 - Must happen before - Must happen before
4706 the following the following
4707 atomicrmw. atomicrmw.
4708 - Ensures that all - Ensures that all
4709 memory operations memory operations
4710 to local have have
4711 completed before completed before
4712 performing the performing the
4713 atomicrmw that is atomicrmw that is
4714 being released. being released.
4716 2. flat_atomic 2. flat_atomic
4717 3. s_waitcnt lgkmcnt(0) 3. s_waitcnt lgkmcnt(0) &
4718 vm/vscnt(0)
4720 - If CU wavefront execution mode, omit vm/vscnt.
4721 - If OpenCL, omit. - If OpenCL, omit
4722 waitcnt lgkmcnt(0).
4723 - Must happen before - Must happen before
4724 any following the following
4725 global/generic buffer_gl0_inv.
4726 load/load
4727 atomic/store/store
4728 atomic/atomicrmw.
4729 - Ensures any - Ensures any
4730 following global following global
4731 data read is no data read is no
4732 older than the load older than the load
4733 atomic value being atomic value being
4734 acquired. acquired.
4736 3. buffer_gl0_inv
4738 - If CU wavefront execution mode, omit.
4739 - Ensures that
4740 following
4741 loads will not see
4742 stale data.
4744 atomicrmw acq_rel - agent - global 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
4745 - system vmcnt(0) vmcnt(0) & vscnt(0)
4747 - If OpenCL, omit - If OpenCL, omit
4748 lgkmcnt(0). lgkmcnt(0).
4749 - Could be split into - Could be split into
4750 separate s_waitcnt separate s_waitcnt
4751 vmcnt(0) and vmcnt(0), s_waitcnt
4752 s_waitcnt vscnt(0) and s_waitcnt
4753 lgkmcnt(0) to allow lgkmcnt(0) to allow
4754 them to be them to be
4755 independently moved independently moved
4756 according to the according to the
4757 following rules. following rules.
4758 - s_waitcnt vmcnt(0) - s_waitcnt vmcnt(0)
4759 must happen after must happen after
4760 any preceding any preceding
4761 global/generic global/generic
4762 load/store/load load/load atomic/
4763 atomic/store atomicrmw-with-return-value.
4764 atomic/atomicrmw.
4765 - s_waitcnt vscnt(0)
4766 must happen after
4767 any preceding
4768 global/generic
4769 store/store atomic/
4770 atomicrmw-no-return-value.
4771 - s_waitcnt lgkmcnt(0) - s_waitcnt lgkmcnt(0)
4772 must happen after must happen after
4773 any preceding any preceding
4774 local/generic local/generic
4775 load/store/load load/store/load
4776 atomic/store atomic/store
4777 atomic/atomicrmw. atomic/atomicrmw.
4778 - Must happen before - Must happen before
4779 the following the following
4780 atomicrmw. atomicrmw.
4781 - Ensures that all - Ensures that all
4782 memory operations memory operations
4783 to global have to global have
4784 completed before completed before
4785 performing the performing the
4786 atomicrmw that is atomicrmw that is
4787 being released. being released.
4789 2. buffer/global/flat_atomic 2. buffer/global_atomic
4790 3. s_waitcnt vmcnt(0) 3. s_waitcnt vm/vscnt(0)
4792 - Use vmcnt if atomic with
4793 return and vscnt if atomic
4794 with no-return.
4795 waitcnt lgkmcnt(0).
4796 - Must happen before - Must happen before
4797 following following
4798 buffer_wbinvl1_vol. buffer_gl*_inv.
4799 - Ensures the - Ensures the
4800 atomicrmw has atomicrmw has
4801 completed before completed before
4802 invalidating the invalidating the
4803 cache. caches.
4805 4. buffer_wbinvl1_vol 4. buffer_gl0_inv;
4806 buffer_gl1_inv
4808 - Must happen before - Must happen before
4809 any following any following
4810 global/generic global/generic
4811 load/load load/load
4812 atomic/atomicrmw. atomic/atomicrmw.
4813 - Ensures that - Ensures that
4814 following loads following loads
4815 will not see stale will not see stale
4816 global data. global data.
4818 atomicrmw acq_rel - agent - generic 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
4819 - system vmcnt(0) vmcnt(0) & vscnt(0)
4821 - If OpenCL, omit - If OpenCL, omit
4822 lgkmcnt(0). lgkmcnt(0).
4823 - Could be split into - Could be split into
4824 separate s_waitcnt separate s_waitcnt
4825 vmcnt(0) and vmcnt(0), s_waitcnt
4826 s_waitcnt vscnt(0) and s_waitcnt
4827 lgkmcnt(0) to allow lgkmcnt(0) to allow
4828 them to be them to be
4829 independently moved independently moved
4830 according to the according to the
4831 following rules. following rules.
4832 - s_waitcnt vmcnt(0) - s_waitcnt vmcnt(0)
4833 must happen after must happen after
4834 any preceding any preceding
4835 global/generic global/generic
4836 load/store/load load/load atomic
4837 atomic/store atomicrmw-with-return-value.
4838 atomic/atomicrmw.
4839 - s_waitcnt vscnt(0)
4840 must happen after
4841 any preceding
4842 global/generic
4843 store/store atomic/
4844 atomicrmw-no-return-value.
4845 - s_waitcnt lgkmcnt(0) - s_waitcnt lgkmcnt(0)
4846 must happen after must happen after
4847 any preceding any preceding
4848 local/generic local/generic
4849 load/store/load load/store/load
4850 atomic/store atomic/store
4851 atomic/atomicrmw. atomic/atomicrmw.
4852 - Must happen before - Must happen before
4853 the following the following
4854 atomicrmw. atomicrmw.
4855 - Ensures that all - Ensures that all
4856 memory operations memory operations
4857 to global have have
4858 completed before completed before
4859 performing the performing the
4860 atomicrmw that is atomicrmw that is
4861 being released. being released.
4863 2. flat_atomic 2. flat_atomic
4864 3. s_waitcnt vmcnt(0) & 3. s_waitcnt vm/vscnt(0) &
4865 lgkmcnt(0) lgkmcnt(0)
4867 - If OpenCL, omit - If OpenCL, omit
4868 lgkmcnt(0). lgkmcnt(0).
4869 - Use vmcnt if atomic with
4870 return and vscnt if atomic
4871 with no-return.
4872 - Must happen before - Must happen before
4873 following following
4874 buffer_wbinvl1_vol. buffer_gl*_inv.
4875 - Ensures the - Ensures the
4876 atomicrmw has atomicrmw has
4877 completed before completed before
4878 invalidating the invalidating the
4879 cache. caches.
4881 4. buffer_wbinvl1_vol 4. buffer_gl0_inv;
4882 buffer_gl1_inv
4884 - Must happen before - Must happen before
4885 any following any following
4886 global/generic global/generic
4887 load/load load/load
4888 atomic/atomicrmw. atomic/atomicrmw.
4889 - Ensures that - Ensures that
4890 following loads following loads
4891 will not see stale will not see stale
4892 global data. global data.
4894 fence acq_rel - singlethread *none* *none* *none*
4895 - wavefront
4896 fence acq_rel - workgroup *none* 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
4897 vmcnt(0) & vscnt(0)
4899 - If CU wavefront execution mode, omit vmcnt and
4900 vscnt.
4901 - If OpenCL and - If OpenCL and
4902 address space is address space is
4903 not generic, omit. not generic, omit
4904 lgkmcnt(0).
4905 - If OpenCL and
4906 address space is
4907 local, omit
4908 vmcnt(0) and vscnt(0).
4909 - However, - However,
4910 since LLVM since LLVM
4911 currently has no currently has no
4912 address space on address space on
4913 the fence need to the fence need to
4914 conservatively conservatively
4915 always generate always generate
4916 (see comment for (see comment for
4917 previous fence). previous fence).
4918 - Must happen after
4919 any preceding
4920 local/generic
4921 load/load
4922 atomic/store/store
4923 atomic/atomicrmw.
4924 - Could be split into
4925 separate s_waitcnt
4926 vmcnt(0), s_waitcnt
4927 vscnt(0) and s_waitcnt
4928 lgkmcnt(0) to allow
4929 them to be
4930 independently moved
4931 according to the
4932 following rules.
4933 - s_waitcnt vmcnt(0)
4934 must happen after
4935 any preceding
4936 global/generic
4937 load/load
4938 atomic/
4939 atomicrmw-with-return-value.
4940 - s_waitcnt vscnt(0)
4941 must happen after
4942 any preceding
4943 global/generic
4944 store/store atomic/
4945 atomicrmw-no-return-value.
4946 - s_waitcnt lgkmcnt(0)
4947 must happen after
4948 any preceding
4949 local/generic
4950 load/store/load
4951 atomic/store atomic/
4952 atomicrmw.
4953 - Must happen before - Must happen before
4954 any following any following
4955 global/generic global/generic
4956 load/load load/load
4957 atomic/store/store atomic/store/store
4958 atomic/atomicrmw. atomic/atomicrmw.
4959 - Ensures that all - Ensures that all
4960 memory operations memory operations
4961 to local have have
4962 completed before completed before
4963 performing any performing any
4964 following global following global
4965 memory operations. memory operations.
4966 - Ensures that the - Ensures that the
4967 preceding preceding
4968 local/generic load local/generic load
4969 atomic/atomicrmw atomic/atomicrmw
4970 with an equal or with an equal or
4971 wider sync scope wider sync scope
4972 and memory ordering and memory ordering
4973 stronger than stronger than
4974 unordered (this is unordered (this is
4975 termed the termed the
4976 acquire-fence-paired-atomic acquire-fence-paired-atomic
4977 ) has completed ) has completed
4978 before following before following
4979 global memory global memory
4980 operations. This operations. This
4981 satisfies the satisfies the
4982 requirements of requirements of
4983 acquire. acquire.
4984 - Ensures that all - Ensures that all
4985 previous memory previous memory
4986 operations have operations have
4987 completed before a completed before a
4988 following following
4989 local/generic store local/generic store
4990 atomic/atomicrmw atomic/atomicrmw
4991 with an equal or with an equal or
4992 wider sync scope wider sync scope
4993 and memory ordering and memory ordering
4994 stronger than stronger than
4995 unordered (this is unordered (this is
4996 termed the termed the
4997 release-fence-paired-atomic release-fence-paired-atomic
4998 ). This satisfies the ). This satisfies the
4999 requirements of requirements of
5000 release. release.
5001 - Must happen before
5002 the following
5003 buffer_gl0_inv.
5004 - Ensures that the
5005 acquire-fence-paired
5006 atomic has completed
5007 before invalidating
5008 the
5009 cache. Therefore
5010 any following
5011 locations read must
5012 be no older than
5013 the value read by
5014 the
5015 acquire-fence-paired-atomic.
5017 3. buffer_gl0_inv
5019 - If CU wavefront execution mode, omit.
5020 - Ensures that
5021 following
5022 loads will not see
5023 stale data.
5025 fence acq_rel - agent *none* 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
5026 - system vmcnt(0) vmcnt(0) & vscnt(0)
5028 - If OpenCL and - If OpenCL and
5029 address space is address space is
5030 not generic, omit not generic, omit
5031 lgkmcnt(0). lgkmcnt(0).
5032 - If OpenCL and
5033 address space is
5034 local, omit
5035 vmcnt(0) and vscnt(0).
5036 - However, since LLVM - However, since LLVM
5037 currently has no currently has no
5038 address space on address space on
5039 the fence need to the fence need to
5040 conservatively conservatively
5041 always generate always generate
5042 (see comment for (see comment for
5043 previous fence). previous fence).
5044 - Could be split into - Could be split into
5045 separate s_waitcnt separate s_waitcnt
5046 vmcnt(0) and vmcnt(0), s_waitcnt
5047 s_waitcnt vscnt(0) and s_waitcnt
5048 lgkmcnt(0) to allow lgkmcnt(0) to allow
5049 them to be them to be
5050 independently moved independently moved
5051 according to the according to the
5052 following rules. following rules.
5053 - s_waitcnt vmcnt(0) - s_waitcnt vmcnt(0)
5054 must happen after must happen after
5055 any preceding any preceding
5056 global/generic global/generic
5057 load/store/load load/load
5058 atomic/store atomic/
5059 atomic/atomicrmw. atomicrmw-with-return-value.
5060 - s_waitcnt vscnt(0)
5061 must happen after
5062 any preceding
5063 global/generic
5064 store/store atomic/
5065 atomicrmw-no-return-value.
5066 - s_waitcnt lgkmcnt(0) - s_waitcnt lgkmcnt(0)
5067 must happen after must happen after
5068 any preceding any preceding
5069 local/generic local/generic
5070 load/store/load load/store/load
5071 atomic/store atomic/store
5072 atomic/atomicrmw. atomic/atomicrmw.
5073 - Must happen before - Must happen before
5074 the following the following
5075 buffer_wbinvl1_vol. buffer_gl*_inv.
5076 - Ensures that the - Ensures that the
5077 preceding preceding
5078 global/local/generic global/local/generic
5079 load load
5080 atomic/atomicrmw atomic/atomicrmw
5081 with an equal or with an equal or
5082 wider sync scope wider sync scope
5083 and memory ordering and memory ordering
5084 stronger than stronger than
5085 unordered (this is unordered (this is
5086 termed the termed the
5087 acquire-fence-paired-atomic acquire-fence-paired-atomic
5088 ) has completed ) has completed
5089 before invalidating before invalidating
5090 the cache. This the caches. This
5091 satisfies the satisfies the
5092 requirements of requirements of
5093 acquire. acquire.
5094 - Ensures that all - Ensures that all
5095 previous memory previous memory
5096 operations have operations have
5097 completed before a completed before a
5098 following following
5099 global/local/generic global/local/generic
5100 store store
5101 atomic/atomicrmw atomic/atomicrmw
5102 with an equal or with an equal or
5103 wider sync scope wider sync scope
5104 and memory ordering and memory ordering
5105 stronger than stronger than
5106 unordered (this is unordered (this is
5107 termed the termed the
5108 release-fence-paired-atomic release-fence-paired-atomic
5109 ). This satisfies the ). This satisfies the
5110 requirements of requirements of
5111 release. release.
5113 2. buffer_wbinvl1_vol 2. buffer_gl0_inv;
5114 buffer_gl1_inv
5116 - Must happen before - Must happen before
5117 any following any following
5118 global/generic global/generic
5119 load/load load/load
5120 atomic/store/store atomic/store/store
5121 atomic/atomicrmw. atomic/atomicrmw.
5122 - Ensures that - Ensures that
5123 following loads following loads
5124 will not see stale will not see stale
5125 global data. This global data. This
5126 satisfies the satisfies the
5127 requirements of requirements of
5128 acquire. acquire.
5130 **Sequential Consistent Atomic**
5131 ----------------------------------------------------------------------------------------------------------------------
5132 load atomic seq_cst - singlethread - global *Same as corresponding *Same as corresponding
5133 - wavefront - local load atomic acquire, load atomic acquire,
5134 - generic except must generated except must generated
5135 all instructions even all instructions even
5136 for OpenCL.* for OpenCL.*
5137 load atomic seq_cst - workgroup - global 1. s_waitcnt lgkmcnt(0) 1. s_waitcnt lgkmcnt(0) &
5138 - generic vmcnt(0) & vscnt(0)
5140 - If CU wavefront execution mode, omit vmcnt and
5141 vscnt.
5142 - Could be split into
5143 separate s_waitcnt
5144 vmcnt(0), s_waitcnt
5145 vscnt(0) and s_waitcnt
5146 lgkmcnt(0) to allow
5147 them to be
5148 independently moved
5149 according to the
5150 following rules.
5151 - Must - waitcnt lgkmcnt(0) must
5152 happen after happen after
5153 preceding preceding
5154 global/generic load local load
5155 atomic/store atomic/store
5156 atomic/atomicrmw atomic/atomicrmw
5157 with memory with memory
5158 ordering of seq_cst ordering of seq_cst
5159 and with equal or and with equal or
5160 wider sync scope. wider sync scope.
5161 (Note that seq_cst (Note that seq_cst
5162 fences have their fences have their
5163 own s_waitcnt own s_waitcnt
5164 lgkmcnt(0) and so do lgkmcnt(0) and so do
5165 not need to be not need to be
5166 considered.) considered.)
5167 - waitcnt vmcnt(0)
5168 Must happen after
5169 preceding
5170 global/generic load
5171 atomic/
5172 atomicrmw-with-return-value
5173 with memory
5174 ordering of seq_cst
5175 and with equal or
5176 wider sync scope.
5177 (Note that seq_cst
5178 fences have their
5179 own s_waitcnt
5180 vmcnt(0) and so do
5181 not need to be
5182 considered.)
5183 - waitcnt vscnt(0)
5184 Must happen after
5185 preceding
5186 global/generic store
5187 atomic/
5188 atomicrmw-no-return-value
5189 with memory
5190 ordering of seq_cst
5191 and with equal or
5192 wider sync scope.
5193 (Note that seq_cst
5194 fences have their
5195 own s_waitcnt
5196 vscnt(0) and so do
5197 not need to be
5198 considered.)
5199 - Ensures any - Ensures any
5200 preceding preceding
5201 sequential sequential
5202 consistent local consistent global/local
5203 memory instructions memory instructions
5204 have completed have completed
5205 before executing before executing
5206 this sequentially this sequentially
5207 consistent consistent
5208 instruction. This instruction. This
5209 prevents reordering prevents reordering
5210 a seq_cst store a seq_cst store
5211 followed by a followed by a
5212 seq_cst load. (Note seq_cst load. (Note
5213 that seq_cst is that seq_cst is
5214 stronger than stronger than
5215 acquire/release as acquire/release as
5216 the reordering of the reordering of
5217 load acquire load acquire
5218 followed by a store followed by a store
5219 release is release is
5220 prevented by the prevented by the
5221 waitcnt of waitcnt of
5222 the release, but the release, but
5223 there is nothing there is nothing
5224 preventing a store preventing a store
5225 release followed by release followed by
5226 load acquire from load acquire from
5227 competing out of competing out of
5228 order.) order.)
5230 2. *Following 2. *Following
5231 instructions same as instructions same as
5232 corresponding load corresponding load
5233 atomic acquire, atomic acquire,
5234 except must generated except must generated
5235 all instructions even all instructions even
5236 for OpenCL.* for OpenCL.*
5237 load atomic seq_cst - workgroup - local *Same as corresponding
5238 load atomic acquire,
5239 except must generated
5240 all instructions even
5241 for OpenCL.*
5243 1. s_waitcnt vmcnt(0) & vscnt(0)
5245 - If CU wavefront execution mode, omit.
5246 - Could be split into
5247 separate s_waitcnt
5248 vmcnt(0) and s_waitcnt
5249 vscnt(0) to allow
5250 them to be
5251 independently moved
5252 according to the
5253 following rules.
5254 - waitcnt vmcnt(0)
5255 Must happen after
5256 preceding
5257 global/generic load
5258 atomic/
5259 atomicrmw-with-return-value
5260 with memory
5261 ordering of seq_cst
5262 and with equal or
5263 wider sync scope.
5264 (Note that seq_cst
5265 fences have their
5266 own s_waitcnt
5267 vmcnt(0) and so do
5268 not need to be
5269 considered.)
5270 - waitcnt vscnt(0)
5271 Must happen after
5272 preceding
5273 global/generic store
5274 atomic/
5275 atomicrmw-no-return-value
5276 with memory
5277 ordering of seq_cst
5278 and with equal or
5279 wider sync scope.
5280 (Note that seq_cst
5281 fences have their
5282 own s_waitcnt
5283 vscnt(0) and so do
5284 not need to be
5285 considered.)
5286 - Ensures any
5287 preceding
5288 sequential
5289 consistent global
5290 memory instructions
5291 have completed
5292 before executing
5293 this sequentially
5294 consistent
5295 instruction. This
5296 prevents reordering
5297 a seq_cst store
5298 followed by a
5299 seq_cst load. (Note
5300 that seq_cst is
5301 stronger than
5302 acquire/release as
5303 the reordering of
5304 load acquire
5305 followed by a store
5306 release is
5307 prevented by the
5308 waitcnt of
5309 the release, but
5310 there is nothing
5311 preventing a store
5312 release followed by
5313 load acquire from
5314 competing out of
5315 order.)
5317 2. *Following
5318 instructions same as
5319 corresponding load
5320 atomic acquire,
5321 except must generated
5322 all instructions even
5323 for OpenCL.*
5325 load atomic seq_cst - agent - global 1. s_waitcnt lgkmcnt(0) & 1. s_waitcnt lgkmcnt(0) &
5326 - system - generic vmcnt(0) vmcnt(0) & vscnt(0)
5328 - Could be split into - Could be split into
5329 separate s_waitcnt separate s_waitcnt
5330 vmcnt(0) vmcnt(0), s_waitcnt
5331 and s_waitcnt vscnt(0) and s_waitcnt
5332 lgkmcnt(0) to allow lgkmcnt(0) to allow
5333 them to be them to be
5334 independently moved independently moved
5335 according to the according to the
5336 following rules. following rules.
5337 - waitcnt lgkmcnt(0) - waitcnt lgkmcnt(0)
5338 must happen after must happen after
5339 preceding preceding
5340 global/generic load local load
5341 atomic/store atomic/store
5342 atomic/atomicrmw atomic/atomicrmw
5343 with memory with memory
5344 ordering of seq_cst ordering of seq_cst
5345 and with equal or and with equal or
5346 wider sync scope. wider sync scope.
5347 (Note that seq_cst (Note that seq_cst
5348 fences have their fences have their
5349 own s_waitcnt own s_waitcnt
5350 lgkmcnt(0) and so do lgkmcnt(0) and so do
5351 not need to be not need to be
5352 considered.) considered.)
5353 - waitcnt vmcnt(0) - waitcnt vmcnt(0)
5354 must happen after must happen after
5355 preceding preceding
5356 global/generic load global/generic load
5357 atomic/store atomic/
5358 atomic/atomicrmw atomicrmw-with-return-value
5359 with memory with memory
5360 ordering of seq_cst ordering of seq_cst
5361 and with equal or and with equal or
5362 wider sync scope. wider sync scope.
5363 (Note that seq_cst (Note that seq_cst
5364 fences have their fences have their
5365 own s_waitcnt own s_waitcnt
5366 vmcnt(0) and so do vmcnt(0) and so do
5367 not need to be not need to be
5368 considered.) considered.)
5369 - waitcnt vscnt(0)
5370 Must happen after
5371 preceding
5372 global/generic store
5373 atomic/
5374 atomicrmw-no-return-value
5375 with memory
5376 ordering of seq_cst
5377 and with equal or
5378 wider sync scope.
5379 (Note that seq_cst
5380 fences have their
5381 own s_waitcnt
5382 vscnt(0) and so do
5383 not need to be
5384 considered.)
5385 - Ensures any - Ensures any
5386 preceding preceding
5387 sequential sequential
5388 consistent global consistent global
5389 memory instructions memory instructions
5390 have completed have completed
5391 before executing before executing
5392 this sequentially this sequentially
5393 consistent consistent
5394 instruction. This instruction. This
5395 prevents reordering prevents reordering
5396 a seq_cst store a seq_cst store
5397 followed by a followed by a
5398 seq_cst load. (Note seq_cst load. (Note
5399 that seq_cst is that seq_cst is
5400 stronger than stronger than
5401 acquire/release as acquire/release as
5402 the reordering of the reordering of
5403 load acquire load acquire
5404 followed by a store followed by a store
5405 release is release is
5406 prevented by the prevented by the
5407 waitcnt of waitcnt of
5408 the release, but the release, but
5409 there is nothing there is nothing
5410 preventing a store preventing a store
5411 release followed by release followed by
5412 load acquire from load acquire from
5413 competing out of competing out of
5414 order.) order.)
5416 2. *Following 2. *Following
5417 instructions same as instructions same as
5418 corresponding load corresponding load
5419 atomic acquire, atomic acquire,
5420 except must generated except must generated
5421 all instructions even all instructions even
5422 for OpenCL.* for OpenCL.*
5423 store atomic seq_cst - singlethread - global *Same as corresponding *Same as corresponding
5424 - wavefront - local store atomic release, store atomic release,
5425 - workgroup - generic except must generated except must generated
5426 all instructions even all instructions even
5427 for OpenCL.* for OpenCL.*
5428 store atomic seq_cst - agent - global *Same as corresponding *Same as corresponding
5429 - system - generic store atomic release, store atomic release,
5430 except must generated except must generated
5431 all instructions even all instructions even
5432 for OpenCL.* for OpenCL.*
5433 atomicrmw seq_cst - singlethread - global *Same as corresponding *Same as corresponding
5434 - wavefront - local atomicrmw acq_rel, atomicrmw acq_rel,
5435 - workgroup - generic except must generated except must generated
5436 all instructions even all instructions even
5437 for OpenCL.* for OpenCL.*
5438 atomicrmw seq_cst - agent - global *Same as corresponding *Same as corresponding
5439 - system - generic atomicrmw acq_rel, atomicrmw acq_rel,
5440 except must generated except must generated
5441 all instructions even all instructions even
5442 for OpenCL.* for OpenCL.*
5443 fence seq_cst - singlethread *none* *Same as corresponding *Same as corresponding
5444 - wavefront fence acq_rel, fence acq_rel,
5445 - workgroup except must generated except must generated
5446 - agent all instructions even all instructions even
5447 - system for OpenCL.* for OpenCL.*
5448 ============ ============ ============== ========== =============================== ==================================
5450 The memory order also adds the single thread optimization constrains defined in
5451 table
5452 :ref:`amdgpu-amdhsa-memory-model-single-thread-optimization-constraints-gfx6-gfx10-table`.
5454 .. table:: AMDHSA Memory Model Single Thread Optimization Constraints GFX6-GFX10
5455 :name: amdgpu-amdhsa-memory-model-single-thread-optimization-constraints-gfx6-gfx10-table
5457 ============ ==============================================================
5458 LLVM Memory Optimization Constraints
5459 Ordering
5460 ============ ==============================================================
5461 unordered *none*
5462 monotonic *none*
5463 acquire - If a load atomic/atomicrmw then no following load/load
5464 atomic/store/ store atomic/atomicrmw/fence instruction can
5465 be moved before the acquire.
5466 - If a fence then same as load atomic, plus no preceding
5467 associated fence-paired-atomic can be moved after the fence.
5468 release - If a store atomic/atomicrmw then no preceding load/load
5469 atomic/store/ store atomic/atomicrmw/fence instruction can
5470 be moved after the release.
5471 - If a fence then same as store atomic, plus no following
5472 associated fence-paired-atomic can be moved before the
5473 fence.
5474 acq_rel Same constraints as both acquire and release.
5475 seq_cst - If a load atomic then same constraints as acquire, plus no
5476 preceding sequentially consistent load atomic/store
5477 atomic/atomicrmw/fence instruction can be moved after the
5478 seq_cst.
5479 - If a store atomic then the same constraints as release, plus
5480 no following sequentially consistent load atomic/store
5481 atomic/atomicrmw/fence instruction can be moved before the
5482 seq_cst.
5483 - If an atomicrmw/fence then same constraints as acq_rel.
5484 ============ ==============================================================
5486 Trap Handler ABI
5487 ~~~~~~~~~~~~~~~~
5489 For code objects generated by AMDGPU backend for HSA [HSA]_ compatible runtimes
5490 (such as ROCm [AMD-ROCm]_), the runtime installs a trap handler that supports
5491 the ``s_trap`` instruction with the following usage:
5493 .. table:: AMDGPU Trap Handler for AMDHSA OS
5494 :name: amdgpu-trap-handler-for-amdhsa-os-table
5496 =================== =============== =============== =======================
5497 Usage Code Sequence Trap Handler Description
5498 Inputs
5499 =================== =============== =============== =======================
5500 reserved ``s_trap 0x00`` Reserved by hardware.
5501 ``debugtrap(arg)`` ``s_trap 0x01`` ``SGPR0-1``: Reserved for HSA
5502 ``queue_ptr`` ``debugtrap``
5503 ``VGPR0``: intrinsic (not
5504 ``arg`` implemented).
5505 ``llvm.trap`` ``s_trap 0x02`` ``SGPR0-1``: Causes dispatch to be
5506 ``queue_ptr`` terminated and its
5507 associated queue put
5508 into the error state.
5509 ``llvm.debugtrap`` ``s_trap 0x03`` - If debugger not
5510 installed then
5511 behaves as a
5512 no-operation. The
5513 trap handler is
5514 entered and
5515 immediately returns
5516 to continue
5517 execution of the
5518 wavefront.
5519 - If the debugger is
5520 installed, causes
5521 the debug trap to be
5522 reported by the
5523 debugger and the
5524 wavefront is put in
5525 the halt state until
5526 resumed by the
5527 debugger.
5528 reserved ``s_trap 0x04`` Reserved.
5529 reserved ``s_trap 0x05`` Reserved.
5530 reserved ``s_trap 0x06`` Reserved.
5531 debugger breakpoint ``s_trap 0x07`` Reserved for debugger
5532 breakpoints.
5533 reserved ``s_trap 0x08`` Reserved.
5534 reserved ``s_trap 0xfe`` Reserved.
5535 reserved ``s_trap 0xff`` Reserved.
5536 =================== =============== =============== =======================
5538 AMDPAL
5539 ------
5541 This section provides code conventions used when the target triple OS is
5542 ``amdpal`` (see :ref:`amdgpu-target-triples`) for passing runtime parameters
5543 from the application/runtime to each invocation of a hardware shader. These
5544 parameters include both generic, application-controlled parameters called
5545 *user data* as well as system-generated parameters that are a product of the
5546 draw or dispatch execution.
5548 User Data
5549 ~~~~~~~~~
5551 Each hardware stage has a set of 32-bit *user data registers* which can be
5552 written from a command buffer and then loaded into SGPRs when waves are launched
5553 via a subsequent dispatch or draw operation. This is the way most arguments are
5554 passed from the application/runtime to a hardware shader.
5556 Compute User Data
5557 ~~~~~~~~~~~~~~~~~
5559 Compute shader user data mappings are simpler than graphics shaders, and have a
5560 fixed mapping.
5562 Note that there are always 10 available *user data entries* in registers -
5563 entries beyond that limit must be fetched from memory (via the spill table
5564 pointer) by the shader.
5566 .. table:: PAL Compute Shader User Data Registers
5567 :name: pal-compute-user-data-registers
5569 ============= ================================
5570 User Register Description
5571 ============= ================================
5572 0 Global Internal Table (32-bit pointer)
5573 1 Per-Shader Internal Table (32-bit pointer)
5574 2 - 11 Application-Controlled User Data (10 32-bit values)
5575 12 Spill Table (32-bit pointer)
5576 13 - 14 Thread Group Count (64-bit pointer)
5577 15 GDS Range
5578 ============= ================================
5580 Graphics User Data
5581 ~~~~~~~~~~~~~~~~~~
5583 Graphics pipelines support a much more flexible user data mapping:
5585 .. table:: PAL Graphics Shader User Data Registers
5586 :name: pal-graphics-user-data-registers
5588 ============= ================================
5589 User Register Description
5590 ============= ================================
5591 0 Global Internal Table (32-bit pointer)
5592 + Per-Shader Internal Table (32-bit pointer)
5593 + 1-15 Application Controlled User Data
5594 (1-15 Contiguous 32-bit Values in Registers)
5595 + Spill Table (32-bit pointer)
5596 + Draw Index (First Stage Only)
5597 + Vertex Offset (First Stage Only)
5598 + Instance Offset (First Stage Only)
5599 ============= ================================
5601 The placement of the global internal table remains fixed in the first *user
5602 data SGPR register*. Otherwise all parameters are optional, and can be mapped
5603 to any desired *user data SGPR register*, with the following regstrictions:
5605 * Draw Index, Vertex Offset, and Instance Offset can only be used by the first
5606 activehardware stage in a graphics pipeline (i.e. where the API vertex
5607 shader runs).
5609 * Application-controlled user data must be mapped into a contiguous range of
5610 user data registers.
5612 * The application-controlled user data range supports compaction remapping, so
5613 only *entries* that are actually consumed by the shader must be assigned to
5614 corresponding *registers*. Note that in order to support an efficient runtime
5615 implementation, the remapping must pack *registers* in the same order as
5616 *entries*, with unused *entries* removed.
5618 .. _pal_global_internal_table:
5620 Global Internal Table
5621 ~~~~~~~~~~~~~~~~~~~~~
5623 The global internal table is a table of *shader resource descriptors* (SRDs) that
5624 define how certain engine-wide, runtime-managed resources should be accessed
5625 from a shader. The majority of these resources have HW-defined formats, and it
5626 is up to the compiler to write/read data as required by the target hardware.
5628 The following table illustrates the required format:
5630 .. table:: PAL Global Internal Table
5631 :name: pal-git-table
5633 ============= ================================
5634 Offset Description
5635 ============= ================================
5636 0-3 Graphics Scratch SRD
5637 4-7 Compute Scratch SRD
5638 8-11 ES/GS Ring Output SRD
5639 12-15 ES/GS Ring Input SRD
5640 16-19 GS/VS Ring Output #0
5641 20-23 GS/VS Ring Output #1
5642 24-27 GS/VS Ring Output #2
5643 28-31 GS/VS Ring Output #3
5644 32-35 GS/VS Ring Input SRD
5645 36-39 Tessellation Factor Buffer SRD
5646 40-43 Off-Chip LDS Buffer SRD
5647 44-47 Off-Chip Param Cache Buffer SRD
5648 48-51 Sample Position Buffer SRD
5649 52 vaRange::ShadowDescriptorTable High Bits
5650 ============= ================================
5652 The pointer to the global internal table passed to the shader as user data
5653 is a 32-bit pointer. The top 32 bits should be assumed to be the same as
5654 the top 32 bits of the pipeline, so the shader may use the program
5655 counter's top 32 bits.
5657 Unspecified OS
5658 --------------
5660 This section provides code conventions used when the target triple OS is
5661 empty (see :ref:`amdgpu-target-triples`).
5663 Trap Handler ABI
5664 ~~~~~~~~~~~~~~~~
5666 For code objects generated by AMDGPU backend for non-amdhsa OS, the runtime does
5667 not install a trap handler. The ``llvm.trap`` and ``llvm.debugtrap``
5668 instructions are handled as follows:
5670 .. table:: AMDGPU Trap Handler for Non-AMDHSA OS
5671 :name: amdgpu-trap-handler-for-non-amdhsa-os-table
5673 =============== =============== ===========================================
5674 Usage Code Sequence Description
5675 =============== =============== ===========================================
5676 llvm.trap s_endpgm Causes wavefront to be terminated.
5677 llvm.debugtrap *none* Compiler warning given that there is no
5678 trap handler installed.
5679 =============== =============== ===========================================
5681 Source Languages
5682 ================
5684 .. _amdgpu-opencl:
5686 OpenCL
5687 ------
5689 When the language is OpenCL the following differences occur:
5691 1. The OpenCL memory model is used (see :ref:`amdgpu-amdhsa-memory-model`).
5692 2. The AMDGPU backend appends additional arguments to the kernel's explicit
5693 arguments for the AMDHSA OS (see
5694 :ref:`opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table`).
5695 3. Additional metadata is generated
5696 (see :ref:`amdgpu-amdhsa-code-object-metadata`).
5698 .. table:: OpenCL kernel implicit arguments appended for AMDHSA OS
5699 :name: opencl-kernel-implicit-arguments-appended-for-amdhsa-os-table
5701 ======== ==== ========= ===========================================
5702 Position Byte Byte Description
5703 Size Alignment
5704 ======== ==== ========= ===========================================
5705 1 8 8 OpenCL Global Offset X
5706 2 8 8 OpenCL Global Offset Y
5707 3 8 8 OpenCL Global Offset Z
5708 4 8 8 OpenCL address of printf buffer
5709 5 8 8 OpenCL address of virtual queue used by
5710 enqueue_kernel.
5711 6 8 8 OpenCL address of AqlWrap struct used by
5712 enqueue_kernel.
5713 7 8 8 Pointer argument used for Multi-gird
5714 synchronization.
5715 ======== ==== ========= ===========================================
5717 .. _amdgpu-hcc:
5719 HCC
5720 ---
5722 When the language is HCC the following differences occur:
5724 1. The HSA memory model is used (see :ref:`amdgpu-amdhsa-memory-model`).
5726 .. _amdgpu-assembler:
5728 Assembler
5729 ---------
5731 AMDGPU backend has LLVM-MC based assembler which is currently in development.
5732 It supports AMDGCN GFX6-GFX10.
5734 This section describes general syntax for instructions and operands.
5736 Instructions
5737 ~~~~~~~~~~~~
5739 .. toctree::
5740 :hidden:
5742 AMDGPU/AMDGPUAsmGFX7
5743 AMDGPU/AMDGPUAsmGFX8
5744 AMDGPU/AMDGPUAsmGFX9
5745 AMDGPU/AMDGPUAsmGFX10
5746 AMDGPUModifierSyntax
5747 AMDGPUOperandSyntax
5748 AMDGPUInstructionSyntax
5749 AMDGPUInstructionNotation
5751 An instruction has the following :doc:`syntax<AMDGPUInstructionSyntax>`:
5753 ``<``\ *opcode*\ ``> <``\ *operand0*\ ``>, <``\ *operand1*\ ``>,... <``\ *modifier0*\ ``> <``\ *modifier1*\ ``>...``
5755 :doc:`Operands<AMDGPUOperandSyntax>` are normally comma-separated while
5756 :doc:`modifiers<AMDGPUModifierSyntax>` are space-separated.
5758 The order of *operands* and *modifiers* is fixed.
5759 Most *modifiers* are optional and may be omitted.
5761 See detailed instruction syntax description for :doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`,
5762 :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`, :doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`
5763 and :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>`.
5765 Note that features under development are not included in this description.
5767 For more information about instructions, their semantics and supported combinations of
5768 operands, refer to one of instruction set architecture manuals
5769 [AMD-GCN-GFX6]_, [AMD-GCN-GFX7]_, [AMD-GCN-GFX8]_, [AMD-GCN-GFX9]_ and
5770 [AMD-GCN-GFX10]_.
5772 Operands
5773 ~~~~~~~~
5775 Detailed description of operands may be found :doc:`here<AMDGPUOperandSyntax>`.
5777 Modifiers
5778 ~~~~~~~~~
5780 Detailed description of modifiers may be found :doc:`here<AMDGPUModifierSyntax>`.
5782 Instruction Examples
5783 ~~~~~~~~~~~~~~~~~~~~
5785 DS
5786 ++
5788 .. code-block:: nasm
5790 ds_add_u32 v2, v4 offset:16
5791 ds_write_src2_b64 v2 offset0:4 offset1:8
5792 ds_cmpst_f32 v2, v4, v6
5793 ds_min_rtn_f64 v[8:9], v2, v[4:5]
5796 For full list of supported instructions, refer to "LDS/GDS instructions" in ISA Manual.
5798 FLAT
5799 ++++
5801 .. code-block:: nasm
5803 flat_load_dword v1, v[3:4]
5804 flat_store_dwordx3 v[3:4], v[5:7]
5805 flat_atomic_swap v1, v[3:4], v5 glc
5806 flat_atomic_cmpswap v1, v[3:4], v[5:6] glc slc
5807 flat_atomic_fmax_x2 v[1:2], v[3:4], v[5:6] glc
5809 For full list of supported instructions, refer to "FLAT instructions" in ISA Manual.
5811 MUBUF
5812 +++++
5814 .. code-block:: nasm
5816 buffer_load_dword v1, off, s[4:7], s1
5817 buffer_store_dwordx4 v[1:4], v2, ttmp[4:7], s1 offen offset:4 glc tfe
5818 buffer_store_format_xy v[1:2], off, s[4:7], s1
5819 buffer_wbinvl1
5820 buffer_atomic_inc v1, v2, s[8:11], s4 idxen offset:4 slc
5822 For full list of supported instructions, refer to "MUBUF Instructions" in ISA Manual.
5824 SMRD/SMEM
5825 +++++++++
5827 .. code-block:: nasm
5829 s_load_dword s1, s[2:3], 0xfc
5830 s_load_dwordx8 s[8:15], s[2:3], s4
5831 s_load_dwordx16 s[88:103], s[2:3], s4
5832 s_dcache_inv_vol
5833 s_memtime s[4:5]
5835 For full list of supported instructions, refer to "Scalar Memory Operations" in ISA Manual.
5837 SOP1
5838 ++++
5840 .. code-block:: nasm
5842 s_mov_b32 s1, s2
5843 s_mov_b64 s[0:1], 0x80000000
5844 s_cmov_b32 s1, 200
5845 s_wqm_b64 s[2:3], s[4:5]
5846 s_bcnt0_i32_b64 s1, s[2:3]
5847 s_swappc_b64 s[2:3], s[4:5]
5848 s_cbranch_join s[4:5]
5850 For full list of supported instructions, refer to "SOP1 Instructions" in ISA Manual.
5852 SOP2
5853 ++++
5855 .. code-block:: nasm
5857 s_add_u32 s1, s2, s3
5858 s_and_b64 s[2:3], s[4:5], s[6:7]
5859 s_cselect_b32 s1, s2, s3
5860 s_andn2_b32 s2, s4, s6
5861 s_lshr_b64 s[2:3], s[4:5], s6
5862 s_ashr_i32 s2, s4, s6
5863 s_bfm_b64 s[2:3], s4, s6
5864 s_bfe_i64 s[2:3], s[4:5], s6
5865 s_cbranch_g_fork s[4:5], s[6:7]
5867 For full list of supported instructions, refer to "SOP2 Instructions" in ISA Manual.
5869 SOPC
5870 ++++
5872 .. code-block:: nasm
5874 s_cmp_eq_i32 s1, s2
5875 s_bitcmp1_b32 s1, s2
5876 s_bitcmp0_b64 s[2:3], s4
5877 s_setvskip s3, s5
5879 For full list of supported instructions, refer to "SOPC Instructions" in ISA Manual.
5881 SOPP
5882 ++++
5884 .. code-block:: nasm
5886 s_barrier
5887 s_nop 2
5888 s_endpgm
5889 s_waitcnt 0 ; Wait for all counters to be 0
5890 s_waitcnt vmcnt(0) & expcnt(0) & lgkmcnt(0) ; Equivalent to above
5891 s_waitcnt vmcnt(1) ; Wait for vmcnt counter to be 1.
5892 s_sethalt 9
5893 s_sleep 10
5894 s_sendmsg 0x1
5895 s_sendmsg sendmsg(MSG_INTERRUPT)
5896 s_trap 1
5898 For full list of supported instructions, refer to "SOPP Instructions" in ISA Manual.
5900 Unless otherwise mentioned, little verification is performed on the operands
5901 of SOPP Instructions, so it is up to the programmer to be familiar with the
5902 range or acceptable values.
5904 VALU
5905 ++++
5907 For vector ALU instruction opcodes (VOP1, VOP2, VOP3, VOPC, VOP_DPP, VOP_SDWA),
5908 the assembler will automatically use optimal encoding based on its operands.
5909 To force specific encoding, one can add a suffix to the opcode of the instruction:
5911 * _e32 for 32-bit VOP1/VOP2/VOPC
5912 * _e64 for 64-bit VOP3
5913 * _dpp for VOP_DPP
5914 * _sdwa for VOP_SDWA
5916 VOP1/VOP2/VOP3/VOPC examples:
5918 .. code-block:: nasm
5920 v_mov_b32 v1, v2
5921 v_mov_b32_e32 v1, v2
5922 v_nop
5923 v_cvt_f64_i32_e32 v[1:2], v2
5924 v_floor_f32_e32 v1, v2
5925 v_bfrev_b32_e32 v1, v2
5926 v_add_f32_e32 v1, v2, v3
5927 v_mul_i32_i24_e64 v1, v2, 3
5928 v_mul_i32_i24_e32 v1, -3, v3
5929 v_mul_i32_i24_e32 v1, -100, v3
5930 v_addc_u32 v1, s[0:1], v2, v3, s[2:3]
5931 v_max_f16_e32 v1, v2, v3
5933 VOP_DPP examples:
5935 .. code-block:: nasm
5937 v_mov_b32 v0, v0 quad_perm:[0,2,1,1]
5938 v_sin_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
5939 v_mov_b32 v0, v0 wave_shl:1
5940 v_mov_b32 v0, v0 row_mirror
5941 v_mov_b32 v0, v0 row_bcast:31
5942 v_mov_b32 v0, v0 quad_perm:[1,3,0,1] row_mask:0xa bank_mask:0x1 bound_ctrl:0
5943 v_add_f32 v0, v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
5944 v_max_f16 v1, v2, v3 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0
5946 VOP_SDWA examples:
5948 .. code-block:: nasm
5950 v_mov_b32 v1, v2 dst_sel:BYTE_0 dst_unused:UNUSED_PRESERVE src0_sel:DWORD
5951 v_min_u32 v200, v200, v1 dst_sel:WORD_1 dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD
5952 v_sin_f32 v0, v0 dst_unused:UNUSED_PAD src0_sel:WORD_1
5953 v_fract_f32 v0, |v0| dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1
5954 v_cmpx_le_u32 vcc, v1, v2 src0_sel:BYTE_2 src1_sel:WORD_0
5956 For full list of supported instructions, refer to "Vector ALU instructions".
5958 .. TODO
5959 Remove once we switch to code object v3 by default.
5961 .. _amdgpu-amdhsa-assembler-predefined-symbols-v2:
5963 Code Object V2 Predefined Symbols (-mattr=-code-object-v3)
5964 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
5966 .. warning:: Code Object V2 is not the default code object version emitted by
5967 this version of LLVM. For a description of the predefined symbols available
5968 with the default configuration (Code Object V3) see
5969 :ref:`amdgpu-amdhsa-assembler-predefined-symbols-v3`.
5971 The AMDGPU assembler defines and updates some symbols automatically. These
5972 symbols do not affect code generation.
5974 .option.machine_version_major
5975 +++++++++++++++++++++++++++++
5977 Set to the GFX major generation number of the target being assembled for. For
5978 example, when assembling for a "GFX9" target this will be set to the integer
5979 value "9". The possible GFX major generation numbers are presented in
5980 :ref:`amdgpu-processors`.
5982 .option.machine_version_minor
5983 +++++++++++++++++++++++++++++
5985 Set to the GFX minor generation number of the target being assembled for. For
5986 example, when assembling for a "GFX810" target this will be set to the integer
5987 value "1". The possible GFX minor generation numbers are presented in
5988 :ref:`amdgpu-processors`.
5990 .option.machine_version_stepping
5991 ++++++++++++++++++++++++++++++++
5993 Set to the GFX stepping generation number of the target being assembled for.
5994 For example, when assembling for a "GFX704" target this will be set to the
5995 integer value "4". The possible GFX stepping generation numbers are presented
5996 in :ref:`amdgpu-processors`.
5998 .kernel.vgpr_count
5999 ++++++++++++++++++
6001 Set to zero each time a
6002 :ref:`amdgpu-amdhsa-assembler-directive-amdgpu_hsa_kernel` directive is
6003 encountered. At each instruction, if the current value of this symbol is less
6004 than or equal to the maximum VPGR number explicitly referenced within that
6005 instruction then the symbol value is updated to equal that VGPR number plus
6006 one.
6008 .kernel.sgpr_count
6009 ++++++++++++++++++
6011 Set to zero each time a
6012 :ref:`amdgpu-amdhsa-assembler-directive-amdgpu_hsa_kernel` directive is
6013 encountered. At each instruction, if the current value of this symbol is less
6014 than or equal to the maximum VPGR number explicitly referenced within that
6015 instruction then the symbol value is updated to equal that SGPR number plus
6016 one.
6018 .. _amdgpu-amdhsa-assembler-directives-v2:
6020 Code Object V2 Directives (-mattr=-code-object-v3)
6021 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
6023 .. warning:: Code Object V2 is not the default code object version emitted by
6024 this version of LLVM. For a description of the directives supported with
6025 the default configuration (Code Object V3) see
6026 :ref:`amdgpu-amdhsa-assembler-directives-v3`.
6028 AMDGPU ABI defines auxiliary data in output code object. In assembly source,
6029 one can specify them with assembler directives.
6031 .hsa_code_object_version major, minor
6032 +++++++++++++++++++++++++++++++++++++
6034 *major* and *minor* are integers that specify the version of the HSA code
6035 object that will be generated by the assembler.
6037 .hsa_code_object_isa [major, minor, stepping, vendor, arch]
6038 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
6041 *major*, *minor*, and *stepping* are all integers that describe the instruction
6042 set architecture (ISA) version of the assembly program.
6044 *vendor* and *arch* are quoted strings. *vendor* should always be equal to
6045 "AMD" and *arch* should always be equal to "AMDGPU".
6047 By default, the assembler will derive the ISA version, *vendor*, and *arch*
6048 from the value of the -mcpu option that is passed to the assembler.
6050 .. _amdgpu-amdhsa-assembler-directive-amdgpu_hsa_kernel:
6052 .amdgpu_hsa_kernel (name)
6053 +++++++++++++++++++++++++
6055 This directives specifies that the symbol with given name is a kernel entry point
6056 (label) and the object should contain corresponding symbol of type STT_AMDGPU_HSA_KERNEL.
6058 .amd_kernel_code_t
6059 ++++++++++++++++++
6061 This directive marks the beginning of a list of key / value pairs that are used
6062 to specify the amd_kernel_code_t object that will be emitted by the assembler.
6063 The list must be terminated by the *.end_amd_kernel_code_t* directive. For
6064 any amd_kernel_code_t values that are unspecified a default value will be
6065 used. The default value for all keys is 0, with the following exceptions:
6067 - *amd_code_version_major* defaults to 1.
6068 - *amd_kernel_code_version_minor* defaults to 2.
6069 - *amd_machine_kind* defaults to 1.
6070 - *amd_machine_version_major*, *machine_version_minor*, and
6071 *amd_machine_version_stepping* are derived from the value of the -mcpu option
6072 that is passed to the assembler.
6073 - *kernel_code_entry_byte_offset* defaults to 256.
6074 - *wavefront_size* defaults 6 for all targets before GFX10. For GFX10 onwards
6075 defaults to 6 if target feature ``wavefrontsize64`` is enabled, otherwise 5.
6076 Note that wavefront size is specified as a power of two, so a value of **n**
6077 means a size of 2^ **n**.
6078 - *call_convention* defaults to -1.
6079 - *kernarg_segment_alignment*, *group_segment_alignment*, and
6080 *private_segment_alignment* default to 4. Note that alignments are specified
6081 as a power of 2, so a value of **n** means an alignment of 2^ **n**.
6082 - *enable_wgp_mode* defaults to 1 if target feature ``cumode`` is disabled for
6083 GFX10 onwards.
6084 - *enable_mem_ordered* defaults to 1 for GFX10 onwards.
6086 The *.amd_kernel_code_t* directive must be placed immediately after the
6087 function label and before any instructions.
6089 For a full list of amd_kernel_code_t keys, refer to AMDGPU ABI document,
6090 comments in lib/Target/AMDGPU/AmdKernelCodeT.h and test/CodeGen/AMDGPU/hsa.s.
6092 .. _amdgpu-amdhsa-assembler-example-v2:
6094 Code Object V2 Example Source Code (-mattr=-code-object-v3)
6095 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
6097 .. warning:: Code Object V2 is not the default code object version emitted by
6098 this version of LLVM. For a description of the directives supported with
6099 the default configuration (Code Object V3) see
6100 :ref:`amdgpu-amdhsa-assembler-example-v3`.
6102 Here is an example of a minimal assembly source file, defining one HSA kernel:
6104 .. code-block:: none
6106 .hsa_code_object_version 1,0
6107 .hsa_code_object_isa
6109 .hsatext
6110 .globl hello_world
6111 .p2align 8
6112 .amdgpu_hsa_kernel hello_world
6114 hello_world:
6116 .amd_kernel_code_t
6117 enable_sgpr_kernarg_segment_ptr = 1
6118 is_ptr64 = 1
6119 compute_pgm_rsrc1_vgprs = 0
6120 compute_pgm_rsrc1_sgprs = 0
6121 compute_pgm_rsrc2_user_sgpr = 2
6122 compute_pgm_rsrc1_wgp_mode = 0
6123 compute_pgm_rsrc1_mem_ordered = 0
6124 compute_pgm_rsrc1_fwd_progress = 1
6125 .end_amd_kernel_code_t
6127 s_load_dwordx2 s[0:1], s[0:1] 0x0
6128 v_mov_b32 v0, 3.14159
6129 s_waitcnt lgkmcnt(0)
6130 v_mov_b32 v1, s0
6131 v_mov_b32 v2, s1
6132 flat_store_dword v[1:2], v0
6133 s_endpgm
6134 .Lfunc_end0:
6135 .size hello_world, .Lfunc_end0-hello_world
6137 .. _amdgpu-amdhsa-assembler-predefined-symbols-v3:
6139 Code Object V3 Predefined Symbols (-mattr=+code-object-v3)
6140 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
6142 The AMDGPU assembler defines and updates some symbols automatically. These
6143 symbols do not affect code generation.
6145 .amdgcn.gfx_generation_number
6146 +++++++++++++++++++++++++++++
6148 Set to the GFX major generation number of the target being assembled for. For
6149 example, when assembling for a "GFX9" target this will be set to the integer
6150 value "9". The possible GFX major generation numbers are presented in
6151 :ref:`amdgpu-processors`.
6153 .amdgcn.gfx_generation_minor
6154 ++++++++++++++++++++++++++++
6156 Set to the GFX minor generation number of the target being assembled for. For
6157 example, when assembling for a "GFX810" target this will be set to the integer
6158 value "1". The possible GFX minor generation numbers are presented in
6159 :ref:`amdgpu-processors`.
6161 .amdgcn.gfx_generation_stepping
6162 +++++++++++++++++++++++++++++++
6164 Set to the GFX stepping generation number of the target being assembled for.
6165 For example, when assembling for a "GFX704" target this will be set to the
6166 integer value "4". The possible GFX stepping generation numbers are presented
6167 in :ref:`amdgpu-processors`.
6169 .. _amdgpu-amdhsa-assembler-symbol-next_free_vgpr:
6171 .amdgcn.next_free_vgpr
6172 ++++++++++++++++++++++
6174 Set to zero before assembly begins. At each instruction, if the current value
6175 of this symbol is less than or equal to the maximum VGPR number explicitly
6176 referenced within that instruction then the symbol value is updated to equal
6177 that VGPR number plus one.
6179 May be used to set the `.amdhsa_next_free_vpgr` directive in
6180 :ref:`amdhsa-kernel-directives-table`.
6182 May be set at any time, e.g. manually set to zero at the start of each kernel.
6184 .. _amdgpu-amdhsa-assembler-symbol-next_free_sgpr:
6186 .amdgcn.next_free_sgpr
6187 ++++++++++++++++++++++
6189 Set to zero before assembly begins. At each instruction, if the current value
6190 of this symbol is less than or equal the maximum SGPR number explicitly
6191 referenced within that instruction then the symbol value is updated to equal
6192 that SGPR number plus one.
6194 May be used to set the `.amdhsa_next_free_spgr` directive in
6195 :ref:`amdhsa-kernel-directives-table`.
6197 May be set at any time, e.g. manually set to zero at the start of each kernel.
6199 .. _amdgpu-amdhsa-assembler-directives-v3:
6201 Code Object V3 Directives (-mattr=+code-object-v3)
6202 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
6204 Directives which begin with ``.amdgcn`` are valid for all ``amdgcn``
6205 architecture processors, and are not OS-specific. Directives which begin with
6206 ``.amdhsa`` are specific to ``amdgcn`` architecture processors when the
6207 ``amdhsa`` OS is specified. See :ref:`amdgpu-target-triples` and
6208 :ref:`amdgpu-processors`.
6210 .amdgcn_target <target>
6211 +++++++++++++++++++++++
6213 Optional directive which declares the target supported by the containing
6214 assembler source file. Valid values are described in
6215 :ref:`amdgpu-amdhsa-code-object-target-identification`. Used by the assembler
6216 to validate command-line options such as ``-triple``, ``-mcpu``, and those
6217 which specify target features.
6219 .amdhsa_kernel <name>
6220 +++++++++++++++++++++
6222 Creates a correctly aligned AMDHSA kernel descriptor and a symbol,
6223 ``<name>.kd``, in the current location of the current section. Only valid when
6224 the OS is ``amdhsa``. ``<name>`` must be a symbol that labels the first
6225 instruction to execute, and does not need to be previously defined.
6227 Marks the beginning of a list of directives used to generate the bytes of a
6228 kernel descriptor, as described in :ref:`amdgpu-amdhsa-kernel-descriptor`.
6229 Directives which may appear in this list are described in
6230 :ref:`amdhsa-kernel-directives-table`. Directives may appear in any order, must
6231 be valid for the target being assembled for, and cannot be repeated. Directives
6232 support the range of values specified by the field they reference in
6233 :ref:`amdgpu-amdhsa-kernel-descriptor`. If a directive is not specified, it is
6234 assumed to have its default value, unless it is marked as "Required", in which
6235 case it is an error to omit the directive. This list of directives is
6236 terminated by an ``.end_amdhsa_kernel`` directive.
6238 .. table:: AMDHSA Kernel Assembler Directives
6239 :name: amdhsa-kernel-directives-table
6241 ======================================================== =================== ============ ===================
6242 Directive Default Supported On Description
6243 ======================================================== =================== ============ ===================
6244 ``.amdhsa_group_segment_fixed_size`` 0 GFX6-GFX10 Controls GROUP_SEGMENT_FIXED_SIZE in
6245 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6246 ``.amdhsa_private_segment_fixed_size`` 0 GFX6-GFX10 Controls PRIVATE_SEGMENT_FIXED_SIZE in
6247 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6248 ``.amdhsa_user_sgpr_private_segment_buffer`` 0 GFX6-GFX10 Controls ENABLE_SGPR_PRIVATE_SEGMENT_BUFFER in
6249 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6250 ``.amdhsa_user_sgpr_dispatch_ptr`` 0 GFX6-GFX10 Controls ENABLE_SGPR_DISPATCH_PTR in
6251 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6252 ``.amdhsa_user_sgpr_queue_ptr`` 0 GFX6-GFX10 Controls ENABLE_SGPR_QUEUE_PTR in
6253 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6254 ``.amdhsa_user_sgpr_kernarg_segment_ptr`` 0 GFX6-GFX10 Controls ENABLE_SGPR_KERNARG_SEGMENT_PTR in
6255 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6256 ``.amdhsa_user_sgpr_dispatch_id`` 0 GFX6-GFX10 Controls ENABLE_SGPR_DISPATCH_ID in
6257 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6258 ``.amdhsa_user_sgpr_flat_scratch_init`` 0 GFX6-GFX10 Controls ENABLE_SGPR_FLAT_SCRATCH_INIT in
6259 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6260 ``.amdhsa_user_sgpr_private_segment_size`` 0 GFX6-GFX10 Controls ENABLE_SGPR_PRIVATE_SEGMENT_SIZE in
6261 :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6262 ``.amdhsa_wavefront_size32`` Target GFX10 Controls ENABLE_WAVEFRONT_SIZE32 in
6263 Feature :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6264 Specific
6265 (-wavefrontsize64)
6266 ``.amdhsa_system_sgpr_private_segment_wavefront_offset`` 0 GFX6-GFX10 Controls ENABLE_SGPR_PRIVATE_SEGMENT_WAVEFRONT_OFFSET in
6267 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6268 ``.amdhsa_system_sgpr_workgroup_id_x`` 1 GFX6-GFX10 Controls ENABLE_SGPR_WORKGROUP_ID_X in
6269 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6270 ``.amdhsa_system_sgpr_workgroup_id_y`` 0 GFX6-GFX10 Controls ENABLE_SGPR_WORKGROUP_ID_Y in
6271 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6272 ``.amdhsa_system_sgpr_workgroup_id_z`` 0 GFX6-GFX10 Controls ENABLE_SGPR_WORKGROUP_ID_Z in
6273 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6274 ``.amdhsa_system_sgpr_workgroup_info`` 0 GFX6-GFX10 Controls ENABLE_SGPR_WORKGROUP_INFO in
6275 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6276 ``.amdhsa_system_vgpr_workitem_id`` 0 GFX6-GFX10 Controls ENABLE_VGPR_WORKITEM_ID in
6277 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6278 Possible values are defined in
6279 :ref:`amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table`.
6280 ``.amdhsa_next_free_vgpr`` Required GFX6-GFX10 Maximum VGPR number explicitly referenced, plus one.
6281 Used to calculate GRANULATED_WORKITEM_VGPR_COUNT in
6282 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6283 ``.amdhsa_next_free_sgpr`` Required GFX6-GFX10 Maximum SGPR number explicitly referenced, plus one.
6284 Used to calculate GRANULATED_WAVEFRONT_SGPR_COUNT in
6285 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6286 ``.amdhsa_reserve_vcc`` 1 GFX6-GFX10 Whether the kernel may use the special VCC SGPR.
6287 Used to calculate GRANULATED_WAVEFRONT_SGPR_COUNT in
6288 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6289 ``.amdhsa_reserve_flat_scratch`` 1 GFX7-GFX10 Whether the kernel may use flat instructions to access
6290 scratch memory. Used to calculate
6291 GRANULATED_WAVEFRONT_SGPR_COUNT in
6292 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6293 ``.amdhsa_reserve_xnack_mask`` Target GFX8-GFX10 Whether the kernel may trigger XNACK replay.
6294 Feature Used to calculate GRANULATED_WAVEFRONT_SGPR_COUNT in
6295 Specific :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6296 (+xnack)
6297 ``.amdhsa_float_round_mode_32`` 0 GFX6-GFX10 Controls FLOAT_ROUND_MODE_32 in
6298 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6299 Possible values are defined in
6300 :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
6301 ``.amdhsa_float_round_mode_16_64`` 0 GFX6-GFX10 Controls FLOAT_ROUND_MODE_16_64 in
6302 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6303 Possible values are defined in
6304 :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`.
6305 ``.amdhsa_float_denorm_mode_32`` 0 GFX6-GFX10 Controls FLOAT_DENORM_MODE_32 in
6306 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6307 Possible values are defined in
6308 :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
6309 ``.amdhsa_float_denorm_mode_16_64`` 3 GFX6-GFX10 Controls FLOAT_DENORM_MODE_16_64 in
6310 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6311 Possible values are defined in
6312 :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`.
6313 ``.amdhsa_dx10_clamp`` 1 GFX6-GFX10 Controls ENABLE_DX10_CLAMP in
6314 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6315 ``.amdhsa_ieee_mode`` 1 GFX6-GFX10 Controls ENABLE_IEEE_MODE in
6316 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6317 ``.amdhsa_fp16_overflow`` 0 GFX9-GFX10 Controls FP16_OVFL in
6318 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6319 ``.amdhsa_workgroup_processor_mode`` Target GFX10 Controls ENABLE_WGP_MODE in
6320 Feature :ref:`amdgpu-amdhsa-kernel-descriptor-gfx6-gfx10-table`.
6321 Specific
6322 (-cumode)
6323 ``.amdhsa_memory_ordered`` 1 GFX10 Controls MEM_ORDERED in
6324 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6325 ``.amdhsa_forward_progress`` 0 GFX10 Controls FWD_PROGRESS in
6326 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1-gfx6-gfx10-table`.
6327 ``.amdhsa_exception_fp_ieee_invalid_op`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_IEEE_754_FP_INVALID_OPERATION in
6328 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6329 ``.amdhsa_exception_fp_denorm_src`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_FP_DENORMAL_SOURCE in
6330 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6331 ``.amdhsa_exception_fp_ieee_div_zero`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_IEEE_754_FP_DIVISION_BY_ZERO in
6332 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6333 ``.amdhsa_exception_fp_ieee_overflow`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_IEEE_754_FP_OVERFLOW in
6334 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6335 ``.amdhsa_exception_fp_ieee_underflow`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_IEEE_754_FP_UNDERFLOW in
6336 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6337 ``.amdhsa_exception_fp_ieee_inexact`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_IEEE_754_FP_INEXACT in
6338 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6339 ``.amdhsa_exception_int_div_zero`` 0 GFX6-GFX10 Controls ENABLE_EXCEPTION_INT_DIVIDE_BY_ZERO in
6340 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx10-table`.
6341 ======================================================== =================== ============ ===================
6343 .amdgpu_metadata
6344 ++++++++++++++++
6346 Optional directive which declares the contents of the ``NT_AMDGPU_METADATA``
6347 note record (see :ref:`amdgpu-elf-note-records-table-v3`).
6349 The contents must be in the [YAML]_ markup format, with the same structure and
6350 semantics described in :ref:`amdgpu-amdhsa-code-object-metadata-v3`.
6352 This directive is terminated by an ``.end_amdgpu_metadata`` directive.
6354 .. _amdgpu-amdhsa-assembler-example-v3:
6356 Code Object V3 Example Source Code (-mattr=+code-object-v3)
6357 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
6359 Here is an example of a minimal assembly source file, defining one HSA kernel:
6361 .. code-block:: none
6363 .amdgcn_target "amdgcn-amd-amdhsa--gfx900+xnack" // optional
6365 .text
6366 .globl hello_world
6367 .p2align 8
6368 .type hello_world,@function
6369 hello_world:
6370 s_load_dwordx2 s[0:1], s[0:1] 0x0
6371 v_mov_b32 v0, 3.14159
6372 s_waitcnt lgkmcnt(0)
6373 v_mov_b32 v1, s0
6374 v_mov_b32 v2, s1
6375 flat_store_dword v[1:2], v0
6376 s_endpgm
6377 .Lfunc_end0:
6378 .size hello_world, .Lfunc_end0-hello_world
6380 .rodata
6381 .p2align 6
6382 .amdhsa_kernel hello_world
6383 .amdhsa_user_sgpr_kernarg_segment_ptr 1
6384 .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
6385 .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
6386 .end_amdhsa_kernel
6388 .amdgpu_metadata
6389 ---
6390 amdhsa.version:
6391 - 1
6392 - 0
6393 amdhsa.kernels:
6394 - .name: hello_world
6395 .symbol: hello_world.kd
6396 .kernarg_segment_size: 48
6397 .group_segment_fixed_size: 0
6398 .private_segment_fixed_size: 0
6399 .kernarg_segment_align: 4
6400 .wavefront_size: 64
6401 .sgpr_count: 2
6402 .vgpr_count: 3
6403 .max_flat_workgroup_size: 256
6404 ...
6405 .end_amdgpu_metadata
6407 If an assembly source file contains multiple kernels and/or functions, the
6408 :ref:`amdgpu-amdhsa-assembler-symbol-next_free_vgpr` and
6409 :ref:`amdgpu-amdhsa-assembler-symbol-next_free_sgpr` symbols may be reset using
6410 the ``.set <symbol>, <expression>`` directive. For example, in the case of two
6411 kernels, where ``function1`` is only called from ``kernel1`` it is sufficient
6412 to group the function with the kernel that calls it and reset the symbols
6413 between the two connected components:
6415 .. code-block:: none
6417 .amdgcn_target "amdgcn-amd-amdhsa--gfx900+xnack" // optional
6419 // gpr tracking symbols are implicitly set to zero
6421 .text
6422 .globl kern0
6423 .p2align 8
6424 .type kern0,@function
6425 kern0:
6426 // ...
6427 s_endpgm
6428 .Lkern0_end:
6429 .size kern0, .Lkern0_end-kern0
6431 .rodata
6432 .p2align 6
6433 .amdhsa_kernel kern0
6434 // ...
6435 .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
6436 .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
6437 .end_amdhsa_kernel
6439 // reset symbols to begin tracking usage in func1 and kern1
6440 .set .amdgcn.next_free_vgpr, 0
6441 .set .amdgcn.next_free_sgpr, 0
6443 .text
6444 .hidden func1
6445 .global func1
6446 .p2align 2
6447 .type func1,@function
6448 func1:
6449 // ...
6450 s_setpc_b64 s[30:31]
6451 .Lfunc1_end:
6452 .size func1, .Lfunc1_end-func1
6454 .globl kern1
6455 .p2align 8
6456 .type kern1,@function
6457 kern1:
6458 // ...
6459 s_getpc_b64 s[4:5]
6460 s_add_u32 s4, s4, func1@rel32@lo+4
6461 s_addc_u32 s5, s5, func1@rel32@lo+4
6462 s_swappc_b64 s[30:31], s[4:5]
6463 // ...
6464 s_endpgm
6465 .Lkern1_end:
6466 .size kern1, .Lkern1_end-kern1
6468 .rodata
6469 .p2align 6
6470 .amdhsa_kernel kern1
6471 // ...
6472 .amdhsa_next_free_vgpr .amdgcn.next_free_vgpr
6473 .amdhsa_next_free_sgpr .amdgcn.next_free_sgpr
6474 .end_amdhsa_kernel
6476 These symbols cannot identify connected components in order to automatically
6477 track the usage for each kernel. However, in some cases careful organization of
6478 the kernels and functions in the source file means there is minimal additional
6479 effort required to accurately calculate GPR usage.
6481 Additional Documentation
6482 ========================
6484 .. [AMD-RADEON-HD-2000-3000] `AMD R6xx shader ISA <http://developer.amd.com/wordpress/media/2012/10/R600_Instruction_Set_Architecture.pdf>`__
6485 .. [AMD-RADEON-HD-4000] `AMD R7xx shader ISA <http://developer.amd.com/wordpress/media/2012/10/R700-Family_Instruction_Set_Architecture.pdf>`__
6486 .. [AMD-RADEON-HD-5000] `AMD Evergreen shader ISA <http://developer.amd.com/wordpress/media/2012/10/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf>`__
6487 .. [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>`__
6488 .. [AMD-GCN-GFX6] `AMD Southern Islands Series ISA <http://developer.amd.com/wordpress/media/2012/12/AMD_Southern_Islands_Instruction_Set_Architecture.pdf>`__
6489 .. [AMD-GCN-GFX7] `AMD Sea Islands Series ISA <http://developer.amd.com/wordpress/media/2013/07/AMD_Sea_Islands_Instruction_Set_Architecture.pdf>`_
6490 .. [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>`__
6491 .. [AMD-GCN-GFX9] `AMD "Vega" Instruction Set Architecture <http://developer.amd.com/wordpress/media/2013/12/Vega_Shader_ISA_28July2017.pdf>`__
6492 .. [AMD-GCN-GFX10] AMD "Navi" Instruction Set Architecture *TBA*
6493 .. TODO
6494 ttye Add link when made public.
6495 .. [AMD-ROCm] `ROCm: Open Platform for Development, Discovery and Education Around GPU Computing <http://gpuopen.com/compute-product/rocm/>`__
6496 .. [AMD-ROCm-github] `ROCm github <http://github.com/RadeonOpenCompute>`__
6497 .. [HSA] `Heterogeneous System Architecture (HSA) Foundation <http://www.hsafoundation.com/>`__
6498 .. [ELF] `Executable and Linkable Format (ELF) <http://www.sco.com/developers/gabi/>`__
6499 .. [DWARF] `DWARF Debugging Information Format <http://dwarfstd.org/>`__
6500 .. [YAML] `YAML Ain't Markup Language (YAML™) Version 1.2 <http://www.yaml.org/spec/1.2/spec.html>`__
6501 .. [MsgPack] `Message Pack <http://www.msgpack.org/>`__
6502 .. [OpenCL] `The OpenCL Specification Version 2.0 <http://www.khronos.org/registry/cl/specs/opencl-2.0.pdf>`__
6503 .. [HRF] `Heterogeneous-race-free Memory Models <http://benedictgaster.org/wp-content/uploads/2014/01/asplos269-FINAL.pdf>`__
6504 .. [CLANG-ATTR] `Attributes in Clang <http://clang.llvm.org/docs/AttributeReference.html>`__