Lines Matching refs:the
11 This document describes the Clang driver and code generation steps for creating
14 document is to illustrate the steps necessary to create an offloading
22 Clang and then executed using the ``libomptarget`` runtime and the associated
23 plugin for the target architecture, e.g. ``libomptarget.rtl.cuda``. This section
24 describes the steps necessary to create a functioning device image that can be
25 loaded by the OpenMP runtime. More information on the OpenMP runtimes can be
26 found at the `OpenMP documentation page <https://openmp.llvm.org>`__.
34 can be run on the target device. OpenMP offloading creates executable images by
35 compiling the input file for both the host and the target device. The output
36 from the device phase then needs to be embedded into the host to create a fat
37 object. A special tool then needs to extract the device code from the fat
38 objects, run the device linking step, and embed the final image in a symbol the
39 host runtime library can use to register the library and access the symbols on
40 the device.
45 The compiler performs the following high-level actions to generate OpenMP
48 * Compile the input file for the host to produce a bitcode file. Lower ``#pragma
50 Entries>` and create metadata to indicate which entries are on the device.
51 * Compile the input file for the target :ref:`device <Device Compilation>` using
52 the :ref:`offloading entry <Generating Offloading Entries>` metadata created
53 by the host.
54 * Link the OpenMP device runtime library and run the backend to create a device
56 * Run the backend on the host bitcode file and create a :ref:`fat object file
57 <Creating Fat Objects>` using the device object file.
58 * Pass the fat object file to the :ref:`linker wrapper tool <Device Linking>`
59 and extract the device objects. Run the device linking action on the extracted
61 * :ref:`Wrap <Device Binary Wrapping>` the :ref:`device images <Device linking>`
63 can be accessed by the host.
64 * Add the :ref:`wrapped binary <Device Binary Wrapping>` to the linker input and
65 run the host linking action. Link with ``libomptarget`` to register and
66 execute the images.
73 The first step in compilation is to generate offloading entries for the host.
75 be provided by the device. Blocks contained in a ``#pragma omp target`` or
77 entries generated. The following table shows the :ref:`offload entry structure
88 … | char* | name | Name of the symbol |
90 … | size_t | size | Size of the entry info (0 if it is a function) |
92 … | int32_t | flags | Flags associated with the entry (see :ref:`table-offload_entry_flags`) |
94 … | int32_t | reserved | Reserved, to be used by the runtime library. |
97 The address of the global symbol will be set to the device pointer value by the
98 runtime once the device image is loaded. The flags are set to indicate the
99 handling required for the offloading entry. If the offloading entry is an entry
100 to a target region it can have one of the following :ref:`entry flags
109 | OMPTargetRegionEntryTargetRegion | 0x00 | Mark the entry as generic target region |
111 | OMPTargetRegionEntryCtor | 0x02 | Mark the entry as a global constructor |
113 | OMPTargetRegionEntryDtor | 0x04 | Mark the entry as a global destructor |
116 If the offloading entry is a global variable, indicated by a non-zero size, it
117 will instead have one of the following :ref:`global
126 …| OMPTargetGlobalVarEntryTo | 0x00 | Mark the entry as a 'to' attribute (w.r.t. the to clause) …
128 …| OMPTargetGlobalVarEntryLink | 0x01 | Mark the entry as a 'link' attribute (w.r.t. the link clau…
131 The target offload entries are used by the runtime to access the device kernels
132 and globals that will be provided by the final device image. Each offloading
133 entry is set to use the ``omp_offloading_entries`` section. When the final
134 application is created the linker will provide the
136 which are used to create the :ref:`final image <Device Binary Wrapping>`.
138 This information is used by the device compilation stage to determine which
139 symbols need to be exported from the device. We use the ``omp_offload.info``
142 Accessing Entries on the Device
145 Accessing the entries in the device is done using the address field in the
147 the address to the pointer associated with the device image during runtime
148 initialization. This is used to call the corresponding kernel function when
149 entering a ``#pragma omp target`` region. For variables, the runtime maintains a
151 ``#pragma omp target declare`` directive are first initialized to the host's
152 address. Once the device address is initialized we insert it into the table to
153 map the host address to the device address.
159 ``libomptarget``. This allows the front-end to generate information the runtime
160 library uses for more informative error messages. This is done using the
163 locations to the runtime.
171 …| int32_t | reserved | Reserved, to be used by the runtime library. …
175 …| int32_t | reserved | Reserved, to be used by the runtime library. …
177 …| int32_t | reserved | Reserved, to be used by the runtime library. …
182 If debugging information is enabled, we will also create strings to indicate the
183 names and declarations of variables mapped in target regions. These have the
184 same format as the source location in the :ref:`identifier structure
185 <table-ident_t_structure>`, but the function name is replaced with the variable
194 compilation stage is performed differently from the host stage. Namely, we do
195 not generate any offloading entries. This is set by passing the
196 ``-fopenmp-is-device`` flag to the front-end. We use the host bitcode to
197 determine which symbols to export from the device. The bitcode file is passed in
198 from the previous stage using the ``-fopenmp-host-ir-file-path`` flag.
201 When compiling for the OpenMP device, we set the visibility of all device
203 class of errors where a symbol in the target device could preempt a host
206 The OpenMP runtime library is linked in during compilation to provide the
209 ``libomptarget-nvptx64-sm_70.bc``) using the ``-mlink-builtin-bitcode`` flag.
211 the target is a standard architecture with an existing ``libomp``
213 create a relocatable device object file that can be embedded in the host.
221 device. We create a fat object by embedding the output of the device compilation
222 stage into the host as a named section. The output from the device compilation
223 is passed to the host backend using the ``-fembed-offload-object`` flag. This
224 embeds the device image into the ``.llvm.offloading`` section using a special
226 bundle metadata about the image so the linker can associate the proper device
227 linking action with the image. Each device image will start with the magic bytes
234 The device code will then be placed in the corresponding section one the backend
235 is run on the host, creating a fat object. Using fat objects allows us to treat
237 contain the following :ref:`offloading sections <table-offloading_sections>`. We
248 …| .llvm.offloading | Embedded device object file for the target device and archite…
259 over the host linking job. It scans the input object files for the offloading
261 extracted and passed to the appropriate linking job. The linked device image is
262 then :ref:`wrapped <Device Binary Wrapping>` to create the symbols used to load
263 the device image and link it with the host.
266 optimization (LTO). This is used whenever the object files embedded in the host
268 support a relocatable object format, such as AMDGPU or SPIR-V, or if the user
269 requested it using the ``-foffload-lto`` flag.
276 Various structures and functions are used to create the information necessary to
277 offload code on the device. We use the :ref:`linked device executable <Device
278 Linking>` with the corresponding offloading entries to create the symbols
279 necessary to load and execute the device image.
287 entries are stored using the ``__start_omp_offloading_entries`` and
288 ``__stop_omp_offloading_entries`` symbols generated by the linker using the
297 | void* | ImageStart | Pointer to the target code start |
299 | void* | ImageEnd | Pointer to the target code end |
337 …entries | __tgt_offload_entry | .omp_offloading_entries | Begin symbol for the offload entries tab…
339 …_entries | __tgt_offload_entry | .omp_offloading_entries | End symbol for the offload entries tab…
341 …__tgt_offload_entry | .omp_offloading_entries | Dummy zero-sized object in the offload entries …
345 … __tgt_device_image | .omp_offloading_entries | ELF device code object of the first image. …
347 … __tgt_device_image | .omp_offloading_entries | ELF device code object of the (N+1)th image. …
359 This object is passed to the offloading runtime at program startup and it
360 describes all device images available in the executable or shared library. It
398 The global constructor (``.omp_offloading.descriptor_reg()``) registers the
399 device images with the runtime by calling the ``__tgt_register_lib()`` runtime
401 is run once when the program starts. Similarly, the global destructor
403 the destructor and is also defined in ``.text.startup`` section and run when the
431 This code is compiled using the following Clang flags.
437 The output section in the object file can be seen using the ``readelf`` utility.
438 The ``.llvm.offloading`` section has the ``SHF_EXCLUDE`` flag so it will be
439 removed from the final executable or shared library by the linker.
450 Compiling this file again will invoke the ``clang-linker-wrapper`` utility to
451 extract and link the device code stored at the section named
453 the section named ``omp_offloading_entries`` to create the symbols necessary for
454 ``libomptarget`` to register the device image and call the entry function.
461 We can see the steps created by clang to generate the offloading code using the
462 ``-ccc-print-phases`` option in Clang. This matches the description in