| Age | Commit message (Collapse) | Author | Files | Lines |
|---|
|
Fix test initialization
|
|
This reverts commit d41368134478d1d41726aa85ba82f49b5bce130c.
|
|
Fix test array initialization.
|
|
Just add
// REQUIRES: libomptarget-debug
So that test will not run with release compiler.
|
|
Currently we are missing set up-boundary address for FinalArraySection
as highests elements in partial struct data.
Currently for:
\#pragma omp target map(D.a) map(D.b[:2])
The size is:
%a = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 0
%b = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 1
%arrayidx = getelementptr inbounds [2 x float], ptr %b, i64 0, i64 0
%2 = getelementptr float, ptr %arrayidx, i32 1
%3 = ptrtoint ptr %2 to i64
%4 = ptrtoint ptr %a to i64
%5 = sub i64 %3, %4
%6 = sdiv exact i64 %5, ptrtoint (ptr getelementptr (i8, ptr null, i32
1) to i64)
Where %2 is wrong for (D.b[:2]) is pointer to first element of array
section. It should pointe to last element of array section.
The fix is to emit the pointer to the last element of array section and
use this pointer as the highest element in partial struct data.
After change IR:
%a = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 0
%b = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 1
%arrayidx = getelementptr inbounds [2 x float], ptr %b, i64 0, i64 0
%b1 = getelementptr inbounds %struct.DataTy, ptr %D, i32 0, i32 1
%arrayidx2 = getelementptr inbounds [2 x float], ptr %b1, i64 0, i64 1
%1 = getelementptr float, ptr %arrayidx2, i32 1
%2 = ptrtoint ptr %1 to i64
%3 = ptrtoint ptr %a to i64
%4 = sub i64 %2, %3
%5 = sdiv exact i64 %4, ptrtoint (ptr getelementptr (i8, ptr null, i32
1) to i64)
|
|
Summary:
This was added in a previous patch to update how we export the static
library used for OpenMP offloading. By mistake this if-else was using
the output incorrectly.
Fixes https://github.com/llvm/llvm-project/issues/74079
|
|
Reverts llvm/llvm-project#74360
As I wrote in the analysis of #74360:
Since
https://github.com/llvm/llvm-project/commit/bc4e0c048aa3cd940b0cea787014c7e8680e5040
we will not add PluginAdaptors into the container of all plugin adaptors
before the plugin is not ready. The error is thereby gone. When and old
HSA loads other libraries they can call register_image but that will
simply not register the image with the plugin we are currently
initializing. That seems like reasonable behavior, thought it is good to
keep in mind if we ever want a kernel library (@jhuber6 @mjklemm). We
can still have a standalone kernel library though or load it late after
all plugins are setup (which seems reasonable).
I did not expect one our tests actually doing exactly what this will not
allow anymore, at least when you use rocm <5.5.0. Need to figure out if
we want this behavior (for rocm <5.5.0).
|
|
|
|
Before we expected all symbols in the device image to be backed up with
data that we could read. However, uninitialized values are not. We now
check for this case and avoid reading random memory.
This also replaces the correct readGlobalFromImage call with a
isSymbolInImage check after
https://github.com/llvm/llvm-project/pull/74550 picked the wrong one.
Fixes: https://github.com/llvm/llvm-project/issues/74582
|
|
We probably should use a temporary name, but having stable names helps
debugging.
|
|
Remove `DelayedBinDesc` as it is not necessary since
https://github.com/llvm/llvm-project/commit/bc4e0c048aa3cd940b0cea787014c7e8680e5040.
See
https://github.com/llvm/llvm-project/pull/74360#issuecomment-1843603736
for details.
|
|
By default we now only look for the plugins we build, but the user can
overwrite that with `LIBOMPTARGET_PLUGINS_TO_LOAD="cuda,amdgpu,x86_64"`
|
|
Summary:
There are now a few cases that check if a symbol is present before
continuing, effectively making them optional features if present in the
image. This was done in at least three locations and required an ugly
operation to consume the error. This patch makes a utility function to
handle that instead.
|
|
Unblock build bot, while investigating. Issue is tracked under llvm
https://github.com/llvm/llvm-project/issues/74582
|
|
|
|
This introduces checked errors into the creation and initialization of
`PluginAdaptorTy`. We also allow the adaptor to "hide" devices from the
user if the initialization failed. The new organization avoids the
"initOnce" stuff but we still do not eagerly initialize the plugin
devices (I think we should merge `PluginAdaptorTy::initDevices` into
`PluginAdaptorTy::init`)
|
|
This fixes two bugs and adds a test for them:
- A shared library with declare target functions but without kernels
should not error out due to missing globals.
- Enabling LIBOMPTARGET_INFO=32 should not deadlock in the presence of
indirect declare targets.
|
|
We accessed the `Devices` container most of the time while holding the
RTLsMtx, but not always. Sometimes we used the mutex for the size query,
but then accessed Devices again unguarded. From now we properly
encapsulate the container in a ProtectedObj which ensures exclusive
accesses. We also hide the "isReady" part in the `getDevice` accessor
and use an `llvm::Expected` to allow to return errors.
|
|
- `nothing` directive was effecting the `if` block structure which it
should not. So return an empty statement instead of an error statement
while parsing to avoid this.
|
|
|
|
|
|
This moves the offload entry logic into classes and provides convenient
accessors. No functional change intended but we can now print all
offload entries (and later look them up), tested via
`OMPTARGET_DUMP_OFFLOAD_ENTRIES=<device_no>`.
|
|
|
|
This is not completely NFC since we now check all 4 requirements and the
test is checking the good and the bad case for combining flags.
|
|
|
|
|
|
OpenMP allows 3 different offload policies, handling of which we want to
encapsulate.
|
|
This basically moves code around again, but this time to provide cleaner
interfaces and remove duplication. PluginAdaptorManagerTy is almost all
gone after this.
|
|
This patch addresses the concern about multiple devices and also adds
more tests for `map(to:)`, `map(from:)` and named common blocks.
|
|
The plugin checks the values of HSA_AMD_INTERFACE_VERSION_* so we now
set them to something safe in the header.
|
|
Libomptarget cannot be build because of the recent refactoring
introduced in patch 148dec9fa43b :
[OpenMP][NFC] Separate Envar (environment variable) handling (#73994)
That patch moved handling of environment variables from libomptarget
library. That's why we don't need usage of "llvm::omp::target" namespace
if we handle environment variables.
|
|
|
|
|
|
This simply puts the profiling logic into the `Profiler` class and
allows non-RAII profiling via `beginSection` and `endSection`.
|
|
|
|
|
|
Summary:
Currently, the `libomp.so` and `libomptarget.so` are emitted in the
`./lib` build directory generally. This logic is internal to the
`add_llvm_library` function we use to build `libomptarget`. The
DeviceRTl static library however is in the middle of the OpenMP runtime
build, which can vary depending on if this is a runtimes or projects
build. This patch changes this to install the DeviceRTL static library
alongside the other OpenMP libraries so they are easier to find.
|
|
|
|
This reverts commit 97e16da450e94c92456fa5a74768ec1b22fe6b63 because it
causes build error on i386 system.
|
|
Summary:
We use `size_t` internally in the omp.h header, which is normally
provided by `stdlib.h` which is already included. Howevever, some cases
when using `-ffreestanding` can result in this not being defined via
`stdlib.h`. This patch simply adds an explicit inclusion of this header,
which is provided by the `clang` resource directory, to resolve this in
all cases.
|
|
|
|
If we don't have a team reduction we don't need a kernel launch
environment (for now). In that case we can avoid the cost.
|
|
While this does not really encapsulate the mapping code, it at least
moves most of the declarations out of the way.
|
|
|
|
"private.h" will go.
|
|
Also revert the ifdef OMPT_SUPPORT order to have the short fallback
first and not after 400 lines.
|
|
|
|
|
|
For historic reasons we had it setup that there was
` plugin-nextgen/common/PluginInterface/<sources + headers>`
which is not what we do anywhere else.
Now it looks like the rest:
```
plugin-nextgen/common/include/<headers>
plugin-nextgen/common/src/<sources>
```
As part of this, `dlwrap.h` was moved into common/include (as
`DLWrap.h`)
since it is exclusively used by the plugins.
|
|
|
