Commit Graph

1647 Commits

Author SHA1 Message Date
Vyacheslav Zakharin f2f88f3e7a An attempt to abandon omptarget out-of-tree builds.
I want to start using LLVM component libraries in libomptarget
to stop duplicating implementations already available in LLVM
(e.g. LLVMObject, LLVMSupport, etc.). Without relying on LLVM
in all libomptarget builds one has to provide fallback implementation
for each used LLVM feature.

This is an attempt to stop supporting out-of-llvm-tree builds of libomptarget.

I understand that I may need to revert this,
if this affects downstream projects in a bad way.

Differential Revision: https://reviews.llvm.org/D101509
2021-05-07 12:43:50 -07:00
Joseph Huber a15f8589f4 [libomptarget] Add support for target memory allocators to cuda RTL
Summary:
The allocator interface added in D97883 allows the RTL to allocate shared and
host-pinned memory from the cuda plugin. This patch adds support for these to
the runtime.

Reviewed By: grokos

Differential Revision: https://reviews.llvm.org/D102000
2021-05-07 10:27:02 -04:00
Jon Chesterfield 44ee974e2f [libomptarget][nfc] Refactor amdgpu partial barrier to simplify adding a second one
[libomptarget][nfc] Refactor amdgpu partial barrier to simplify adding a second one

D101976 would require a second barrier instance. This NFC to amdgpu makes it
simpler to add one (an extra global, one more line in init). Also renames the
current barrier to L0.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D102016
2021-05-06 23:52:19 +01:00
Jon Chesterfield 7e9351b9de [libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin
[libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin

Drops an enum that was identical to a HSA one, localises some functions where
they were only called from one TU. Covers everything internalize + adce can
identify as dead, except for msgpack::dump which is useful when debugging.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D102014
2021-05-06 23:16:32 +01:00
Jon Chesterfield 25fe17d3c1 [libomptarget] Initial documentation on amdgpu offload
[libomptarget] Initial documentation on amdgpu offload

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D101927
2021-05-05 19:58:52 +01:00
Peyton, Jonathan L 9982f33e2c [OpenMP] Refactor/Rework topology discovery code
This patch does the following:

1) Introduce kmp_topology_t as the runtime-friendly structure (the
corresponding global variable is __kmp_topology) to determine the
exact machine topology which can vary widely among current and future
architectures. The current design is not easy to expand beyond the assumed
three layer topology: sockets, cores, and threads so a rework capable of
using the existing KMP_AFFINITY mechanisms is required.

This new topology structure has:
* The depth and types of the topology
* Ratio count for each consecutive level (e.g., number of cores per
   socket, number of threads per core)
* Absolute count for each level (e.g., 2 sockets, 16 cores, 32 threads)
* Equivalent topology layer map (e.g., Numa domain is equivalent to
   socket, L1/L2 cache equivalent to core)
* Whether it is uniform or not

The hardware threads are represented with the kmp_hw_thread_t
structure. This structure contains the ids (e.g., socket 0, core 1,
thread 0) and other information grabbed from the previous Address
structure. The kmp_topology_t structure contains an array of these.

2) Generalize the KMP_HW_SUBSET envirable for the new
kmp_topology_t structure. The algorithm doesn't assume any order with
tiles,numa domains,sockets,cores,threads. Instead it just parses the
envirable, makes sure it is consistent with the detected topology
(including taking into account equivalent layers) and then trims away
the unneeded subset of hardware threads. To enable this, a new
kmp_hw_subset_t structure is introduced which contains a vector of
items (hardware type, number user wants, offset). Any keyword within
__kmp_hw_get_keyword() can be used as a name and can be shortened as
well. e.g.,
KMP_HW_SUBSET=1s,2numa,4tile,2c,3t can be used on the KNL SNC-4 machine.

3) Simplify topology detection functions so they only do the singular
task of detecting the machine's topology. Printing, and all
canonicalizing functionality is now done afterwards. So many lines of
duplicated code are eliminated.

4) Add new ll_caches and numa_domains to OMP_PLACES, and
consequently, KMP_AFFINITY's granularity setting. All the names within
__kmp_hw_get_keyword() are available for use in OMP_PLACES or
KMP_AFFINITY's granularity setting.

5) Simplify and future-proof code where explicit lists of allowed
affinity settings keywords inside if() conditions.

6) Add x86 CPUID leaf 4 cache detection to existing x2apic id method
so equivalent caches could be detected (in particular for the ll_caches
place).

