OpenMP 5.2, sec. 10.2 "teams Construct", p. 232, L9-12 restricts what
regions can be strictly nested within a `teams` construct. This patch
relaxes Clang's enforcement of this restriction in the case of nested
`atomic` constructs unless `-fno-openmp-extensions` is specified.
Cases like the following then seem to work fine with no additional
implementation changes:
```
#pragma omp target teams map(tofrom:x)
#pragma omp atomic update
x++;
```
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D126323
Summary:
We usually used the `OMP_LIKELY` and `OMP_UNLIKELY` macros to add branch
prediction intrinsics to help the optimizer ignore unlikely loops. This
wasn't applied to this one loop so add that in.
Summary:
This patch adds the `leaf` attribute to the `vprintf` declaration in the
OpenMP runtime. This attribute allows us to determine that the `vprintf`
function will not call any functions within the translation unit,
allowing us to deduce `norecurse` attributes on the caller.
Currently the library ignores requested wait policy in the presence
of tasking. Threads always actively spin. The patch fixes this problem
making the wait policy passive if this explicitly requested by user.
Differential Revision: https://reviews.llvm.org/D123044
The OpenMP device offloading library is a bitcode library and thus only
expect to build and linked with the same version of clang that was used
to create it. This somewhat copmlicates the building process as we
require the Clang that was just built to be used to create the library.
This is either done with a two-step build, where OpenMP is built with
the Clang that was just installed, or through the
`-DLLLVM_ENABLE_RUNTIMES=openmp` option. This has always been the case,
but recent changes have caused this to make it difficult to build the
rest of OpenMP. This patchs adds a check to not build the OpenMP device
runtime if the current compiler is not Clang with the same version as
the LLVM installation. This should allow users to build OpenMP as a
project using any compiler without it erroring out due to the bitcode
library, but if users require it they will need to use the above methods
to compile it.
Reviewed By: jdoerfert, tianshilei1992, ye-luo
Differential Revision: https://reviews.llvm.org/D125698
Summary:
This patch allows users to compile the static library without CUDA
installed on the system. This requires the new flag `--cuda-feature` to
indicate that we need `+ptx61` in order to compile the runtime.
This patch adds the necessary CMake configuration to build a static
library version of the device runtime, `libomptarget.devicertl.a`.
Various improvements in how we handle static libraries and generating
offloading code should allow us to treat the device library as a regular
project without needing to invoke the clang front-end directly. Here we
generate a job for each offloading architecture supported. Each
offloading architecture will be embedded into the static library and
used as-needed by the host.
This library will primarily be used to replace the bitcode library when
performing LTO. Currently, we need to manually pass in the bitcode
library which requires foreknowledge of the offloading architecture.
This approach lets us handle that in the linker wrapper instead.
Furthermore this should improve our interface to the device runtime. We
can now build it fully under a release build and have all the expected
entry points, as well as supporting debug builds.
Depends on D125265 D125256 D125260 D125314 D125563
Reviewed By: tianshilei1992
Differential Revision: https://reviews.llvm.org/D125315
We used to globally include the libomptarget include directory for all
projects. This caused some conflicts with the other files named
"Debug.h". This patch changes the cmake to include these files via the
target include instead.
Reviewed By: tianshilei1992
Differential Revision: https://reviews.llvm.org/D125563
This patche attemps to address the current warnings in the OpenMP
offloading device runtime. Previously we did not see these because we
compiled the runtime without the standard warning flags enabled.
However, these warnings are used when we now build the static library
version of this runtime. This became extremely noisy when coupled with
the fact the we compile each file roughly 32 times when all the
architectures are considered. So it would be ideal to not have all these
warnings show up when building.
Most of these errors were simply implicit switch-case fallthroughs,
which can be addressed using C++17's fallthrough attribute. Additionally
there was a volatile variable that was being casted away. This is most
likely safe to remove because we cast it away before its even used and
didn't seem to affect anything in testing.
Depends on D125260
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D125339
Currently the OpenMP offloading device runtime is only expected to be
compiled for the specific architecture it's targeting. This is
problematic if we want to make compiling the device runtime more general
via the standar `clang` driver rather than invoking the clang front-end
directly. This patch addresses this by primarily changing the declare
type to `nohost` so the host will not contain any of this code.
Additionally we forward declare the functions that are defined via
variants, otherwise these would cause problems on the host.
Reviewed By: jdoerfert, tianshilei1992
Differential Revision: https://reviews.llvm.org/D125260
Intel Inspector uses itt notifications to analyze code execution, and it
reports race conditions in dependent tasks.
This patch fixes the issue notifying Inspector on tasks dependency
synchronizations.
Differential Revision: https://reviews.llvm.org/D123042
Currently, globals on the device will have an infinite reference count
and an unknown name when using `LIBOMPTARGET_INFO` to print the mapping
table. We already store the name of the global in the offloading entry
so we should be able to use it, although there will be no source
location. To do this we need to create a valid `ident_t` string from a
name only.
Reviewed By: tianshilei1992
Differential Revision: https://reviews.llvm.org/D124381
Consider checking whether a pointer has been mapped can be achieved via omp_get_mapped_ptr.
omp_target_is_present is more needed to check whether the storage being pointed is mapped.
This restore the old behavior of omp_target_is_present before D123093
Fixes https://github.com/llvm/llvm-project/issues/54899
Reviewed By: jdenny
Differential Revision: https://reviews.llvm.org/D123891
When hwloc is used and is installed outside of the default paths, the omp CMake target
needs to provide the needed include path thru the CMake target by adding it with
target_include_directories to it, so libompd gets it as well when it defines it's cmake
target using target_link_libraries.
As suggested in D122667
Reviewed By: ye-luo
Differential Revision: https://reviews.llvm.org/D123888
The remote offloading server and plugin rely on OpenMP, so this needs to be added as a linker flag. Without this, applications segfault.
Differential Revision: https://reviews.llvm.org/D124200
This fixes a compile-time error recently introduced within the remote
offloading plugin. This patch also removes some extra linker flags that are unnecessary, and adds an explicit abseil linker flag without which we occasionally get problems.
Differential Revision: https://reviews.llvm.org/D119984
Summary:
One test had an old "unsupported" string that used the old `newDriver`
string which was removed. This test should be updated to use the
`oldDriver` one instead.
Previously an opt-in flag `-fopenmp-new-driver` was used to enable the
new offloading driver. After passing tests for a few months it should be
sufficiently mature to flip the switch and make it the default. The new
offloading driver is now enabled if there is OpenMP and OpenMP
offloading present and the new `-fno-openmp-new-driver` is not present.
The new offloading driver has three main benefits over the old method:
- Static library support
- Device-side LTO
- Unified clang driver stages
Depends on D122683
Differential Revision: https://reviews.llvm.org/D122831
Summary:
The changes made in D123177 added new targets to the
`LIBOMPTARGET_TESTED_PLUGINS` variable which was linked against when
building the `llvm-omp-target-info` tool. This caused linker errors on
the export scripts. This patch removes that dependency, it still builds
and runs as expected so I will assume it's correct.
As described in 5.1 spec
2.21.7.2 Pointer Initialization for Device Data Environments
Reviewed By: RaviNarayanaswamy
Differential Revision: https://reviews.llvm.org/D123093
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
The target allocators have been supported for NVPTX offloading for
awhile. The tests should use the allocators instead of calling the
functions manually. Also the comments indicating these being a preview
should be removed.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D123242
In a clean build directory, `check-openmp` or `check-libomptarget` will fail because of missing device RTL .bc files. Ensure that the new targets new custom targets `omptarget.devicertl.nvptx` and `omptarget.devicertl.amdgpu` (corresponding to the plugin rtl targets `omptarget.rtl.cuda`, respectively `omptarget.rlt.amdgpu` ) are dependencies of the regression tests.
Reviewed By: JonChesterfield
Differential Revision: https://reviews.llvm.org/D123177
This change adds support for ompt_callback_dispatch with the new
dispatch chunk type introduced in 5.2. Definitions of the new
ompt_work_loop types were also added in the header file.
Differential Revision: https://reviews.llvm.org/D122107
Latest OpenMP spec says parallel_data is NULL for initial/implicit-task-end.
We nevertheless need to cleanup the ParallelData here, as there is no other
callback for the end of the implicit parallel region. We can use the reference
stored in the TaskData.
Reviewed By: dreachem
Differential Revision: https://reviews.llvm.org/D114005
If we decided to delete a mapping entry we did not act on it right away
but first issued and waited for memory copies. In the meantime some
other thread might reuse the entry. While there was some logic to avoid
colliding on the actual "deletion" part, there were two races happening:
1) The data transfer back of the thread deleting the entry and
the data transfer back of the thread taking over the entry raced.
2) The update to the shadow map happened regardless if the entry was
actually reused by another thread which left the shadow map in a
inconsistent state.
To fix both issues we will now update the shadow map and delete the
entry only if we are sure the thread is responsible for deletion, hence
no other thread took over the entry and reused it. We also wait for a
potential former data transfer from the device to finish before we issue
another one that would race with it.
Fixes https://github.com/llvm/llvm-project/issues/54216
Differential Revision: https://reviews.llvm.org/D121058
Inline assembly is scary but we need to support it for the OpenMP GPU
device runtime. The new assumption expresses the fact that it may not
have call semantics, that is, it will not call another function but
simply perform an operation or side-effect. This is important for
reachability in the presence of inline assembly.
Differential Revision: https://reviews.llvm.org/D109986
As we mentioned in the code comments for function `ResourcePoolTy::release`,
at some point there could be two identical resources on the two sides of `Next`
mark. It is usually not an issue, unless the following case:
1. Some resources are not returned.
2. We need to iterate the pool and free the element.
That will cause double free, which is the case for event pool. Since we don't release
events hold by the data map, it can happen that the `Next` mark is not reset, and
we have two identical items in the pool. When the pool is destroyed, we will call
`cuEventDestroy` twice on the same event. In the best case, we can only observe
CUDA errors. In the worst case, it can cause internal failures in CUDART and further
crash.
This patch fixes the issue by tracking all resources that have been given using
an `unordered_set`. We don't remove it when a resource is returned. When the pool
is destroyed, we merge the pool (a `vector`) and the set. In this way, we can make
sure that the set contains all resources allocated from the device. We just need
to iterate the set and free the resource accordingly.
For now, only event pool is set to use it. Stream pool is not because we can make
sure all streams are returned when the plugin is destroyed.
Someone might be wondering, why don't we release all events hold in the data map.
That is because, plugins are determined to be destroyed *before* `libomptarget`.
If we can somehow make the plugin outlast `libomptarget`, life will be much
easier.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D122014
This patch adds the necessary AMDGPU calling convention to the ctor /
dtor kernels. These are fundamentally device kenels called by the host
on image load. Without this calling convention information the AMDGPU
plugin is unable to identify them.
Depends on D122504
Fixes#54091
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D122515
Joel E. Denny