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