mirrors/llvm-project
0
0
Fork 0
mirror of https://github.com/llvm/llvm-project.git synced 2025-04-16 03:16:54 +00:00
Code Issues Projects Releases Wiki Activity
llvm-project/offload/libomptarget/exports

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

79 lines
2.2 KiB
Plaintext
Raw Normal View History

[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
VERS1.0 {
global:
[Libomptarget] Remove global ctor and use reference counting (#80499) Summary: Currently we rely on global constructors to initialize and shut down the OpenMP runtime library and plugin manager. This causes some issues because we do not have a defined lifetime that we can rely on to release and allocate resources. This patch instead adds some simple reference counted initialization and deinitialization function. A future patch will use the `deinit` interface to more intelligently handle plugin deinitilization. Right now we do nothing and rely on `atexit` inside of the plugins to tear them down. This isn't great because it limits our ability to control these things. Note that I made the `__tgt_register_lib` functions do the initialization instead of adding calls to the new runtime functions in the linker wrapper. The reason for this is because in the past it's been easier to not introduce a new function call, since sometimes the user's compiler will link against an older `libomptarget`. Maybe if we change the name with offloading in the future we can simplify this. Depends on https://github.com/llvm/llvm-project/pull/80460
2024-02-22 12:01:52 -06:00
__tgt_rtl_init;
__tgt_rtl_deinit;
[OpenMP][libomptarget] Enable requires flags for target libraries. Summary: Target link variables are currently implemented by creating a copy of the variables on the device side and unified memory never gets exploited. When the prgram uses the: ``` #pragma omp requires unified_shared_memory ``` directive in conjunction with a declare target link, the linked variable is no longer allocated on the device and the host version is used instead. This behavior is overridden by performing an explicit mapping. A Clang side patch is required. Reviewers: ABataev, AlexEichenberger, grokos, Hahnfeld Reviewed By: AlexEichenberger, grokos, Hahnfeld Subscribers: Hahnfeld, jfb, guansong, jdoerfert, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D60223 llvm-svn: 361294
2019-05-21 19:35:02 +00:00
__tgt_register_requires;
[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
__tgt_register_lib;
__tgt_unregister_lib;
[OpenMP][Tool] Introducing the `llvm-omp-device-info` tool This patch introduces the `llvm-omp-device-info` tool, which uses the omptarget library and interface to query the device info from all the available devices as seen by OpenMP. This is inspired by PGI's `pgaccelinfo` Since omptarget usually requires a description structure with executable kernels, I split the initialization of the RTLs and Devices to be able to initialize all possible devices and query each of them. This revision relies on the patch that introduces the print device info. A limitation is that the order in which the devices are initialized, and the corresponding device ID is not necesarily the one seen by OpenMP. The changes are as follows: 1. Separate the RTL initialization that was performed in `RegisterLib` to its own `initRTLonce` function 2. Create an `initAllRTLs` method that initializes all available RTLs at runtime 3. Created the `llvm-deviceinfo.cpp` tool that uses `omptarget` to query each device and prints its information. Example Output: ``` Device (0): print_device_info not implemented Device (1): print_device_info not implemented Device (2): print_device_info not implemented Device (3): print_device_info not implemented Device (4): CUDA Driver Version: 11000 CUDA Device Number: 0 Device Name: Quadro P1000 Global Memory Size: 4236312576 bytes Number of Multiprocessors: 5 Concurrent Copy and Execution: Yes Total Constant Memory: 65536 bytes Max Shared Memory per Block: 49152 bytes Registers per Block: 65536 Warp Size: 32 Threads Maximum Threads per Block: 1024 Maximum Block Dimensions: 1024, 1024, 64 Maximum Grid Dimensions: 2147483647 x 65535 x 65535 Maximum Memory Pitch: 2147483647 bytes Texture Alignment: 512 bytes Clock Rate: 1480500 kHz Execution Timeout: Yes Integrated Device: No Can Map Host Memory: Yes Compute Mode: DEFAULT Concurrent Kernels: Yes ECC Enabled: No Memory Clock Rate: 2505000 kHz Memory Bus Width: 128 bits L2 Cache Size: 1048576 bytes Max Threads Per SMP: 2048 Async Engines: Yes (2) Unified Addressing: Yes Managed Memory: Yes Concurrent Managed Memory: Yes Preemption Supported: Yes Cooperative Launch: Yes Multi-Device Boars: No Compute Capabilities: 61 ``` Reviewed By: tianshilei1992 Differential Revision: https://reviews.llvm.org/D106752
2021-07-27 22:38:27 -04:00
__tgt_init_all_rtls;
[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
__tgt_target_data_begin;
__tgt_target_data_end;
__tgt_target_data_update;
__tgt_target;
__tgt_target_teams;
__tgt_target_data_begin_nowait;
__tgt_target_data_end_nowait;
__tgt_target_data_update_nowait;
__tgt_target_nowait;
__tgt_target_teams_nowait;
[OpenMP][Offload] Declare mapper runtime implementation Libomptarget patch adding runtime support for "declare mapper". Patch co-developed by Lingda Li and George Rokos. Differential revision: https://reviews.llvm.org/D68100
2020-07-15 13:24:03 -07:00
__tgt_target_data_begin_mapper;
__tgt_target_data_end_mapper;
__tgt_target_data_update_mapper;
__tgt_target_mapper;
__tgt_target_teams_mapper;
__tgt_target_data_begin_nowait_mapper;
__tgt_target_data_end_nowait_mapper;
__tgt_target_data_update_nowait_mapper;
__tgt_target_nowait_mapper;
__tgt_target_teams_nowait_mapper;
[Libomptarget] Implement a unified kernel entry function This patch implements a unified kernel entry function that will be targeted from both teams and non-teams clauses. We introduce a new interface and make the old functions call in using the new one. A following patch will include the necessary changes to Clang to call these new functions instead. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D128549
2022-06-23 14:57:59 -04:00
__tgt_target_kernel;
__tgt_target_kernel_nowait;
[OpenMP] Add non-blocking support for target nowait regions This patch better integrates the target nowait functions with the tasking runtime. It splits the nowait execution into two stages: a dispatch stage, which triggers all the necessary asynchronous device operations and stores a set of post-processing procedures that must be executed after said ops; and a synchronization stage, responsible for synchronizing the previous operations in a non-blocking manner and running the appropriate post-processing functions. Suppose during the synchronization stage the operations are not completed. In that case, the attached hidden helper task is re-enqueued to any hidden helper thread to be later synchronized, allowing other target nowait regions to be concurrently dispatched. Reviewed By: jdoerfert, tianshilei1992 Differential Revision: https://reviews.llvm.org/D132005
2022-12-14 13:46:23 -03:00
__tgt_target_nowait_query;
[OpenMP] Support kernel record and replay This patch adds functionality for recording and replaying the execution of OpenMP offload kernels, based on an original implementation by Steve Rangel. The patch extends libomptarget to extract a json description of the kernel, the device image binary, and a device memory snapshot before and after the execution of a recorded kernel. Kernel recording/replaying in libomptarget is controlled through env vars (LIBOMPTARGET_RECORD, LIBOMPTARGET_REPLAY). It provides a tool, llvm-omp-kernel-replay, for replaying a kernel using the extracted information with the ability to verify replayed execution using the post-execution device memory snapshot, also supporting changing the number of teams/threads for replaying. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D138931
2023-01-17 15:35:44 -08:00
__tgt_target_kernel_replay;
[OpenMP] Support for OpenMP-Offload Record Replay Enable record-replay for OpenMP offload kernels. On recording the initialization is performed on device initialization by reading env variables. (This is similar to the way rr used to operate). The primary change takes place in the replay phase with the replay tool explicitly initializing the record-replay functionality. Differential Revision: https://reviews.llvm.org/D156174 Fix
2023-07-19 10:32:07 -07:00
__tgt_activate_record_replay;
[OpenMP 5.0] libomptarget interface for declare mapper functions. This patch implements the libomptarget runtime interface for OpenMP 5.0 declare mapper functions. The declare mapper functions generated by Clang will call them to complete the mapping of members. kmpc_mapper_num_components gets the current number of components for a user-defined mapper; kmpc_push_mapper_component pushes back one component for a user-defined mapper. The design slides can be found at https://github.com/lingda-li/public-sharing/blob/master/mapper_runtime_design.pptx Patch by Lingda Li <lildmh@gmail.com> Differential Revision: https://reviews.llvm.org/D60972 llvm-svn: 367772
2019-08-04 04:18:28 +00:00
__tgt_mapper_num_components;
__tgt_push_mapper_component;
[OpenMP] Restore backwards compatibility for libomptarget Summary: The changes introduced in D87946 changed the API for libomptarget functions. `__kmpc_push_target_tripcount` was a function in Clang 11.x but was not given a backward-compatible interface. This change will require people using Clang 13.x or 12.x to recompile their offloading programs. Reviewed By: jdoerfert cchen Differential Revision: https://reviews.llvm.org/D98358
2021-03-10 13:25:33 -05:00
__kmpc_push_target_tripcount;
__kmpc_push_target_tripcount_mapper;
[OpenMP] Add OpenMP extension API to dump mapping tables (#85381) This adds an API call ompx_dump_mapping_tables. This allows users to debug the mapping tables and can be especially useful for unified shared memory applications to check if the code behaves in the way it should. The implementation reuses code already present to dump mapping tables (in a debug setting). --------- Co-authored-by: Joseph Huber <huberjn@outlook.com>
2024-03-18 20:09:20 +01:00
ompx_dump_mapping_tables;
[OpenMP] Implement `omp_get_mapped_ptr` This patch implements the function `omp_get_mapped_ptr`. Fix #59945. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D141545
2023-01-11 22:05:33 -05:00
omp_get_mapped_ptr;
[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
omp_get_num_devices;
[OpenMP] Implementing omp_get_device_num() This patch implements omp_get_device_num() in the host and the device. It uses the already existing getDeviceNum in the device config for the device. And in the host it uses the omp_get_num_devices(). Two simple tests added Differential Revision: https://reviews.llvm.org/D128347
2022-06-22 10:05:34 -05:00
omp_get_device_num;
[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
omp_get_initial_device;
omp_target_alloc;
omp_target_free;
omp_target_is_present;
omp_target_memcpy;
omp_target_memcpy_rect;
[OpenMP] asynchronous memory copy support We introduced the implementation of supporting asynchronous routines with depend objects specified in Version 5.1 of the OpenMP Application Programming Interface. In brief, these routines omp_target_memcpy_async and omp_target_memcpy_rect_async perform asynchronous (nonblocking) memory copies between any combination of host and device pointers. The basic idea is to create the implicit tasks to carry the memory copy calls and handle dependencies specified by depend objects. The implicit tasks are executed via hidden helper thread in OpenMP runtime. Reviewed By: jdoerfert, tianshilei1992 Committed By: jplehr Differential Revision: https://reviews.llvm.org/D136103
2023-03-28 10:32:11 -04:00
omp_target_memcpy_async;
omp_target_memcpy_rect_async;
[libomptarget][OpenMP] Initial implementation of omp_target_memset() and omp_target_memset_async() (#68706) Implement a slow-path version of omp_target_memset*() There is a TODO to implement a fast path that uses an on-device kernel instead of the host-based memory fill operation. This may require some additional plumbing to have kernels in libomptarget.so
2023-10-19 15:29:36 +02:00
omp_target_memset;
omp_target_memset_async;
[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
omp_target_associate_ptr;
omp_target_disassociate_ptr;
[libomptarget] Add allocator support for target memory This patch adds the infrastructure for allocator support for target memory. Three allocators are introduced for device, host and shared memory. The corresponding API functions have the llvm_ prefix temporarily, until they become part of the OpenMP standard. Differential Revision: https://reviews.llvm.org/D97883
2021-03-03 11:48:32 -08:00
llvm_omp_target_alloc_host;
llvm_omp_target_alloc_shared;
llvm_omp_target_alloc_device;
[Libomptarget] Change device free routines to accept the allocation kind Previous support for device memory allocators used a single free routine and did not provide the original kind of the allocation. This is problematic as some of these memory types required different handling. Previously this was worked around using a map in runtime to record the original kind of each pointer. Instead, this patch introduces new free routines similar to the existing allocation routines. This allows us to avoid a map traversal every time we free a device pointer. The only interfaces defined by the standard are `omp_target_alloc` and `omp_target_free`, these do not take a kind as `omp_alloc` does. The standard dictates the following: "The omp_target_alloc routine returns a device pointer that references the device address of a storage location of size bytes. The storage location is dynamically allocated in the device data environment of the device specified by device_num." Which suggests that these routines only allocate the default device memory for the kind. So this has been changed to reflect this. This change is somewhat breaking if users were using `omp_target_free` as previously shown in the tests. Reviewed By: JonChesterfield, tianshilei1992 Differential Revision: https://reviews.llvm.org/D133053
2022-08-31 15:55:14 -05:00
llvm_omp_target_free_host;
llvm_omp_target_free_shared;
llvm_omp_target_free_device;
[OpenMP] Add dynamic memory function to omp.h and add documentation This patch adds the `llvm_omp_target_dynamic_shared_alloc` function to the `omp.h` header file so users can access it by default. Also changed the name to keep it consistent with the other target allocators. Added some documentation so users know how to use it. Didn't add the interface for Fortran since there's no way to test it right now. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D123246
2022-04-06 15:56:19 -04:00
llvm_omp_target_dynamic_shared_alloc;
[OpenMP][libomptarget][AMDGPU] lock/unlock (pin/unpin) mechanism in libomptarget amdgpu plugin (API and implementation) The current only way to obtain pinned memory with libomptarget is to use a custom allocator llvm_omp_target_alloc_host. This reflects well the CUDA implementation of libomptarget, but it does not correctly expose the AMDGPU runtime API, where any system allocated page can be locked/unlocked through a call to hsa_amd_memory_lock/unlock. This patch enables users to allocate memory through malloc (mmap, sbreak) and then pin the related memory pages with a libomptarget special call. It is a base support in the amdgpu libomptarget plugin to enable users to prelock their host memory pages so that the runtime doesn't need to lock them itself for asynchronous memory transfers. Reviewed By: jdoerfert, ye-luo Differential Revision: https://reviews.llvm.org/D139208
2023-01-13 12:18:49 -06:00
llvm_omp_target_lock_mem;
llvm_omp_target_unlock_mem;
[OpenMP] Add function for setting LIBOMPTARGET_INFO at runtime Summary: This patch adds a new runtime function __tgt_set_info_flag that allows the user to set the information level at runtime without using the environment variable. Using this will require an extern function, but will eventually be added into an auxilliary library for OpenMP support functions. This patch required moving the current InfoLevel to a global variable which must be instantiated by each plugin. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D100774
2021-04-21 17:31:09 -04:00
__tgt_set_info_flag;
[OpenMP][Libomptarget] Adding `print_device_info` to RTL and `omptarget` This patch introduces a function in the device's plugin to print the device information. This patch relates to another patch that introduces a CLI tool to obtain the device information from the omplibrary directly. It is inspired by PGI's pgaccelinfo. The modifications are as follows: 1. Introduce the optional `void __tgt_rtl_print_device_info(RTLdevID)` function into the RTL. 2. Introduce the `bool __tgt_print_device_info(devID)` function into `omptarget` interface. Returns false if the RTL is not implemented 3. Added `bool printDeviceInfo(RTLDevID)` to the `DeviceTy` 4. Implement the `__tgt_rtl_print_device_info` for CUDA. Added additional CUDA Runtime calls. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D106751
2021-07-27 21:47:40 -04:00
__tgt_print_device_info;
Runtime for Interop directive This implements the runtime portion of the interop directive. It expects the frontend and IRBuilder portions to be in place for proper execution. It currently works only for GPUs and has several TODOs that should be addressed going forward. Reviewed By: RaviNarayanaswamy Differential Revision: https://reviews.llvm.org/D106674
2022-01-27 15:15:25 -05:00
omp_get_interop_ptr;
omp_get_interop_str;
omp_get_interop_int;
omp_get_interop_name;
omp_get_interop_type_desc;
__tgt_interop_init;
__tgt_interop_use;
__tgt_interop_destroy;
[Offload][CUDA] Allow CUDA kernels to use LLVM/Offload (#94549) Through the new `-foffload-via-llvm` flag, CUDA kernels can now be lowered to the LLVM/Offload API. On the Clang side, this is simply done by using the OpenMP offload toolchain and emitting calls to `llvm*` functions to orchestrate the kernel launch rather than `cuda*` functions. These `llvm*` functions are implemented on top of the existing LLVM/Offload API. As we are about to redefine the Offload API, this wil help us in the design process as a second offload language. We do not support any CUDA APIs yet, however, we could: https://www.osti.gov/servlets/purl/1892137 For proper host execution we need to resurrect/rebase https://tianshilei.me/wp-content/uploads/2021/12/llpp-2021.pdf (which was designed for debugging). ``` ❯❯❯ cat test.cu extern "C" { void *llvm_omp_target_alloc_shared(size_t Size, int DeviceNum); void llvm_omp_target_free_shared(void *DevicePtr, int DeviceNum); } __global__ void square(int *A) { *A = 42; } int main(int argc, char **argv) { int DevNo = 0; int *Ptr = reinterpret_cast<int *>(llvm_omp_target_alloc_shared(4, DevNo)); *Ptr = 7; printf("Ptr %p, *Ptr %i\n", Ptr, *Ptr); square<<<1, 1>>>(Ptr); printf("Ptr %p, *Ptr %i\n", Ptr, *Ptr); llvm_omp_target_free_shared(Ptr, DevNo); } ❯❯❯ clang++ test.cu -O3 -o test123 -foffload-via-llvm --offload-arch=native ❯❯❯ llvm-objdump --offloading test123 test123: file format elf64-x86-64 OFFLOADING IMAGE [0]: kind elf arch gfx90a triple amdgcn-amd-amdhsa producer openmp ❯❯❯ LIBOMPTARGET_INFO=16 ./test123 Ptr 0x155448ac8000, *Ptr 7 Ptr 0x155448ac8000, *Ptr 42 ```
2024-08-12 17:44:58 -07:00
__llvmPushCallConfiguration;
__llvmPopCallConfiguration;
llvmLaunchKernel;
[OpenMP] Initial implementation of OpenMP offloading library - libomptarget. This is the patch upstreaming the device-agnostic part of libomptarget. Differential Revision: https://reviews.llvm.org/D14031 llvm-svn: 293094
2017-01-25 21:27:24 +00:00
local:
*;
};
Reference in New Issue Copy Permalink