838 lines · plain
1.. raw:: html2 3 <style type="text/css">4 .none { background-color: #FFCCCC }5 .part { background-color: #FFFF99 }6 .good { background-color: #CCFF99 }7 </style>8 9.. role:: none10.. role:: part11.. role:: good12 13.. contents::14 :local:15 16=============17HIP Support18=============19 20HIP (Heterogeneous-Compute Interface for Portability) `<https://github.com/ROCm/HIP>`_ is21a C++ Runtime API and Kernel Language. It enables developers to create portable applications for22offloading computation to different hardware platforms from a single source code.23 24AMD GPU Support25===============26 27Clang provides HIP support on AMD GPUs via the ROCm platform `<https://rocm.docs.amd.com/en/latest/#>`_.28The ROCm runtime forms the base for HIP host APIs, while HIP device APIs are realized through HIP header29files and the ROCm device library. The Clang driver uses the HIPAMD toolchain to compile HIP device code30to AMDGPU ISA via the AMDGPU backend, or SPIR-V via the workflow outlined below.31The compiled code is then bundled and embedded in the host executables.32 33Intel GPU Support34=================35 36Clang provides partial HIP support on Intel GPUs using the CHIP-Star project `<https://github.com/CHIP-SPV/chipStar>`_.37CHIP-Star implements the HIP runtime over oneAPI Level Zero or OpenCL runtime. The Clang driver uses the HIPSPV38toolchain to compile HIP device code into LLVM IR, which is subsequently translated to SPIR-V via the SPIR-V39backend or the out-of-tree LLVM-SPIRV translator. The SPIR-V is then bundled and embedded into the host executables.40 41.. note::42 While Clang does not directly provide HIP support for NVIDIA GPUs and CPUs, these platforms are supported via other means:43 44 - NVIDIA GPUs: HIP support is offered through the HIP project `<https://github.com/ROCm/HIP>`_, which provides a header-only library for translating HIP runtime APIs into CUDA runtime APIs. The code is subsequently compiled using NVIDIA's `nvcc`.45 46 - CPUs: HIP support is available through the HIP-CPU runtime library `<https://github.com/ROCm/HIP-CPU>`_. This header-only library enables CPUs to execute unmodified HIP code.47 48 49Example Usage50=============51 52To compile a HIP program, use the following command:53 54.. code-block:: shell55 56 clang++ -c --offload-arch=gfx906 -xhip sample.cpp -o sample.o57 58The ``-xhip`` option indicates that the source is a HIP program. If the file has a ``.hip`` extension,59Clang will automatically recognize it as a HIP program:60 61.. code-block:: shell62 63 clang++ -c --offload-arch=gfx906 sample.hip -o sample.o64 65To link a HIP program, use this command:66 67.. code-block:: shell68 69 clang++ --hip-link --offload-arch=gfx906 sample.o -o sample70 71In the above command, the ``--hip-link`` flag instructs Clang to link the HIP runtime library. However,72the use of this flag is unnecessary if a HIP input file is already present in your program.73 74For convenience, Clang also supports compiling and linking in a single step:75 76.. code-block:: shell77 78 clang++ --offload-arch=gfx906 -xhip sample.cpp -o sample79 80In the above commands, ``gfx906`` is the GPU architecture that the code is being compiled for. The supported GPU81architectures can be found in the `AMDGPU Processor Table <https://llvm.org/docs/AMDGPUUsage.html#processors>`_.82Alternatively, you can use the ``amdgpu-arch`` tool that comes with Clang to list the GPU architecture on your system:83 84.. code-block:: shell85 86 amdgpu-arch87 88You can use ``--offload-arch=native`` to automatically detect the GPU architectures on your system:89 90.. code-block:: shell91 92 clang++ --offload-arch=native -xhip sample.cpp -o sample93 94 95Path Setting for Dependencies96=============================97 98Compiling a HIP program depends on the HIP runtime and device library. The paths to the HIP runtime and device libraries99can be specified either using compiler options or environment variables. The paths can also be set through the ROCm path100if they follow the ROCm installation directory structure.101 102Order of Precedence for HIP Path103--------------------------------104 1051. ``--hip-path`` compiler option1062. ``HIP_PATH`` environment variable *(use with caution)*1073. ``--rocm-path`` compiler option1084. ``ROCM_PATH`` environment variable *(use with caution)*1095. Default automatic detection (relative to Clang or at the default ROCm installation location)110 111Order of Precedence for Device Library Path112-------------------------------------------113 1141. ``--hip-device-lib-path`` compiler option1152. ``HIP_DEVICE_LIB_PATH`` environment variable *(use with caution)*1163. ``--rocm-path`` compiler option1174. ``ROCM_PATH`` environment variable *(use with caution)*1185. Default automatic detection (relative to Clang or at the default ROCm installation location)119 120.. list-table::121 :header-rows: 1122 123 * - Compiler Option124 - Environment Variable125 - Description126 - Default Value127 * - ``--rocm-path=<path>``128 - ``ROCM_PATH``129 - Specifies the ROCm installation path.130 - Automatic detection131 * - ``--hip-path=<path>``132 - ``HIP_PATH``133 - Specifies the HIP runtime installation path.134 - Determined by ROCm directory structure135 * - ``--hip-device-lib-path=<path>``136 - ``HIP_DEVICE_LIB_PATH``137 - Specifies the HIP device library installation path.138 - Determined by ROCm directory structure139 140.. note::141 142 We recommend using the compiler options as the primary method for specifying these paths. While the environment variables ``ROCM_PATH``, ``HIP_PATH``, and ``HIP_DEVICE_LIB_PATH`` are supported, their use can lead to implicit dependencies that might cause issues in the long run. Use them with caution.143 144 145Predefined Macros146=================147 148.. list-table::149 :header-rows: 1150 151 * - Macro152 - Description153 * - ``__CLANG_RDC__``154 - Defined when Clang is compiling code in Relocatable Device Code (RDC) mode. RDC, enabled with the ``-fgpu-rdc`` compiler option, is necessary for linking device codes across translation units.155 * - ``__HIP__``156 - Defined when compiling with HIP language support, indicating that the code targets the HIP environment.157 * - ``__HIPCC__``158 - Alias to ``__HIP__``.159 * - ``__HIP_DEVICE_COMPILE__``160 - Defined during device code compilation in Clang's separate compilation process for the host and each offloading GPU architecture.161 * - ``__HIP_MEMORY_SCOPE_SINGLETHREAD``162 - Represents single-thread memory scope in HIP (value is 1).163 * - ``__HIP_MEMORY_SCOPE_WAVEFRONT``164 - Represents wavefront memory scope in HIP (value is 2).165 * - ``__HIP_MEMORY_SCOPE_WORKGROUP``166 - Represents workgroup memory scope in HIP (value is 3).167 * - ``__HIP_MEMORY_SCOPE_CLUSTER``168 - Represents cluster memory scope in HIP (value is 6).169 * - ``__HIP_MEMORY_SCOPE_AGENT``170 - Represents agent memory scope in HIP (value is 4).171 * - ``__HIP_MEMORY_SCOPE_SYSTEM``172 - Represents system-wide memory scope in HIP (value is 5).173 * - ``__HIP_NO_IMAGE_SUPPORT__``174 - Defined with a value of 1 when the target device lacks support for HIP image functions.175 * - ``__HIP_NO_IMAGE_SUPPORT``176 - Alias to ``__HIP_NO_IMAGE_SUPPORT__``. Deprecated.177 * - ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``178 - Defined when the GPU default stream is set to per-thread mode.179 * - ``HIP_API_PER_THREAD_DEFAULT_STREAM``180 - Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.181 182Note that some architecture specific AMDGPU macros will have default values when183used from the HIP host compilation.184 185Compilation Modes186=================187 188Each HIP source file contains intertwined device and host code. Depending on the chosen compilation mode by the compiler options ``-fno-gpu-rdc`` and ``-fgpu-rdc``, these portions of code are compiled differently.189 190Device Code Compilation191-----------------------192 193**``-fno-gpu-rdc`` Mode (default)**:194 195- Compiles to a self-contained, fully linked offloading device binary for each offloading device architecture.196- Device code within a Translation Unit (TU) cannot call functions located in another TU.197 198**``-fgpu-rdc`` Mode**:199 200- Compiles to a bitcode for each GPU architecture.201- For each offloading device architecture, the bitcode from different TUs are linked together to create a single offloading device binary.202- Device code in one TU can call functions located in another TU.203 204Host Code Compilation205---------------------206 207**Both Modes**:208 209- Compiles to a relocatable object for each TU.210- These relocatable objects are then linked together.211- Host code within a TU can call host functions and launch kernels from another TU.212 213HIP Fat Binary Registration and Unregistration214==============================================215 216When compiling HIP for AMD GPUs, Clang embeds device code into HIP "fat217binaries" and generates host-side helper functions that register these218fat binaries with the HIP runtime at program start and unregister them at219program exit. In non-RDC mode (``-fno-gpu-rdc``), each compilation unit220typically produces its own HIP fat binary: a container that holds, for every221enabled GPU architecture, a fully linked offloading device image (for example,222a GPU code object) that can be loaded directly by the HIP runtime. In RDC mode223(``-fgpu-rdc``), each compilation unit contributes device code in a relocatable224form (for example, GPU object files or LLVM IR). A later device-link step links225those relocatable inputs into fully linked device images per GPU architecture226and then packages those images into a HIP fat binary container.227 228Registering a HIP fat binary allows the runtime to discover the kernels and229device variables defined in that container and to associate host-side addresses230and symbols with the corresponding GPU-side entities. For example, when a231host-side kernel launch stub is called, the HIP runtime uses information232established during registration (and the fat binary handle it returned) to233identify which GPU kernel symbol to launch from which device image.234 235At the LLVM IR level, Clang/LLVM typically create an internal module236constructor (for example ``__hip_module_ctor`` or a ``.hip.fatbin_reg``237function) and add it to ``@llvm.global_ctors``. This constructor is called by238the C runtime before ``main`` and it:239 240* calls ``__hipRegisterFatBinary`` with a pointer to an internal wrapper241 object that describes the HIP fat binary;242* stores the returned handle in an internal global variable;243* calls an internal helper such as ``__hip_register_globals`` to register244 kernels, device variables and other metadata associated with the fat binary;245* registers a corresponding module destructor with ``atexit`` so it will run246 during program termination and use the stored handle to unregister the fat247 binary from the HIP runtime.248 249The module destructor (for example ``__hip_module_dtor`` or a250``.hip.fatbin_unreg`` function) loads the stored handle, checks that it is251non-null, calls ``__hipUnregisterFatBinary`` to unregister the fat binary from252the HIP runtime, and then clears the handle. This ensures that the HIP runtime253sees each fat binary registered exactly once and that it is unregistered once254at exit, even when multiple translation units contribute HIP kernels to the255same host program.256 257These registration/unregistration helpers are implementation details of Clang's258HIP code generation; user code should not call ``__hipRegisterFatBinary`` or259``__hipUnregisterFatBinary`` directly.260 261Implications for HIP Application Developers262-------------------------------------------263 264From the point of view of HIP application code, Clang and the HIP runtime265provide the following guarantees:266 267* Kernels and device variables defined in HIP code will be registered with the268 HIP runtime before ``main`` begins execution.269* Fat binaries will be unregistered via an ``atexit``-registered module270 destructor after ``main`` returns (or after ``exit`` is called).271 272Beyond these points, the detailed ordering of fat binary registration and273unregistration relative to user-defined global constructors, destructors and274other ``atexit`` handlers is not specified and should not be relied upon.275Applications should avoid depending on HIP kernels or device variables being276usable from global constructors or destructors, and instead perform HIP277initialization and teardown that touches device state in ``main`` (or in278functions called from ``main``).279 280Implications for HIP Runtime Developers281---------------------------------------282 283HIP runtime implementations that are linked with Clang-generated host code284must handle registration and unregistration in the presence of uncertain285global ctor/dtor ordering:286 287* ``__hipRegisterFatBinary`` must accept a pointer to the compiler-generated288 wrapper object and return an opaque handle that remains valid for as long as289 the fat binary may be used.290* ``__hipUnregisterFatBinary`` must accept the handle previously returned by291 ``__hipRegisterFatBinary`` and perform any necessary cleanup. It may be292 called late in process teardown, after other parts of the runtime have293 started shutting down, so it should be robust in the presence of partially294 torn-down state.295* Runtimes should use appropriate synchronization and guards so that fat296 binary registration does not observe uninitialized resources and297 unregistration does not release resources that are still required by other298 runtime components. In particular, registration and unregistration routines299 should be written to be safe under repeated calls and in the presence of300 concurrent or overlapping initialization/teardown logic.301 302Syntax Difference with CUDA303===========================304 305Clang's front end, used for both CUDA and HIP programming models, shares the same parsing and semantic analysis mechanisms. This includes the resolution of overloads concerning device and host functions. While there exists a comprehensive documentation on the syntax differences between Clang and NVCC for CUDA at `Dialect Differences Between Clang and NVCC <https://llvm.org/docs/CompileCudaWithLLVM.html#dialect-differences-between-clang-and-nvcc>`_, it is important to note that these differences also apply to HIP code compilation.306 307Predefined Macros for Differentiation308-------------------------------------309 310To facilitate differentiation between HIP and CUDA code, as well as between device and host compilations within HIP, Clang defines specific macros:311 312- ``__HIP__`` : This macro is defined only when compiling HIP code. It can be used to conditionally compile code specific to HIP, enabling developers to write portable code that can be compiled for both CUDA and HIP.313 314- ``__HIP_DEVICE_COMPILE__`` : Defined exclusively during HIP device compilation, this macro allows for conditional compilation of device-specific code. It provides a mechanism to segregate device and host code, ensuring that each can be optimized for their respective execution environments.315 316Function Pointers Support317=========================318 319Function pointers' support varies with the usage mode in Clang with HIP. The following table provides an overview of the support status across different use-cases and modes.320 321.. list-table:: Function Pointers Support Overview322 :widths: 25 25 25323 :header-rows: 1324 325 * - Use Case326 - ``-fno-gpu-rdc`` Mode (default)327 - ``-fgpu-rdc`` Mode328 * - Defined and used in the same TU329 - Supported330 - Supported331 * - Defined in one TU and used in another TU332 - Not Supported333 - Supported334 335In the ``-fno-gpu-rdc`` mode, the compiler calculates the resource usage of kernels based only on functions present within the same TU. This mode does not support the use of function pointers defined in a different TU due to the possibility of incorrect resource usage calculations, leading to undefined behavior.336 337On the other hand, the ``-fgpu-rdc`` mode allows the definition and use of function pointers across different TUs, as resource usage calculations can accommodate functions from disparate TUs.338 339Virtual Function Support340========================341 342In Clang with HIP, support for calling virtual functions of an object in device or host code is contingent on where the object is constructed.343 344- **Constructed in Device Code**: Virtual functions of an object can be called in device code on a specific offloading device if the object is constructed in device code on an offloading device with the same architecture.345- **Constructed in Host Code**: Virtual functions of an object can be called in host code if the object is constructed in host code.346 347In other scenarios, calling virtual functions is not allowed.348 349Explanation350-----------351 352An object constructed on the device side contains a pointer to the virtual function table on the device side, which is not accessible in host code, and vice versa. Thus, trying to invoke virtual functions from a context different from where the object was constructed will be disallowed because the appropriate virtual table cannot be accessed. The virtual function tables for offloading devices with different architectures are different, therefore trying to invoke virtual functions from an offloading device with a different architecture than where the object is constructed is also disallowed.353 354Example Usage355-------------356 357.. code-block:: c++358 359 class Base {360 public:361 __device__ virtual void virtualFunction() {362 // Base virtual function implementation363 }364 };365 366 class Derived : public Base {367 public:368 __device__ void virtualFunction() override {369 // Derived virtual function implementation370 }371 };372 373 __global__ void kernel() {374 Derived obj;375 Base* basePtr = &obj;376 basePtr->virtualFunction(); // Allowed since obj is constructed in device code377 }378 379Alias Attribute Support380=======================381 382Clang supports alias attributes in HIP code, allowing creation of alternative names for functions and variables. 383 - Aliases work with ``__host__``, ``__device__``, and ``__host__ __device__`` functions and variables.384 - The alias attribute uses the syntax ``__attribute__((alias("target_name")))``. Both weak and strong aliases are supported.385 - Outside of ``extern "C"``, the alias target must use the mangled name of the aliasee386 - The alias is only emitted if the aliasee is emitted on the same side (ie __host__ or __device__), otherwise it is ignored.387 388Example Usage389-------------390 391.. code-block:: c++392 393 extern "C" {394 // Host function alias395 int __HostFunc(void) { return 0; }396 int HostFunc(void) __attribute__((weak, alias("__HostFunc")));397 398 // Device function alias399 __device__ int __DeviceFunc(void) { return 1; }400 __device__ int DeviceFunc(void) __attribute__((weak, alias("__DeviceFunc")));401 402 // Host-device function alias403 __host__ __device__ int __BothFunc(void) { return 2; }404 __host__ __device__ int BothFunc(void) __attribute__((alias("__BothFunc")));405 406 // Variable alias407 int __host_var = 3;408 extern int __attribute__((weak, alias("__host_var"))) host_var;409 }410 // Mangled / overload alias411 __host__ __device__ float __Four(float f) { return 2.0f * f; }412 __host__ __device__ int Four(void) __attribute__((weak, alias("_Z6__Fourv")));413 __host__ __device__ float Four(float f) __attribute__((weak, alias("_Z6__Fourf")));414 415 416Host and Device Attributes of Default Destructors417===================================================418 419If a default destructor does not have explicit host or device attributes,420clang infers these attributes based on the destructors of its data members421and base classes. If any conflicts are detected among these destructors,422clang diagnoses the issue. Otherwise, clang adds an implicit host or device423attribute according to whether the data members's and base classes's424destructors can execute on the host or device side.425 426For explicit template classes with virtual destructors, which must be emitted,427the inference adopts a conservative approach. In this case, implicit host or428device attributes from member and base class destructors are ignored. This429precaution is necessary because, although a constexpr destructor carries430implicit host or device attributes, a constexpr function may call a431non-constexpr function, which is by default a host function.432 433Users can override the inferred host and device attributes of default434destructors by adding explicit host and device attributes to them.435 436C++ Standard Parallelism Offload Support: Compiler And Runtime437==============================================================438 439Introduction440============441 442This section describes the implementation of support for offloading the443execution of standard C++ algorithms to accelerators that can be targeted via444HIP. Furthermore, it enumerates restrictions on user defined code, as well as445the interactions with runtimes.446 447Algorithm Offload: What, Why, Where448===================================449 450C++17 introduced overloads451`for most algorithms in the standard library <https://www.open-std.org/jtc1/sc22/wg21/docs/papers/2016/p0024r2.html>`_452which allow the user to specify a desired453`execution policy <https://en.cppreference.com/w/cpp/algorithm#Execution_policies>`_.454The `parallel_unsequenced_policy <https://en.cppreference.com/w/cpp/algorithm/execution_policy_tag_t>`_455maps relatively well to the execution model of AMD GPUs. This, coupled with the456the availability and maturity of GPU accelerated algorithm libraries that457implement most / all corresponding algorithms in the standard library458(e.g. `rocThrust <https://github.com/ROCm/rocm-libraries/tree/develop/projects/rocthrust>`__), makes459it feasible to provide seamless accelerator offload for supported algorithms,460when an accelerated version exists. Thus, it becomes possible to easily access461the computational resources of an AMD accelerator, via a well specified,462familiar, algorithmic interface, without having to delve into low-level hardware463specific details. Putting it all together:464 465- **What**: standard library algorithms, when invoked with the466 ``parallel_unsequenced_policy``467- **Why**: democratise AMDGPU accelerator programming, without loss of user468 familiarity469- **Where**: only AMDGPU accelerators targeted by Clang/LLVM via HIP470 471Small Example472=============473 474Given the following C++ code:475 476.. code-block:: C++477 478 bool has_the_answer(const std::vector<int>& v) {479 return std::find(std::execution::par_unseq, std::cbegin(v), std::cend(v), 42) != std::cend(v);480 }481 482if Clang is invoked with the ``--hipstdpar --offload-arch=foo`` flags, the call483to ``find`` will be offloaded to an accelerator that is part of the ``foo``484target family. If either ``foo`` or its runtime environment do not support485transparent on-demand paging (such as e.g. that provided in Linux via486`HMM <https://docs.kernel.org/mm/hmm.html>`_), it is necessary to also include487the ``--hipstdpar-interpose-alloc`` flag. If the accelerator specific algorithm488library ``foo`` uses doesn't have an implementation of a particular algorithm,489execution seamlessly falls back to the host CPU. It is legal to specify multiple490``--offload-arch``\s. All the flags we introduce, as well as a thorough view of491various restrictions an their implementations, will be provided below.492 493Implementation - General View494=============================495 496We built support for Algorithm Offload support atop the pre-existing HIP497infrastructure. More specifically, when one requests offload via ``--hipstdpar``,498compilation is switched to HIP compilation, as if ``-x hip`` was specified.499Similarly, linking is also switched to HIP linking, as if ``--hip-link`` was500specified. Note that these are implicit, and one should not assume that any501interop with HIP specific language constructs is available e.g. ``__device__``502annotations are neither necessary nor guaranteed to work.503 504Since there are no language restriction mechanisms in place, it is necessary to505relax HIP language specific semantic checks performed by the FE; they would506identify otherwise valid, offloadable code, as invalid HIP code. Given that we507know that the user intended only for certain algorithms to be offloaded, and508encoded this by specifying the ``parallel_unsequenced_policy``, we rely on a509pass over IR to clean up any and all code that was not "meant" for offload. If510requested, allocation interposition is also handled via a separate pass over IR.511 512To interface with the client HIP runtime, and to forward offloaded algorithm513invocations to the corresponding accelerator specific library implementation, an514implementation detail forwarding header is implicitly included by the driver,515when compiling with ``--hipstdpar``. In what follows, we will delve into each516component that contributes to implementing Algorithm Offload support.517 518Implementation - Driver519=======================520 521We augment the ``clang`` driver with the following flags:522 523- ``--hipstdpar`` enables algorithm offload, which depending on phase, has the524 following effects:525 526 - when compiling:527 528 - ``-x hip`` gets prepended to enable HIP support;529 - the ``ROCmToolchain`` component checks for the ``hipstdpar_lib.hpp``530 forwarding header,531 `rocThrust <https://rocm.docs.amd.com/projects/rocThrust/en/latest/>`_ and532 `rocPrim <https://rocm.docs.amd.com/projects/rocPRIM/en/latest/>`_ in533 their canonical locations, which can be overriden via flags found below;534 if all are found, the forwarding header gets implicitly included,535 otherwise an error listing the missing component is generated;536 - the ``LangOpts.HIPStdPar`` member is set.537 538 - when linking:539 540 - ``--hip-link`` and ``-frtlib-add-rpath`` gets appended to enable HIP541 support.542 543- ``--hipstdpar-interpose-alloc`` enables the interposition of standard544 allocation / deallocation functions with accelerator aware equivalents; the545 ``LangOpts.HIPStdParInterposeAlloc`` member is set;546- ``--hipstdpar-path=`` specifies a non-canonical path for the forwarding547 header; it must point to the folder where the header is located and not to the548 header itself;549- ``--hipstdpar-thrust-path=`` specifies a non-canonical path for550 `rocThrust <https://rocm.docs.amd.com/projects/rocThrust/en/latest/>`_; it551 must point to the folder where the library is installed / built under a552 ``/thrust`` subfolder;553- ``--hipstdpar-prim-path=`` specifies a non-canonical path for554 `rocPrim <https://rocm.docs.amd.com/projects/rocPRIM/en/latest/>`_; it must555 point to the folder where the library is installed / built under a556 ``/rocprim`` subfolder;557 558The `--offload-arch <https://llvm.org/docs/AMDGPUUsage.html#amdgpu-processors>`_559flag can be used to specify the accelerator for which offload code is to be560generated.561 562Implementation - Front-End563==========================564 565When ``LangOpts.HIPStdPar`` is set, we relax some of the HIP language specific566``Sema`` checks to account for the fact that we want to consume pure unannotated567C++ code:568 5691. ``__device__`` / ``__host__ __device__`` functions (which would originate in570 the accelerator specific algorithm library) are allowed to call implicitly571 ``__host__`` functions;5722. ``__global__`` functions (which would originate in the accelerator specific573 algorithm library) are allowed to call implicitly ``__host__`` functions;5743. resolving ``__builtin`` availability is deferred, because it is possible that575 a ``__builtin`` that is unavailable on the target accelerator is not576 reachable from any offloaded algorithm, and thus will be safely removed in577 the middle-end;5784. ASM parsing / checking is deferred, because it is possible that an ASM block579 that e.g. uses some constraints that are incompatible with the target580 accelerator is not reachable from any offloaded algorithm, and thus will be581 safely removed in the middle-end.582 583``CodeGen`` is similarly relaxed, with implicitly ``__host__`` functions being584emitted as well.585 586Implementation - Middle-End587===========================588 589We add two ``opt`` passes:590 5911. ``HipStdParAcceleratorCodeSelectionPass``592 593 - For all kernels in a ``Module``, compute reachability, where a function594 ``F`` is reachable from a kernel ``K`` if and only if there exists a direct595 call-chain rooted in ``F`` that includes ``K``;596 - Remove all functions that are not reachable from kernels;597 - This pass is only run when compiling for the accelerator.598 599The first pass assumes that the only code that the user intended to offload was600that which was directly or transitively invocable as part of an algorithm601execution. It also assumes that an accelerator aware algorithm implementation602would rely on accelerator specific special functions (kernels), and that these603effectively constitute the only roots for accelerator execution graphs. Both of604these assumptions are based on observing how widespread accelerators,605such as GPUs, work.606 6071. ``HipStdParAllocationInterpositionPass``608 609 - Iterate through all functions in a ``Module``, and replace standard610 allocation / deallocation functions with accelerator-aware equivalents,611 based on a pre-established table; the list of functions that can be612 interposed is available613 `here <https://github.com/ROCm/roc-stdpar#allocation--deallocation-interposition-status>`__;614 - This is only run when compiling for the host.615 616The second pass is optional.617 618Implementation - Forwarding Header619==================================620 621The forwarding header implements two pieces of functionality:622 6231. It forwards algorithms to a target accelerator, which is done by relying on624 C++ language rules around overloading:625 626 - overloads taking an explicit argument of type627 ``parallel_unsequenced_policy`` are introduced into the ``std`` namespace;628 - these will get preferentially selected versus the master template;629 - the body forwards to the equivalent algorithm from the accelerator specific630 library631 6322. It provides allocation / deallocation functions that are equivalent to the633 standard ones, but obtain memory by invoking634 `hipMallocManaged <https://rocm.docs.amd.com/projects/HIP/en/latest/.doxygen/docBin/html/group___memory_m.html#gab8cfa0e292193fa37e0cc2e4911fa90a>`_635 and release it via `hipFree <https://rocm.docs.amd.com/projects/HIP/en/latest/.doxygen/docBin/html/group___memory.html#ga740d08da65cae1441ba32f8fedb863d1>`_.636 637Predefined Macros638=================639 640.. list-table::641 :header-rows: 1642 643 * - Macro644 - Description645 * - ``__HIPSTDPAR__``646 - Defined when Clang is compiling code in algorithm offload mode, enabled647 with the ``--hipstdpar`` compiler option.648 * - ``__HIPSTDPAR_INTERPOSE_ALLOC__`` / ``__HIPSTDPAR_INTERPOSE_ALLOC_V1__``649 - Defined only when compiling in algorithm offload mode, when the user650 enables interposition mode with the ``--hipstdpar-interpose-alloc``651 compiler option, indicating that all dynamic memory allocation /652 deallocation functions should be replaced with accelerator aware653 variants.654 655Restrictions656============657 658We define two modes in which runtime execution can occur:659 6601. **HMM Mode** - this assumes that the661 `HMM <https://docs.kernel.org/mm/hmm.html>`_ subsystem of the Linux kernel662 is used to provide transparent on-demand paging i.e. memory obtained from a663 system / OS allocator such as via a call to ``malloc`` or ``operator new`` is664 directly accessible to the accelerator and it follows the C++ memory model;6652. **Interposition Mode** - this is a fallback mode for cases where transparent666 on-demand paging is unavailable (e.g. in the Windows OS), which means that667 memory must be allocated via an accelerator aware mechanism, and system668 allocated memory is inaccessible for the accelerator.669 670The following restrictions imposed on user code apply to both modes:671 6721. Pointers to function, and all associated features, such as e.g. dynamic673 polymorphism, cannot be used (directly or transitively) by the user provided674 callable passed to an algorithm invocation;6752. ``static`` (except for program-wide unique ones) / ``thread`` storage676 duration variables cannot be used (directly or transitively) in name by the677 user provided callable;6783. User code must be compiled in ``-fgpu-rdc`` mode in order for global /679 namespace scope variables / program-wide unique ``static`` storage duration680 variables to be usable in name by the user provided callable;6814. Only algorithms that are invoked with the ``parallel_unsequenced_policy`` are682 candidates for offload;6835. Only algorithms that are invoked with iterator arguments that model684 `random_access_iterator <https://en.cppreference.com/w/cpp/iterator/random_access_iterator>`_685 are candidates for offload;6866. `Exceptions <https://en.cppreference.com/w/cpp/language/exceptions>`_ cannot687 be used by the user provided callable;6887. Dynamic memory allocation (e.g. ``operator new``) cannot be used by the user689 provided callable;6908. Selective offload is not possible i.e. it is not possible to indicate that691 only some algorithms invoked with the ``parallel_unsequenced_policy`` are to692 be executed on the accelerator.693 694In addition to the above, using **Interposition Mode** imposes the following695additional restrictions:696 6971. All code that is expected to interoperate has to be recompiled with the698 ``--hipstdpar-interpose-alloc`` flag i.e. it is not safe to compose libraries699 that have been independently compiled;700 701Current Support702===============703 704At the moment, C++ Standard Parallelism Offload is only available for AMD GPUs,705when the `ROCm <https://rocm.docs.amd.com/en/latest/>`_ stack is used, on the706Linux operating system. Support is synthesised in the following table:707 708.. list-table::709 :header-rows: 1710 711 * - `Processor <https://llvm.org/docs/AMDGPUUsage.html#amdgpu-processors>`_712 - HMM Mode713 - Interposition Mode714 * - GCN GFX9 (Vega)715 - YES716 - YES717 * - GCN GFX10.1 (RDNA 1)718 - *NO*719 - YES720 * - GCN GFX10.3 (RDNA 2)721 - *NO*722 - YES723 * - GCN GFX11 (RDNA 3)724 - *NO*725 - YES726 * - GCN GFX12 (RDNA 4)727 - *NO*728 - YES729 730The minimum Linux kernel version for running in HMM mode is 6.4.731 732The forwarding header is packaged by733`ROCm <https://rocm.docs.amd.com/en/latest/>`_, and is obtainable by installing734the `hipstdpar` packege. The list algorithms that can be offloaded is available735`here <https://github.com/ROCm/roc-stdpar#algorithm-support-status>`_. More736details are available via the dedicated blog737`<https://rocm.blogs.amd.com/software-tools-optimization/hipstdpar/README.html>`_.738 739HIP Specific Elements740---------------------741 7421. There is no defined interop with the743 `HIP kernel language <https://rocm.docs.amd.com/projects/HIP/en/latest/reference/kernel_language.html>`_;744 whilst things like using `__device__` annotations might accidentally "work",745 they are not guaranteed to, and thus cannot be relied upon by user code;746 747 - A consequence of the above is that both bitcode linking and linking748 relocatable object files will "work", but it is not guaranteed to remain749 working or actively tested at the moment; this restriction might be relaxed750 in the future.751 7522. Combining explicit HIP, CUDA or OpenMP Offload compilation with753 ``--hipstdpar`` based offloading is not allowed or supported in any way.7543. There is no way to target different accelerators via a standard algorithm755 invocation (`this might be addressed in future C++ standards <https://www.open-std.org/jtc1/sc22/wg21/docs/papers/2023/p2500r1.html>`_);756 an unsafe (per the point above) way of achieving this is to spawn new threads757 and invoke the `hipSetDevice <https://rocm.docs.amd.com/projects/HIP/en/latest/.doxygen/docBin/html/group___device.html#ga43c1e7f15925eeb762195ccb5e063eae>`_758 interface e.g.:759 760 .. code-block:: c++761 762 int accelerator_0 = ...;763 int accelerator_1 = ...;764 765 bool multiple_accelerators(const std::vector<int>& u, const std::vector<int>& v) {766 std::atomic<unsigned int> r{0u};767 768 thread t0{[&]() {769 hipSetDevice(accelerator_0);770 771 r += std::count(std::execution::par_unseq, std::cbegin(u), std::cend(u), 42);772 }};773 thread t1{[&]() {774 hitSetDevice(accelerator_1);775 776 r += std::count(std::execution::par_unseq, std::cbegin(v), std::cend(v), 314152)777 }};778 779 t0.join();780 t1.join();781 782 return r;783 }784 785 Note that this is a temporary, unsafe workaround for a deficiency in the C++786 Standard.787 788Open Questions / Future Developments789====================================790 7911. The restriction on the use of ``static`` / ``thread`` storage duration792 variables in offloaded algorithms might be lifted;7932. The restriction on the use of dynamic memory allocation in offloaded794 algorithms will be lifted in the future.7953. The restriction on the use of pointers to function, and associated features796 such as dynamic polymorphism might be lifted in the future, when running in797 **HMM Mode**;7984. Offload support might be extended to cases where the ``parallel_policy`` is799 used for some or all targets.800 801SPIR-V Support on HIPAMD ToolChain802==================================803 804The HIPAMD ToolChain supports targeting805`AMDGCN Flavoured SPIR-V <https://llvm.org/docs/SPIRVUsage.html#target-triples>`_.806The support for SPIR-V in the ROCm and HIPAMD ToolChain is under active807development.808 809Compilation Process810-------------------811 812When compiling HIP programs with the intent of utilizing SPIR-V, the process813diverges from the traditional compilation flow:814 815Using ``--offload-arch=amdgcnspirv``816^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^817 818- **Target Triple**: The ``--offload-arch=amdgcnspirv`` flag instructs the819 compiler to use the target triple ``spirv64-amd-amdhsa``. This approach does820 generates generic AMDGCN SPIR-V which retains architecture specific elements821 without hardcoding them, thus allowing for optimal target specific code to be822 generated at run time, when the concrete target is known.823 824- **LLVM IR Translation**: The program is compiled to LLVM Intermediate825 Representation (IR), which is subsequently translated into SPIR-V. In the826 future, this translation step will be replaced by direct SPIR-V emission via827 the SPIR-V Back-end.828 829- **Clang Offload Bundler**: The resulting SPIR-V is embedded in the Clang830 offload bundler with the bundle ID ``hip-spirv64-amd-amdhsa--amdgcnspirv``.831 832Architecture Specific Macros833----------------------------834 835None of the architecture specific :doc:`AMDGPU macros <AMDGPUSupport>` are836defined when targeting SPIR-V. An alternative, more flexible mechanism to enable837doing per target / per feature code selection will be added in the future.838