| ============================= |
| Offloading Design & Internals |
| ============================= |
| |
| .. contents:: |
| :local: |
| |
| Introduction |
| ============ |
| |
| This document describes the Clang driver and code generation steps for creating |
| offloading applications. Clang supports offloading to various architectures |
| using programming models like CUDA, HIP, and OpenMP. The purpose of this |
| document is to illustrate the steps necessary to create an offloading |
| application using Clang. |
| |
| OpenMP Offloading |
| ================= |
| |
| Clang supports OpenMP target offloading to several different architectures such |
| as NVPTX, AMDGPU, X86_64, Arm, and PowerPC. Offloading code is generated by |
| Clang and then executed using the ``libomptarget`` runtime and the associated |
| plugin for the target architecture, e.g. ``libomptarget.rtl.cuda``. This section |
| describes the steps necessary to create a functioning device image that can be |
| loaded by the OpenMP runtime. More information on the OpenMP runtimes can be |
| found at the `OpenMP documentation page <https://openmp.llvm.org>`__. |
| |
| .. _Offloading Overview: |
| |
| Offloading Overview |
| ------------------- |
| |
| The goal of offloading compilation is to create an executable device image that |
| can be run on the target device. OpenMP offloading creates executable images by |
| 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 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 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 |
| Entries>` and create metadata to indicate which entries are on the device. |
| * Compile the input file for the target :ref:`device <Device Compilation>` using |
| the :ref:`offloading entry <Generating Offloading Entries>` metadata created |
| by the host. |
| * Link the OpenMP device runtime library and run the backend to create a device |
| object file. |
| * Run the backend on the host bitcode file and create a :ref:`fat object file |
| <Creating Fat Objects>` using the device object file. |
| * Pass the fat object file to the :ref:`linker wrapper tool <Device Linking>` |
| and extract the device objects. Run the device linking action on the extracted |
| objects. |
| * :ref:`Wrap <Device Binary Wrapping>` the :ref:`device images <Device linking>` |
| and :ref:`offload entries <Generating Offloading Entries>` in a symbol that |
| can be accessed by the host. |
| * Add the :ref:`wrapped binary <Device Binary Wrapping>` to the linker input and |
| run the host linking action. Link with ``libomptarget`` to register and |
| execute the images. |
| |
| .. _Generating Offloading Entries: |
| |
| Generating Offloading Entries |
| ----------------------------- |
| |
| The first step in compilation is to generate offloading entries for the host. |
| This information is used to identify function kernels or global values that will |
| be provided by the device. Blocks contained in a ``#pragma omp target`` or |
| symbols inside a ``#pragma omp declare target`` directive will have offloading |
| entries generated. The following table shows the :ref:`offload entry structure |
| <table-tgt_offload_entry_structure>`. |
| |
| .. table:: __tgt_offload_entry Structure |
| :name: table-tgt_offload_entry_structure |
| |
| +---------+------------+------------------------------------------------------------------------+ |
| | Type | Identifier | Description | |
| +=========+============+========================================================================+ |
| | void* | addr | Address of global symbol within device image (function or global) | |
| +---------+------------+------------------------------------------------------------------------+ |
| | char* | name | Name of the symbol | |
| +---------+------------+------------------------------------------------------------------------+ |
| | size_t | size | Size of the entry info (0 if it is a function) | |
| +---------+------------+------------------------------------------------------------------------+ |
| | int32_t | flags | Flags associated with the entry (see :ref:`table-offload_entry_flags`) | |
| +---------+------------+------------------------------------------------------------------------+ |
| | int32_t | reserved | Reserved, to be used by the runtime library. | |
| +---------+------------+------------------------------------------------------------------------+ |
| |
| 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>`. |
| |
| .. table:: Target Region Entry Flags |
| :name: table-offload_entry_flags |
| |
| +----------------------------------+-------+-----------------------------------------+ |
| | Name | Value | Description | |
| +==================================+=======+=========================================+ |
| | OMPTargetRegionEntryTargetRegion | 0x00 | Mark the entry as generic target region | |
| +----------------------------------+-------+-----------------------------------------+ |
| | OMPTargetRegionEntryCtor | 0x02 | Mark the entry as a global constructor | |
| +----------------------------------+-------+-----------------------------------------+ |
| | OMPTargetRegionEntryDtor | 0x04 | Mark the entry as a global destructor | |
| +----------------------------------+-------+-----------------------------------------+ |
| |
| If the offloading entry is a global variable, indicated by a non-zero size, it |
| will instead have one of the following :ref:`global |
| <table-offload_global_flags>` flags. |
| |
| .. table:: Target Region Global |
| :name: table-offload_global_flags |
| |
| +-----------------------------+-------+---------------------------------------------------------------+ |
| | Name | Value | Description | |
| +=============================+=======+===============================================================+ |
| | OMPTargetGlobalVarEntryTo | 0x00 | Mark the entry as a 'to' attribute (w.r.t. the to clause) | |
| +-----------------------------+-------+---------------------------------------------------------------+ |
| | OMPTargetGlobalVarEntryLink | 0x01 | Mark the entry as a 'link' attribute (w.r.t. the link clause) | |
| +-----------------------------+-------+---------------------------------------------------------------+ |
| |
| The target offload entries are used by the runtime to access the device kernels |
| and globals that will be provided by the final device image. Each offloading |
| entry is set to use the ``omp_offloading_entries`` section. When the final |
| 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 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 |
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| Accessing the entries in the device is done using the address field in the |
| :ref:`offload entry<table-tgt_offload_entry_structure>`. The runtime will set |
| 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 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. |
| |
| Debugging Information |
| ^^^^^^^^^^^^^^^^^^^^^ |
| |
| We generate structures to hold debugging information that is passed to |
| ``libomptarget``. This allows the front-end to generate information the runtime |
| library uses for more informative error messages. This is done using the |
| standard :ref:`identifier structure <table-ident_t_structure>` used in |
| ``libomp`` and ``libomptarget``. This is used to pass information and source |
| locations to the runtime. |
| |
| .. table:: ident_t Structure |
| :name: table-ident_t_structure |
| |
| +---------+------------+-----------------------------------------------------------------------------+ |
| | Type | Identifier | Description | |
| +=========+============+=============================================================================+ |
| | int32_t | reserved | Reserved, to be used by the runtime library. | |
| +---------+------------+-----------------------------------------------------------------------------+ |
| | int32_t | flags | Flags used to indicate some features, mostly unused. | |
| +---------+------------+-----------------------------------------------------------------------------+ |
| | int32_t | reserved | Reserved, to be used by the runtime library. | |
| +---------+------------+-----------------------------------------------------------------------------+ |
| | int32_t | reserved | Reserved, to be used by the runtime library. | |
| +---------+------------+-----------------------------------------------------------------------------+ |
| | char* | psource | Program source information, stored as ";filename;function;line;column;;\\0" | |
| +---------+------------+-----------------------------------------------------------------------------+ |
| |
| 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 function name is replaced with the variable |
| name. |
| |
| .. _Device Compilation: |
| |
| Offload Device Compilation |
| -------------------------- |
| |
| The input file is compiled for each active device toolchain. The device |
| compilation stage is performed differently from the host stage. Namely, we do |
| not generate any offloading entries. This is set by passing the |
| ``-fopenmp-is-target-device`` flag to the front-end. We use the host bitcode to |
| determine which symbols to export from the device. The bitcode file is passed in |
| from the previous stage using the ``-fopenmp-host-ir-file-path`` flag. |
| Compilation is otherwise performed as it would be for any other target triple. |
| |
| When compiling for the OpenMP device, we set the visibility of all device |
| symbols to be ``protected`` by default. This improves performance and prevents a |
| class of errors where a symbol in the target device could preempt a host |
| library. |
| |
| The OpenMP runtime library is linked in during compilation to provide the |
| implementations for standard OpenMP functionality. For GPU targets this is done |
| by linking in a special bitcode library during compilation, (e.g. |
| ``libomptarget-nvptx64-sm_70.bc``) using the ``-mlink-builtin-bitcode`` flag. |
| Other device libraries, such as CUDA's libdevice, are also linked this way. If |
| the target is a standard architecture with an existing ``libomp`` |
| implementation, that will be linked instead. Finally, device tools are used to |
| create a relocatable device object file that can be embedded in the host. |
| |
| .. _Creating Fat Objects: |
| |
| Creating Fat Objects |
| -------------------- |
| |
| A fat binary is a binary file that contains information intended for another |
| device. We create a fat object by embedding the output of the device compilation |
| stage into the host as a named section. The output from the device compilation |
| is passed to the host backend using the ``-fembed-offload-object`` flag. This |
| embeds the device image into the ``.llvm.offloading`` section using a special |
| binary format that behaves like a string map. This binary format is used to |
| bundle metadata about the image so the linker can associate the proper device |
| linking action with the image. Each device image will start with the magic bytes |
| ``0x10FF10AD``. |
| |
| .. code-block:: llvm |
| |
| @llvm.embedded.object = private constant [1 x i8] c"\00", section ".llvm.offloading" |
| |
| The device code will then be placed in the corresponding section one the backend |
| is run on the host, creating a fat object. Using fat objects allows us to treat |
| offloading objects as standard host objects. The final object file should |
| contain the following :ref:`offloading sections <table-offloading_sections>`. We |
| will use this information when :ref:`Device Linking`. |
| |
| .. table:: Offloading Sections |
| :name: table-offloading_sections |
| |
| +----------------------------------+------------------------------------------------------------------------------+ |
| | Section | Description | |
| +==================================+==============================================================================+ |
| | omp_offloading_entries | Offloading entry information (see :ref:`table-tgt_offload_entry_structure`) | |
| +----------------------------------+------------------------------------------------------------------------------+ |
| | .llvm.offloading | Embedded device object file for the target device and architecture | |
| +----------------------------------+------------------------------------------------------------------------------+ |
| |
| .. _Device Linking: |
| |
| Linking Target Device Code |
| -------------------------- |
| |
| Objects containing :ref:`table-offloading_sections` require special handling to |
| create an executable device image. This is done using a Clang tool, see |
| :doc:`ClangLinkerWrapper` for more information. This tool works as a wrapper |
| over the host linking job. It scans the input object files for the offloading |
| section ``.llvm.offloading``. The device files stored in this section are then |
| extracted and passed to the appropriate linking job. The linked device image is |
| then :ref:`wrapped <Device Binary Wrapping>` to create the symbols used to load |
| the device image and link it with the host. |
| |
| 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 |
| requested it using the ``-foffload-lto`` flag. |
| |
| .. _Device Binary Wrapping: |
| |
| Device Binary Wrapping |
| ---------------------- |
| |
| Various structures and functions are used to create the information necessary to |
| offload code on the device. We use the :ref:`linked device executable <Device |
| Linking>` with the corresponding offloading entries to create the symbols |
| necessary to load and execute the device image. |
| |
| Structure Types |
| ^^^^^^^^^^^^^^^ |
| |
| Several different structures are used to store offloading information. The |
| :ref:`device image structure <table-device_image_structure>` stores a single |
| linked device image and its associated offloading entries. The offloading |
| entries are stored using the ``__start_omp_offloading_entries`` and |
| ``__stop_omp_offloading_entries`` symbols generated by the linker using the |
| :ref:`table-tgt_offload_entry_structure`. |
| |
| .. table:: __tgt_device_image Structure |
| :name: table-device_image_structure |
| |
| +----------------------+--------------+----------------------------------------+ |
| | Type | Identifier | Description | |
| +======================+==============+========================================+ |
| | void* | ImageStart | Pointer to the target code start | |
| +----------------------+--------------+----------------------------------------+ |
| | void* | ImageEnd | Pointer to the target code end | |
| +----------------------+--------------+----------------------------------------+ |
| | __tgt_offload_entry* | EntriesBegin | Begin of table with all target entries | |
| +----------------------+--------------+----------------------------------------+ |
| | __tgt_offload_entry* | EntriesEnd | End of table (non inclusive) | |
| +----------------------+--------------+----------------------------------------+ |
| |
| The target :ref:`target binary descriptor <table-target_binary_descriptor>` is |
| used to store all binary images and offloading entries in an array. |
| |
| .. table:: __tgt_bin_desc Structure |
| :name: table-target_binary_descriptor |
| |
| +----------------------+------------------+------------------------------------------+ |
| | Type | Identifier | Description | |
| +======================+==================+==========================================+ |
| | int32_t | NumDeviceImages | Number of device types supported | |
| +----------------------+------------------+------------------------------------------+ |
| | __tgt_device_image* | DeviceImages | Array of device images (1 per dev. type) | |
| +----------------------+------------------+------------------------------------------+ |
| | __tgt_offload_entry* | HostEntriesBegin | Begin of table with all host entries | |
| +----------------------+------------------+------------------------------------------+ |
| | __tgt_offload_entry* | HostEntriesEnd | End of table (non inclusive) | |
| +----------------------+------------------+------------------------------------------+ |
| |
| Global Variables |
| ---------------- |
| |
| :ref:`table-global_variables` lists various global variables, along with their |
| type and their explicit ELF sections, which are used to store device images and |
| related symbols. |
| |
| .. table:: Global Variables |
| :name: table-global_variables |
| |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | Variable | Type | ELF Section | Description | |
| +================================+=====================+=========================+=========================================================+ |
| | __start_omp_offloading_entries | __tgt_offload_entry | .omp_offloading_entries | Begin symbol for the offload entries table. | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | __stop_omp_offloading_entries | __tgt_offload_entry | .omp_offloading_entries | End symbol for the offload entries table. | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | __dummy.omp_offloading.entry | __tgt_offload_entry | .omp_offloading_entries | Dummy zero-sized object in the offload entries | |
| | | | | section to force linker to define begin/end | |
| | | | | symbols defined above. | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | .omp_offloading.device_image | __tgt_device_image | .omp_offloading_entries | ELF device code object of the first image. | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | .omp_offloading.device_image.N | __tgt_device_image | .omp_offloading_entries | ELF device code object of the (N+1)th image. | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | .omp_offloading.device_images | __tgt_device_image | .omp_offloading_entries | Array of images. | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| | .omp_offloading.descriptor | __tgt_bin_desc | .omp_offloading_entries | Binary descriptor object (see :ref:`binary_descriptor`) | |
| +--------------------------------+---------------------+-------------------------+---------------------------------------------------------+ |
| |
| .. _binary_descriptor: |
| |
| Binary Descriptor for Device Images |
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| This object is passed to the offloading runtime at program startup and it |
| describes all device images available in the executable or shared library. It |
| is defined as follows: |
| |
| .. code-block:: c |
| |
| __attribute__((visibility("hidden"))) |
| extern __tgt_offload_entry *__start_omp_offloading_entries; |
| __attribute__((visibility("hidden"))) |
| extern __tgt_offload_entry *__stop_omp_offloading_entries; |
| static const char Image0[] = { <Bufs.front() contents> }; |
| ... |
| static const char ImageN[] = { <Bufs.back() contents> }; |
| static const __tgt_device_image Images[] = { |
| { |
| Image0, /*ImageStart*/ |
| Image0 + sizeof(Image0), /*ImageEnd*/ |
| __start_omp_offloading_entries, /*EntriesBegin*/ |
| __stop_omp_offloading_entries /*EntriesEnd*/ |
| }, |
| ... |
| { |
| ImageN, /*ImageStart*/ |
| ImageN + sizeof(ImageN), /*ImageEnd*/ |
| __start_omp_offloading_entries, /*EntriesBegin*/ |
| __stop_omp_offloading_entries /*EntriesEnd*/ |
| } |
| }; |
| static const __tgt_bin_desc BinDesc = { |
| sizeof(Images) / sizeof(Images[0]), /*NumDeviceImages*/ |
| Images, /*DeviceImages*/ |
| __start_omp_offloading_entries, /*HostEntriesBegin*/ |
| __stop_omp_offloading_entries /*HostEntriesEnd*/ |
| }; |
| |
| |
| Global Constructor and Destructor |
| ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ |
| |
| The global constructor (``.omp_offloading.descriptor_reg()``) registers the |
| device images with the runtime by calling the ``__tgt_register_lib()`` runtime |
| function. The constructor is explicitly defined in ``.text.startup`` section and |
| is run once when the program starts. Similarly, the global destructor |
| (``.omp_offloading.descriptor_unreg()``) calls ``__tgt_unregister_lib()`` for |
| the destructor and is also defined in ``.text.startup`` section and run when the |
| program exits. |
| |
| Offloading Example |
| ------------------ |
| |
| This section contains a simple example of generating offloading code using |
| OpenMP offloading. We will use a simple ``ZAXPY`` BLAS routine. |
| |
| .. code-block:: c++ |
| |
| #include <complex> |
| |
| using complex = std::complex<double>; |
| |
| void zaxpy(complex *X, complex *Y, complex D, std::size_t N) { |
| #pragma omp target teams distribute parallel for |
| for (std::size_t i = 0; i < N; ++i) |
| Y[i] = D * X[i] + Y[i]; |
| } |
| |
| int main() { |
| const std::size_t N = 1024; |
| complex X[N], Y[N], D; |
| #pragma omp target data map(to:X[0 : N]) map(tofrom:Y[0 : N]) |
| zaxpy(X, Y, D, N); |
| } |
| |
| This code is compiled using the following Clang flags. |
| |
| .. code-block:: console |
| |
| $ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 zaxpy.cpp -c |
| |
| The output section in the object file can be seen using the ``readelf`` utility. |
| The ``.llvm.offloading`` section has the ``SHF_EXCLUDE`` flag so it will be |
| removed from the final executable or shared library by the linker. |
| |
| .. code-block:: text |
| |
| $ llvm-readelf -WS zaxpy.o |
| Section Headers: |
| [Nr] Name Type Address Off Size ES Flg Lk Inf Al |
| [11] omp_offloading_entries PROGBITS 0000000000000000 0001f0 000040 00 A 0 0 1 |
| [12] .llvm.offloading PROGBITS 0000000000000000 000260 030950 00 E 0 0 8 |
| |
| |
| Compiling this file again will invoke the ``clang-linker-wrapper`` utility to |
| extract and link the device code stored at the section named |
| ``.llvm.offloading`` and then use entries stored in |
| the section named ``omp_offloading_entries`` to create the symbols necessary for |
| ``libomptarget`` to register the device image and call the entry function. |
| |
| .. code-block:: console |
| |
| $ clang++ -fopenmp -fopenmp-targets=nvptx64 zaxpy.o -o zaxpy |
| $ ./zaxpy |
| |
| We can see the steps created by clang to generate the offloading code using the |
| ``-ccc-print-phases`` option in Clang. This matches the description in |
| :ref:`Offloading Overview`. |
| |
| .. code-block:: console |
| |
| $ clang++ -fopenmp -fopenmp-targets=nvptx64 -ccc-print-phases zaxpy.cpp |
| # "x86_64-unknown-linux-gnu" - "clang", inputs: ["zaxpy.cpp"], output: "/tmp/zaxpy-host.bc" |
| # "nvptx64-nvidia-cuda" - "clang", inputs: ["zaxpy.cpp", "/tmp/zaxpy-e6a41b.bc"], output: "/tmp/zaxpy-07f434.s" |
| # "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["/tmp/zaxpy-07f434.s"], output: "/tmp/zaxpy-0af7b7.o" |
| # "x86_64-unknown-linux-gnu" - "clang", inputs: ["/tmp/zaxpy-e6a41b.bc", "/tmp/zaxpy-0af7b7.o"], output: "/tmp/zaxpy-416cad.o" |
| # "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["/tmp/zaxpy-416cad.o"], output: "a.out" |
| |
| Relocatable Linking |
| ------------------- |
| |
| The offloading compilation pipeline normally will defer the final device linking |
| and runtime registration until the ``clang-linker-wrapper`` is run to create the |
| executable. This is the standard behaviour when compiling for OpenMP offloading |
| or CUDA and HIP in ``-fgpu-rdc`` mode. However, there are some cases where the |
| user may wish to perform this device handling prematurely. This is described in |
| the :doc:`linker wrapper documentation<ClangLinkerWrapper>`. |
| |
| Effectively, this allows the user to handle offloading specific linking ahead of |
| time when shipping objects or static libraries. This can be thought of as |
| performing a standard ``-fno-gpu-rdc`` compilation on a subset of object files. |
| This can be useful to reduce link time, prevent users from interacting with the |
| library's device code, or for shipping libraries to incompatible compilers. |
| |
| Normally, if a relocatable link is done using ``clang -r`` it will simply merge |
| the ``.llvm.offloading`` sections which will then be linked later when the |
| executable is created. However, if the ``-r`` flag is used with the offloading |
| toolchain, it will perform the device linking and registration phases and then |
| merge the registration code into the final relocatable object file. |
| |
| The following example shows how using the relocatable link with the offloading |
| pipeline can create a static library with offloading code that can be |
| redistributed without requiring any additional handling. |
| |
| .. code-block:: console |
| |
| $ clang++ -fopenmp -fopenmp-targets=nvptx64 foo.cpp -c |
| $ clang++ -lomptarget.devicertl --offload-link -r foo.o -o merged.o |
| $ llvm-ar rcs libfoo.a merged.o |
| # g++ app.cpp -L. -lfoo |
| |
