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