This PR adds support in the Level Zero plugin to execute
constructors/destructors on the device code. As spirv-link has some
limitations, it mimics the CUDA plugin behavior where the RTL constructs
the device side tables before invoking the kernel that will execute
them.
The kernel and other necessary symbols to create the device tables are
created by the SPIRVCtorDtorLowering pass to be added in #187509
Additional cleanup improvements on error conditions (in addition to
those in #187597):
* Fixed incomplete cleanup in L0Context::init()
* Fixed build log leak in addModule()
* Fixed context inconsistent state in findDevices()
Disclaimer: The base of this PR was generated by Claude and adjusted by
me afterwards.
Summary:
We create the RPC doorbell signal lazily and destroy it at the plugin
level. This means that we can't rely on the normal 'per-device' handling
so this needs to be called unconditionally. We only create the signal if
a device is registered, but deinit is called unconditionally. Just check
the handle.
Summary:
One of the main disadvantages to using the RPC interface is that it
requires a server thread to spin on the mailboxes checking for work.
The vast majority of the time, there will be no work and work will come
in large bursts.
The HSA / KFD interface supports device-side interrupts and already has
handling for binding these events to an HSA signal. This means that we
can send interrupts from the GPU to wake a sleeping thread on the CPU.
The sleeping thread will be descheduled with a blocking HSA wait call
and woken up when its event ID is raised through the kernel driver's
interrupt.
This is very target-specific handling, but I believe it is valuable
enough to warrant it being in the protocol. It is completely optional,
as it is ignored if uninitialized. This should bring this support at
parity with the interface HIP expects.
Summary:
This PR changes the handling of the emitted kernels when targeting a CPU
to be a pointer struct.
The old handling emitted a standard function prototype, this
necessitated a target specific ABI to call it because the signature
differed with the number of arguments. Instead, this PR emits a void
pointer to a naturally aligned struct, this is what APIs like `pthreads`
assert.
This allows us to remove all the complexity around launching host
kernels and just pass the argument list.
Summary:
The changes in https://www.github.com/llvm/llvm-project/pull/185552
allowed us to
start building the standard `libclang_rt.profile.a` for GPU targets.
This PR expands this by adding an optimized GPU routine for counter
increment and removing the special-case handling of these functions in
the OpenMP runtime.
Vast majority of these functions are boilerplate, but we should be able
to do more interesting things with this in the future, like value or
memory profiling.
Summary:
Follow up on removal of OPENMP_STANDALONE_BUILD in openmp (#149878).
This
build method is redundant and can be accomplished via runtimes.
Removes support for:
`cmake -S <llvm-project>/offload ...`
Switches over to:
`make -S <llvm-project>/runtimes -DLLVM_ENABLE_RUNTIMES=openmp;offload
...`
Libomptarget has a dependency on libomp.so and requires the omp cmake
target to exist at build time, which is why both runtimes are listed.
Updates cmake compiler logic in offload/CMakeLists.txt to mirror openmp
changes:
[openmp] Allow testing OpenMP without a full clang build tree (#182470)
User will still need to have a separate invocation to build openmp
DeviceRTL via:
`-DLLVM_ENABLE_RUNTIMES=openmp`
`-DLLVM_DEFAULT_TARGET_TRIPLE=<amdgcn-amd-amdhsa|nvptx64-nvidia-cuda>`
In the pre-merge CI we need a top-level visible target that can be used
to build offload, i.e., libomptarget and LLVMOffload.
The related PR to include offload into pre-merge CI is here:
https://github.com/llvm/llvm-project/pull/174955
This change improves handling of errors during synchronization in Level
Zero plugin by ensuring cleanup of queues and events in case of an
synchronization error. As a result multiple tests stopped hanging.
---------
Co-authored-by: Duran, Alex <alejandro.duran@intel.com>
Queue's WaitEvent collection wasn't being cleared after synchronization
and resetting of the events. This led to hangs on subsequent host
synchronizations if not preceeded by any other operation.
This commit removes the `LIBOMPTARGET_SHARED_MEMORY_SIZE` envar and
outputs a runtime warning if it is defined. Access to dynamic shared memory
should be obtained through the `dyn_groupprivate` clause (OpenMP 6.1) or
the launch arguments in liboffload kernel launch.
Summary:
We use this `dyn_ptr` argument in Clang/OpenMP to handle the
`KernelLaunchEnvironment`. This is a per-kernel argument used to share
some information. Currenetly, it's prepended to the argument list and we
generate storage for it in the runtime.
This is bad for a few reasons:
1. It changes the ABI by shifting user arguments
2. It cannot be trivially be left uninitialized if unused
3. The runtime must allocate its own memory for it
This PR changes it to be appended instead. Additionally, space for this
is always emitted. This means the OMPIRBuilder itself will provide the
storage, we simply need to populate it in the runtime if it is used.
This means that if it's unused we don't always pay the cost and it's
easier for non-OpenMP users to ignore it.
Backward compatibility is maintained by auto-upgrading the kernel
arguments. In `libomptarget` we completely allocate a new buffer to
store this in the new format. The plugins still need to respect the old
ABI of the called device object, so we simply rotate it if it's the old
version.
Converting the current bot config to use the CMake cache file that we
use in other bots (offload/cmake/caches/AMDGPUBot.cmake). This PR
removes all CMake settings that the cache file already sets and only
leaves those that were either not set explicitly or which differ.
Thus, first load the cache file and then adjust the settings to override
existing values.
### Issue: Dimension override missing
When variable count expressions were used with stride, the constant
subsection path computed size first. This marked `ArgSizes` with byte
size semantics. Variable expression logic later triggered, but reused
`ArgSizes` assuming "bytes" semantics
`OMPIRBuilder.cpp` didn't handle dimension count for
`OMP_MAP_NON_CONTIG` flag
**Result**: `ArgSizes` wasn't overwritten with dimension count, breaking
non-contiguous mapping.
**Fixes**:
`llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp` - Expression semantics for
non-contiguous.
stride/count.
Generate 3D descriptor structures with runtime dimensions.
Fix dimension override to use dimension count instead of byte size.
Added testcases to cover stack arrays, heap pointers, struct members,
etc.
- Accept OffloadBinaries as valid images by plugins that support them in
the PluginInterface.
- Add support in L0 plugin to extract SPIRV images and their associated
metadata from an OffloadBinary image.
Depends on:
- #185663
Follow-up PRs:
- #185413 (Changes SPIRV wrapper generation to use OffloadBinary)
- #185425 (Adjusts llvm-objdump)
- #184774 (Adjusts llvm-offload-binary)
Summary:
As suggested in https://github.com/llvm/llvm-project/pull/177665, we
should build a GPU version of the compiler-rt profile library instead of
writing it in-line in the lowering. This PR does not define anything GPU
specific, it simply re-uses the baremetal handling. Later PRs will
prevent the GPU specific handling we would want to do to optimize
counter handling on the GPU.
Note that this will require using the cache file, or setting these
options
manually for existing users. Hopefully if people are using the cache
file
as they should it won't break anything.
As discussed in #185404 we might want to provide a way for plugins to
validate images not recognized by the common layer.
This PR adds such extension and uses it to validate pure SPIRV images by
the Level Zero plugin.
Summary:
This was a regression from the original LLVM-gpu-loader. We used to
handle `-mwavefrontsize64` correctly in the loader by over-allocating
memory and just leaving the upper 32-bits masked off. In order to handle
this in offload we need to scan loaded kernels to see how much memory we
need to allocate. This should be safe, the protocol is designed to
handle an arbitrary size and worst-case this just wastes space.
This fixes a bug in USM mode where the `close` map type modifer was
attached to some `map.info.op`'s corresponding to user-defined type
members while the parent type instance itself is not marked as `close`.
This fix ensures that if a parent record type map does not have the
'close' flag, it is cleared from its members as well, maintaining
consistency.
Gemini was used to create tests. AI generated test code was reviewed
line-by-line by me. Which were derived from a reproducer I was working
with to debug the issue.
Assisted-by: Gemini <gemini@google.com>
This patch implements the CodeGen logic for calling __llvm_omp_indirect_call_lookup
on the device when an indirect function call or a virtual function call is made
within an OpenMP target region.
---------
Co-authored-by: Youngsuk Kim
This patch implements the CodeGen logic for calling __llvm_omp_indirect_call_lookup
on the device when an indirect function call or a virtual function call is made
within an OpenMP target region.
---------
Co-authored-by: Youngsuk Kim
Some tests that were checking for prints inside/outside `target` regions
needed to be updated to work on systems where the ordering wasn't
deterministic.
Reverts llvm/llvm-project#184240
Original description from #165494:
-----
OpenMP allows cases like the following:
```c
int *p1, *p2, x;
p1 = p2 = &x;
...
#pragma omp target_exit_data map(delete: p1[:]) from(p2[0])
```
Which means, when the runtime encounters the `from` entry, the ref-count
may
not be zero, but it will go down to zero at the end of the current
construct,
which should cause the "from" transfer to happen.
Similarly, a user may have:
```c
struct S {
int *p;
};
#pragma omp declare_mapper (id1: S s) map(s.p) map(present, alloc: s.p[0:10])
#pragma omp declare_mapper (id2: S s) map(s.p, s.p[0:10])
S s1;
// present-check should fail here
#pragma omp target_enter_data map(alloc: s.p[0:10]) map(mapper(id1), to: s)
// "to" should be honored here
#pragma omp target_enter_data map(alloc: s.p[0:10]) map(mapper(id2), to: s)
```
Where the allocation happens before the "to" entry is encountered by the
runtime. Or, an allocation happens before a "present" entry is
encountered.
To handle cases like this, we need to use the state information of
previously
seen new allocations, deletions, "from" entries, when honoring
`to`/`from`/`present` map entries.
-----
OpenMP allows cases like the following:
```c
int *p1, *p2, x;
p1 = p2 = &x;
...
#pragma omp target_exit_data map(delete: p1[:]) from(p2[0])
```
Which means, when the runtime encounters the `from` entry, the ref-count
may not be zero, but it will go down to zero at the end of the current
construct, which should cause the "from" transfer to happen.
Similarly, a user may have:
```c
struct S {
int *p;
};
#pragma omp declare_mapper (id1: S s) map(s.p) map(present, alloc: s.p[0:10])
#pragma omp declare_mapper (id2: S s) map(s.p, s.p[0:10])
S s1;
// present-check should fail here
#pragma omp target_enter_data map(alloc: s.p[0:10]) map(mapper(id1), to: s)
// "to" should be honored here
#pragma omp target_enter_data map(alloc: s.p[0:10]) map(mapper(id2), to: s)
```
Where the allocation happens before the "to" entry is encountered by the
runtime. Or, an allocation happens before a "present" entry is
encountered.
To handle cases like this, we need to use the state information of
previously seen new allocations, deletions, "from" entries, when
honoring `to`/`from`/`present` map entries.
Implement support for the OpenMP `mapper` modifier on `target update` `to` and
`from` clauses in Flang.
Semantic name resolution is extended to bind the mapper symbol for
`OmpClause::To` and `OmpClause::From` via a shared `ResolveMapperModifier`
helper. Lowering is extended in `ClauseProcessor` with a `getMapperIdentifier`
template helper to extract the mapper name for both `map` and `target update`
clauses and forward it to `omp.map_info`.
Fixes#168701.
Reviewed By: TIFitis, kparzysz
Assited By: Copilot( For review and articulations of messages)
Summary:
This enables primarily `stop.cpp` and `descriptor.cpp`. Requires a
little bit of wrangling to get it to compile. Unlike the CUDA build,
this build uses an in-tree libc++ configured for the GPU. This is
configured without thread support, environment, or filesystem, and it is
not POSIX at all. So, no mutexes, pthreads, or get/setenv.
I tested stop, but i don't know if it's actually legal to exit from
OpenMP offloading.
`spirv-link` seems to internalize all symbols, which ends up causing the
OpenMP Device Environment global generated by the OMP FE to get
optimized out which causes `liboffload` to run in the wrong
parallelization mode which breaks at least one liboffload lit test.
Pass `--create-library` to tell it not to do that.
```
--create-library
Link the binaries into a library, keeping all exported symbols.
```
This fixes the test.
Closes: https://github.com/llvm/llvm-project/issues/182901
Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
Summary:
Expands on the previous support to enable formatted output, characters,
and checking basic iostat. We intentionally do not handle cases where
the descriptor is non-null as this is a non-trivial class that cannot
easily be shepherded across the wire.
Summary:
This PR provides the minimal support for Fortran I/O coming from a GPU
in OpenMP offloading. We use the same support the `libc` uses for its
printing through the RPC server. The helper functions `rpc::dispatch`
and `rpc::invoke` help make this mostly automatic.
Becaus Fortran I/O is not reentrant, the vast majority of complexity
comes from needing to stitch together calls from the GPU until they can
be executed all at once. This is needed not only because of the
limitations of recursive I/O, but without this the output would all be
interleaved because of the GPU's lock-step execution.
As such, the return values from the intermediate functions are
meaningless, all returning true. The final value is correct however. For
cookies we create a context pointer on the server to chain these
together.
Works on both my AMD and NVIDIA GPUs.
```fortran
program hello_gpu
implicit none
!$omp target teams num_teams(1)
!$omp parallel num_threads(2)
! Print strings
print *, "Hello from GPU"
!$omp end parallel
!$omp end target teams
end program hello_gpu
```
```console
> flang hello.f90 -O2 -fopenmp --offload-arch=gfx1030
> ./a.out
Hello from GPU
Hello from GPU
> flang hello.f90 -O2 -fopenmp --offload-arch=sm_89
> ./a.out
Hello from GPU
Hello from GPU
```
