maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[Clang][Docs] Fix some typos in offloading design documentation

This commit is contained in:
Joseph Huber 2022-02-07 15:23:50 -05:00
parent 1237c1496f
commit 5c9ee35138
1 changed files with 18 additions and 14 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

32
clang/docs/OffloadingDesign.rst
View File

@ -18,8 +18,9 @@ OpenMP Offloading
=================
.. 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
@ -40,13 +41,14 @@ compiling the input file for both the host and the target device. The output
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
@ -97,11 +99,11 @@ entries generated. The following table shows the :ref:`offload entry structure
| 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
@ -138,9 +140,9 @@ application is created the linker will provide the
``__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
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@ -151,7 +153,7 @@ the address to the pointer associated with the device image during runtime
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.
@ -185,7 +187,8 @@ locations to the runtime.
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:
@ -267,7 +270,7 @@ The linker wrapper tool supports linking bitcode files through link time
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:
@ -392,6 +395,7 @@ is defined as follows:
__stop_omp_offloading_entries /*HostEntriesEnd*/
};
Global Constructor and Destructor
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^