=================
.. note::
- This documentation describes Clang's behavior using the new offloading driver
- which. This currently must be enabled manually using ``-fopenmp-new-driver``.
+ This documentation describes Clang's behavior using the new offloading
+ driver. This currently must be enabled manually using
+ ``-fopenmp-new-driver``.
Clang supports OpenMP target offloading to several different architectures such
as NVPTX, AMDGPU, X86_64, Arm, and PowerPC. Offloading code is generated by
from the device phase then needs to be embedded into the host to create a fat
object. A special tool then needs to extract the device code from the fat
objects, run the device linking step, and embed the final image in a symbol the
-host can use to register the library and access the symbols on the device.
+host runtime library can use to register the library and access the symbols on
+the device.
Compilation Process
^^^^^^^^^^^^^^^^^^^
-The compiler performs the following high-level actions to generate offloading
-code:
+The compiler performs the following high-level actions to generate OpenMP
+offloading code:
* Compile the input file for the host to produce a bitcode file. Lower ``#pragma
omp target`` declarations to :ref:`offloading entries <Generating Offloading
| int32_t | reserved | Reserved, to be used by the runtime library. |
+---------+------------+------------------------------------------------------------------------+
-The address of the global symbol will be set to the appropriate value by the
+The address of the global symbol will be set to the device pointer value by the
runtime once the device image is loaded. The flags are set to indicate the
handling required for the offloading entry. If the offloading entry is an entry
-to a target region it can have one of the following
-:ref:`entry flags <table-offload_entry_flags>`.
+to a target region it can have one of the following :ref:`entry flags
+<table-offload_entry_flags>`.
.. table:: Target Region Entry Flags
:name: table-offload_entry_flags
``__start_omp_offloading_entries`` and ``__stop_omp_offloading_entries`` symbols
which are used to create the :ref:`final image <Device Binary Wrapping>`.
-This information is by the device compilation stage to determine which symbols
-need to be exported from the device. We use the ``omp_offload.info`` metadata
-node to pass this information device compilation stage.
+This information is used by the device compilation stage to determine which
+symbols need to be exported from the device. We use the ``omp_offload.info``
+metadata node to pass this information device compilation stage.
Accessing Entries on the Device
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
initialization. This is used to call the corresponding kernel function when
entering a ``#pragma omp target`` region. For variables, the runtime maintains a
table mapping host pointers to device pointers. Global variables inside a
-``#pragma omp target reclare`` directive are first initialized to the host's
+``#pragma omp target declare`` directive are first initialized to the host's
address. Once the device address is initialized we insert it into the table to
map the host address to the device address.
If debugging information is enabled, we will also create strings to indicate the
names and declarations of variables mapped in target regions. These have the
same format as the source location in the :ref:`identifier structure
-<table-ident_t_structure>`, but the filename is replaced with the variable name.
+<table-ident_t_structure>`, but the function name is replaced with the variable
+name.
.. _Device Compilation:
optimization (LTO). This is used whenever the object files embedded in the host
contain LLVM bitcode. Bitcode will be embedded for architectures that do not
support a relocatable object format, such as AMDGPU or SPIR-V, or if the user
-passed in ``-foffload-lto``.
+requested it using the ``-foffload-lto`` flag.
.. _Device Binary Wrapping:
__stop_omp_offloading_entries /*HostEntriesEnd*/
};
+
Global Constructor and Destructor
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^