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