Differential Revision: https://reviews.llvm.org/D100997
2021-05-03 18:00:24 -05:00
Pushpinder Singh ae845d6426 [AMDGPU][OpenMP] Enable Libomptarget runtime tests
This enables the runtime tests on amdgpu targets.
10 tests have been marked as XFAIL on amdgcn currently mostly due to
missing printf.

Reviewed By: protze.joachim

Differential Revision: https://reviews.llvm.org/D99656
2021-05-03 05:56:42 +00:00
Martin Storsjö 01d27fc408 [OpenMP] Fix warnings due to redundant semicolons. NFC. 2021-05-02 21:51:06 +03:00
Kevin Athey bc9120047b Correct tiny misspelling (readlef -> readelf).
Getting my feet wet here as a new committer.

Correct misspelling in check-depends.pl.

Reviewed By: vitalybuka

Differential Revision: https://reviews.llvm.org/D101552
2021-04-30 17:20:35 -07:00
Michael Kruse 7308862ff5 [OpenMP][CMake] Use in-project clang as CUDA->IR compiler.
If available, use the clang that is already built in the same project as
CUDA compiler unless another executable is explicitly defined. This also
ensures the generated deviceRTL IR will be consistent with the version
of Clang.

This patch is required to reliably test OpenMP offloading in a buildbot
without either a two-stage build (e.g. with LLVM_ENABLE_RUNTIMES) or a
separately installed clang on the worker that will eventually become
outdated.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D101265
2021-04-30 12:45:52 -05:00
Michael Kruse 3244a8b536 [OpenMP][CMake] Pass --cuda-path to regression tests.
The OpenMP runtime can be compiled using a CUDA installed at non-default
location with the -DCUDA_TOOLKIT_ROOT_DIR setting. However, check-openmp
will fail afterwards because Clang needs to know where to find the CUDA
headers.

Fix by passing -cuda-path to Clang using the value of
CUDA_TOOLKIT_ROOT_DIR which has been determined by CMake. Also set
LD_LIBRARY_PATH such that it can find the cuda runtime when executing.
This will ensure that the regression test do not depend on the current
environment, but use the environment it was configured for.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D101266
2021-04-27 16:27:40 -05:00
Joachim Protze 24f836e8fd [OpenMP][libomptarget] Separate lit tests for different offloading targets (2/2)
This patch fuses the RUN lines for most libomptarget tests. The previous patch
D101315 created separate test targets for each supported offloading triple.

This patch updates the RUN lines in libomptarget tests to use a generic run
line independent of the offloading target selected for the lit instance.

In cases, where no RUN line was defined for a specific offloading target,
the corresponding target is declared as XFAIL. If it turns out that a test
actually supports the target, the XFAIL line can be removed.

Differential Revision: https://reviews.llvm.org/D101326
2021-04-27 15:54:32 +02:00
Joachim Protze b845217b1d [OpenMP][libomptarget] Separate lit tests for different offloading targets (1/2)
This patch creates a separate test directory for each offloading target to be
tested. This allows to test multiple architectures in one configuration, while
still see all failing tests separately. The lit test names include the target
triple, so that it will be easier to spot the failing target.

This patch also allows to mark expected failing tests based on the
target-triple, as the currently used triple is added to the lit "features":
```
// XFAIL: nvptx64-nvidia-cuda
```

Differential Revision: https://reviews.llvm.org/D101315
2021-04-27 12:30:01 +02:00
Joseph Huber 077fe0f739 [OpenMP][Documentation] Add FAQ entry for dynamically linked libraries
Summary:
Add an FAW entry detailing the support for using dynamically linked libraries
with OpenMP Offloading
2021-04-26 14:21:17 -04:00
Jon Chesterfield 58f125493d [libomptarget] Enable AMDGPU devicertl
[libomptarget] Enable AMDGPU devicertl

The amdgpu devicertl is written in freestanding openmp and compiles to a
bitcode library (per listed gfx arch) with no unresolved symbols. It requires
a recent clang, preferably the one from the same monorepo checkout.

This is D98658, with printf explicitly stubbed out, after patching clang to no
longer require an llvm with the amdgpu target enabled.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D101213
2021-04-24 02:24:44 +01:00
Johannes Doerfert 17330a3cb1 [OpenMP] Avoid reading uninitialized parallel level values
In a last minute change request for a2dbfb6b72 we introduced a
read of the uninitialized parallel level value in SPMD-mode.
We go back to initializing the array early and checking for an
adjusted level.

