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