1============================= 2User Guide for AMDGPU Backend 3============================= 4 5.. contents:: 6 :local: 7 8Introduction 9============ 10 11The AMDGPU backend provides ISA code generation for AMD GPUs, starting with the 12R600 family up until the current GCN families. It lives in the 13``lib/Target/AMDGPU`` directory. 14 15LLVM 16==== 17 18.. _amdgpu-target-triples: 19 20Target Triples 21-------------- 22 23Use the ``clang -target <Architecture>-<Vendor>-<OS>-<Environment>`` option to 24specify the target triple: 25 26 .. table:: AMDGPU Target Triples 27 :name: amdgpu-target-triples-table 28 29 ============ ======== ========= =========== 30 Architecture Vendor OS Environment 31 ============ ======== ========= =========== 32 r600 amd <empty> <empty> 33 amdgcn amd <empty> <empty> 34 amdgcn amd amdhsa <empty> 35 amdgcn amd amdhsa opencl 36 amdgcn amd amdhsa amdgizcl 37 amdgcn amd amdhsa amdgiz 38 amdgcn amd amdhsa hcc 39 ============ ======== ========= =========== 40 41``r600-amd--`` 42 Supports AMD GPUs HD2XXX-HD6XXX for graphics and compute shaders executed on 43 the MESA runtime. 44 45``amdgcn-amd--`` 46 Supports AMD GPUs GCN GFX6 onwards for graphics and compute shaders executed on 47 the MESA runtime. 48 49``amdgcn-amd-amdhsa-`` 50 Supports AMD GCN GPUs GFX6 onwards for compute kernels executed on HSA [HSA]_ 51 compatible runtimes such as AMD's ROCm [AMD-ROCm]_. 52 53``amdgcn-amd-amdhsa-opencl`` 54 Supports AMD GCN GPUs GFX6 onwards for OpenCL compute kernels executed on HSA 55 [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. See 56 :ref:`amdgpu-opencl`. 57 58``amdgcn-amd-amdhsa-amdgizcl`` 59 Same as ``amdgcn-amd-amdhsa-opencl`` except a different address space mapping 60 is used (see :ref:`amdgpu-address-spaces`). 61 62``amdgcn-amd-amdhsa-amdgiz`` 63 Same as ``amdgcn-amd-amdhsa-`` except a different address space mapping is 64 used (see :ref:`amdgpu-address-spaces`). 65 66``amdgcn-amd-amdhsa-hcc`` 67 Supports AMD GCN GPUs GFX6 onwards for AMD HC language compute kernels 68 executed on HSA [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. See 69 :ref:`amdgpu-hcc`. 70 71.. _amdgpu-processors: 72 73Processors 74---------- 75 76Use the ``clang -mcpu <Processor>`` option to specify the AMD GPU processor. The 77names from both the *Processor* and *Alternative Processor* can be used. 78 79 .. table:: AMDGPU Processors 80 :name: amdgpu-processors-table 81 82 ========== =========== ============ ===== ======= ================== 83 Processor Alternative Target dGPU/ Runtime Example 84 Processor Triple APU Support Products 85 Architecture 86 ========== =========== ============ ===== ======= ================== 87 **R600** [AMD-R6xx]_ 88 -------------------------------------------------------------------- 89 r600 r600 dGPU 90 r630 r600 dGPU 91 rs880 r600 dGPU 92 rv670 r600 dGPU 93 **R700** [AMD-R7xx]_ 94 -------------------------------------------------------------------- 95 rv710 r600 dGPU 96 rv730 r600 dGPU 97 rv770 r600 dGPU 98 **Evergreen** [AMD-Evergreen]_ 99 -------------------------------------------------------------------- 100 cedar r600 dGPU 101 redwood r600 dGPU 102 sumo r600 dGPU 103 juniper r600 dGPU 104 cypress r600 dGPU 105 **Northern Islands** [AMD-Cayman-Trinity]_ 106 -------------------------------------------------------------------- 107 barts r600 dGPU 108 turks r600 dGPU 109 caicos r600 dGPU 110 cayman r600 dGPU 111 **GCN GFX6 (Southern Islands (SI))** [AMD-Souther-Islands]_ 112 -------------------------------------------------------------------- 113 gfx600 - tahiti amdgcn dGPU 114 gfx601 - pitcairn amdgcn dGPU 115 - verde 116 - oland 117 - hainan 118 **GCN GFX7 (Sea Islands (CI))** [AMD-Sea-Islands]_ 119 -------------------------------------------------------------------- 120 gfx700 - bonaire amdgcn dGPU - Radeon HD 7790 121 - Radeon HD 8770 122 - R7 260 123 - R7 260X 124 \ - kaveri amdgcn APU - A6-7000 125 - A6 Pro-7050B 126 - A8-7100 127 - A8 Pro-7150B 128 - A10-7300 129 - A10 Pro-7350B 130 - FX-7500 131 - A8-7200P 132 - A10-7400P 133 - FX-7600P 134 gfx701 - hawaii amdgcn dGPU ROCm - FirePro W8100 135 - FirePro W9100 136 - FirePro S9150 137 - FirePro S9170 138 gfx702 dGPU ROCm - Radeon R9 290 139 - Radeon R9 290x 140 - Radeon R390 141 - Radeon R390x 142 gfx703 - kabini amdgcn APU - E1-2100 143 - mullins - E1-2200 144 - E1-2500 145 - E2-3000 146 - E2-3800 147 - A4-5000 148 - A4-5100 149 - A6-5200 150 - A4 Pro-3340B 151 **GCN GFX8 (Volcanic Islands (VI))** [AMD-Volcanic-Islands]_ 152 -------------------------------------------------------------------- 153 gfx800 - iceland amdgcn dGPU - FirePro S7150 154 - FirePro S7100 155 - FirePro W7100 156 - Radeon R285 157 - Radeon R9 380 158 - Radeon R9 385 159 - Mobile FirePro 160 M7170 161 gfx801 - carrizo amdgcn APU - A6-8500P 162 - Pro A6-8500B 163 - A8-8600P 164 - Pro A8-8600B 165 - FX-8800P 166 - Pro A12-8800B 167 \ amdgcn APU ROCm - A10-8700P 168 - Pro A10-8700B 169 - A10-8780P 170 \ amdgcn APU - A10-9600P 171 - A10-9630P 172 - A12-9700P 173 - A12-9730P 174 - FX-9800P 175 - FX-9830P 176 \ amdgcn APU - E2-9010 177 - A6-9210 178 - A9-9410 179 gfx802 - tonga amdgcn dGPU ROCm Same as gfx800 180 gfx803 - fiji amdgcn dGPU ROCm - Radeon R9 Nano 181 - Radeon R9 Fury 182 - Radeon R9 FuryX 183 - Radeon Pro Duo 184 - FirePro S9300x2 185 - Radeon Instinct MI8 186 \ - polaris10 amdgcn dGPU ROCm - Radeon RX 470 187 - Radeon RX 480 188 - Radeon Instinct MI6 189 \ - polaris11 amdgcn dGPU ROCm - Radeon RX 460 190 gfx804 amdgcn dGPU Same as gfx803 191 gfx810 - stoney amdgcn APU 192 **GCN GFX9** [AMD-Vega]_ 193 -------------------------------------------------------------------- 194 gfx900 amdgcn dGPU - Radeon Vega 195 Frontier Edition 196 - Radeon RX Vega 56 197 - Radeon RX Vega 64 198 - Radeon RX Vega 64 199 Liquid 200 - Radeon Instinct MI25 201 gfx901 amdgcn dGPU ROCm Same as gfx900 202 except XNACK is 203 enabled 204 gfx902 amdgcn APU *TBA* 205 206 .. TODO 207 Add product 208 names. 209 gfx903 amdgcn APU Same as gfx902 210 except XNACK is 211 enabled 212 ========== =========== ============ ===== ======= ================== 213 214.. _amdgpu-address-spaces: 215 216Address Spaces 217-------------- 218 219The AMDGPU backend uses the following address space mappings. 220 221The memory space names used in the table, aside from the region memory space, is 222from the OpenCL standard. 223 224LLVM Address Space number is used throughout LLVM (for example, in LLVM IR). 225 226 .. table:: Address Space Mapping 227 :name: amdgpu-address-space-mapping-table 228 229 ================== ================= ================= ================= ================= 230 LLVM Address Space Memory Space 231 ------------------ ----------------------------------------------------------------------- 232 \ Current Default amdgiz/amdgizcl hcc Future Default 233 ================== ================= ================= ================= ================= 234 0 Private (Scratch) Generic (Flat) Generic (Flat) Generic (Flat) 235 1 Global Global Global Global 236 2 Constant Constant Constant Region (GDS) 237 3 Local (group/LDS) Local (group/LDS) Local (group/LDS) Local (group/LDS) 238 4 Generic (Flat) Region (GDS) Region (GDS) Constant 239 5 Region (GDS) Private (Scratch) Private (Scratch) Private (Scratch) 240 ================== ================= ================= ================= ================= 241 242Current Default 243 This is the current default address space mapping used for all languages 244 except hcc. This will shortly be deprecated. 245 246amdgiz/amdgizcl 247 This is the current address space mapping used when ``amdgiz`` or ``amdgizcl`` 248 is specified as the target triple environment value. 249 250hcc 251 This is the current address space mapping used when ``hcc`` is specified as 252 the target triple environment value.This will shortly be deprecated. 253 254Future Default 255 This will shortly be the only address space mapping for all languages using 256 AMDGPU backend. 257 258.. _amdgpu-memory-scopes: 259 260Memory Scopes 261------------- 262 263This section provides LLVM memory synchronization scopes supported by the AMDGPU 264backend memory model when the target triple OS is ``amdhsa`` (see 265:ref:`amdgpu-amdhsa-memory-model` and :ref:`amdgpu-target-triples`). 266 267The memory model supported is based on the HSA memory model [HSA]_ which is 268based in turn on HRF-indirect with scope inclusion [HRF]_. The happens-before 269relation is transitive over the synchonizes-with relation independent of scope, 270and synchonizes-with allows the memory scope instances to be inclusive (see 271table :ref:`amdgpu-amdhsa-llvm-sync-scopes-amdhsa-table`). 272 273This is different to the OpenCL [OpenCL]_ memory model which does not have scope 274inclusion and requires the memory scopes to exactly match. However, this 275is conservatively correct for OpenCL. 276 277 .. table:: AMDHSA LLVM Sync Scopes for AMDHSA 278 :name: amdgpu-amdhsa-llvm-sync-scopes-amdhsa-table 279 280 ================ ========================================================== 281 LLVM Sync Scope Description 282 ================ ========================================================== 283 *none* The default: ``system``. 284 285 Synchronizes with, and participates in modification and 286 seq_cst total orderings with, other operations (except 287 image operations) for all address spaces (except private, 288 or generic that accesses private) provided the other 289 operation's sync scope is: 290 291 - ``system``. 292 - ``agent`` and executed by a thread on the same agent. 293 - ``workgroup`` and executed by a thread in the same 294 workgroup. 295 - ``wavefront`` and executed by a thread in the same 296 wavefront. 297 298 ``agent`` Synchronizes with, and participates in modification and 299 seq_cst total orderings with, other operations (except 300 image operations) for all address spaces (except private, 301 or generic that accesses private) provided the other 302 operation's sync scope is: 303 304 - ``system`` or ``agent`` and executed by a thread on the 305 same agent. 306 - ``workgroup`` and executed by a thread in the same 307 workgroup. 308 - ``wavefront`` and executed by a thread in the same 309 wavefront. 310 311 ``workgroup`` Synchronizes with, and participates in modification and 312 seq_cst total orderings with, other operations (except 313 image operations) for all address spaces (except private, 314 or generic that accesses private) provided the other 315 operation's sync scope is: 316 317 - ``system``, ``agent`` or ``workgroup`` and executed by a 318 thread in the same workgroup. 319 - ``wavefront`` and executed by a thread in the same 320 wavefront. 321 322 ``wavefront`` Synchronizes with, and participates in modification and 323 seq_cst total orderings with, other operations (except 324 image operations) for all address spaces (except private, 325 or generic that accesses private) provided the other 326 operation's sync scope is: 327 328 - ``system``, ``agent``, ``workgroup`` or ``wavefront`` 329 and executed by a thread in the same wavefront. 330 331 ``singlethread`` Only synchronizes with, and participates in modification 332 and seq_cst total orderings with, other operations (except 333 image operations) running in the same thread for all 334 address spaces (for example, in signal handlers). 335 ================ ========================================================== 336 337AMDGPU Intrinsics 338----------------- 339 340The AMDGPU backend implements the following intrinsics. 341 342*This section is WIP.* 343 344.. TODO 345 List AMDGPU intrinsics 346 347Code Object 348=========== 349 350The AMDGPU backend generates a standard ELF [ELF]_ relocatable code object that 351can be linked by ``lld`` to produce a standard ELF shared code object which can 352be loaded and executed on an AMDGPU target. 353 354Header 355------ 356 357The AMDGPU backend uses the following ELF header: 358 359 .. table:: AMDGPU ELF Header 360 :name: amdgpu-elf-header-table 361 362 ========================== =============================== 363 Field Value 364 ========================== =============================== 365 ``e_ident[EI_CLASS]`` ``ELFCLASS64`` 366 ``e_ident[EI_DATA]`` ``ELFDATA2LSB`` 367 ``e_ident[EI_OSABI]`` ``ELFOSABI_AMDGPU_HSA``, 368 ``ELFOSABI_AMDGPU_PAL`` or 369 ``ELFOSABI_AMDGPU_MESA3D`` 370 ``e_ident[EI_ABIVERSION]`` ``ELFABIVERSION_AMDGPU_HSA``, 371 ``ELFABIVERSION_AMDGPU_PAL`` or 372 ``ELFABIVERSION_AMDGPU_MESA3D`` 373 ``e_type`` ``ET_REL`` or ``ET_DYN`` 374 ``e_machine`` ``EM_AMDGPU`` 375 ``e_entry`` 0 376 ``e_flags`` 0 377 ========================== =============================== 378 379.. 380 381 .. table:: AMDGPU ELF Header Enumeration Values 382 :name: amdgpu-elf-header-enumeration-values-table 383 384 =============================== ===== 385 Name Value 386 =============================== ===== 387 ``EM_AMDGPU`` 224 388 ``ELFOSABI_AMDGPU_HSA`` 64 389 ``ELFOSABI_AMDGPU_PAL`` 65 390 ``ELFOSABI_AMDGPU_MESA3D`` 66 391 ``ELFABIVERSION_AMDGPU_HSA`` 1 392 ``ELFABIVERSION_AMDGPU_PAL`` 0 393 ``ELFABIVERSION_AMDGPU_MESA3D`` 0 394 =============================== ===== 395 396``e_ident[EI_CLASS]`` 397 The ELF class is always ``ELFCLASS64``. The AMDGPU backend only supports 64 398 bit applications. 399 400``e_ident[EI_DATA]`` 401 All AMDGPU targets use ELFDATA2LSB for little-endian byte ordering. 402 403``e_ident[EI_OSABI]`` 404 One of the following AMD GPU architecture specific OS ABIs: 405 406 * ``ELFOSABI_AMDGPU_HSA`` is used to specify that the code object conforms to 407 the AMD HSA runtime ABI [HSA]_. 408 409 * ``ELFOSABI_AMDGPU_PAL`` is used to specify that the code object conforms to 410 the AMD PAL runtime ABI. 411 412 * ``ELFOSABI_AMDGPU_MESA3D`` is used to specify that the code object conforms 413 to the AMD MESA runtime ABI. 414 415``e_ident[EI_ABIVERSION]`` 416 The ABI version of the AMD GPU architecture specific OS ABI to which the code 417 object conforms: 418 419 * ``ELFABIVERSION_AMDGPU_HSA`` is used to specify the version of AMD HSA 420 runtime ABI. 421 422 * ``ELFABIVERSION_AMDGPU_PAL`` is used to specify the version of AMD PAL 423 runtime ABI. 424 425 * ``ELFABIVERSION_AMDGPU_MESA3D`` is used to specify the version of AMD MESA 426 runtime ABI. 427 428``e_type`` 429 Can be one of the following values: 430 431 432 ``ET_REL`` 433 The type produced by the AMD GPU backend compiler as it is relocatable code 434 object. 435 436 ``ET_DYN`` 437 The type produced by the linker as it is a shared code object. 438 439 The AMD HSA runtime loader requires a ``ET_DYN`` code object. 440 441``e_machine`` 442 The value ``EM_AMDGPU`` is used for the machine for all members of the AMD GPU 443 architecture family. The specific member is specified in the 444 ``NT_AMD_AMDGPU_ISA`` entry in the ``.note`` section (see 445 :ref:`amdgpu-note-records`). 446 447``e_entry`` 448 The entry point is 0 as the entry points for individual kernels must be 449 selected in order to invoke them through AQL packets. 450 451``e_flags`` 452 The value is 0 as no flags are used. 453 454Sections 455-------- 456 457An AMDGPU target ELF code object has the standard ELF sections which include: 458 459 .. table:: AMDGPU ELF Sections 460 :name: amdgpu-elf-sections-table 461 462 ================== ================ ================================= 463 Name Type Attributes 464 ================== ================ ================================= 465 ``.bss`` ``SHT_NOBITS`` ``SHF_ALLOC`` + ``SHF_WRITE`` 466 ``.data`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_WRITE`` 467 ``.debug_``\ *\** ``SHT_PROGBITS`` *none* 468 ``.dynamic`` ``SHT_DYNAMIC`` ``SHF_ALLOC`` 469 ``.dynstr`` ``SHT_PROGBITS`` ``SHF_ALLOC`` 470 ``.dynsym`` ``SHT_PROGBITS`` ``SHF_ALLOC`` 471 ``.got`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_WRITE`` 472 ``.hash`` ``SHT_HASH`` ``SHF_ALLOC`` 473 ``.note`` ``SHT_NOTE`` *none* 474 ``.rela``\ *name* ``SHT_RELA`` *none* 475 ``.rela.dyn`` ``SHT_RELA`` *none* 476 ``.rodata`` ``SHT_PROGBITS`` ``SHF_ALLOC`` 477 ``.shstrtab`` ``SHT_STRTAB`` *none* 478 ``.strtab`` ``SHT_STRTAB`` *none* 479 ``.symtab`` ``SHT_SYMTAB`` *none* 480 ``.text`` ``SHT_PROGBITS`` ``SHF_ALLOC`` + ``SHF_EXECINSTR`` 481 ================== ================ ================================= 482 483These sections have their standard meanings (see [ELF]_) and are only generated 484if needed. 485 486``.debug``\ *\** 487 The standard DWARF sections. See :ref:`amdgpu-dwarf` for information on the 488 DWARF produced by the AMDGPU backend. 489 490``.dynamic``, ``.dynstr``, ``.dynsym``, ``.hash`` 491 The standard sections used by a dynamic loader. 492 493``.note`` 494 See :ref:`amdgpu-note-records` for the note records supported by the AMDGPU 495 backend. 496 497``.rela``\ *name*, ``.rela.dyn`` 498 For relocatable code objects, *name* is the name of the section that the 499 relocation records apply. For example, ``.rela.text`` is the section name for 500 relocation records associated with the ``.text`` section. 501 502 For linked shared code objects, ``.rela.dyn`` contains all the relocation 503 records from each of the relocatable code object's ``.rela``\ *name* sections. 504 505 See :ref:`amdgpu-relocation-records` for the relocation records supported by 506 the AMDGPU backend. 507 508``.text`` 509 The executable machine code for the kernels and functions they call. Generated 510 as position independent code. See :ref:`amdgpu-code-conventions` for 511 information on conventions used in the isa generation. 512 513.. _amdgpu-note-records: 514 515Note Records 516------------ 517 518As required by ``ELFCLASS64``, minimal zero byte padding must be generated after 519the ``name`` field to ensure the ``desc`` field is 4 byte aligned. In addition, 520minimal zero byte padding must be generated to ensure the ``desc`` field size is 521a multiple of 4 bytes. The ``sh_addralign`` field of the ``.note`` section must 522be at least 4 to indicate at least 8 byte alignment. 523 524The AMDGPU backend code object uses the following ELF note records in the 525``.note`` section. The *Description* column specifies the layout of the note 526record’s ``desc`` field. All fields are consecutive bytes. Note records with 527variable size strings have a corresponding ``*_size`` field that specifies the 528number of bytes, including the terminating null character, in the string. The 529string(s) come immediately after the preceding fields. 530 531Additional note records can be present. 532 533 .. table:: AMDGPU ELF Note Records 534 :name: amdgpu-elf-note-records-table 535 536 ===== ============================== ====================================== 537 Name Type Description 538 ===== ============================== ====================================== 539 "AMD" ``NT_AMD_AMDGPU_HSA_METADATA`` <metadata null terminated string> 540 "AMD" ``NT_AMD_AMDGPU_ISA`` <isa name null terminated string> 541 ===== ============================== ====================================== 542 543.. 544 545 .. table:: AMDGPU ELF Note Record Enumeration Values 546 :name: amdgpu-elf-note-record-enumeration-values-table 547 548 ============================== ===== 549 Name Value 550 ============================== ===== 551 *reserved* 0-9 552 ``NT_AMD_AMDGPU_HSA_METADATA`` 10 553 ``NT_AMD_AMDGPU_ISA`` 11 554 ============================== ===== 555 556``NT_AMD_AMDGPU_ISA`` 557 Specifies the instruction set architecture used by the machine code contained 558 in the code object. 559 560 This note record is required for code objects containing machine code for 561 processors matching the ``amdgcn`` architecture in table 562 :ref:`amdgpu-processors`. 563 564 The null terminated string has the following syntax: 565 566 *architecture*\ ``-``\ *vendor*\ ``-``\ *os*\ ``-``\ *environment*\ ``-``\ *processor* 567 568 where: 569 570 *architecture* 571 The architecture from table :ref:`amdgpu-target-triples-table`. 572 573 This is always ``amdgcn`` when the target triple OS is ``amdhsa`` (see 574 :ref:`amdgpu-target-triples`). 575 576 *vendor* 577 The vendor from table :ref:`amdgpu-target-triples-table`. 578 579 For the AMDGPU backend this is always ``amd``. 580 581 *os* 582 The OS from table :ref:`amdgpu-target-triples-table`. 583 584 *environment* 585 An environment from table :ref:`amdgpu-target-triples-table`, or blank if 586 the environment has no affect on the execution of the code object. 587 588 For the AMDGPU backend this is currently always blank. 589 *processor* 590 The processor from table :ref:`amdgpu-processors-table`. 591 592 For example: 593 594 ``amdgcn-amd-amdhsa--gfx901`` 595 596``NT_AMD_AMDGPU_HSA_METADATA`` 597 Specifies extensible metadata associated with the code objects executed on HSA 598 [HSA]_ compatible runtimes such as AMD's ROCm [AMD-ROCm]_. It is required when 599 the target triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`). See 600 :ref:`amdgpu-amdhsa-hsa-code-object-metadata` for the syntax of the code 601 object metadata string. 602 603.. _amdgpu-symbols: 604 605Symbols 606------- 607 608Symbols include the following: 609 610 .. table:: AMDGPU ELF Symbols 611 :name: amdgpu-elf-symbols-table 612 613 ===================== ============== ============= ================== 614 Name Type Section Description 615 ===================== ============== ============= ================== 616 *link-name* ``STT_OBJECT`` - ``.data`` Global variable 617 - ``.rodata`` 618 - ``.bss`` 619 *link-name*\ ``@kd`` ``STT_OBJECT`` - ``.rodata`` Kernel descriptor 620 *link-name* ``STT_FUNC`` - ``.text`` Kernel entry point 621 ===================== ============== ============= ================== 622 623Global variable 624 Global variables both used and defined by the compilation unit. 625 626 If the symbol is defined in the compilation unit then it is allocated in the 627 appropriate section according to if it has initialized data or is readonly. 628 629 If the symbol is external then its section is ``STN_UNDEF`` and the loader 630 will resolve relocations using the definition provided by another code object 631 or explicitly defined by the runtime. 632 633 All global symbols, whether defined in the compilation unit or external, are 634 accessed by the machine code indirectly through a GOT table entry. This 635 allows them to be preemptable. The GOT table is only supported when the target 636 triple OS is ``amdhsa`` (see :ref:`amdgpu-target-triples`). 637 638 .. TODO 639 Add description of linked shared object symbols. Seems undefined symbols 640 are marked as STT_NOTYPE. 641 642Kernel descriptor 643 Every HSA kernel has an associated kernel descriptor. It is the address of the 644 kernel descriptor that is used in the AQL dispatch packet used to invoke the 645 kernel, not the kernel entry point. The layout of the HSA kernel descriptor is 646 defined in :ref:`amdgpu-amdhsa-kernel-descriptor`. 647 648Kernel entry point 649 Every HSA kernel also has a symbol for its machine code entry point. 650 651.. _amdgpu-relocation-records: 652 653Relocation Records 654------------------ 655 656AMDGPU backend generates ``Elf64_Rela`` relocation records. Supported 657relocatable fields are: 658 659``word32`` 660 This specifies a 32-bit field occupying 4 bytes with arbitrary byte 661 alignment. These values use the same byte order as other word values in the 662 AMD GPU architecture. 663 664``word64`` 665 This specifies a 64-bit field occupying 8 bytes with arbitrary byte 666 alignment. These values use the same byte order as other word values in the 667 AMD GPU architecture. 668 669Following notations are used for specifying relocation calculations: 670 671**A** 672 Represents the addend used to compute the value of the relocatable field. 673 674**G** 675 Represents the offset into the global offset table at which the relocation 676 entry’s symbol will reside during execution. 677 678**GOT** 679 Represents the address of the global offset table. 680 681**P** 682 Represents the place (section offset for ``et_rel`` or address for ``et_dyn``) 683 of the storage unit being relocated (computed using ``r_offset``). 684 685**S** 686 Represents the value of the symbol whose index resides in the relocation 687 entry. 688 689The following relocation types are supported: 690 691 .. table:: AMDGPU ELF Relocation Records 692 :name: amdgpu-elf-relocation-records-table 693 694 ========================== ===== ========== ============================== 695 Relocation Type Value Field Calculation 696 ========================== ===== ========== ============================== 697 ``R_AMDGPU_NONE`` 0 *none* *none* 698 ``R_AMDGPU_ABS32_LO`` 1 ``word32`` (S + A) & 0xFFFFFFFF 699 ``R_AMDGPU_ABS32_HI`` 2 ``word32`` (S + A) >> 32 700 ``R_AMDGPU_ABS64`` 3 ``word64`` S + A 701 ``R_AMDGPU_REL32`` 4 ``word32`` S + A - P 702 ``R_AMDGPU_REL64`` 5 ``word64`` S + A - P 703 ``R_AMDGPU_ABS32`` 6 ``word32`` S + A 704 ``R_AMDGPU_GOTPCREL`` 7 ``word32`` G + GOT + A - P 705 ``R_AMDGPU_GOTPCREL32_LO`` 8 ``word32`` (G + GOT + A - P) & 0xFFFFFFFF 706 ``R_AMDGPU_GOTPCREL32_HI`` 9 ``word32`` (G + GOT + A - P) >> 32 707 ``R_AMDGPU_REL32_LO`` 10 ``word32`` (S + A - P) & 0xFFFFFFFF 708 ``R_AMDGPU_REL32_HI`` 11 ``word32`` (S + A - P) >> 32 709 ========================== ===== ========== ============================== 710 711.. _amdgpu-dwarf: 712 713DWARF 714----- 715 716Standard DWARF [DWARF]_ Version 2 sections can be generated. These contain 717information that maps the code object executable code and data to the source 718language constructs. It can be used by tools such as debuggers and profilers. 719 720Address Space Mapping 721~~~~~~~~~~~~~~~~~~~~~ 722 723The following address space mapping is used: 724 725 .. table:: AMDGPU DWARF Address Space Mapping 726 :name: amdgpu-dwarf-address-space-mapping-table 727 728 =================== ================= 729 DWARF Address Space Memory Space 730 =================== ================= 731 1 Private (Scratch) 732 2 Local (group/LDS) 733 *omitted* Global 734 *omitted* Constant 735 *omitted* Generic (Flat) 736 *not supported* Region (GDS) 737 =================== ================= 738 739See :ref:`amdgpu-address-spaces` for information on the memory space terminology 740used in the table. 741 742An ``address_class`` attribute is generated on pointer type DIEs to specify the 743DWARF address space of the value of the pointer when it is in the *private* or 744*local* address space. Otherwise the attribute is omitted. 745 746An ``XDEREF`` operation is generated in location list expressions for variables 747that are allocated in the *private* and *local* address space. Otherwise no 748``XDREF`` is omitted. 749 750Register Mapping 751~~~~~~~~~~~~~~~~ 752 753*This section is WIP.* 754 755.. TODO 756 Define DWARF register enumeration. 757 758 If want to present a wavefront state then should expose vector registers as 759 64 wide (rather than per work-item view that LLVM uses). Either as separate 760 registers, or a 64x4 byte single register. In either case use a new LANE op 761 (akin to XDREF) to select the current lane usage in a location 762 expression. This would also allow scalar register spilling to vector register 763 lanes to be expressed (currently no debug information is being generated for 764 spilling). If choose a wide single register approach then use LANE in 765 conjunction with PIECE operation to select the dword part of the register for 766 the current lane. If the separate register approach then use LANE to select 767 the register. 768 769Source Text 770~~~~~~~~~~~ 771 772*This section is WIP.* 773 774.. TODO 775 DWARF extension to include runtime generated source text. 776 777.. _amdgpu-code-conventions: 778 779Code Conventions 780================ 781 782This section provides code conventions used for each supported target triple OS 783(see :ref:`amdgpu-target-triples`). 784 785AMDHSA 786------ 787 788This section provides code conventions used when the target triple OS is 789``amdhsa`` (see :ref:`amdgpu-target-triples`). 790 791.. _amdgpu-amdhsa-hsa-code-object-metadata: 792 793Code Object Metadata 794~~~~~~~~~~~~~~~~~~~~ 795 796The code object metadata specifies extensible metadata associated with the code 797objects executed on HSA [HSA]_ compatible runtimes such as AMD's ROCm 798[AMD-ROCm]_. It is specified by the ``NT_AMD_AMDGPU_HSA_METADATA`` note record 799(see :ref:`amdgpu-note-records`) and is required when the target triple OS is 800``amdhsa`` (see :ref:`amdgpu-target-triples`). It must contain the minimum 801information necessary to support the ROCM kernel queries. For example, the 802segment sizes needed in a dispatch packet. In addition, a high level language 803runtime may require other information to be included. For example, the AMD 804OpenCL runtime records kernel argument information. 805 806The metadata is specified as a YAML formatted string (see [YAML]_ and 807:doc:`YamlIO`). 808 809.. TODO 810 Is the string null terminated? It probably should not if YAML allows it to 811 contain null characters, otherwise it should be. 812 813The metadata is represented as a single YAML document comprised of the mapping 814defined in table :ref:`amdgpu-amdhsa-code-object-metadata-mapping-table` and 815referenced tables. 816 817For boolean values, the string values of ``false`` and ``true`` are used for 818false and true respectively. 819 820Additional information can be added to the mappings. To avoid conflicts, any 821non-AMD key names should be prefixed by "*vendor-name*.". 822 823 .. table:: AMDHSA Code Object Metadata Mapping 824 :name: amdgpu-amdhsa-code-object-metadata-mapping-table 825 826 ========== ============== ========= ======================================= 827 String Key Value Type Required? Description 828 ========== ============== ========= ======================================= 829 "Version" sequence of Required - The first integer is the major 830 2 integers version. Currently 1. 831 - The second integer is the minor 832 version. Currently 0. 833 "Printf" sequence of Each string is encoded information 834 strings about a printf function call. The 835 encoded information is organized as 836 fields separated by colon (':'): 837 838 ``ID:N:S[0]:S[1]:...:S[N-1]:FormatString`` 839 840 where: 841 842 ``ID`` 843 A 32 bit integer as a unique id for 844 each printf function call 845 846 ``N`` 847 A 32 bit integer equal to the number 848 of arguments of printf function call 849 minus 1 850 851 ``S[i]`` (where i = 0, 1, ... , N-1) 852 32 bit integers for the size in bytes 853 of the i-th FormatString argument of 854 the printf function call 855 856 FormatString 857 The format string passed to the 858 printf function call. 859 "Kernels" sequence of Required Sequence of the mappings for each 860 mapping kernel in the code object. See 861 :ref:`amdgpu-amdhsa-code-object-kernel-metadata-mapping-table` 862 for the definition of the mapping. 863 ========== ============== ========= ======================================= 864 865.. 866 867 .. table:: AMDHSA Code Object Kernel Metadata Mapping 868 :name: amdgpu-amdhsa-code-object-kernel-metadata-mapping-table 869 870 ================= ============== ========= ================================ 871 String Key Value Type Required? Description 872 ================= ============== ========= ================================ 873 "Name" string Required Source name of the kernel. 874 "SymbolName" string Required Name of the kernel 875 descriptor ELF symbol. 876 "Language" string Source language of the kernel. 877 Values include: 878 879 - "OpenCL C" 880 - "OpenCL C++" 881 - "HCC" 882 - "OpenMP" 883 884 "LanguageVersion" sequence of - The first integer is the major 885 2 integers version. 886 - The second integer is the 887 minor version. 888 "Attrs" mapping Mapping of kernel attributes. 889 See 890 :ref:`amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table` 891 for the mapping definition. 892 "Arguments" sequence of Sequence of mappings of the 893 mapping kernel arguments. See 894 :ref:`amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table` 895 for the definition of the mapping. 896 "CodeProps" mapping Mapping of properties related to 897 the kernel code. See 898 :ref:`amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table` 899 for the mapping definition. 900 "DebugProps" mapping Mapping of properties related to 901 the kernel debugging. See 902 :ref:`amdgpu-amdhsa-code-object-kernel-debug-properties-metadata-mapping-table` 903 for the mapping definition. 904 ================= ============== ========= ================================ 905 906.. 907 908 .. table:: AMDHSA Code Object Kernel Attribute Metadata Mapping 909 :name: amdgpu-amdhsa-code-object-kernel-attribute-metadata-mapping-table 910 911 =================== ============== ========= ============================== 912 String Key Value Type Required? Description 913 =================== ============== ========= ============================== 914 "ReqdWorkGroupSize" sequence of The dispatch work-group size 915 3 integers X, Y, Z must correspond to the 916 specified values. 917 918 Corresponds to the OpenCL 919 ``reqd_work_group_size`` 920 attribute. 921 "WorkGroupSizeHint" sequence of The dispatch work-group size 922 3 integers X, Y, Z is likely to be the 923 specified values. 924 925 Corresponds to the OpenCL 926 ``work_group_size_hint`` 927 attribute. 928 "VecTypeHint" string The name of a scalar or vector 929 type. 930 931 Corresponds to the OpenCL 932 ``vec_type_hint`` attribute. 933 934 "RuntimeHandle" string The external symbol name 935 associated with a kernel. 936 OpenCL runtime allocates a 937 global buffer for the symbol 938 and saves the kernel's address 939 to it, which is used for 940 device side enqueueing. Only 941 available for device side 942 enqueued kernels. 943 =================== ============== ========= ============================== 944 945.. 946 947 .. table:: AMDHSA Code Object Kernel Argument Metadata Mapping 948 :name: amdgpu-amdhsa-code-object-kernel-argument-metadata-mapping-table 949 950 ================= ============== ========= ================================ 951 String Key Value Type Required? Description 952 ================= ============== ========= ================================ 953 "Name" string Kernel argument name. 954 "TypeName" string Kernel argument type name. 955 "Size" integer Required Kernel argument size in bytes. 956 "Align" integer Required Kernel argument alignment in 957 bytes. Must be a power of two. 958 "ValueKind" string Required Kernel argument kind that 959 specifies how to set up the 960 corresponding argument. 961 Values include: 962 963 "ByValue" 964 The argument is copied 965 directly into the kernarg. 966 967 "GlobalBuffer" 968 A global address space pointer 969 to the buffer data is passed 970 in the kernarg. 971 972 "DynamicSharedPointer" 973 A group address space pointer 974 to dynamically allocated LDS 975 is passed in the kernarg. 976 977 "Sampler" 978 A global address space 979 pointer to a S# is passed in 980 the kernarg. 981 982 "Image" 983 A global address space 984 pointer to a T# is passed in 985 the kernarg. 986 987 "Pipe" 988 A global address space pointer 989 to an OpenCL pipe is passed in 990 the kernarg. 991 992 "Queue" 993 A global address space pointer 994 to an OpenCL device enqueue 995 queue is passed in the 996 kernarg. 997 998 "HiddenGlobalOffsetX" 999 The OpenCL grid dispatch 1000 global offset for the X 1001 dimension is passed in the 1002 kernarg. 1003 1004 "HiddenGlobalOffsetY" 1005 The OpenCL grid dispatch 1006 global offset for the Y 1007 dimension is passed in the 1008 kernarg. 1009 1010 "HiddenGlobalOffsetZ" 1011 The OpenCL grid dispatch 1012 global offset for the Z 1013 dimension is passed in the 1014 kernarg. 1015 1016 "HiddenNone" 1017 An argument that is not used 1018 by the kernel. Space needs to 1019 be left for it, but it does 1020 not need to be set up. 1021 1022 "HiddenPrintfBuffer" 1023 A global address space pointer 1024 to the runtime printf buffer 1025 is passed in kernarg. 1026 1027 "HiddenDefaultQueue" 1028 A global address space pointer 1029 to the OpenCL device enqueue 1030 queue that should be used by 1031 the kernel by default is 1032 passed in the kernarg. 1033 1034 "HiddenCompletionAction" 1035 *TBD* 1036 1037 .. TODO 1038 Add description. 1039 1040 "ValueType" string Required Kernel argument value type. Only 1041 present if "ValueKind" is 1042 "ByValue". For vector data 1043 types, the value is for the 1044 element type. Values include: 1045 1046 - "Struct" 1047 - "I8" 1048 - "U8" 1049 - "I16" 1050 - "U16" 1051 - "F16" 1052 - "I32" 1053 - "U32" 1054 - "F32" 1055 - "I64" 1056 - "U64" 1057 - "F64" 1058 1059 .. TODO 1060 How can it be determined if a 1061 vector type, and what size 1062 vector? 1063 "PointeeAlign" integer Alignment in bytes of pointee 1064 type for pointer type kernel 1065 argument. Must be a power 1066 of 2. Only present if 1067 "ValueKind" is 1068 "DynamicSharedPointer". 1069 "AddrSpaceQual" string Kernel argument address space 1070 qualifier. Only present if 1071 "ValueKind" is "GlobalBuffer" or 1072 "DynamicSharedPointer". Values 1073 are: 1074 1075 - "Private" 1076 - "Global" 1077 - "Constant" 1078 - "Local" 1079 - "Generic" 1080 - "Region" 1081 1082 .. TODO 1083 Is GlobalBuffer only Global 1084 or Constant? Is 1085 DynamicSharedPointer always 1086 Local? Can HCC allow Generic? 1087 How can Private or Region 1088 ever happen? 1089 "AccQual" string Kernel argument access 1090 qualifier. Only present if 1091 "ValueKind" is "Image" or 1092 "Pipe". Values 1093 are: 1094 1095 - "ReadOnly" 1096 - "WriteOnly" 1097 - "ReadWrite" 1098 1099 .. TODO 1100 Does this apply to 1101 GlobalBuffer? 1102 "ActualAcc" string The actual memory accesses 1103 performed by the kernel on the 1104 kernel argument. Only present if 1105 "ValueKind" is "GlobalBuffer", 1106 "Image", or "Pipe". This may be 1107 more restrictive than indicated 1108 by "AccQual" to reflect what the 1109 kernel actual does. If not 1110 present then the runtime must 1111 assume what is implied by 1112 "AccQual" and "IsConst". Values 1113 are: 1114 1115 - "ReadOnly" 1116 - "WriteOnly" 1117 - "ReadWrite" 1118 1119 "IsConst" boolean Indicates if the kernel argument 1120 is const qualified. Only present 1121 if "ValueKind" is 1122 "GlobalBuffer". 1123 1124 "IsRestrict" boolean Indicates if the kernel argument 1125 is restrict qualified. Only 1126 present if "ValueKind" is 1127 "GlobalBuffer". 1128 1129 "IsVolatile" boolean Indicates if the kernel argument 1130 is volatile qualified. Only 1131 present if "ValueKind" is 1132 "GlobalBuffer". 1133 1134 "IsPipe" boolean Indicates if the kernel argument 1135 is pipe qualified. Only present 1136 if "ValueKind" is "Pipe". 1137 1138 .. TODO 1139 Can GlobalBuffer be pipe 1140 qualified? 1141 ================= ============== ========= ================================ 1142 1143.. 1144 1145 .. table:: AMDHSA Code Object Kernel Code Properties Metadata Mapping 1146 :name: amdgpu-amdhsa-code-object-kernel-code-properties-metadata-mapping-table 1147 1148 ============================ ============== ========= ===================== 1149 String Key Value Type Required? Description 1150 ============================ ============== ========= ===================== 1151 "KernargSegmentSize" integer Required The size in bytes of 1152 the kernarg segment 1153 that holds the values 1154 of the arguments to 1155 the kernel. 1156 "GroupSegmentFixedSize" integer Required The amount of group 1157 segment memory 1158 required by a 1159 work-group in 1160 bytes. This does not 1161 include any 1162 dynamically allocated 1163 group segment memory 1164 that may be added 1165 when the kernel is 1166 dispatched. 1167 "PrivateSegmentFixedSize" integer Required The amount of fixed 1168 private address space 1169 memory required for a 1170 work-item in 1171 bytes. If 1172 IsDynamicCallstack 1173 is 1 then additional 1174 space must be added 1175 to this value for the 1176 call stack. 1177 "KernargSegmentAlign" integer Required The maximum byte 1178 alignment of 1179 arguments in the 1180 kernarg segment. Must 1181 be a power of 2. 1182 "WavefrontSize" integer Required Wavefront size. Must 1183 be a power of 2. 1184 "NumSGPRs" integer Number of scalar 1185 registers used by a 1186 wavefront for 1187 GFX6-GFX9. This 1188 includes the special 1189 SGPRs for VCC, Flat 1190 Scratch (GFX7-GFX9) 1191 and XNACK (for 1192 GFX8-GFX9). It does 1193 not include the 16 1194 SGPR added if a trap 1195 handler is 1196 enabled. It is not 1197 rounded up to the 1198 allocation 1199 granularity. 1200 "NumVGPRs" integer Number of vector 1201 registers used by 1202 each work-item for 1203 GFX6-GFX9 1204 "MaxFlatWorkgroupSize" integer Maximum flat 1205 work-group size 1206 supported by the 1207 kernel in work-items. 1208 "IsDynamicCallStack" boolean Indicates if the 1209 generated machine 1210 code is using a 1211 dynamically sized 1212 call stack. 1213 "IsXNACKEnabled" boolean Indicates if the 1214 generated machine 1215 code is capable of 1216 supporting XNACK. 1217 ============================ ============== ========= ===================== 1218 1219.. 1220 1221 .. table:: AMDHSA Code Object Kernel Debug Properties Metadata Mapping 1222 :name: amdgpu-amdhsa-code-object-kernel-debug-properties-metadata-mapping-table 1223 1224 =================================== ============== ========= ============== 1225 String Key Value Type Required? Description 1226 =================================== ============== ========= ============== 1227 "DebuggerABIVersion" string 1228 "ReservedNumVGPRs" integer 1229 "ReservedFirstVGPR" integer 1230 "PrivateSegmentBufferSGPR" integer 1231 "WavefrontPrivateSegmentOffsetSGPR" integer 1232 =================================== ============== ========= ============== 1233 1234.. TODO 1235 Plan to remove the debug properties metadata. 1236 1237Kernel Dispatch 1238~~~~~~~~~~~~~~~ 1239 1240The HSA architected queuing language (AQL) defines a user space memory interface 1241that can be used to control the dispatch of kernels, in an agent independent 1242way. An agent can have zero or more AQL queues created for it using the ROCm 1243runtime, in which AQL packets (all of which are 64 bytes) can be placed. See the 1244*HSA Platform System Architecture Specification* [HSA]_ for the AQL queue 1245mechanics and packet layouts. 1246 1247The packet processor of a kernel agent is responsible for detecting and 1248dispatching HSA kernels from the AQL queues associated with it. For AMD GPUs the 1249packet processor is implemented by the hardware command processor (CP), 1250asynchronous dispatch controller (ADC) and shader processor input controller 1251(SPI). 1252 1253The ROCm runtime can be used to allocate an AQL queue object. It uses the kernel 1254mode driver to initialize and register the AQL queue with CP. 1255 1256To dispatch a kernel the following actions are performed. This can occur in the 1257CPU host program, or from an HSA kernel executing on a GPU. 1258 12591. A pointer to an AQL queue for the kernel agent on which the kernel is to be 1260 executed is obtained. 12612. A pointer to the kernel descriptor (see 1262 :ref:`amdgpu-amdhsa-kernel-descriptor`) of the kernel to execute is 1263 obtained. It must be for a kernel that is contained in a code object that that 1264 was loaded by the ROCm runtime on the kernel agent with which the AQL queue is 1265 associated. 12663. Space is allocated for the kernel arguments using the ROCm runtime allocator 1267 for a memory region with the kernarg property for the kernel agent that will 1268 execute the kernel. It must be at least 16 byte aligned. 12694. Kernel argument values are assigned to the kernel argument memory 1270 allocation. The layout is defined in the *HSA Programmer’s Language Reference* 1271 [HSA]_. For AMDGPU the kernel execution directly accesses the kernel argument 1272 memory in the same way constant memory is accessed. (Note that the HSA 1273 specification allows an implementation to copy the kernel argument contents to 1274 another location that is accessed by the kernel.) 12755. An AQL kernel dispatch packet is created on the AQL queue. The ROCm runtime 1276 api uses 64 bit atomic operations to reserve space in the AQL queue for the 1277 packet. The packet must be set up, and the final write must use an atomic 1278 store release to set the packet kind to ensure the packet contents are 1279 visible to the kernel agent. AQL defines a doorbell signal mechanism to 1280 notify the kernel agent that the AQL queue has been updated. These rules, and 1281 the layout of the AQL queue and kernel dispatch packet is defined in the *HSA 1282 System Architecture Specification* [HSA]_. 12836. A kernel dispatch packet includes information about the actual dispatch, 1284 such as grid and work-group size, together with information from the code 1285 object about the kernel, such as segment sizes. The ROCm runtime queries on 1286 the kernel symbol can be used to obtain the code object values which are 1287 recorded in the :ref:`amdgpu-amdhsa-hsa-code-object-metadata`. 12887. CP executes micro-code and is responsible for detecting and setting up the 1289 GPU to execute the wavefronts of a kernel dispatch. 12908. CP ensures that when the a wavefront starts executing the kernel machine 1291 code, the scalar general purpose registers (SGPR) and vector general purpose 1292 registers (VGPR) are set up as required by the machine code. The required 1293 setup is defined in the :ref:`amdgpu-amdhsa-kernel-descriptor`. The initial 1294 register state is defined in 1295 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`. 12969. The prolog of the kernel machine code (see 1297 :ref:`amdgpu-amdhsa-kernel-prolog`) sets up the machine state as necessary 1298 before continuing executing the machine code that corresponds to the kernel. 129910. When the kernel dispatch has completed execution, CP signals the completion 1300 signal specified in the kernel dispatch packet if not 0. 1301 1302.. _amdgpu-amdhsa-memory-spaces: 1303 1304Memory Spaces 1305~~~~~~~~~~~~~ 1306 1307The memory space properties are: 1308 1309 .. table:: AMDHSA Memory Spaces 1310 :name: amdgpu-amdhsa-memory-spaces-table 1311 1312 ================= =========== ======== ======= ================== 1313 Memory Space Name HSA Segment Hardware Address NULL Value 1314 Name Name Size 1315 ================= =========== ======== ======= ================== 1316 Private private scratch 32 0x00000000 1317 Local group LDS 32 0xFFFFFFFF 1318 Global global global 64 0x0000000000000000 1319 Constant constant *same as 64 0x0000000000000000 1320 global* 1321 Generic flat flat 64 0x0000000000000000 1322 Region N/A GDS 32 *not implemented 1323 for AMDHSA* 1324 ================= =========== ======== ======= ================== 1325 1326The global and constant memory spaces both use global virtual addresses, which 1327are the same virtual address space used by the CPU. However, some virtual 1328addresses may only be accessible to the CPU, some only accessible by the GPU, 1329and some by both. 1330 1331Using the constant memory space indicates that the data will not change during 1332the execution of the kernel. This allows scalar read instructions to be 1333used. The vector and scalar L1 caches are invalidated of volatile data before 1334each kernel dispatch execution to allow constant memory to change values between 1335kernel dispatches. 1336 1337The local memory space uses the hardware Local Data Store (LDS) which is 1338automatically allocated when the hardware creates work-groups of wavefronts, and 1339freed when all the wavefronts of a work-group have terminated. The data store 1340(DS) instructions can be used to access it. 1341 1342The private memory space uses the hardware scratch memory support. If the kernel 1343uses scratch, then the hardware allocates memory that is accessed using 1344wavefront lane dword (4 byte) interleaving. The mapping used from private 1345address to physical address is: 1346 1347 ``wavefront-scratch-base + 1348 (private-address * wavefront-size * 4) + 1349 (wavefront-lane-id * 4)`` 1350 1351There are different ways that the wavefront scratch base address is determined 1352by a wavefront (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). This 1353memory can be accessed in an interleaved manner using buffer instruction with 1354the scratch buffer descriptor and per wave scratch offset, by the scratch 1355instructions, or by flat instructions. If each lane of a wavefront accesses the 1356same private address, the interleaving results in adjacent dwords being accessed 1357and hence requires fewer cache lines to be fetched. Multi-dword access is not 1358supported except by flat and scratch instructions in GFX9. 1359 1360The generic address space uses the hardware flat address support available in 1361GFX7-GFX9. This uses two fixed ranges of virtual addresses (the private and 1362local appertures), that are outside the range of addressible global memory, to 1363map from a flat address to a private or local address. 1364 1365FLAT instructions can take a flat address and access global, private (scratch) 1366and group (LDS) memory depending in if the address is within one of the 1367apperture ranges. Flat access to scratch requires hardware aperture setup and 1368setup in the kernel prologue (see :ref:`amdgpu-amdhsa-flat-scratch`). Flat 1369access to LDS requires hardware aperture setup and M0 (GFX7-GFX8) register setup 1370(see :ref:`amdgpu-amdhsa-m0`). 1371 1372To convert between a segment address and a flat address the base address of the 1373appertures address can be used. For GFX7-GFX8 these are available in the 1374:ref:`amdgpu-amdhsa-hsa-aql-queue` the address of which can be obtained with 1375Queue Ptr SGPR (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). For 1376GFX9 the appature base addresses are directly available as inline constant 1377registers ``SRC_SHARED_BASE/LIMIT`` and ``SRC_PRIVATE_BASE/LIMIT``. In 64 bit 1378address mode the apperture sizes are 2^32 bytes and the base is aligned to 2^32 1379which makes it easier to convert from flat to segment or segment to flat. 1380 1381Image and Samplers 1382~~~~~~~~~~~~~~~~~~ 1383 1384Image and sample handles created by the ROCm runtime are 64 bit addresses of a 1385hardware 32 byte V# and 48 byte S# object respectively. In order to support the 1386HSA ``query_sampler`` operations two extra dwords are used to store the HSA BRIG 1387enumeration values for the queries that are not trivially deducible from the S# 1388representation. 1389 1390HSA Signals 1391~~~~~~~~~~~ 1392 1393HSA signal handles created by the ROCm runtime are 64 bit addresses of a 1394structure allocated in memory accessible from both the CPU and GPU. The 1395structure is defined by the ROCm runtime and subject to change between releases 1396(see [AMD-ROCm-github]_). 1397 1398.. _amdgpu-amdhsa-hsa-aql-queue: 1399 1400HSA AQL Queue 1401~~~~~~~~~~~~~ 1402 1403The HSA AQL queue structure is defined by the ROCm runtime and subject to change 1404between releases (see [AMD-ROCm-github]_). For some processors it contains 1405fields needed to implement certain language features such as the flat address 1406aperture bases. It also contains fields used by CP such as managing the 1407allocation of scratch memory. 1408 1409.. _amdgpu-amdhsa-kernel-descriptor: 1410 1411Kernel Descriptor 1412~~~~~~~~~~~~~~~~~ 1413 1414A kernel descriptor consists of the information needed by CP to initiate the 1415execution of a kernel, including the entry point address of the machine code 1416that implements the kernel. 1417 1418Kernel Descriptor for GFX6-GFX9 1419+++++++++++++++++++++++++++++++ 1420 1421CP microcode requires the Kernel descritor to be allocated on 64 byte alignment. 1422 1423 .. table:: Kernel Descriptor for GFX6-GFX9 1424 :name: amdgpu-amdhsa-kernel-descriptor-gfx6-gfx9-table 1425 1426 ======= ======= =============================== =========================== 1427 Bits Size Field Name Description 1428 ======= ======= =============================== =========================== 1429 31:0 4 bytes group_segment_fixed_size The amount of fixed local 1430 address space memory 1431 required for a work-group 1432 in bytes. This does not 1433 include any dynamically 1434 allocated local address 1435 space memory that may be 1436 added when the kernel is 1437 dispatched. 1438 63:32 4 bytes private_segment_fixed_size The amount of fixed 1439 private address space 1440 memory required for a 1441 work-item in bytes. If 1442 is_dynamic_callstack is 1 1443 then additional space must 1444 be added to this value for 1445 the call stack. 1446 95:64 4 bytes max_flat_workgroup_size Maximum flat work-group 1447 size supported by the 1448 kernel in work-items. 1449 96 1 bit is_dynamic_call_stack Indicates if the generated 1450 machine code is using a 1451 dynamically sized call 1452 stack. 1453 97 1 bit is_xnack_enabled Indicates if the generated 1454 machine code is capable of 1455 suppoting XNACK. 1456 127:98 30 bits Reserved. Must be 0. 1457 191:128 8 bytes kernel_code_entry_byte_offset Byte offset (possibly 1458 negative) from base 1459 address of kernel 1460 descriptor to kernel's 1461 entry point instruction 1462 which must be 256 byte 1463 aligned. 1464 383:192 24 Reserved. Must be 0. 1465 bytes 1466 415:384 4 bytes compute_pgm_rsrc1 Compute Shader (CS) 1467 program settings used by 1468 CP to set up 1469 ``COMPUTE_PGM_RSRC1`` 1470 configuration 1471 register. See 1472 :ref:`amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table`. 1473 447:416 4 bytes compute_pgm_rsrc2 Compute Shader (CS) 1474 program settings used by 1475 CP to set up 1476 ``COMPUTE_PGM_RSRC2`` 1477 configuration 1478 register. See 1479 :ref:`amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table`. 1480 448 1 bit enable_sgpr_private_segment Enable the setup of the 1481 _buffer SGPR user data registers 1482 (see 1483 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). 1484 1485 The total number of SGPR 1486 user data registers 1487 requested must not exceed 1488 16 and match value in 1489 ``compute_pgm_rsrc2.user_sgpr.user_sgpr_count``. 1490 Any requests beyond 16 1491 will be ignored. 1492 449 1 bit enable_sgpr_dispatch_ptr *see above* 1493 450 1 bit enable_sgpr_queue_ptr *see above* 1494 451 1 bit enable_sgpr_kernarg_segment_ptr *see above* 1495 452 1 bit enable_sgpr_dispatch_id *see above* 1496 453 1 bit enable_sgpr_flat_scratch_init *see above* 1497 454 1 bit enable_sgpr_private_segment *see above* 1498 _size 1499 455 1 bit enable_sgpr_grid_workgroup Not implemented in CP and 1500 _count_X should always be 0. 1501 456 1 bit enable_sgpr_grid_workgroup Not implemented in CP and 1502 _count_Y should always be 0. 1503 457 1 bit enable_sgpr_grid_workgroup Not implemented in CP and 1504 _count_Z should always be 0. 1505 463:458 6 bits Reserved. Must be 0. 1506 511:464 4 Reserved. Must be 0. 1507 bytes 1508 512 **Total size 64 bytes.** 1509 ======= =================================================================== 1510 1511.. 1512 1513 .. table:: compute_pgm_rsrc1 for GFX6-GFX9 1514 :name: amdgpu-amdhsa-compute_pgm_rsrc1_t-gfx6-gfx9-table 1515 1516 ======= ======= =============================== =========================================================================== 1517 Bits Size Field Name Description 1518 ======= ======= =============================== =========================================================================== 1519 5:0 6 bits granulated_workitem_vgpr_count Number of vector registers 1520 used by each work-item, 1521 granularity is device 1522 specific: 1523 1524 GFX6-9 1525 roundup((max-vgpg + 1) 1526 / 4) - 1 1527 1528 Used by CP to set up 1529 ``COMPUTE_PGM_RSRC1.VGPRS``. 1530 9:6 4 bits granulated_wavefront_sgpr_count Number of scalar registers 1531 used by a wavefront, 1532 granularity is device 1533 specific: 1534 1535 GFX6-8 1536 roundup((max-sgpg + 1) 1537 / 8) - 1 1538 GFX9 1539 roundup((max-sgpg + 1) 1540 / 16) - 1 1541 1542 Includes the special SGPRs 1543 for VCC, Flat Scratch (for 1544 GFX7 onwards) and XNACK 1545 (for GFX8 onwards). It does 1546 not include the 16 SGPR 1547 added if a trap handler is 1548 enabled. 1549 1550 Used by CP to set up 1551 ``COMPUTE_PGM_RSRC1.SGPRS``. 1552 11:10 2 bits priority Must be 0. 1553 1554 Start executing wavefront 1555 at the specified priority. 1556 1557 CP is responsible for 1558 filling in 1559 ``COMPUTE_PGM_RSRC1.PRIORITY``. 1560 13:12 2 bits float_mode_round_32 Wavefront starts execution 1561 with specified rounding 1562 mode for single (32 1563 bit) floating point 1564 precision floating point 1565 operations. 1566 1567 Floating point rounding 1568 mode values are defined in 1569 :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`. 1570 1571 Used by CP to set up 1572 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``. 1573 15:14 2 bits float_mode_round_16_64 Wavefront starts execution 1574 with specified rounding 1575 denorm mode for half/double (16 1576 and 64 bit) floating point 1577 precision floating point 1578 operations. 1579 1580 Floating point rounding 1581 mode values are defined in 1582 :ref:`amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table`. 1583 1584 Used by CP to set up 1585 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``. 1586 17:16 2 bits float_mode_denorm_32 Wavefront starts execution 1587 with specified denorm mode 1588 for single (32 1589 bit) floating point 1590 precision floating point 1591 operations. 1592 1593 Floating point denorm mode 1594 values are defined in 1595 :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`. 1596 1597 Used by CP to set up 1598 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``. 1599 19:18 2 bits float_mode_denorm_16_64 Wavefront starts execution 1600 with specified denorm mode 1601 for half/double (16 1602 and 64 bit) floating point 1603 precision floating point 1604 operations. 1605 1606 Floating point denorm mode 1607 values are defined in 1608 :ref:`amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table`. 1609 1610 Used by CP to set up 1611 ``COMPUTE_PGM_RSRC1.FLOAT_MODE``. 1612 20 1 bit priv Must be 0. 1613 1614 Start executing wavefront 1615 in privilege trap handler 1616 mode. 1617 1618 CP is responsible for 1619 filling in 1620 ``COMPUTE_PGM_RSRC1.PRIV``. 1621 21 1 bit enable_dx10_clamp Wavefront starts execution 1622 with DX10 clamp mode 1623 enabled. Used by the vector 1624 ALU to force DX-10 style 1625 treatment of NaN's (when 1626 set, clamp NaN to zero, 1627 otherwise pass NaN 1628 through). 1629 1630 Used by CP to set up 1631 ``COMPUTE_PGM_RSRC1.DX10_CLAMP``. 1632 22 1 bit debug_mode Must be 0. 1633 1634 Start executing wavefront 1635 in single step mode. 1636 1637 CP is responsible for 1638 filling in 1639 ``COMPUTE_PGM_RSRC1.DEBUG_MODE``. 1640 23 1 bit enable_ieee_mode Wavefront starts execution 1641 with IEEE mode 1642 enabled. Floating point 1643 opcodes that support 1644 exception flag gathering 1645 will quiet and propagate 1646 signaling-NaN inputs per 1647 IEEE 754-2008. Min_dx10 and 1648 max_dx10 become IEEE 1649 754-2008 compliant due to 1650 signaling-NaN propagation 1651 and quieting. 1652 1653 Used by CP to set up 1654 ``COMPUTE_PGM_RSRC1.IEEE_MODE``. 1655 24 1 bit bulky Must be 0. 1656 1657 Only one work-group allowed 1658 to execute on a compute 1659 unit. 1660 1661 CP is responsible for 1662 filling in 1663 ``COMPUTE_PGM_RSRC1.BULKY``. 1664 25 1 bit cdbg_user Must be 0. 1665 1666 Flag that can be used to 1667 control debugging code. 1668 1669 CP is responsible for 1670 filling in 1671 ``COMPUTE_PGM_RSRC1.CDBG_USER``. 1672 31:26 6 bits Reserved. Must be 0. 1673 32 **Total size 4 bytes** 1674 ======= =================================================================================================================== 1675 1676.. 1677 1678 .. table:: compute_pgm_rsrc2 for GFX6-GFX9 1679 :name: amdgpu-amdhsa-compute_pgm_rsrc2-gfx6-gfx9-table 1680 1681 ======= ======= =============================== =========================================================================== 1682 Bits Size Field Name Description 1683 ======= ======= =============================== =========================================================================== 1684 0 1 bit enable_sgpr_private_segment Enable the setup of the 1685 _wave_offset SGPR wave scratch offset 1686 system register (see 1687 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). 1688 1689 Used by CP to set up 1690 ``COMPUTE_PGM_RSRC2.SCRATCH_EN``. 1691 5:1 5 bits user_sgpr_count The total number of SGPR 1692 user data registers 1693 requested. This number must 1694 match the number of user 1695 data registers enabled. 1696 1697 Used by CP to set up 1698 ``COMPUTE_PGM_RSRC2.USER_SGPR``. 1699 6 1 bit enable_trap_handler Set to 1 if code contains a 1700 TRAP instruction which 1701 requires a trap handler to 1702 be enabled. 1703 1704 CP sets 1705 ``COMPUTE_PGM_RSRC2.TRAP_PRESENT`` 1706 if the runtime has 1707 installed a trap handler 1708 regardless of the setting 1709 of this field. 1710 7 1 bit enable_sgpr_workgroup_id_x Enable the setup of the 1711 system SGPR register for 1712 the work-group id in the X 1713 dimension (see 1714 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). 1715 1716 Used by CP to set up 1717 ``COMPUTE_PGM_RSRC2.TGID_X_EN``. 1718 8 1 bit enable_sgpr_workgroup_id_y Enable the setup of the 1719 system SGPR register for 1720 the work-group id in the Y 1721 dimension (see 1722 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). 1723 1724 Used by CP to set up 1725 ``COMPUTE_PGM_RSRC2.TGID_Y_EN``. 1726 9 1 bit enable_sgpr_workgroup_id_z Enable the setup of the 1727 system SGPR register for 1728 the work-group id in the Z 1729 dimension (see 1730 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). 1731 1732 Used by CP to set up 1733 ``COMPUTE_PGM_RSRC2.TGID_Z_EN``. 1734 10 1 bit enable_sgpr_workgroup_info Enable the setup of the 1735 system SGPR register for 1736 work-group information (see 1737 :ref:`amdgpu-amdhsa-initial-kernel-execution-state`). 1738 1739 Used by CP to set up 1740 ``COMPUTE_PGM_RSRC2.TGID_SIZE_EN``. 1741 12:11 2 bits enable_vgpr_workitem_id Enable the setup of the 1742 VGPR system registers used 1743 for the work-item ID. 1744 :ref:`amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table` 1745 defines the values. 1746 1747 Used by CP to set up 1748 ``COMPUTE_PGM_RSRC2.TIDIG_CMP_CNT``. 1749 13 1 bit enable_exception_address_watch Must be 0. 1750 1751 Wavefront starts execution 1752 with address watch 1753 exceptions enabled which 1754 are generated when L1 has 1755 witnessed a thread access 1756 an *address of 1757 interest*. 1758 1759 CP is responsible for 1760 filling in the address 1761 watch bit in 1762 ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB`` 1763 according to what the 1764 runtime requests. 1765 14 1 bit enable_exception_memory Must be 0. 1766 1767 Wavefront starts execution 1768 with memory violation 1769 exceptions exceptions 1770 enabled which are generated 1771 when a memory violation has 1772 occurred for this wave from 1773 L1 or LDS 1774 (write-to-read-only-memory, 1775 mis-aligned atomic, LDS 1776 address out of range, 1777 illegal address, etc.). 1778 1779 CP sets the memory 1780 violation bit in 1781 ``COMPUTE_PGM_RSRC2.EXCP_EN_MSB`` 1782 according to what the 1783 runtime requests. 1784 23:15 9 bits granulated_lds_size Must be 0. 1785 1786 CP uses the rounded value 1787 from the dispatch packet, 1788 not this value, as the 1789 dispatch may contain 1790 dynamically allocated group 1791 segment memory. CP writes 1792 directly to 1793 ``COMPUTE_PGM_RSRC2.LDS_SIZE``. 1794 1795 Amount of group segment 1796 (LDS) to allocate for each 1797 work-group. Granularity is 1798 device specific: 1799 1800 GFX6: 1801 roundup(lds-size / (64 * 4)) 1802 GFX7-GFX9: 1803 roundup(lds-size / (128 * 4)) 1804 1805 24 1 bit enable_exception_ieee_754_fp Wavefront starts execution 1806 _invalid_operation with specified exceptions 1807 enabled. 1808 1809 Used by CP to set up 1810 ``COMPUTE_PGM_RSRC2.EXCP_EN`` 1811 (set from bits 0..6). 1812 1813 IEEE 754 FP Invalid 1814 Operation 1815 25 1 bit enable_exception_fp_denormal FP Denormal one or more 1816 _source input operands is a 1817 denormal number 1818 26 1 bit enable_exception_ieee_754_fp IEEE 754 FP Division by 1819 _division_by_zero Zero 1820 27 1 bit enable_exception_ieee_754_fp IEEE 754 FP FP Overflow 1821 _overflow 1822 28 1 bit enable_exception_ieee_754_fp IEEE 754 FP Underflow 1823 _underflow 1824 29 1 bit enable_exception_ieee_754_fp IEEE 754 FP Inexact 1825 _inexact 1826 30 1 bit enable_exception_int_divide_by Integer Division by Zero 1827 _zero (rcp_iflag_f32 instruction 1828 only) 1829 31 1 bit Reserved. Must be 0. 1830 32 **Total size 4 bytes.** 1831 ======= =================================================================================================================== 1832 1833.. 1834 1835 .. table:: Floating Point Rounding Mode Enumeration Values 1836 :name: amdgpu-amdhsa-floating-point-rounding-mode-enumeration-values-table 1837 1838 ===================================== ===== =============================== 1839 Enumeration Name Value Description 1840 ===================================== ===== =============================== 1841 AMD_FLOAT_ROUND_MODE_NEAR_EVEN 0 Round Ties To Even 1842 AMD_FLOAT_ROUND_MODE_PLUS_INFINITY 1 Round Toward +infinity 1843 AMD_FLOAT_ROUND_MODE_MINUS_INFINITY 2 Round Toward -infinity 1844 AMD_FLOAT_ROUND_MODE_ZERO 3 Round Toward 0 1845 ===================================== ===== =============================== 1846 1847.. 1848 1849 .. table:: Floating Point Denorm Mode Enumeration Values 1850 :name: amdgpu-amdhsa-floating-point-denorm-mode-enumeration-values-table 1851 1852 ===================================== ===== =============================== 1853 Enumeration Name Value Description 1854 ===================================== ===== =============================== 1855 AMD_FLOAT_DENORM_MODE_FLUSH_SRC_DST 0 Flush Source and Destination 1856 Denorms 1857 AMD_FLOAT_DENORM_MODE_FLUSH_DST 1 Flush Output Denorms 1858 AMD_FLOAT_DENORM_MODE_FLUSH_SRC 2 Flush Source Denorms 1859 AMD_FLOAT_DENORM_MODE_FLUSH_NONE 3 No Flush 1860 ===================================== ===== =============================== 1861 1862.. 1863 1864 .. table:: System VGPR Work-Item ID Enumeration Values 1865 :name: amdgpu-amdhsa-system-vgpr-work-item-id-enumeration-values-table 1866 1867 ===================================== ===== =============================== 1868 Enumeration Name Value Description 1869 ===================================== ===== =============================== 1870 AMD_SYSTEM_VGPR_WORKITEM_ID_X 0 Set work-item X dimension ID. 1871 AMD_SYSTEM_VGPR_WORKITEM_ID_X_Y 1 Set work-item X and Y 1872 dimensions ID. 1873 AMD_SYSTEM_VGPR_WORKITEM_ID_X_Y_Z 2 Set work-item X, Y and Z 1874 dimensions ID. 1875 AMD_SYSTEM_VGPR_WORKITEM_ID_UNDEFINED 3 Undefined. 1876 ===================================== ===== =============================== 1877 1878.. _amdgpu-amdhsa-initial-kernel-execution-state: 1879 1880Initial Kernel Execution State 1881~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ 1882 1883This section defines the register state that will be set up by the packet 1884processor prior to the start of execution of every wavefront. This is limited by 1885the constraints of the hardware controllers of CP/ADC/SPI. 1886 1887The order of the SGPR registers is defined, but the compiler can specify which 1888ones are actually setup in the kernel descriptor using the ``enable_sgpr_*`` bit 1889fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used 1890for enabled registers are dense starting at SGPR0: the first enabled register is 1891SGPR0, the next enabled register is SGPR1 etc.; disabled registers do not have 1892an SGPR number. 1893 1894The initial SGPRs comprise up to 16 User SRGPs that are set by CP and apply to 1895all waves of the grid. It is possible to specify more than 16 User SGPRs using 1896the ``enable_sgpr_*`` bit fields, in which case only the first 16 are actually 1897initialized. These are then immediately followed by the System SGPRs that are 1898set up by ADC/SPI and can have different values for each wave of the grid 1899dispatch. 1900 1901SGPR register initial state is defined in 1902:ref:`amdgpu-amdhsa-sgpr-register-set-up-order-table`. 1903 1904 .. table:: SGPR Register Set Up Order 1905 :name: amdgpu-amdhsa-sgpr-register-set-up-order-table 1906 1907 ========== ========================== ====== ============================== 1908 SGPR Order Name Number Description 1909 (kernel descriptor enable of 1910 field) SGPRs 1911 ========== ========================== ====== ============================== 1912 First Private Segment Buffer 4 V# that can be used, together 1913 (enable_sgpr_private with Scratch Wave Offset as an 1914 _segment_buffer) offset, to access the private 1915 memory space using a segment 1916 address. 1917 1918 CP uses the value provided by 1919 the runtime. 1920 then Dispatch Ptr 2 64 bit address of AQL dispatch 1921 (enable_sgpr_dispatch_ptr) packet for kernel dispatch 1922 actually executing. 1923 then Queue Ptr 2 64 bit address of amd_queue_t 1924 (enable_sgpr_queue_ptr) object for AQL queue on which 1925 the dispatch packet was 1926 queued. 1927 then Kernarg Segment Ptr 2 64 bit address of Kernarg 1928 (enable_sgpr_kernarg segment. This is directly 1929 _segment_ptr) copied from the 1930 kernarg_address in the kernel 1931 dispatch packet. 1932 1933 Having CP load it once avoids 1934 loading it at the beginning of 1935 every wavefront. 1936 then Dispatch Id 2 64 bit Dispatch ID of the 1937 (enable_sgpr_dispatch_id) dispatch packet being 1938 executed. 1939 then Flat Scratch Init 2 This is 2 SGPRs: 1940 (enable_sgpr_flat_scratch 1941 _init) GFX6 1942 Not supported. 1943 GFX7-GFX8 1944 The first SGPR is a 32 bit 1945 byte offset from 1946 ``SH_HIDDEN_PRIVATE_BASE_VIMID`` 1947 to per SPI base of memory 1948 for scratch for the queue 1949 executing the kernel 1950 dispatch. CP obtains this 1951 from the runtime. (The 1952 Scratch Segment Buffer base 1953 address is 1954 ``SH_HIDDEN_PRIVATE_BASE_VIMID`` 1955 plus this offset.) The value 1956 of Scratch Wave Offset must 1957 be added to this offset by 1958 the kernel machine code, 1959 right shifted by 8, and 1960 moved to the FLAT_SCRATCH_HI 1961 SGPR register. 1962 FLAT_SCRATCH_HI corresponds 1963 to SGPRn-4 on GFX7, and 1964 SGPRn-6 on GFX8 (where SGPRn 1965 is the highest numbered SGPR 1966 allocated to the wave). 1967 FLAT_SCRATCH_HI is 1968 multiplied by 256 (as it is 1969 in units of 256 bytes) and 1970 added to 1971 ``SH_HIDDEN_PRIVATE_BASE_VIMID`` 1972 to calculate the per wave 1973 FLAT SCRATCH BASE in flat 1974 memory instructions that 1975 access the scratch 1976 apperture. 1977 1978 The second SGPR is 32 bit 1979 byte size of a single 1980 work-item’s scratch memory 1981 usage. CP obtains this from 1982 the runtime, and it is 1983 always a multiple of DWORD. 1984 CP checks that the value in 1985 the kernel dispatch packet 1986 Private Segment Byte Size is 1987 not larger, and requests the 1988 runtime to increase the 1989 queue's scratch size if 1990 necessary. The kernel code 1991 must move it to 1992 FLAT_SCRATCH_LO which is 1993 SGPRn-3 on GFX7 and SGPRn-5 1994 on GFX8. FLAT_SCRATCH_LO is 1995 used as the FLAT SCRATCH 1996 SIZE in flat memory 1997 instructions. Having CP load 1998 it once avoids loading it at 1999 the beginning of every 2000 wavefront. GFX9 This is the 2001 64 bit base address of the 2002 per SPI scratch backing 2003 memory managed by SPI for 2004 the queue executing the 2005 kernel dispatch. CP obtains 2006 this from the runtime (and 2007 divides it if there are 2008 multiple Shader Arrays each 2009 with its own SPI). The value 2010 of Scratch Wave Offset must 2011 be added by the kernel 2012 machine code and the result 2013 moved to the FLAT_SCRATCH 2014 SGPR which is SGPRn-6 and 2015 SGPRn-5. It is used as the 2016 FLAT SCRATCH BASE in flat 2017 memory instructions. then 2018 Private Segment Size 1 The 2019 32 bit byte size of a 2020 (enable_sgpr_private single 2021 work-item's 2022 scratch_segment_size) memory 2023 allocation. This is the 2024 value from the kernel 2025 dispatch packet Private 2026 Segment Byte Size rounded up 2027 by CP to a multiple of 2028 DWORD. 2029 2030 Having CP load it once avoids 2031 loading it at the beginning of 2032 every wavefront. 2033 2034 This is not used for 2035 GFX7-GFX8 since it is the same 2036 value as the second SGPR of 2037 Flat Scratch Init. However, it 2038 may be needed for GFX9 which 2039 changes the meaning of the 2040 Flat Scratch Init value. 2041 then Grid Work-Group Count X 1 32 bit count of the number of 2042 (enable_sgpr_grid work-groups in the X dimension 2043 _workgroup_count_X) for the grid being 2044 executed. Computed from the 2045 fields in the kernel dispatch 2046 packet as ((grid_size.x + 2047 workgroup_size.x - 1) / 2048 workgroup_size.x). 2049 then Grid Work-Group Count Y 1 32 bit count of the number of 2050 (enable_sgpr_grid work-groups in the Y dimension 2051 _workgroup_count_Y && for the grid being 2052 less than 16 previous executed. Computed from the 2053 SGPRs) fields in the kernel dispatch 2054 packet as ((grid_size.y + 2055 workgroup_size.y - 1) / 2056 workgroupSize.y). 2057 2058 Only initialized if <16 2059 previous SGPRs initialized. 2060 then Grid Work-Group Count Z 1 32 bit count of the number of 2061 (enable_sgpr_grid work-groups in the Z dimension 2062 _workgroup_count_Z && for the grid being 2063 less than 16 previous executed. Computed from the 2064 SGPRs) fields in the kernel dispatch 2065 packet as ((grid_size.z + 2066 workgroup_size.z - 1) / 2067 workgroupSize.z). 2068 2069 Only initialized if <16 2070 previous SGPRs initialized. 2071 then Work-Group Id X 1 32 bit work-group id in X 2072 (enable_sgpr_workgroup_id dimension of grid for 2073 _X) wavefront. 2074 then Work-Group Id Y 1 32 bit work-group id in Y 2075 (enable_sgpr_workgroup_id dimension of grid for 2076 _Y) wavefront. 2077 then Work-Group Id Z 1 32 bit work-group id in Z 2078 (enable_sgpr_workgroup_id dimension of grid for 2079 _Z) wavefront. 2080 then Work-Group Info 1 {first_wave, 14’b0000, 2081 (enable_sgpr_workgroup ordered_append_term[10:0], 2082 _info) threadgroup_size_in_waves[5:0]} 2083 then Scratch Wave Offset 1 32 bit byte offset from base 2084 (enable_sgpr_private of scratch base of queue 2085 _segment_wave_offset) executing the kernel 2086 dispatch. Must be used as an 2087 offset with Private 2088 segment address when using 2089 Scratch Segment Buffer. It 2090 must be used to set up FLAT 2091 SCRATCH for flat addressing 2092 (see 2093 :ref:`amdgpu-amdhsa-flat-scratch`). 2094 ========== ========================== ====== ============================== 2095 2096The order of the VGPR registers is defined, but the compiler can specify which 2097ones are actually setup in the kernel descriptor using the ``enable_vgpr*`` bit 2098fields (see :ref:`amdgpu-amdhsa-kernel-descriptor`). The register numbers used 2099for enabled registers are dense starting at VGPR0: the first enabled register is 2100VGPR0, the next enabled register is VGPR1 etc.; disabled registers do not have a 2101VGPR number. 2102 2103VGPR register initial state is defined in 2104:ref:`amdgpu-amdhsa-vgpr-register-set-up-order-table`. 2105 2106 .. table:: VGPR Register Set Up Order 2107 :name: amdgpu-amdhsa-vgpr-register-set-up-order-table 2108 2109 ========== ========================== ====== ============================== 2110 VGPR Order Name Number Description 2111 (kernel descriptor enable of 2112 field) VGPRs 2113 ========== ========================== ====== ============================== 2114 First Work-Item Id X 1 32 bit work item id in X 2115 (Always initialized) dimension of work-group for 2116 wavefront lane. 2117 then Work-Item Id Y 1 32 bit work item id in Y 2118 (enable_vgpr_workitem_id dimension of work-group for 2119 > 0) wavefront lane. 2120 then Work-Item Id Z 1 32 bit work item id in Z 2121 (enable_vgpr_workitem_id dimension of work-group for 2122 > 1) wavefront lane. 2123 ========== ========================== ====== ============================== 2124 2125The setting of registers is is done by GPU CP/ADC/SPI hardware as follows: 2126 21271. SGPRs before the Work-Group Ids are set by CP using the 16 User Data 2128 registers. 21292. Work-group Id registers X, Y, Z are set by ADC which supports any 2130 combination including none. 21313. Scratch Wave Offset is set by SPI in a per wave basis which is why its value 2132 cannot included with the flat scratch init value which is per queue. 21334. The VGPRs are set by SPI which only supports specifying either (X), (X, Y) 2134 or (X, Y, Z). 2135 2136Flat Scratch register pair are adjacent SGRRs so they can be moved as a 64 bit 2137value to the hardware required SGPRn-3 and SGPRn-4 respectively. 2138 2139The global segment can be accessed either using buffer instructions (GFX6 which 2140has V# 64 bit address support), flat instructions (GFX7-9), or global 2141instructions (GFX9). 2142 2143If buffer operations are used then the compiler can generate a V# with the 2144following properties: 2145 2146* base address of 0 2147* no swizzle 2148* ATC: 1 if IOMMU present (such as APU) 2149* ptr64: 1 2150* MTYPE set to support memory coherence that matches the runtime (such as CC for 2151 APU and NC for dGPU). 2152 2153.. _amdgpu-amdhsa-kernel-prolog: 2154 2155Kernel Prolog 2156~~~~~~~~~~~~~ 2157 2158.. _amdgpu-amdhsa-m0: 2159 2160M0 2161++ 2162 2163GFX6-GFX8 2164 The M0 register must be initialized with a value at least the total LDS size 2165 if the kernel may access LDS via DS or flat operations. Total LDS size is 2166 available in dispatch packet. For M0, it is also possible to use maximum 2167 possible value of LDS for given target (0x7FFF for GFX6 and 0xFFFF for 2168 GFX7-GFX8). 2169GFX9 2170 The M0 register is not used for range checking LDS accesses and so does not 2171 need to be initialized in the prolog. 2172 2173.. _amdgpu-amdhsa-flat-scratch: 2174 2175Flat Scratch 2176++++++++++++ 2177 2178If the kernel may use flat operations to access scratch memory, the prolog code 2179must set up FLAT_SCRATCH register pair (FLAT_SCRATCH_LO/FLAT_SCRATCH_HI which 2180are in SGPRn-4/SGPRn-3). Initialization uses Flat Scratch Init and Scratch Wave 2181Offset SGPR registers (see :ref:`amdgpu-amdhsa-initial-kernel-execution-state`): 2182 2183GFX6 2184 Flat scratch is not supported. 2185 2186GFX7-8 2187 1. The low word of Flat Scratch Init is 32 bit byte offset from 2188 ``SH_HIDDEN_PRIVATE_BASE_VIMID`` to the base of scratch backing memory 2189 being managed by SPI for the queue executing the kernel dispatch. This is 2190 the same value used in the Scratch Segment Buffer V# base address. The 2191 prolog must add the value of Scratch Wave Offset to get the wave's byte 2192 scratch backing memory offset from ``SH_HIDDEN_PRIVATE_BASE_VIMID``. Since 2193 FLAT_SCRATCH_LO is in units of 256 bytes, the offset must be right shifted 2194 by 8 before moving into FLAT_SCRATCH_LO. 2195 2. The second word of Flat Scratch Init is 32 bit byte size of a single 2196 work-items scratch memory usage. This is directly loaded from the kernel 2197 dispatch packet Private Segment Byte Size and rounded up to a multiple of 2198 DWORD. Having CP load it once avoids loading it at the beginning of every 2199 wavefront. The prolog must move it to FLAT_SCRATCH_LO for use as FLAT SCRATCH 2200 SIZE. 2201GFX9 2202 The Flat Scratch Init is the 64 bit address of the base of scratch backing 2203 memory being managed by SPI for the queue executing the kernel dispatch. The 2204 prolog must add the value of Scratch Wave Offset and moved to the FLAT_SCRATCH 2205 pair for use as the flat scratch base in flat memory instructions. 2206 2207.. _amdgpu-amdhsa-memory-model: 2208 2209Memory Model 2210~~~~~~~~~~~~ 2211 2212This section describes the mapping of LLVM memory model onto AMDGPU machine code 2213(see :ref:`memmodel`). *The implementation is WIP.* 2214 2215.. TODO 2216 Update when implementation complete. 2217 2218 Support more relaxed OpenCL memory model to be controlled by environment 2219 component of target triple. 2220 2221The AMDGPU backend supports the memory synchronization scopes specified in 2222:ref:`amdgpu-memory-scopes`. 2223 2224The code sequences used to implement the memory model are defined in table 2225:ref:`amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx9-table`. 2226 2227The sequences specify the order of instructions that a single thread must 2228execute. The ``s_waitcnt`` and ``buffer_wbinvl1_vol`` are defined with respect 2229to other memory instructions executed by the same thread. This allows them to be 2230moved earlier or later which can allow them to be combined with other instances 2231of the same instruction, or hoisted/sunk out of loops to improve 2232performance. Only the instructions related to the memory model are given; 2233additional ``s_waitcnt`` instructions are required to ensure registers are 2234defined before being used. These may be able to be combined with the memory 2235model ``s_waitcnt`` instructions as described above. 2236 2237The AMDGPU memory model supports both the HSA [HSA]_ memory model, and the 2238OpenCL [OpenCL]_ memory model. The HSA memory model uses a single happens-before 2239relation for all address spaces (see :ref:`amdgpu-address-spaces`). The OpenCL 2240memory model which has separate happens-before relations for the global and 2241local address spaces, and only a fence specifying both global and local address 2242space joins the relationships. Since the LLVM ``memfence`` instruction does not 2243allow an address space to be specified the OpenCL fence has to convervatively 2244assume both local and global address space was specified. However, optimizations 2245can often be done to eliminate the additional ``s_waitcnt``instructions when 2246there are no intervening corresponding ``ds/flat_load/store/atomic`` memory 2247instructions. The code sequences in the table indicate what can be omitted for 2248the OpenCL memory. The target triple environment is used to determine if the 2249source language is OpenCL (see :ref:`amdgpu-opencl`). 2250 2251``ds/flat_load/store/atomic`` instructions to local memory are termed LDS 2252operations. 2253 2254``buffer/global/flat_load/store/atomic`` instructions to global memory are 2255termed vector memory operations. 2256 2257For GFX6-GFX9: 2258 2259* Each agent has multiple compute units (CU). 2260* Each CU has multiple SIMDs that execute wavefronts. 2261* The wavefronts for a single work-group are executed in the same CU but may be 2262 executed by different SIMDs. 2263* Each CU has a single LDS memory shared by the wavefronts of the work-groups 2264 executing on it. 2265* All LDS operations of a CU are performed as wavefront wide operations in a 2266 global order and involve no caching. Completion is reported to a wavefront in 2267 execution order. 2268* The LDS memory has multiple request queues shared by the SIMDs of a 2269 CU. Therefore, the LDS operations performed by different waves of a work-group 2270 can be reordered relative to each other, which can result in reordering the 2271 visibility of vector memory operations with respect to LDS operations of other 2272 wavefronts in the same work-group. A ``s_waitcnt lgkmcnt(0)`` is required to 2273 ensure synchronization between LDS operations and vector memory operations 2274 between waves of a work-group, but not between operations performed by the 2275 same wavefront. 2276* The vector memory operations are performed as wavefront wide operations and 2277 completion is reported to a wavefront in execution order. The exception is 2278 that for GFX7-9 ``flat_load/store/atomic`` instructions can report out of 2279 vector memory order if they access LDS memory, and out of LDS operation order 2280 if they access global memory. 2281* The vector memory operations access a vector L1 cache shared by all wavefronts 2282 on a CU. Therefore, no special action is required for coherence between 2283 wavefronts in the same work-group. A ``buffer_wbinvl1_vol`` is required for 2284 coherence between waves executing in different work-groups as they may be 2285 executing on different CUs. 2286* The scalar memory operations access a scalar L1 cache shared by all wavefronts 2287 on a group of CUs. The scalar and vector L1 caches are not coherent. However, 2288 scalar operations are used in a restricted way so do not impact the memory 2289 model. See :ref:`amdgpu-amdhsa-memory-spaces`. 2290* The vector and scalar memory operations use an L2 cache shared by all CUs on 2291 the same agent. 2292* The L2 cache has independent channels to service disjoint ranges of virtual 2293 addresses. 2294* Each CU has a separate request queue per channel. Therefore, the vector and 2295 scalar memory operations performed by waves executing in different work-groups 2296 (which may be executing on different CUs) of an agent can be reordered 2297 relative to each other. A ``s_waitcnt vmcnt(0)`` is required to ensure 2298 synchronization between vector memory operations of different CUs. It ensures a 2299 previous vector memory operation has completed before executing a subsequent 2300 vector memory or LDS operation and so can be used to meet the requirements of 2301 acquire and release. 2302* The L2 cache can be kept coherent with other agents on some targets, or ranges 2303 of virtual addresses can be set up to bypass it to ensure system coherence. 2304 2305Private address space uses ``buffer_load/store`` using the scratch V# (GFX6-8), 2306or ``scratch_load/store`` (GFX9). Since only a single thread is accessing the 2307memory, atomic memory orderings are not meaningful and all accesses are treated 2308as non-atomic. 2309 2310Constant address space uses ``buffer/global_load`` instructions (or equivalent 2311scalar memory instructions). Since the constant address space contents do not 2312change during the execution of a kernel dispatch it is not legal to perform 2313stores, and atomic memory orderings are not meaningful and all access are 2314treated as non-atomic. 2315 2316A memory synchronization scope wider than work-group is not meaningful for the 2317group (LDS) address space and is treated as work-group. 2318 2319The memory model does not support the region address space which is treated as 2320non-atomic. 2321 2322Acquire memory ordering is not meaningful on store atomic instructions and is 2323treated as non-atomic. 2324 2325Release memory ordering is not meaningful on load atomic instructions and is 2326treated a non-atomic. 2327 2328Acquire-release memory ordering is not meaningful on load or store atomic 2329instructions and is treated as acquire and release respectively. 2330 2331AMDGPU backend only uses scalar memory operations to access memory that is 2332proven to not change during the execution of the kernel dispatch. This includes 2333constant address space and global address space for program scope const 2334variables. Therefore the kernel machine code does not have to maintain the 2335scalar L1 cache to ensure it is coherent with the vector L1 cache. The scalar 2336and vector L1 caches are invalidated between kernel dispatches by CP since 2337constant address space data may change between kernel dispatch executions. See 2338:ref:`amdgpu-amdhsa-memory-spaces`. 2339 2340The one execption is if scalar writes are used to spill SGPR registers. In this 2341case the AMDGPU backend ensures the memory location used to spill is never 2342accessed by vector memory operations at the same time. If scalar writes are used 2343then a ``s_dcache_wb`` is inserted before the ``s_endpgm`` and before a function 2344return since the locations may be used for vector memory instructions by a 2345future wave that uses the same scratch area, or a function call that creates a 2346frame at the same address, respectively. There is no need for a ``s_dcache_inv`` 2347as all scalar writes are write-before-read in the same thread. 2348 2349Scratch backing memory (which is used for the private address space) is accessed 2350with MTYPE NC_NV (non-coherenent non-volatile). Since the private address space 2351is only accessed by a single thread, and is always write-before-read, 2352there is never a need to invalidate these entries from the L1 cache. Hence all 2353cache invalidates are done as ``*_vol`` to only invalidate the volatile cache 2354lines. 2355 2356On dGPU the kernarg backing memory is accessed as UC (uncached) to avoid needing 2357to invalidate the L2 cache. This also causes it to be treated as non-volatile 2358and so is not invalidated by ``*_vol``. On APU it is accessed as CC (cache 2359coherent) and so the L2 cache will coherent with the CPU and other agents. 2360 2361 .. table:: AMDHSA Memory Model Code Sequences GFX6-GFX9 2362 :name: amdgpu-amdhsa-memory-model-code-sequences-gfx6-gfx9-table 2363 2364 ============ ============ ============== ========== ======================= 2365 LLVM Instr LLVM Memory LLVM Memory AMDGPU AMDGPU Machine Code 2366 Ordering Sync Scope Address 2367 Space 2368 ============ ============ ============== ========== ======================= 2369 **Non-Atomic** 2370 --------------------------------------------------------------------------- 2371 load *none* *none* - global non-volatile 2372 - generic 1. buffer/global/flat_load 2373 volatile 2374 1. buffer/global/flat_load 2375 glc=1 2376 load *none* *none* - local 1. ds_load 2377 store *none* *none* - global 1. buffer/global/flat_store 2378 - generic 2379 store *none* *none* - local 1. ds_store 2380 **Unordered Atomic** 2381 --------------------------------------------------------------------------- 2382 load atomic unordered *any* *any* *Same as non-atomic*. 2383 store atomic unordered *any* *any* *Same as non-atomic*. 2384 atomicrmw unordered *any* *any* *Same as monotonic 2385 atomic*. 2386 **Monotonic Atomic** 2387 --------------------------------------------------------------------------- 2388 load atomic monotonic - singlethread - global 1. buffer/global/flat_load 2389 - wavefront - generic 2390 - workgroup 2391 load atomic monotonic - singlethread - local 1. ds_load 2392 - wavefront 2393 - workgroup 2394 load atomic monotonic - agent - global 1. buffer/global/flat_load 2395 - system - generic glc=1 2396 store atomic monotonic - singlethread - global 1. buffer/global/flat_store 2397 - wavefront - generic 2398 - workgroup 2399 - agent 2400 - system 2401 store atomic monotonic - singlethread - local 1. ds_store 2402 - wavefront 2403 - workgroup 2404 atomicrmw monotonic - singlethread - global 1. buffer/global/flat_atomic 2405 - wavefront - generic 2406 - workgroup 2407 - agent 2408 - system 2409 atomicrmw monotonic - singlethread - local 1. ds_atomic 2410 - wavefront 2411 - workgroup 2412 **Acquire Atomic** 2413 --------------------------------------------------------------------------- 2414 load atomic acquire - singlethread - global 1. buffer/global/ds/flat_load 2415 - wavefront - local 2416 - generic 2417 load atomic acquire - workgroup - global 1. buffer/global_load 2418 load atomic acquire - workgroup - local 1. ds/flat_load 2419 - generic 2. s_waitcnt lgkmcnt(0) 2420 2421 - If OpenCL, omit 2422 waitcnt. 2423 - Must happen before 2424 any following 2425 global/generic 2426 load/load 2427 atomic/store/store 2428 atomic/atomicrmw. 2429 - Ensures any 2430 following global 2431 data read is no 2432 older than the load 2433 atomic value being 2434 acquired. 2435 2436 load atomic acquire - agent - global 1. buffer/global_load 2437 - system glc=1 2438 2. s_waitcnt vmcnt(0) 2439 2440 - Must happen before 2441 following 2442 buffer_wbinvl1_vol. 2443 - Ensures the load 2444 has completed 2445 before invalidating 2446 the cache. 2447 2448 3. buffer_wbinvl1_vol 2449 2450 - Must happen before 2451 any following 2452 global/generic 2453 load/load 2454 atomic/atomicrmw. 2455 - Ensures that 2456 following 2457 loads will not see 2458 stale global data. 2459 2460 load atomic acquire - agent - generic 1. flat_load glc=1 2461 - system 2. s_waitcnt vmcnt(0) & 2462 lgkmcnt(0) 2463 2464 - If OpenCL omit 2465 lgkmcnt(0). 2466 - Must happen before 2467 following 2468 buffer_wbinvl1_vol. 2469 - Ensures the flat_load 2470 has completed 2471 before invalidating 2472 the cache. 2473 2474 3. buffer_wbinvl1_vol 2475 2476 - Must happen before 2477 any following 2478 global/generic 2479 load/load 2480 atomic/atomicrmw. 2481 - Ensures that 2482 following loads 2483 will not see stale 2484 global data. 2485 2486 atomicrmw acquire - singlethread - global 1. buffer/global/ds/flat_atomic 2487 - wavefront - local 2488 - generic 2489 atomicrmw acquire - workgroup - global 1. buffer/global_atomic 2490 atomicrmw acquire - workgroup - local 1. ds/flat_atomic 2491 - generic 2. waitcnt lgkmcnt(0) 2492 2493 - If OpenCL, omit 2494 waitcnt. 2495 - Must happen before 2496 any following 2497 global/generic 2498 load/load 2499 atomic/store/store 2500 atomic/atomicrmw. 2501 - Ensures any 2502 following global 2503 data read is no 2504 older than the 2505 atomicrmw value 2506 being acquired. 2507 2508 atomicrmw acquire - agent - global 1. buffer/global_atomic 2509 - system 2. s_waitcnt vmcnt(0) 2510 2511 - Must happen before 2512 following 2513 buffer_wbinvl1_vol. 2514 - Ensures the 2515 atomicrmw has 2516 completed before 2517 invalidating the 2518 cache. 2519 2520 3. buffer_wbinvl1_vol 2521 2522 - Must happen before 2523 any following 2524 global/generic 2525 load/load 2526 atomic/atomicrmw. 2527 - Ensures that 2528 following loads 2529 will not see stale 2530 global data. 2531 2532 atomicrmw acquire - agent - generic 1. flat_atomic 2533 - system 2. s_waitcnt vmcnt(0) & 2534 lgkmcnt(0) 2535 2536 - If OpenCL, omit 2537 lgkmcnt(0). 2538 - Must happen before 2539 following 2540 buffer_wbinvl1_vol. 2541 - Ensures the 2542 atomicrmw has 2543 completed before 2544 invalidating the 2545 cache. 2546 2547 3. buffer_wbinvl1_vol 2548 2549 - Must happen before 2550 any following 2551 global/generic 2552 load/load 2553 atomic/atomicrmw. 2554 - Ensures that 2555 following loads 2556 will not see stale 2557 global data. 2558 2559 fence acquire - singlethread *none* *none* 2560 - wavefront 2561 fence acquire - workgroup *none* 1. s_waitcnt lgkmcnt(0) 2562 2563 - If OpenCL and 2564 address space is 2565 not generic, omit 2566 waitcnt. However, 2567 since LLVM 2568 currently has no 2569 address space on 2570 the fence need to 2571 conservatively 2572 always generate. If 2573 fence had an 2574 address space then 2575 set to address 2576 space of OpenCL 2577 fence flag, or to 2578 generic if both 2579 local and global 2580 flags are 2581 specified. 2582 - Must happen after 2583 any preceding 2584 local/generic load 2585 atomic/atomicrmw 2586 with an equal or 2587 wider sync scope 2588 and memory ordering 2589 stronger than 2590 unordered (this is 2591 termed the 2592 fence-paired-atomic). 2593 - Must happen before 2594 any following 2595 global/generic 2596 load/load 2597 atomic/store/store 2598 atomic/atomicrmw. 2599 - Ensures any 2600 following global 2601 data read is no 2602 older than the 2603 value read by the 2604 fence-paired-atomic. 2605 2606 fence acquire - agent *none* 1. s_waitcnt vmcnt(0) & 2607 - system lgkmcnt(0) 2608 2609 - If OpenCL and 2610 address space is 2611 not generic, omit 2612 lgkmcnt(0). 2613 However, since LLVM 2614 currently has no 2615 address space on 2616 the fence need to 2617 conservatively 2618 always generate 2619 (see comment for 2620 previous fence). 2621 - Could be split into 2622 separate s_waitcnt 2623 vmcnt(0) and 2624 s_waitcnt 2625 lgkmcnt(0) to allow 2626 them to be 2627 independently moved 2628 according to the 2629 following rules. 2630 - s_waitcnt vmcnt(0) 2631 must happen after 2632 any preceding 2633 global/generic load 2634 atomic/atomicrmw 2635 with an equal or 2636 wider sync scope 2637 and memory ordering 2638 stronger than 2639 unordered (this is 2640 termed the 2641 fence-paired-atomic). 2642 - s_waitcnt lgkmcnt(0) 2643 must happen after 2644 any preceding 2645 group/generic load 2646 atomic/atomicrmw 2647 with an equal or 2648 wider sync scope 2649 and memory ordering 2650 stronger than 2651 unordered (this is 2652 termed the 2653 fence-paired-atomic). 2654 - Must happen before 2655 the following 2656 buffer_wbinvl1_vol. 2657 - Ensures that the 2658 fence-paired atomic 2659 has completed 2660 before invalidating 2661 the 2662 cache. Therefore 2663 any following 2664 locations read must 2665 be no older than 2666 the value read by 2667 the 2668 fence-paired-atomic. 2669 2670 2. buffer_wbinvl1_vol 2671 2672 - Must happen before 2673 any following global/generic 2674 load/load 2675 atomic/store/store 2676 atomic/atomicrmw. 2677 - Ensures that 2678 following loads 2679 will not see stale 2680 global data. 2681 2682 **Release Atomic** 2683 --------------------------------------------------------------------------- 2684 store atomic release - singlethread - global 1. buffer/global/ds/flat_store 2685 - wavefront - local 2686 - generic 2687 store atomic release - workgroup - global 1. s_waitcnt lgkmcnt(0) 2688 - generic 2689 - If OpenCL, omit 2690 waitcnt. 2691 - Must happen after 2692 any preceding 2693 local/generic 2694 load/store/load 2695 atomic/store 2696 atomic/atomicrmw. 2697 - Must happen before 2698 the following 2699 store. 2700 - Ensures that all 2701 memory operations 2702 to local have 2703 completed before 2704 performing the 2705 store that is being 2706 released. 2707 2708 2. buffer/global/flat_store 2709 store atomic release - workgroup - local 1. ds_store 2710 store atomic release - agent - global 1. s_waitcnt vmcnt(0) & 2711 - system - generic lgkmcnt(0) 2712 2713 - If OpenCL, omit 2714 lgkmcnt(0). 2715 - Could be split into 2716 separate s_waitcnt 2717 vmcnt(0) and 2718 s_waitcnt 2719 lgkmcnt(0) to allow 2720 them to be 2721 independently moved 2722 according to the 2723 following rules. 2724 - s_waitcnt vmcnt(0) 2725 must happen after 2726 any preceding 2727 global/generic 2728 load/store/load 2729 atomic/store 2730 atomic/atomicrmw. 2731 - s_waitcnt lgkmcnt(0) 2732 must happen after 2733 any preceding 2734 local/generic 2735 load/store/load 2736 atomic/store 2737 atomic/atomicrmw. 2738 - Must happen before 2739 the following 2740 store. 2741 - Ensures that all 2742 memory operations 2743 to global have 2744 completed before 2745 performing the 2746 store that is being 2747 released. 2748 2749 2. buffer/global/ds/flat_store 2750 atomicrmw release - singlethread - global 1. buffer/global/ds/flat_atomic 2751 - wavefront - local 2752 - generic 2753 atomicrmw release - workgroup - global 1. s_waitcnt lgkmcnt(0) 2754 - generic 2755 - If OpenCL, omit 2756 waitcnt. 2757 - Must happen after 2758 any preceding 2759 local/generic 2760 load/store/load 2761 atomic/store 2762 atomic/atomicrmw. 2763 - Must happen before 2764 the following 2765 atomicrmw. 2766 - Ensures that all 2767 memory operations 2768 to local have 2769 completed before 2770 performing the 2771 atomicrmw that is 2772 being released. 2773 2774 2. buffer/global/flat_atomic 2775 atomicrmw release - workgroup - local 1. ds_atomic 2776 atomicrmw release - agent - global 1. s_waitcnt vmcnt(0) & 2777 - system - generic lgkmcnt(0) 2778 2779 - If OpenCL, omit 2780 lgkmcnt(0). 2781 - Could be split into 2782 separate s_waitcnt 2783 vmcnt(0) and 2784 s_waitcnt 2785 lgkmcnt(0) to allow 2786 them to be 2787 independently moved 2788 according to the 2789 following rules. 2790 - s_waitcnt vmcnt(0) 2791 must happen after 2792 any preceding 2793 global/generic 2794 load/store/load 2795 atomic/store 2796 atomic/atomicrmw. 2797 - s_waitcnt lgkmcnt(0) 2798 must happen after 2799 any preceding 2800 local/generic 2801 load/store/load 2802 atomic/store 2803 atomic/atomicrmw. 2804 - Must happen before 2805 the following 2806 atomicrmw. 2807 - Ensures that all 2808 memory operations 2809 to global and local 2810 have completed 2811 before performing 2812 the atomicrmw that 2813 is being released. 2814 2815 2. buffer/global/ds/flat_atomic* 2816 fence release - singlethread *none* *none* 2817 - wavefront 2818 fence release - workgroup *none* 1. s_waitcnt lgkmcnt(0) 2819 2820 - If OpenCL and 2821 address space is 2822 not generic, omit 2823 waitcnt. However, 2824 since LLVM 2825 currently has no 2826 address space on 2827 the fence need to 2828 conservatively 2829 always generate 2830 (see comment for 2831 previous fence). 2832 - Must happen after 2833 any preceding 2834 local/generic 2835 load/load 2836 atomic/store/store 2837 atomic/atomicrmw. 2838 - Must happen before 2839 any following store 2840 atomic/atomicrmw 2841 with an equal or 2842 wider sync scope 2843 and memory ordering 2844 stronger than 2845 unordered (this is 2846 termed the 2847 fence-paired-atomic). 2848 - Ensures that all 2849 memory operations 2850 to local have 2851 completed before 2852 performing the 2853 following 2854 fence-paired-atomic. 2855 2856 fence release - agent *none* 1. s_waitcnt vmcnt(0) & 2857 - system lgkmcnt(0) 2858 2859 - If OpenCL and 2860 address space is 2861 not generic, omit 2862 lgkmcnt(0). 2863 However, since LLVM 2864 currently has no 2865 address space on 2866 the fence need to 2867 conservatively 2868 always generate 2869 (see comment for 2870 previous fence). 2871 - Could be split into 2872 separate s_waitcnt 2873 vmcnt(0) and 2874 s_waitcnt 2875 lgkmcnt(0) to allow 2876 them to be 2877 independently moved 2878 according to the 2879 following rules. 2880 - s_waitcnt vmcnt(0) 2881 must happen after 2882 any preceding 2883 global/generic 2884 load/store/load 2885 atomic/store 2886 atomic/atomicrmw. 2887 - s_waitcnt lgkmcnt(0) 2888 must happen after 2889 any preceding 2890 local/generic 2891 load/store/load 2892 atomic/store 2893 atomic/atomicrmw. 2894 - Must happen before 2895 any following store 2896 atomic/atomicrmw 2897 with an equal or 2898 wider sync scope 2899 and memory ordering 2900 stronger than 2901 unordered (this is 2902 termed the 2903 fence-paired-atomic). 2904 - Ensures that all 2905 memory operations 2906 to global have 2907 completed before 2908 performing the 2909 following 2910 fence-paired-atomic. 2911 2912 **Acquire-Release Atomic** 2913 --------------------------------------------------------------------------- 2914 atomicrmw acq_rel - singlethread - global 1. buffer/global/ds/flat_atomic 2915 - wavefront - local 2916 - generic 2917 atomicrmw acq_rel - workgroup - global 1. s_waitcnt lgkmcnt(0) 2918 2919 - If OpenCL, omit 2920 waitcnt. 2921 - Must happen after 2922 any preceding 2923 local/generic 2924 load/store/load 2925 atomic/store 2926 atomic/atomicrmw. 2927 - Must happen before 2928 the following 2929 atomicrmw. 2930 - Ensures that all 2931 memory operations 2932 to local have 2933 completed before 2934 performing the 2935 atomicrmw that is 2936 being released. 2937 2938 2. buffer/global_atomic 2939 atomicrmw acq_rel - workgroup - local 1. ds_atomic 2940 2. s_waitcnt lgkmcnt(0) 2941 2942 - If OpenCL, omit 2943 waitcnt. 2944 - Must happen before 2945 any following 2946 global/generic 2947 load/load 2948 atomic/store/store 2949 atomic/atomicrmw. 2950 - Ensures any 2951 following global 2952 data read is no 2953 older than the load 2954 atomic value being 2955 acquired. 2956 2957 atomicrmw acq_rel - workgroup - generic 1. s_waitcnt lgkmcnt(0) 2958 2959 - If OpenCL, omit 2960 waitcnt. 2961 - Must happen after 2962 any preceding 2963 local/generic 2964 load/store/load 2965 atomic/store 2966 atomic/atomicrmw. 2967 - Must happen before 2968 the following 2969 atomicrmw. 2970 - Ensures that all 2971 memory operations 2972 to local have 2973 completed before 2974 performing the 2975 atomicrmw that is 2976 being released. 2977 2978 2. flat_atomic 2979 3. s_waitcnt lgkmcnt(0) 2980 2981 - If OpenCL, omit 2982 waitcnt. 2983 - Must happen before 2984 any following 2985 global/generic 2986 load/load 2987 atomic/store/store 2988 atomic/atomicrmw. 2989 - Ensures any 2990 following global 2991 data read is no 2992 older than the load 2993 atomic value being 2994 acquired. 2995 atomicrmw acq_rel - agent - global 1. s_waitcnt vmcnt(0) & 2996 - system lgkmcnt(0) 2997 2998 - If OpenCL, omit 2999 lgkmcnt(0). 3000 - Could be split into 3001 separate s_waitcnt 3002 vmcnt(0) and 3003 s_waitcnt 3004 lgkmcnt(0) to allow 3005 them to be 3006 independently moved 3007 according to the 3008 following rules. 3009 - s_waitcnt vmcnt(0) 3010 must happen after 3011 any preceding 3012 global/generic 3013 load/store/load 3014 atomic/store 3015 atomic/atomicrmw. 3016 - s_waitcnt lgkmcnt(0) 3017 must happen after 3018 any preceding 3019 local/generic 3020 load/store/load 3021 atomic/store 3022 atomic/atomicrmw. 3023 - Must happen before 3024 the following 3025 atomicrmw. 3026 - Ensures that all 3027 memory operations 3028 to global have 3029 completed before 3030 performing the 3031 atomicrmw that is 3032 being released. 3033 3034 2. buffer/global_atomic 3035 3. s_waitcnt vmcnt(0) 3036 3037 - Must happen before 3038 following 3039 buffer_wbinvl1_vol. 3040 - Ensures the 3041 atomicrmw has 3042 completed before 3043 invalidating the 3044 cache. 3045 3046 4. buffer_wbinvl1_vol 3047 3048 - Must happen before 3049 any following 3050 global/generic 3051 load/load 3052 atomic/atomicrmw. 3053 - Ensures that 3054 following loads 3055 will not see stale 3056 global data. 3057 3058 atomicrmw acq_rel - agent - generic 1. s_waitcnt vmcnt(0) & 3059 - system lgkmcnt(0) 3060 3061 - If OpenCL, omit 3062 lgkmcnt(0). 3063 - Could be split into 3064 separate s_waitcnt 3065 vmcnt(0) and 3066 s_waitcnt 3067 lgkmcnt(0) to allow 3068 them to be 3069 independently moved 3070 according to the 3071 following rules. 3072 - s_waitcnt vmcnt(0) 3073 must happen after 3074 any preceding 3075 global/generic 3076 load/store/load 3077 atomic/store 3078 atomic/atomicrmw. 3079 - s_waitcnt lgkmcnt(0) 3080 must happen after 3081 any preceding 3082 local/generic 3083 load/store/load 3084 atomic/store 3085 atomic/atomicrmw. 3086 - Must happen before 3087 the following 3088 atomicrmw. 3089 - Ensures that all 3090 memory operations 3091 to global have 3092 completed before 3093 performing the 3094 atomicrmw that is 3095 being released. 3096 3097 2. flat_atomic 3098 3. s_waitcnt vmcnt(0) & 3099 lgkmcnt(0) 3100 3101 - If OpenCL, omit 3102 lgkmcnt(0). 3103 - Must happen before 3104 following 3105 buffer_wbinvl1_vol. 3106 - Ensures the 3107 atomicrmw has 3108 completed before 3109 invalidating the 3110 cache. 3111 3112 4. buffer_wbinvl1_vol 3113 3114 - Must happen before 3115 any following 3116 global/generic 3117 load/load 3118 atomic/atomicrmw. 3119 - Ensures that 3120 following loads 3121 will not see stale 3122 global data. 3123 3124 fence acq_rel - singlethread *none* *none* 3125 - wavefront 3126 fence acq_rel - workgroup *none* 1. s_waitcnt lgkmcnt(0) 3127 3128 - If OpenCL and 3129 address space is 3130 not generic, omit 3131 waitcnt. However, 3132 since LLVM 3133 currently has no 3134 address space on 3135 the fence need to 3136 conservatively 3137 always generate 3138 (see comment for 3139 previous fence). 3140 - Must happen after 3141 any preceding 3142 local/generic 3143 load/load 3144 atomic/store/store 3145 atomic/atomicrmw. 3146 - Must happen before 3147 any following 3148 global/generic 3149 load/load 3150 atomic/store/store 3151 atomic/atomicrmw. 3152 - Ensures that all 3153 memory operations 3154 to local have 3155 completed before 3156 performing any 3157 following global 3158 memory operations. 3159 - Ensures that the 3160 preceding 3161 local/generic load 3162 atomic/atomicrmw 3163 with an equal or 3164 wider sync scope 3165 and memory ordering 3166 stronger than 3167 unordered (this is 3168 termed the 3169 fence-paired-atomic) 3170 has completed 3171 before following 3172 global memory 3173 operations. This 3174 satisfies the 3175 requirements of 3176 acquire. 3177 - Ensures that all 3178 previous memory 3179 operations have 3180 completed before a 3181 following 3182 local/generic store 3183 atomic/atomicrmw 3184 with an equal or 3185 wider sync scope 3186 and memory ordering 3187 stronger than 3188 unordered (this is 3189 termed the 3190 fence-paired-atomic). 3191 This satisfies the 3192 requirements of 3193 release. 3194 3195 fence acq_rel - agent *none* 1. s_waitcnt vmcnt(0) & 3196 - system lgkmcnt(0) 3197 3198 - If OpenCL and 3199 address space is 3200 not generic, omit 3201 lgkmcnt(0). 3202 However, since LLVM 3203 currently has no 3204 address space on 3205 the fence need to 3206 conservatively 3207 always generate 3208 (see comment for 3209 previous fence). 3210 - Could be split into 3211 separate s_waitcnt 3212 vmcnt(0) and 3213 s_waitcnt 3214 lgkmcnt(0) to allow 3215 them to be 3216 independently moved 3217 according to the 3218 following rules. 3219 - s_waitcnt vmcnt(0) 3220 must happen after 3221 any preceding 3222 global/generic 3223 load/store/load 3224 atomic/store 3225 atomic/atomicrmw. 3226 - s_waitcnt lgkmcnt(0) 3227 must happen after 3228 any preceding 3229 local/generic 3230 load/store/load 3231 atomic/store 3232 atomic/atomicrmw. 3233 - Must happen before 3234 the following 3235 buffer_wbinvl1_vol. 3236 - Ensures that the 3237 preceding 3238 global/local/generic 3239 load 3240 atomic/atomicrmw 3241 with an equal or 3242 wider sync scope 3243 and memory ordering 3244 stronger than 3245 unordered (this is 3246 termed the 3247 fence-paired-atomic) 3248 has completed 3249 before invalidating 3250 the cache. This 3251 satisfies the 3252 requirements of 3253 acquire. 3254 - Ensures that all 3255 previous memory 3256 operations have 3257 completed before a 3258 following 3259 global/local/generic 3260 store 3261 atomic/atomicrmw 3262 with an equal or 3263 wider sync scope 3264 and memory ordering 3265 stronger than 3266 unordered (this is 3267 termed the 3268 fence-paired-atomic). 3269 This satisfies the 3270 requirements of 3271 release. 3272 3273 2. buffer_wbinvl1_vol 3274 3275 - Must happen before 3276 any following 3277 global/generic 3278 load/load 3279 atomic/store/store 3280 atomic/atomicrmw. 3281 - Ensures that 3282 following loads 3283 will not see stale 3284 global data. This 3285 satisfies the 3286 requirements of 3287 acquire. 3288 3289 **Sequential Consistent Atomic** 3290 --------------------------------------------------------------------------- 3291 load atomic seq_cst - singlethread - global *Same as corresponding 3292 - wavefront - local load atomic acquire*. 3293 - workgroup - generic 3294 load atomic seq_cst - agent - global 1. s_waitcnt vmcnt(0) 3295 - system - local 3296 - generic - Must happen after 3297 preceding 3298 global/generic load 3299 atomic/store 3300 atomic/atomicrmw 3301 with memory 3302 ordering of seq_cst 3303 and with equal or 3304 wider sync scope. 3305 (Note that seq_cst 3306 fences have their 3307 own s_waitcnt 3308 vmcnt(0) and so do 3309 not need to be 3310 considered.) 3311 - Ensures any 3312 preceding 3313 sequential 3314 consistent global 3315 memory instructions 3316 have completed 3317 before executing 3318 this sequentially 3319 consistent 3320 instruction. This 3321 prevents reordering 3322 a seq_cst store 3323 followed by a 3324 seq_cst load (Note 3325 that seq_cst is 3326 stronger than 3327 acquire/release as 3328 the reordering of 3329 load acquire 3330 followed by a store 3331 release is 3332 prevented by the 3333 waitcnt vmcnt(0) of 3334 the release, but 3335 there is nothing 3336 preventing a store 3337 release followed by 3338 load acquire from 3339 competing out of 3340 order.) 3341 3342 2. *Following 3343 instructions same as 3344 corresponding load 3345 atomic acquire*. 3346 3347 store atomic seq_cst - singlethread - global *Same as corresponding 3348 - wavefront - local store atomic release*. 3349 - workgroup - generic 3350 store atomic seq_cst - agent - global *Same as corresponding 3351 - system - generic store atomic release*. 3352 atomicrmw seq_cst - singlethread - global *Same as corresponding 3353 - wavefront - local atomicrmw acq_rel*. 3354 - workgroup - generic 3355 atomicrmw seq_cst - agent - global *Same as corresponding 3356 - system - generic atomicrmw acq_rel*. 3357 fence seq_cst - singlethread *none* *Same as corresponding 3358 - wavefront fence acq_rel*. 3359 - workgroup 3360 - agent 3361 - system 3362 ============ ============ ============== ========== ======================= 3363 3364The memory order also adds the single thread optimization constrains defined in 3365table 3366:ref:`amdgpu-amdhsa-memory-model-single-thread-optimization-constraints-gfx6-gfx9-table`. 3367 3368 .. table:: AMDHSA Memory Model Single Thread Optimization Constraints GFX6-GFX9 3369 :name: amdgpu-amdhsa-memory-model-single-thread-optimization-constraints-gfx6-gfx9-table 3370 3371 ============ ============================================================== 3372 LLVM Memory Optimization Constraints 3373 Ordering 3374 ============ ============================================================== 3375 unordered *none* 3376 monotonic *none* 3377 acquire - If a load atomic/atomicrmw then no following load/load 3378 atomic/store/ store atomic/atomicrmw/fence instruction can 3379 be moved before the acquire. 3380 - If a fence then same as load atomic, plus no preceding 3381 associated fence-paired-atomic can be moved after the fence. 3382 release - If a store atomic/atomicrmw then no preceding load/load 3383 atomic/store/ store atomic/atomicrmw/fence instruction can 3384 be moved after the release. 3385 - If a fence then same as store atomic, plus no following 3386 associated fence-paired-atomic can be moved before the 3387 fence. 3388 acq_rel Same constraints as both acquire and release. 3389 seq_cst - If a load atomic then same constraints as acquire, plus no 3390 preceding sequentially consistent load atomic/store 3391 atomic/atomicrmw/fence instruction can be moved after the 3392 seq_cst. 3393 - If a store atomic then the same constraints as release, plus 3394 no following sequentially consistent load atomic/store 3395 atomic/atomicrmw/fence instruction can be moved before the 3396 seq_cst. 3397 - If an atomicrmw/fence then same constraints as acq_rel. 3398 ============ ============================================================== 3399 3400Trap Handler ABI 3401~~~~~~~~~~~~~~~~ 3402 3403For code objects generated by AMDGPU backend for HSA [HSA]_ compatible runtimes 3404(such as ROCm [AMD-ROCm]_), the runtime installs a trap handler that supports 3405the ``s_trap`` instruction with the following usage: 3406 3407 .. table:: AMDGPU Trap Handler for AMDHSA OS 3408 :name: amdgpu-trap-handler-for-amdhsa-os-table 3409 3410 =================== =============== =============== ======================= 3411 Usage Code Sequence Trap Handler Description 3412 Inputs 3413 =================== =============== =============== ======================= 3414 reserved ``s_trap 0x00`` Reserved by hardware. 3415 ``debugtrap(arg)`` ``s_trap 0x01`` ``SGPR0-1``: Reserved for HSA 3416 ``queue_ptr`` ``debugtrap`` 3417 ``VGPR0``: intrinsic (not 3418 ``arg`` implemented). 3419 ``llvm.trap`` ``s_trap 0x02`` ``SGPR0-1``: Causes dispatch to be 3420 ``queue_ptr`` terminated and its 3421 associated queue put 3422 into the error state. 3423 ``llvm.debugtrap`` ``s_trap 0x03`` ``SGPR0-1``: If debugger not 3424 ``queue_ptr`` installed handled 3425 same as ``llvm.trap``. 3426 debugger breakpoint ``s_trap 0x07`` Reserved for debugger 3427 breakpoints. 3428 debugger ``s_trap 0x08`` Reserved for debugger. 3429 debugger ``s_trap 0xfe`` Reserved for debugger. 3430 debugger ``s_trap 0xff`` Reserved for debugger. 3431 =================== =============== =============== ======================= 3432 3433Unspecified OS 3434-------------- 3435 3436This section provides code conventions used when the target triple OS is 3437empty (see :ref:`amdgpu-target-triples`). 3438 3439Trap Handler ABI 3440~~~~~~~~~~~~~~~~ 3441 3442For code objects generated by AMDGPU backend for non-amdhsa OS, the runtime does 3443not install a trap handler. The ``llvm.trap`` and ``llvm.debugtrap`` 3444instructions are handled as follows: 3445 3446 .. table:: AMDGPU Trap Handler for Non-AMDHSA OS 3447 :name: amdgpu-trap-handler-for-non-amdhsa-os-table 3448 3449 =============== =============== =========================================== 3450 Usage Code Sequence Description 3451 =============== =============== =========================================== 3452 llvm.trap s_endpgm Causes wavefront to be terminated. 3453 llvm.debugtrap *none* Compiler warning given that there is no 3454 trap handler installed. 3455 =============== =============== =========================================== 3456 3457Source Languages 3458================ 3459 3460.. _amdgpu-opencl: 3461 3462OpenCL 3463------ 3464 3465When generating code for the OpenCL language the target triple environment 3466should be ``opencl`` or ``amdgizcl`` (see :ref:`amdgpu-target-triples`). 3467 3468When the language is OpenCL the following differences occur: 3469 34701. The OpenCL memory model is used (see :ref:`amdgpu-amdhsa-memory-model`). 34712. The AMDGPU backend adds additional arguments to the kernel. 34723. Additional metadata is generated 3473 (:ref:`amdgpu-amdhsa-hsa-code-object-metadata`). 3474 3475.. TODO 3476 Specify what affect this has. Hidden arguments added. Additional metadata 3477 generated. 3478 3479.. _amdgpu-hcc: 3480 3481HCC 3482--- 3483 3484When generating code for the OpenCL language the target triple environment 3485should be ``hcc`` (see :ref:`amdgpu-target-triples`). 3486 3487When the language is OpenCL the following differences occur: 3488 34891. The HSA memory model is used (see :ref:`amdgpu-amdhsa-memory-model`). 3490 3491.. TODO 3492 Specify what affect this has. 3493 3494Assembler 3495--------- 3496 3497AMDGPU backend has LLVM-MC based assembler which is currently in development. 3498It supports AMDGCN GFX6-GFX8. 3499 3500This section describes general syntax for instructions and operands. For more 3501information about instructions, their semantics and supported combinations of 3502operands, refer to one of instruction set architecture manuals 3503[AMD-Souther-Islands]_, [AMD-Sea-Islands]_, [AMD-Volcanic-Islands]_ and 3504[AMD-Vega]_. 3505 3506An instruction has the following syntax (register operands are normally 3507comma-separated while extra operands are space-separated): 3508 3509*<opcode> <register_operand0>, ... <extra_operand0> ...* 3510 3511Operands 3512~~~~~~~~ 3513 3514The following syntax for register operands is supported: 3515 3516* SGPR registers: s0, ... or s[0], ... 3517* VGPR registers: v0, ... or v[0], ... 3518* TTMP registers: ttmp0, ... or ttmp[0], ... 3519* Special registers: exec (exec_lo, exec_hi), vcc (vcc_lo, vcc_hi), flat_scratch (flat_scratch_lo, flat_scratch_hi) 3520* Special trap registers: tba (tba_lo, tba_hi), tma (tma_lo, tma_hi) 3521* Register pairs, quads, etc: s[2:3], v[10:11], ttmp[5:6], s[4:7], v[12:15], ttmp[4:7], s[8:15], ... 3522* Register lists: [s0, s1], [ttmp0, ttmp1, ttmp2, ttmp3] 3523* Register index expressions: v[2*2], s[1-1:2-1] 3524* 'off' indicates that an operand is not enabled 3525 3526The following extra operands are supported: 3527 3528* offset, offset0, offset1 3529* idxen, offen bits 3530* glc, slc, tfe bits 3531* waitcnt: integer or combination of counter values 3532* VOP3 modifiers: 3533 3534 - abs (\| \|), neg (\-) 3535 3536* DPP modifiers: 3537 3538 - row_shl, row_shr, row_ror, row_rol 3539 - row_mirror, row_half_mirror, row_bcast 3540 - wave_shl, wave_shr, wave_ror, wave_rol, quad_perm 3541 - row_mask, bank_mask, bound_ctrl 3542 3543* SDWA modifiers: 3544 3545 - dst_sel, src0_sel, src1_sel (BYTE_N, WORD_M, DWORD) 3546 - dst_unused (UNUSED_PAD, UNUSED_SEXT, UNUSED_PRESERVE) 3547 - abs, neg, sext 3548 3549Instruction Examples 3550~~~~~~~~~~~~~~~~~~~~ 3551 3552DS 3553~~ 3554 3555.. code-block:: nasm 3556 3557 ds_add_u32 v2, v4 offset:16 3558 ds_write_src2_b64 v2 offset0:4 offset1:8 3559 ds_cmpst_f32 v2, v4, v6 3560 ds_min_rtn_f64 v[8:9], v2, v[4:5] 3561 3562 3563For full list of supported instructions, refer to "LDS/GDS instructions" in ISA Manual. 3564 3565FLAT 3566++++ 3567 3568.. code-block:: nasm 3569 3570 flat_load_dword v1, v[3:4] 3571 flat_store_dwordx3 v[3:4], v[5:7] 3572 flat_atomic_swap v1, v[3:4], v5 glc 3573 flat_atomic_cmpswap v1, v[3:4], v[5:6] glc slc 3574 flat_atomic_fmax_x2 v[1:2], v[3:4], v[5:6] glc 3575 3576For full list of supported instructions, refer to "FLAT instructions" in ISA Manual. 3577 3578MUBUF 3579+++++ 3580 3581.. code-block:: nasm 3582 3583 buffer_load_dword v1, off, s[4:7], s1 3584 buffer_store_dwordx4 v[1:4], v2, ttmp[4:7], s1 offen offset:4 glc tfe 3585 buffer_store_format_xy v[1:2], off, s[4:7], s1 3586 buffer_wbinvl1 3587 buffer_atomic_inc v1, v2, s[8:11], s4 idxen offset:4 slc 3588 3589For full list of supported instructions, refer to "MUBUF Instructions" in ISA Manual. 3590 3591SMRD/SMEM 3592+++++++++ 3593 3594.. code-block:: nasm 3595 3596 s_load_dword s1, s[2:3], 0xfc 3597 s_load_dwordx8 s[8:15], s[2:3], s4 3598 s_load_dwordx16 s[88:103], s[2:3], s4 3599 s_dcache_inv_vol 3600 s_memtime s[4:5] 3601 3602For full list of supported instructions, refer to "Scalar Memory Operations" in ISA Manual. 3603 3604SOP1 3605++++ 3606 3607.. code-block:: nasm 3608 3609 s_mov_b32 s1, s2 3610 s_mov_b64 s[0:1], 0x80000000 3611 s_cmov_b32 s1, 200 3612 s_wqm_b64 s[2:3], s[4:5] 3613 s_bcnt0_i32_b64 s1, s[2:3] 3614 s_swappc_b64 s[2:3], s[4:5] 3615 s_cbranch_join s[4:5] 3616 3617For full list of supported instructions, refer to "SOP1 Instructions" in ISA Manual. 3618 3619SOP2 3620++++ 3621 3622.. code-block:: nasm 3623 3624 s_add_u32 s1, s2, s3 3625 s_and_b64 s[2:3], s[4:5], s[6:7] 3626 s_cselect_b32 s1, s2, s3 3627 s_andn2_b32 s2, s4, s6 3628 s_lshr_b64 s[2:3], s[4:5], s6 3629 s_ashr_i32 s2, s4, s6 3630 s_bfm_b64 s[2:3], s4, s6 3631 s_bfe_i64 s[2:3], s[4:5], s6 3632 s_cbranch_g_fork s[4:5], s[6:7] 3633 3634For full list of supported instructions, refer to "SOP2 Instructions" in ISA Manual. 3635 3636SOPC 3637++++ 3638 3639.. code-block:: nasm 3640 3641 s_cmp_eq_i32 s1, s2 3642 s_bitcmp1_b32 s1, s2 3643 s_bitcmp0_b64 s[2:3], s4 3644 s_setvskip s3, s5 3645 3646For full list of supported instructions, refer to "SOPC Instructions" in ISA Manual. 3647 3648SOPP 3649++++ 3650 3651.. code-block:: nasm 3652 3653 s_barrier 3654 s_nop 2 3655 s_endpgm 3656 s_waitcnt 0 ; Wait for all counters to be 0 3657 s_waitcnt vmcnt(0) & expcnt(0) & lgkmcnt(0) ; Equivalent to above 3658 s_waitcnt vmcnt(1) ; Wait for vmcnt counter to be 1. 3659 s_sethalt 9 3660 s_sleep 10 3661 s_sendmsg 0x1 3662 s_sendmsg sendmsg(MSG_INTERRUPT) 3663 s_trap 1 3664 3665For full list of supported instructions, refer to "SOPP Instructions" in ISA Manual. 3666 3667Unless otherwise mentioned, little verification is performed on the operands 3668of SOPP Instructions, so it is up to the programmer to be familiar with the 3669range or acceptable values. 3670 3671VALU 3672++++ 3673 3674For vector ALU instruction opcodes (VOP1, VOP2, VOP3, VOPC, VOP_DPP, VOP_SDWA), 3675the assembler will automatically use optimal encoding based on its operands. 3676To force specific encoding, one can add a suffix to the opcode of the instruction: 3677 3678* _e32 for 32-bit VOP1/VOP2/VOPC 3679* _e64 for 64-bit VOP3 3680* _dpp for VOP_DPP 3681* _sdwa for VOP_SDWA 3682 3683VOP1/VOP2/VOP3/VOPC examples: 3684 3685.. code-block:: nasm 3686 3687 v_mov_b32 v1, v2 3688 v_mov_b32_e32 v1, v2 3689 v_nop 3690 v_cvt_f64_i32_e32 v[1:2], v2 3691 v_floor_f32_e32 v1, v2 3692 v_bfrev_b32_e32 v1, v2 3693 v_add_f32_e32 v1, v2, v3 3694 v_mul_i32_i24_e64 v1, v2, 3 3695 v_mul_i32_i24_e32 v1, -3, v3 3696 v_mul_i32_i24_e32 v1, -100, v3 3697 v_addc_u32 v1, s[0:1], v2, v3, s[2:3] 3698 v_max_f16_e32 v1, v2, v3 3699 3700VOP_DPP examples: 3701 3702.. code-block:: nasm 3703 3704 v_mov_b32 v0, v0 quad_perm:[0,2,1,1] 3705 v_sin_f32 v0, v0 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0 3706 v_mov_b32 v0, v0 wave_shl:1 3707 v_mov_b32 v0, v0 row_mirror 3708 v_mov_b32 v0, v0 row_bcast:31 3709 v_mov_b32 v0, v0 quad_perm:[1,3,0,1] row_mask:0xa bank_mask:0x1 bound_ctrl:0 3710 v_add_f32 v0, v0, |v0| row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0 3711 v_max_f16 v1, v2, v3 row_shl:1 row_mask:0xa bank_mask:0x1 bound_ctrl:0 3712 3713VOP_SDWA examples: 3714 3715.. code-block:: nasm 3716 3717 v_mov_b32 v1, v2 dst_sel:BYTE_0 dst_unused:UNUSED_PRESERVE src0_sel:DWORD 3718 v_min_u32 v200, v200, v1 dst_sel:WORD_1 dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:DWORD 3719 v_sin_f32 v0, v0 dst_unused:UNUSED_PAD src0_sel:WORD_1 3720 v_fract_f32 v0, |v0| dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:WORD_1 3721 v_cmpx_le_u32 vcc, v1, v2 src0_sel:BYTE_2 src1_sel:WORD_0 3722 3723For full list of supported instructions, refer to "Vector ALU instructions". 3724 3725HSA Code Object Directives 3726~~~~~~~~~~~~~~~~~~~~~~~~~~ 3727 3728AMDGPU ABI defines auxiliary data in output code object. In assembly source, 3729one can specify them with assembler directives. 3730 3731.hsa_code_object_version major, minor 3732+++++++++++++++++++++++++++++++++++++ 3733 3734*major* and *minor* are integers that specify the version of the HSA code 3735object that will be generated by the assembler. 3736 3737.hsa_code_object_isa [major, minor, stepping, vendor, arch] 3738+++++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 3739 3740 3741*major*, *minor*, and *stepping* are all integers that describe the instruction 3742set architecture (ISA) version of the assembly program. 3743 3744*vendor* and *arch* are quoted strings. *vendor* should always be equal to 3745"AMD" and *arch* should always be equal to "AMDGPU". 3746 3747By default, the assembler will derive the ISA version, *vendor*, and *arch* 3748from the value of the -mcpu option that is passed to the assembler. 3749 3750.amdgpu_hsa_kernel (name) 3751+++++++++++++++++++++++++ 3752 3753This directives specifies that the symbol with given name is a kernel entry point 3754(label) and the object should contain corresponding symbol of type STT_AMDGPU_HSA_KERNEL. 3755 3756.amd_kernel_code_t 3757++++++++++++++++++ 3758 3759This directive marks the beginning of a list of key / value pairs that are used 3760to specify the amd_kernel_code_t object that will be emitted by the assembler. 3761The list must be terminated by the *.end_amd_kernel_code_t* directive. For 3762any amd_kernel_code_t values that are unspecified a default value will be 3763used. The default value for all keys is 0, with the following exceptions: 3764 3765- *kernel_code_version_major* defaults to 1. 3766- *machine_kind* defaults to 1. 3767- *machine_version_major*, *machine_version_minor*, and 3768 *machine_version_stepping* are derived from the value of the -mcpu option 3769 that is passed to the assembler. 3770- *kernel_code_entry_byte_offset* defaults to 256. 3771- *wavefront_size* defaults to 6. 3772- *kernarg_segment_alignment*, *group_segment_alignment*, and 3773 *private_segment_alignment* default to 4. Note that alignments are specified 3774 as a power of two, so a value of **n** means an alignment of 2^ **n**. 3775 3776The *.amd_kernel_code_t* directive must be placed immediately after the 3777function label and before any instructions. 3778 3779For a full list of amd_kernel_code_t keys, refer to AMDGPU ABI document, 3780comments in lib/Target/AMDGPU/AmdKernelCodeT.h and test/CodeGen/AMDGPU/hsa.s. 3781 3782Here is an example of a minimal amd_kernel_code_t specification: 3783 3784.. code-block:: none 3785 3786 .hsa_code_object_version 1,0 3787 .hsa_code_object_isa 3788 3789 .hsatext 3790 .globl hello_world 3791 .p2align 8 3792 .amdgpu_hsa_kernel hello_world 3793 3794 hello_world: 3795 3796 .amd_kernel_code_t 3797 enable_sgpr_kernarg_segment_ptr = 1 3798 is_ptr64 = 1 3799 compute_pgm_rsrc1_vgprs = 0 3800 compute_pgm_rsrc1_sgprs = 0 3801 compute_pgm_rsrc2_user_sgpr = 2 3802 kernarg_segment_byte_size = 8 3803 wavefront_sgpr_count = 2 3804 workitem_vgpr_count = 3 3805 .end_amd_kernel_code_t 3806 3807 s_load_dwordx2 s[0:1], s[0:1] 0x0 3808 v_mov_b32 v0, 3.14159 3809 s_waitcnt lgkmcnt(0) 3810 v_mov_b32 v1, s0 3811 v_mov_b32 v2, s1 3812 flat_store_dword v[1:2], v0 3813 s_endpgm 3814 .Lfunc_end0: 3815 .size hello_world, .Lfunc_end0-hello_world 3816 3817Additional Documentation 3818======================== 3819 3820.. [AMD-R6xx] `AMD R6xx shader ISA <http://developer.amd.com/wordpress/media/2012/10/R600_Instruction_Set_Architecture.pdf>`__ 3821.. [AMD-R7xx] `AMD R7xx shader ISA <http://developer.amd.com/wordpress/media/2012/10/R700-Family_Instruction_Set_Architecture.pdf>`__ 3822.. [AMD-Evergreen] `AMD Evergreen shader ISA <http://developer.amd.com/wordpress/media/2012/10/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf>`__ 3823.. [AMD-Cayman-Trinity] `AMD Cayman/Trinity shader ISA <http://developer.amd.com/wordpress/media/2012/10/AMD_HD_6900_Series_Instruction_Set_Architecture.pdf>`__ 3824.. [AMD-Souther-Islands] `AMD Southern Islands Series ISA <http://developer.amd.com/wordpress/media/2012/12/AMD_Southern_Islands_Instruction_Set_Architecture.pdf>`__ 3825.. [AMD-Sea-Islands] `AMD Sea Islands Series ISA <http://developer.amd.com/wordpress/media/2013/07/AMD_Sea_Islands_Instruction_Set_Architecture.pdf>`_ 3826.. [AMD-Volcanic-Islands] `AMD GCN3 Instruction Set Architecture <http://amd-dev.wpengine.netdna-cdn.com/wordpress/media/2013/12/AMD_GCN3_Instruction_Set_Architecture_rev1.1.pdf>`__ 3827.. [AMD-Vega] `AMD "Vega" Instruction Set Architecture <http://developer.amd.com/wordpress/media/2013/12/Vega_Shader_ISA_28July2017.pdf>`__ 3828.. [AMD-OpenCL_Programming-Guide] `AMD Accelerated Parallel Processing OpenCL Programming Guide <http://developer.amd.com/download/AMD_Accelerated_Parallel_Processing_OpenCL_Programming_Guide.pdf>`_ 3829.. [AMD-APP-SDK] `AMD Accelerated Parallel Processing APP SDK Documentation <http://developer.amd.com/tools/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/documentation/>`__ 3830.. [AMD-ROCm] `ROCm: Open Platform for Development, Discovery and Education Around GPU Computing <http://gpuopen.com/compute-product/rocm/>`__ 3831.. [AMD-ROCm-github] `ROCm github <http://github.com/RadeonOpenCompute>`__ 3832.. [HSA] `Heterogeneous System Architecture (HSA) Foundation <http://www.hsafoundation.com/>`__ 3833.. [ELF] `Executable and Linkable Format (ELF) <http://www.sco.com/developers/gabi/>`__ 3834.. [DWARF] `DWARF Debugging Information Format <http://dwarfstd.org/>`__ 3835.. [YAML] `YAML Ain’t Markup Language (YAML™) Version 1.2 <http://www.yaml.org/spec/1.2/spec.html>`__ 3836.. [OpenCL] `The OpenCL Specification Version 2.0 <http://www.khronos.org/registry/cl/specs/opencl-2.0.pdf>`__ 3837.. [HRF] `Heterogeneous-race-free Memory Models <http://benedictgaster.org/wp-content/uploads/2014/01/asplos269-FINAL.pdf>`__ 3838.. [AMD-AMDGPU-Compute-Application-Binary-Interface] `AMDGPU Compute Application Binary Interface <https://github.com/RadeonOpenCompute/ROCm-ComputeABI-Doc/blob/master/AMDGPU-ABI.md>`__ 3839