Found by the miniqmc unit tests:
  https://cdash.qmcpack.org/CDash/viewTest.php?onlyfailed&buildid=203434

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D101123
2021-04-23 11:21:58 -05:00
Joseph Huber 59b6849012 [OpenMP] Replace global InfoLevel with a reference to an internal one.
Summary:
This patch improves the implementation of D100774 by replacing the global
variable introduced with a function that returns a reference to an internal
one. This removes the need to define the variable in every plugin that uses it.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D101102
2021-04-23 09:43:46 -04:00
Joseph Huber 2b6f20082e [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-22 12:48:11 -04:00
Alexey Bataev ca70512099 [OPENMP]Mark test as unsupported to avoid possible unexpected passes,
NFC.
2021-04-22 08:06:25 -07:00
Giorgis Georgakoudis a2dbfb6b72 [OpenMP] Simplify offloading parallel call codegen
This revision simplifies Clang codegen for parallel regions in OpenMP GPU target offloading and corresponding changes in libomptarget: SPMD/non-SPMD parallel calls are unified under a single `kmpc_parallel_51` runtime entry point for parallel regions (which will be commonized between target, host-side parallel regions), data sharing is internalized to the runtime. Tests have been auto-generated using `update_cc_test_checks.py`. Also, the revision contains changes to OpenMPOpt for remark creation on target offloading regions.

Reviewed By: jdoerfert, Meinersbur

Differential Revision: https://reviews.llvm.org/D95976
2021-04-21 18:46:07 -07:00
Alexey Bataev 079884225a [OPENMP]Fix PR49698: OpenMP declare mapper causes segmentation fault.
The implicitly generated mappings for allocation/deallocation in mappers
runtime should be mapped as implicit, also no need to clear member_of
flag to avoid ref counter increment. Also, the ref counter should not be
incremented for the very first element that comes from the mapper
function.

Differential Revision: https://reviews.llvm.org/D100673
2021-04-21 10:38:31 -07:00
Peyton, Jonathan L 4457565757 [OpenMP] Implement GOMP task reductions
Implement the remaining GOMP_* functions to support task reductions
in taskgroup, parallel, loop, and taskloop constructs.  The unused mem
argument to many of the work-sharing constructs has to do with the
scan() directive/ inscan() modifier.  If mem is set, each function
will call KMP_FATAL() and tell the user scan/inscan is unsupported.  The
GOMP reduction implementation is kept separate from our implementation
because of how GOMP presents reduction data and computes the reductions.
GOMP expects the privatized copies to be present even after a #pragma
omp parallel reduction(task:...) region has ended so the data is stored
inside GOMP's uintptr_t* data pseudo-structure.  This style is tightly
coupled with GCC compiler codegen.  There also isn't any init(),
combiner(), fini() functions in GOMP's codegen so the two
implementations were to disparate to try to wrap GOMP's around our own.

Differential Revision: https://reviews.llvm.org/D98806
2021-04-16 16:36:31 -05:00
Peyton, Jonathan L 5ebbb366c4 [OpenMP] Allow affinity to re-detect for child processes
Current atfork() handler for child processes does not reset
the affinity masks array which prevents users from setting their own
affinity in child processes.

Differential Revision: https://reviews.llvm.org/D99218
2021-04-16 16:34:02 -05:00
Hansang Bae 9b98497b44 [OpenMP] Add omp_target_is_accessible() to header files
-- Added omp_target_is_accessible to the header files
-- Added missing const qualifier to device memory routines

Differential Revision: https://reviews.llvm.org/D100420
2021-04-16 07:54:15 -05:00
Joseph Huber 83d4b2e2e0 [OpenMP] Add info for device table changes
Summary:
This patch adds a feature to print information whenever the host-device pointer
mapping table is changed by inserting or removing an entry. This introduces a
new bit field for LIBOMPTARGET_INFO at position 0x8.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D100600
2021-04-15 18:39:48 -04:00
Hansang Bae 77dc7b4653 [OpenMP] Fix printing routine for OMP_TOOL_VERBOSE_INIT
Also fixed typo in the verbose message.

Differential Revision: https://reviews.llvm.org/D100414
2021-04-14 07:55:26 -05:00
Hansang Bae 3da61ddae7 [OpenMP] Define omp_is_initial_device() variants in omp.h
omp_is_initial_device() is marked as a built-in function in the current
compiler, and user code guarded by this call may be optimized away,
resulting in undesired behavior in some cases. This patch provides a
possible fix for such cases by defining the routine as a variant
function and removing it from builtin list.

Differential Revision: https://reviews.llvm.org/D99447
2021-04-06 16:58:01 -05:00
Peyton, Jonathan L 2aebb7cb3c [OpenMP] Fix incorrect KMP_STRLEN() macro
The second argument to the strnlen_s(str, size) function should be
sizeof(str) when str is a true array of characters with known size
(instead of just a char*). Use type traits to determine if first
parameter is a character array and use the correct size based on that
trait.

Differential Revision: https://reviews.llvm.org/D98209
2021-04-05 09:03:09 -05:00
Joseph Huber 0af4e74aef [OpenMP][NFC] Fix typo in libomptarget error message
Summary:
There was a typo suggesting the user to use `LIBOMPTARGET_DEBUG` instead of
`LIBOMPTARGET_INFO`
2021-04-01 12:45:28 -04:00
Joseph Huber 29338459fb [OpenMP] Trim error messages in CUDA plugin
Summary:
Remove some of the error messages printed when the CUDA plugin fails. The current error messages can be confusing because they are the first error messages printed after the async stream finds an error. This means that the printed values aren't related to what caused the issue, but are simply the last asyncronous operation that succeeded on the device. Remove these as they can be misleading.

Reviewers: jdoerfert

Differential Revision: https://reviews.llvm.org/D99510
2021-03-29 12:20:19 -04:00
Alexey Bataev 0411b23319 [OPENMP]Map data field with l-value reference types.
Added initial support dfor the mapping of the data members with l-value
reference types.

Differential Revision: https://reviews.llvm.org/D98812
2021-03-29 07:07:09 -07:00
Joseph Huber 16064e71e9 [OpenMP] Reset async stream properly upon failure
Summary:
If the call to `synchronize` fails, it will currently block the stream indefinitely if execution is continued from this point. Additionally, if the program exits it will trigger an assertion on the non-null value of the async queue and prevent the runtime from printing debugging information.

Reviewers: jdoerfert

Differential Revision: https://reviews.llvm.org/D99443
2021-03-26 19:05:06 -04:00
Hansang Bae 467f39249d [OpenMP] Misc. changes that add or remove pointer/bound checks
-- Added or moved checks to appropriate places.
-- Removed ineffective null check where the pointer is already being
   dereferenced around the code.
-- Initialized variables that can be used without definitions.
-- Added call to dlclose/FreeLibrary in OMPT tool activation.
-- Added a new build compiler definition.

Differential Revision: https://reviews.llvm.org/D98584
2021-03-23 18:55:08 -05:00
Shilei Tian 2df65f87c1 [OpenMP] Fixed a crash in hidden helper thread
It is reported that after enabling hidden helper thread, the program
can hit the assertion `new_gtid < __kmp_threads_capacity` sometimes. The root
cause is explained as follows. Let's say the default `__kmp_threads_capacity` is
`N`. If hidden helper thread is enabled, `__kmp_threads_capacity` will be offset
to `N+8` by default. If the number of threads we need exceeds `N+8`, e.g. via
`num_threads` clause, we need to expand `__kmp_threads`. In
`__kmp_expand_threads`, the expansion starts from `__kmp_threads_capacity`, and
repeatedly doubling it until the new capacity meets the requirement. Let's
assume the new requirement is `Y`.  If `Y` happens to meet the constraint
`(N+8)*2^X=Y` where `X` is the number of iterations, the new capacity is not
enough because we have 8 slots for hidden helper threads.

Here is an example.
```
#include <vector>

int main(int argc, char *argv[]) {
  constexpr const size_t N = 1344;
  std::vector<int> data(N);

#pragma omp parallel for
  for (unsigned i = 0; i < N; ++i) {
    data[i] = i;
  }

#pragma omp parallel for num_threads(N)
  for (unsigned i = 0; i < N; ++i) {
    data[i] += i;
  }

  return 0;
}
```
My CPU is 20C40T, then `__kmp_threads_capacity` is 160. After offset,
`__kmp_threads_capacity` becomes 168. `1344 = (160+8)*2^3`, then the assertions
hit.

Reviewed By: protze.joachim

Differential Revision: https://reviews.llvm.org/D98838
2021-03-18 18:25:36 -04:00
Jon Chesterfield 626a31de15 [libomptarget] Add register usage info to kernel metadata
Add register usage information to the runtime metadata so that it can be used during kernel launch (that change will be in a different commit). Add this information to the kernel trace.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D98829
2021-03-18 17:00:42 +00:00
Jon Chesterfield dbf8f2b089 Revert "[libomptarget] Build amdgcn devicertl by default"
This reverts commit e23f3502d9.
It broke the build of openmp for clang built without amdgcn
support. D98746, under review, would allow this to reland.
2021-03-17 11:34:44 +00:00
Hansang Bae a6f9cb6adc [OpenMP] Add runtime interface for OpenMP 5.1 error directive
The proposed new interface is for supporting `at(execution)` clause in the
error directive.

Differential Revision: https://reviews.llvm.org/D98448
2021-03-16 08:55:25 -05:00
Johannes Doerfert 0a954a528b [OpenMP][FIX] Repair accidental replacement of _shfl_sync with _shfl
This was broken accidentally in D95752.

Reviewed By: ye-luo

Differential Revision: https://reviews.llvm.org/D98677
2021-03-15 22:46:00 -05:00
Jon Chesterfield e23f3502d9 [libomptarget] Build amdgcn devicertl by default
[libomptarget] Build amdgcn devicertl by default

The cmake for this looks for an llvm install and does the right thing when
building as part of enable_runtimes. It will probably do the right thing
in other settings - at least, it won't try to build this with gcc.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D98658
2021-03-15 23:17:50 +00:00
Peyton, Jonathan L 7085f04573 [OpenMP] Remove unused cpu_stackoffset member 2021-03-15 16:52:04 -05:00
Jon Chesterfield bb38d7ff05 [libomptarget][nfc][amdgcn] Use precise triple for devicertl build 2021-03-15 20:24:13 +00:00
Jon Chesterfield d0bc85f04a [libomptarget][nfc] Drop unused DEVICE macro
[libomptarget][nfc] Drop unused DEVICE macro

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D98655
2021-03-15 20:12:50 +00:00
Jon Chesterfield 7da76aaaf4 [libomptarget] Build amdgpu plugin by default
[libomptarget] Build amdgpu plugin by default

This will build the amdgpu plugin if cmake is able to find the hsa
runtime library, which will be the case if rocm is installed or if
the hsa library has been installed somewhere cmake looks.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D98654
2021-03-15 20:12:01 +00:00
Jon Chesterfield bcb3f0f867 [libomptarget] Fix devicertl build
[libomptarget] Fix devicertl build

The target specific functions in target_interface are extern C, but the
implementations for nvptx were mostly C++ mangling. That worked out as
a quirk of DEVICE macro expanding to nothing, except for shuffle.h which
only forward declared the functions with C++ linkage.

Also implements GetWarpSize, as used by shuffle, and includes target_interface
in nvptx target_impl.cu to help catch future divergence between interface and
implementation.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D98651
2021-03-15 19:50:22 +00:00
Jon Chesterfield f675b3df48 [libomptarget] Drop assert.h, use freestanding for amdgcn devicertl
[libomptarget] Drop assert.h, use freestanding for amdgcn devicertl

Promotes the runtime assert to a link time error for the unimplemented
fallback functions. Enables amdgcn to build with only clang provided
headers, which makes it less likely to break other builds when enabled.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D98649
2021-03-15 18:50:09 +00:00
Jon Chesterfield 156842937f [libomptarget][amdgcn] Drop use of inttypes.h, moving closer to freestanding
[libomptarget][amdgcn] Drop use of inttypes.h, moving closer to freestanding

The glibc headers are a periodic source of problems compiling the devicertl.
This patch resolves the following error run into while building llvm on a slightly
different linux system.
```
In file included from .../lib/clang/13.0.0/include/inttypes.h:21:
In file included from /usr/include/inttypes.h:25:
/usr/include/features.h:461:12: fatal error: 'sys/cdefs.h' file not found
#  include <sys/cdefs.h>
           ^~~~~~~~~~~~~
```
As a second patch, removing assert.h from shuffle will let amdgcn build as
-ffreestanding, at which point only the headers that clang itself provides are
used and interactions with the host glibc are eliminated. Doing the same for
nvptx is complicated by printf handling but also seems worthwhile.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D98565
2021-03-15 16:54:58 +00:00
George Rokos 2468fdd9af [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-13 03:47:07 -08:00
Johannes Doerfert 5449fbb5d4 [OpenMP][NFC] Use `AsyncInfo` as the variable name for a `__tgt_async_info`
Reviewed By: grokos, tianshilei1992

Differential Revision: https://reviews.llvm.org/D96444
2021-03-11 23:31:34 -06:00
Johannes Doerfert 66ba494b49 [OpenMP][DeviceRTL] Extract shuffle idiom and port it to declare variant
The shuffle idiom is differently implemented in our supported targets.
To reduce the "target_impl" file we now move the shuffle idiom in it's
own self-contained header that provides the implementation for AMDGPU
and NVPTX. A fallback can be added later on.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D95752
2021-03-11 23:31:30 -06:00
Joseph Huber 807466ef28 [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-11 09:52:11 -05:00