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