Commit Graph

1166 Commits

Author SHA1 Message Date
Jon Chesterfield dbb3fec8ad [libomptarget] Move resource id functions into target specific code, implement for amdgcn
Summary: [libomptarget] Move resource id functions into target specific code, implement for amdgcn

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: jdoerfert

Subscribers: jvesely, mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71382
2019-12-12 22:49:02 +00:00
Jon Chesterfield b399252028 [libomptarget][nfc] Add missing header for amdgcn/target_impl 2019-12-12 09:36:57 +00:00
David Carlier 27535a1449 [OpenMP] Fix linkage issue on FreeBSD
needs kmp_set_thread_affinity_mask_initial implementation.
2019-12-06 15:47:50 +00:00
JonChesterfield 0dd62c5c2e [libomptarget][nfc] Move cuda threadfence functions behind kmpc_impl
Summary:
[libomptarget][nfc] Move cuda threadfence functions behind kmpc_impl

Part of building code under common/ without requiring a cuda compiler

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: ABataev

Subscribers: jvesely, jfb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71102
2019-12-06 15:41:18 +00:00
Jon Chesterfield cd90f49d70 [libomptarget][nfc] Move three more files to common
Summary: [libomptarget][nfc] Move three more files to common

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: ABataev

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71103
2019-12-06 15:29:50 +00:00
Jon Chesterfield 4af84d2686 [libomptarget][nfc] Introduce SHARED, ALIGN macros
Summary:
[libomptarget][nfc] Introduce SHARED, ALIGN macros
Move remaining cuda attributes behind such macros

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: ABataev

Subscribers: openmp-commits, jvesely

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71076
2019-12-05 21:57:58 +00:00
Jon Chesterfield d0b9ed5c49 [libomptarget][nfc] Move omptarget-nvptx under common
Summary:
[libomptarget][nfc] Move omptarget-nvptx under common

Almost all files depend on require omptarget-nvptx, which no longer
contains any obviously architecture dependent code. Moving it under
common unblocks task/loop for amdgcn, and allows moving other code.

At some point there should probably be a widespread symbol renaming to
replace the nvptx string. I'd prefer to get things working first.

Building this (and task.cu, loop.cu) without a cuda library requires
some more refactoring, e.g. wrap threadfence(), use DEVICE macro more
consistently. Patches for that are orthogonal and will be posted shortly.

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: ABataev

Subscribers: mgorny, fedor.sergeev, jfb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D71073
2019-12-05 20:34:15 +00:00
JonChesterfield 3ada8d2a87 [libomptarget] Build a minimal deviceRTL for amdgcn
Summary:
[libomptarget] Build a minimal deviceRTL for amdgcn

Repeat of D70414, with an include path fixed. Diff for sanity checking.

The CMakeLists.txt file is functionally identical to the one used in the aomp fork.
Whitespace changes were made based on nvptx/CMakeLists.txt, plus the
copyright notice updated to match (Greg was the original author so would
like his sign off on that here).

This change will build a small subset of the deviceRTL if an appropriate toolchain is
available, e.g. a local install of rocm. Support.h is moved from nvptx as a dependency
of debug.h.

Reviewers: ABataev, jdoerfert

Reviewed By: ABataev

Subscribers: jvesely, mgorny, jfb, openmp-commits, jdoerfert

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D70971
2019-12-04 16:43:37 +00:00
Alexey Bataev 02b9c5d963 Revert "[libomptarget] Build a minimal deviceRTL for amdgcn"
This reverts commit 877ffa716f because it
breaks the build.
2019-12-03 12:35:08 -05:00
Jon Chesterfield 877ffa716f [libomptarget] Build a minimal deviceRTL for amdgcn
Summary:
[libomptarget] Build a minimal deviceRTL for amdgcn

The CMakeLists.txt file is functionally identical to the one used in the aomp fork.
Whitespace changes were made based on nvptx/CMakeLists.txt, plus the
copyright notice updated to match (Greg was the original author so would
like his sign off on that here).

This change will build a small subset of the deviceRTL if an appropriate toolchain is
available, e.g. a local install of rocm. Support.h is moved from nvptx as a dependency
of debug.h.

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Reviewed By: jdoerfert

Subscribers: jfb, Hahnfeld, jvesely, mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D70414
2019-12-03 15:18:41 +00:00
Bryan Chan 4d3198e243 [OpenMP] build offload plugins before testing them
Summary:
"make check-all" or "make check-libomptarget" would attempt to run offloading
tests before the offload plugins are built. This patch corrects that by adding
dependencies to the libomptarget CMake rules.

Reviewers: jdoerfert

Subscribers: mgorny, guansong, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D70803
2019-11-28 17:43:56 -05:00
AndreyChurbanov bd2fb41c2d [openmp] Fixed nonmonotonic schedule when #threads > #chunks in a loop.
Differential Revision: https://reviews.llvm.org/D70713
2019-11-27 15:26:51 +03:00
AndreyChurbanov 5f8b8d2820 [openmp] Recognise ARMv7ve machine arch.
Patch by raj.khem (Khem Raj)

Differential Revision: https://reviews.llvm.org/D68543
2019-11-26 14:37:24 +03:00
protze@itc.rwth-aachen.de 77ad98c808 [OpenMP][Tool] archer tests require tsan
Testing for tsan capability in the test-compiler in follow-up review
2019-11-22 17:11:16 +01:00
protze@itc.rwth-aachen.de 6b2431e0c2 [OpenMP][Tool] disable archer tests in standalone build
Will be enabled after Build-Bots are fixed
2019-11-22 15:25:43 +01:00
protze@itc.rwth-aachen.de ac21de0d7e [OpenMP][Tool] Fix cmake variable in lit.site.cfg.in
As noted in D45890
2019-11-22 14:31:54 +01:00
JonChesterfield a84b48d01e [nfc][libomptarget] Remove casts of string literals to char* 2019-11-19 19:41:59 +00:00
JonChesterfield 4681e2e434 [nfc][libomptarget] Write amdgcn macros in terms of compiler intrinsics 2019-11-19 17:23:46 +00:00
AndreyChurbanov 3a76b8a538 Fix openmp on PowerPC64-BE-ELFv2 ABI on FreeBSD.
Patch by adalava (Alfredo Dal'Ava J.nior)

Differential Revision: https://reviews.llvm.org/D67190
2019-11-19 19:45:06 +03:00
Aaron Puchert b29c7fdb61 [OpenMP] Remove -Wl,-fini=__kmp_internal_end_fini
Summary:
The termination function duplicated the functionality of the
__attribute((destructor))-annotated function __kmp_internal_end_fini,
and we have no indication that this doesn't work.

The function might cause issues with link-time optimization turned on:
until very recently, none of the usual linkers was reporting functions
named in -Wl,-fini as used to the LTO plugin, so it might be dropped.
If the function is dropped, -Wl,-fini=__kmp_internal_end_fini doesn't
do what we want: with ld.bfd and lld it drops the FINI attribute from
.dynamic and with gold we get FINI = 0x0, which leads to a crash on
cleanup. This can be reproduced by building with

    -DLLVM_ENABLE_PROJECTS="clang;openmp" \
    -DLLVM_ENABLE_LTO=Thin \
    -DLLVM_USE_LINKER=gold

The issue in lld has been fixed in f95273f75a, but gold remains without
fix so far.

Fixes PR43927.

Reviewers: JonChesterfield, jdoerfert, AndreyChurbanov

Reviewed By: AndreyChurbanov

Differential Revision: https://reviews.llvm.org/D69927
2019-11-19 00:54:58 +01:00
Jon Chesterfield 5a4a05d776 [libomptarget][nfc] Move some source into common from nvptx
Summary:
[libomptarget][nfc] Move some source into common from nvptx

Moves some source that compiles cleanly under amdgcn into a common subdirectory
Includes some non-trivial files and some headers. Keeps the cuda file extension.

The build systems for different architectures seem unlikely to have much in
common. The idea is therefore to set include paths such that files under
common/src compile as if they were under arch/src as the mechanism for sharing.
In particular, files under common/src need to be able to include target_impl.h.

The corresponding -Icommon is left out in favour of explicit includes on the
basis that the it makes it clearer which files under common are used by a given
architecture.

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: ABataev

Subscribers: jfb, mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D70328
2019-11-18 18:17:36 +00:00
protze@itc.rwth-aachen.de 2b8115b10b [OpenMP] Add implementation and tests of Archer tool
The tool provides TSAN annotations for OpenMP synchronization. The tool
is activated if no other OMPT tool is loaded.

The tool detects whether the application was built with TSan and rejects
activation according to the OMPT protocol if there is no TSan-rt.

Differential Revision: https://reviews.llvm.org/D45890
2019-11-18 14:45:34 +01:00
Sylvestre Ledru 9b40a7f3bf Remove +x permission on some files 2019-11-16 14:47:20 +01:00
JonChesterfield 32dfbd131d [libomptarget][nfc] Use cuda variable wrappers from support.h
Summary:
[libomptarget][nfc] Use cuda variable wrappers from support.h
Reimplementation of D69693, after the revert of D69885

Use the wrappers in support.h for cuda builtin variables at all call sites.
Localises use of cuda and removes WARPSIZE==32 assumption in debug.h.

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: jdoerfert

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D70186
2019-11-14 12:45:09 +00:00
JonChesterfield fd9fa9995c [libomptarget] Move supporti.h to support.cu
Summary:
[libomptarget] Move supporti.h to support.cu
Reimplementation of D69652, without the unity build and refactors.
Will need a clean build of libomptarget as the cmakelists changed.

Reviewers: ABataev, jdoerfert

Reviewed By: jdoerfert

Subscribers: mgorny, jfb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D70131
2019-11-13 11:36:46 +00:00
Michał Górny 6f8ee2c575 [openmp] [test] Skip one more test that kills NetBSD buildbot 2019-11-07 17:29:57 +01:00
Jon Chesterfield 7cea0cea77 [libomptarget] Revert all improvements to support
Summary:
[libomptarget] Revert all improvements to support

The change to unity build for nvcc has broken the build for some developers.
This patch reverts to a known-working state.

There has been some confusion over exactly how the build broke. I think we
have reached a common understanding that the disappearing symbols are from
the bitcode library built by clang. The static archive built by nvcc may show the
same problem. Some of the confusion arose from building the deviceRTL twice
and using one or the other library based on various environmental factors.

I'm pretty sure the problem is clang expanding `__forceinline__` into both `__inline__`
and `attribute(("always_inline"))`. The `__inline__` attribute resolves to linkonce_odr
which is not safe for exporting symbols from translation units.

"always_inline" is the desired semantic for small functions defined in one translation
unit that are intended to be inlined at link time. "inline" is not.

This therefore reintroduces the dependency hazard of supporti.h and some code
duplication, and blocks progress separating deviceRTL into reusable components.

See also D69857, D69859 for attempts at a fix instead of a revert.

Reviewers: ABataev, jdoerfert, grokos, ikitayama, tianshilei1992

Reviewed By: ABataev

Subscribers: mgorny, jfb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69885
2019-11-06 15:44:10 +00:00
Ron Lieberman dc34b1c94d Test commit: adds a . to comment. NFC 2019-11-04 16:51:03 -06:00
JonChesterfield 94c59ea8dd [libomptarget] Implement target_impl for amdgcn
Summary:
[libomptarget] Implement target_impl for amdgcn

Smallest atomic addition for a new target. Implements enough of the amdgcn
specific code that some of the source files under nvptx/src could be compiled,
without modification, to run on amdgcn.

This foreshadows a work in progress patch to move said source out of nvptx/src.
Patch based on fork at https://github.com/ROCm-Developer-Tools/llvm-project

Reviewers: ABataev, jdoerfert, grokos, ronlieb

Subscribers: jvesely, jfb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69718
2019-11-01 15:46:35 +00:00
Alexey Bataev e57f8ad914 [LIBOMPTARGET]Call GetLaneId function, do not use its address in debug
log functions.
2019-11-01 09:43:47 -04:00
JonChesterfield 9b06ac98d0 [nfc][omptarget] Use builtin var abstraction. Second pass at D69476
Summary:
[nfc][omptarget] Use builtin var abstraction. Second pass at D69476

Use the wrappers in support.h for cuda builtin variables at all call sites.
Localises use of cuda and removes WARPSIZE==32 assumption in debug.h.

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: jdoerfert

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69693
2019-11-01 02:21:44 +00:00
JonChesterfield 764c8420e4 [nfc][libomptarget] Reorganise support header
Summary:
[nfc][libomptarget] Reorganise support header

All functions defined in support implementation are now declared in support.h
Reordered functions in support implementation to match the sequence in support.h
Added include guards to support.h
Added #include interface to support.h to provide kmp_Ident declaration
Move supporti.h to support.cu and s/INLINE/EXTERN/g
Add remaining includes to support.cu

A minor side effect is to change the name mangling of the support functions to
extern "C". If this matters another macro along the lines of INLINE/EXTERN
can be added - perhaps DEVICE as that's the obvious implementation.

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: jdoerfert

Subscribers: mgorny, jfb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69652
2019-10-31 17:15:02 +00:00
Jon Chesterfield e9f9dfab82 [libomptarget] Change nvcc compilation to use a unity build
Summary:
[libomptarget] Change nvcc compilation to use a unity build

This allows nvcc to inline functions between what would otherwise be distinct
translation units, which in turn removes any runtime cost from implementing
functions in source files (as opposed to inline in headers).

This will then allow the circular dependencies in deviceRTL to be readily
broken and individual components more easily shared between architectures.

Reviewers: ABataev, jdoerfert, grokos, RaviNarayanaswamy, hfinkel, ronlieb, gregrodgers

Reviewed By: jdoerfert

Subscribers: mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69489
2019-10-31 01:58:51 +00:00
Jon Chesterfield 8548e2f543 [nfc][libomptarget] Move named_sync() into target_impl
Summary: [nfc][libomptarget] Move named_sync() into target_impl

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: ABataev

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69487
2019-10-30 16:25:05 +00:00
David Carlier 5069928487 [OpenMP] Reset affinity mask in the process child on FreeBSD
Reviewers: dim, chandlerc, jdoerfert

Reviewed By: dim

Differential Revision: https://reviews.llvm.org/D69047
2019-10-30 14:51:22 +00:00
Jon Chesterfield 74bb5ee674 [nfc][libomptarget] Move smid() into target_impl
Summary: [nfc][libomptarget] Move smid() into target_impl

Reviewers: ABataev, jdoerfert, grokos

Reviewed By: ABataev

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69485
2019-10-30 13:39:15 +00:00
Jon Chesterfield 62a161cc00 [libomptarget] Always call malloc, free via SafeMalloc, SafeFree wrapper
Summary:
[libomptarget] Always call malloc, free via SafeMalloc, SafeFree wrapper

NFC for release, adds some verbosity to debug printing. Motivation is to provide
one place where local modifications can be made to the behaviour of all heap
allocation or deallocation while debugging.

Reviewers: jdoerfert, ABataev, grokos

Reviewed By: ABataev

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69492
2019-10-30 13:35:34 +00:00
AndreyChurbanov 27f6eedc57 Enable OpenBSD support.
Patch by devnexen (David CARLIER)

Differential Revision: https://reviews.llvm.org/D69220
2019-10-30 12:37:44 +03:00
Alexey Bataev d7941a6ab9 [LIBOMPTARGET]Fix build, NFC.
Need to include nvptx_interface.h in target_impl.h, otherwise the build
is failed because of missing __kmpc_impl_lanemask_t type.
2019-10-28 10:43:00 -04:00
Jon Chesterfield 174967f153 [nfc][libomptarget] Decrease coupling between files
Summary:
[nfc][libomptarget] Decrease coupling between files

debug.h used the symbol omptarget_device_environment so implicitly required
an include of omptarget-nvptx.h to compile. Similarly interface.h uses size_t.

Moving this declaration to a new header means cancel, critical can now build
without omptarget-nvptx.h. After this change, debug.h, cancel.cu, critical.cu
could move under a common source directory.

Reviewers: ABataev, jdoerfert, grokos

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69473
2019-10-27 14:27:54 +00:00
Jon Chesterfield ad4c42666d [nfc][libomptarget] Inline option into target_impl
Summary:
[nfc][libomptarget] Inline option into target_impl

Subset of D69423. The macros that were in option.h are all target dependent.
Inlining the header simplifies the dependency graph when looking to move code
into a common subdir.

Reviewers: ABataev, jdoerfert, grokos

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69472
2019-10-27 14:26:55 +00:00
Jon Chesterfield f7c3c640af [NFC][libomptarget]Remove TRUE,FALSE macros from option.h
Summary:
[NFC][libomptarget]Remove TRUE,FALSE macros from option.h
Subset of D69423. Patch series ends with removing option.h.

Reviewers: ABataev, jdoerfert, grokos

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69463
2019-10-27 01:31:12 +01:00
Jon Chesterfield 197b7b24c3 [NFC][libomptarget] move remaining device specific code out of omptarget-nvptx.h
Summary:
[NFC][libomptarget] move remaining device specific code out of omptarget-nvptx.h

Strictly there is one remaining difference wrt amdgcn - parallelLevel is
volatile qualified on amdgcn and not on nvptx. Determining whether this is
correct - and how to represent the different semantics of 'volatile' under
various conditions - is beyond the scope of this code motion patch.

Reviewers: ABataev, jdoerfert, grokos

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D69424
2019-10-25 18:58:31 +01:00
AndreyChurbanov be29d92854 OpenMP Tasks dependencies hash re-sizing fixed.
Details:
- nconflicts field initialized;
- formatting fix (moved declaration out of the long line);
- count conflicts in new hash as opposed to old one.

Differential Revision: https://reviews.llvm.org/D68036
2019-10-25 16:04:46 +03:00
Stephan T. Lavavej 2e4f1e112d [www] Change URLs to HTTPS.
This changes most URLs in llvm's html files to HTTPS. Most changes were
search-and-replace with manual verification; some changes were manual.
For a few URLs, the websites were performing redirects or had changed
their anchors; I fixed those up manually. This consistently uses the
official https://wg21.link redirector. This also strips trailing
whitespace and fixes a couple of typos.

Fixes D69363.

There are a very small number of dead links for which I don't know any
replacements (they are equally dead as HTTP or HTTPS):

https://llvm.org/cmds/llvm2cpp.html
https://llvm.org/devmtg/2010-11/videos/Grosser_Polly-desktop.mp4
https://llvm.org/devmtg/2010-11/videos/Grosser_Polly-mobile.mp4
https://llvm.org/devmtg/2011-11/videos/Grosser_PollyOptimizations-desktop.mov
https://llvm.org/devmtg/2011-11/videos/Grosser_PollyOptimizations-mobile.mp4
https://llvm.org/perf/db_default/v4/nts/22463
https://polly.llvm.org/documentation/memaccess.html
2019-10-24 13:25:15 -07:00
Jon Chesterfield d69d1aa131 [libomptarget][nfc] Make interface.h target independent
Summary:
[libomptarget][nfc] Make interface.h target independent

Move interface.h under a top level include directory.
Remove #includes to avoid the interface depending on the implementation.

Reviewers: ABataev, jdoerfert, grokos, ronlieb, RaviNarayanaswamy

Reviewed By: jdoerfert

Subscribers: mgorny, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D68615

llvm-svn: 374919
2019-10-15 17:15:26 +00:00
David Carlier f61f13d4e7 [OpenMP] Enable thread affinity on FreeBSD
Reviewers: chandlerc, jlpeyton, jdoerfert, dim

Reviewed-By: dim

Differential Revision: https://reviews.llvm.org/D68580

llvm-svn: 374118
2019-10-08 21:25:30 +00:00
Andrey Churbanov ca2973bb20 Don't assume Type from `readelf -d` has parentheses
Patch by jbeich (Jan Beich)

Differential Revision: https://reviews.llvm.org/D68053

llvm-svn: 374038
2019-10-08 12:39:04 +00:00
Andrey Churbanov f34271d886 Don't link libm with -Wl,--as-needed on FreeBSD
Patch by jbeich (Jan Beich)

Differential Revision: https://reviews.llvm.org/D68051

llvm-svn: 374037
2019-10-08 12:23:25 +00:00
Jon Chesterfield 58fd6b5b9c [libomptarget][nfc] Update remaining uint32 to use lanemask_t
Summary:
[libomptarget][nfc] Update remaining uint32 to use lanemask_t

Update a few functions in the API to use lanemask_t instead of i32. NFC for
nvptx. Also update the ActiveThreads type in DataSharingStateTy.
This removes a lot of #ifdef from the downsteam amdgcn implementation.

Reviewers: ABataev, jdoerfert, grokos, ronlieb, RaviNarayanaswamy

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D68513

llvm-svn: 373806
2019-10-04 22:30:28 +00:00
Jon Chesterfield 4f75a73796 Use named constant to indicate all lanes, to handle 32 and 64 wide architectures
Summary: Use named constant to indicate all lanes, to handle 32 and 64 wide architectures

Reviewers: ABataev, jdoerfert, grokos, ronlieb

Reviewed By: grokos

Subscribers: ronlieb, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D68369

llvm-svn: 373793
2019-10-04 21:39:22 +00:00
David Carlier fef62e1a68 [OpenMP] FreeBSD address check if mapped more native
/proc unless Linux layer compatibility is activated for CentOS is activated is not present
thus relying on a more native for checking the address.

Reviewers: Hahnfeld, kongyl, jdoerfert, jlpeyton, AndreyChurbanov, emaster, dim

Reviewed By: Hahnfeld

Differential Revision: https://reviews.llvm.org/D67326

llvm-svn: 373152
2019-09-28 19:01:59 +00:00
Sergey Dmitriev 4b343fd84c [Clang][OpenMP Offload] Create start/end symbols for the offloading entry table with a help of a linker
Linker automatically provides __start_<section name> and __stop_<section name> symbols to satisfy unresolved references if <section name> is representable as a C identifier (see https://sourceware.org/binutils/docs/ld/Input-Section-Example.html for details). These symbols indicate the start address and end address of the output section respectively. Therefore, renaming OpenMP offload entries section name from ".omp.offloading_entries" to "omp_offloading_entries" to use this feature.

This is the first part of the patch for eliminating OpenMP linker script (please see https://reviews.llvm.org/D64943).

Differential Revision: https://reviews.llvm.org/D68070

llvm-svn: 373118
2019-09-27 20:00:51 +00:00
Andrey Churbanov de44f434e8 fixed test: eliminated race condition which might cause deadlock
llvm-svn: 372887
2019-09-25 15:25:52 +00:00
Andrey Churbanov a1639b9bba Enable tasks dependencies hashmaps resizing.
Patch by viroulep (Philippe Virouleau)

Differential Revision: https://reviews.llvm.org/D67447

llvm-svn: 372879
2019-09-25 14:40:19 +00:00
Jonas Hahnfeld 673e5476a8 [OpenMP] Change initialization of __kmp_global
There's no need to initialize variables with static storage duration
because they're implicitly initialized to zero. See
https://en.cppreference.com/w/c/language/initialization#Implicit_initialization

I think that's already relied upon because the supplied 0 only sets
'kmp_time_global_t g_time;' in 'struct kmp_base_global'. The other fields
are not set in the code, but implicitly initialized by the compiler.

Differential Revision: https://reviews.llvm.org/D66292

llvm-svn: 370943
2019-09-04 17:47:37 +00:00
Alexey Bataev 4812941776 [OPENMP][NVPTX]Fix parallel level counter in non-SPMD mode.
Summary:
In non-SPMD mode we may end up with the divergent threads when trying to
increment/decrement parallel level counter. It may lead to incorrect
calculations of the parallel level and wrong results when threads are
divergent. We need to reconverge the threads before trying to modify the
parallel level counter.

Reviewers: grokos, jdoerfert

Subscribers: guansong, openmp-commits, caomhin, kkwli0

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66802

llvm-svn: 370803
2019-09-03 18:11:50 +00:00
Jon Chesterfield bbdd282371 [libomptarget] Refactor activemask macro to inline function
Summary:
[libomptarget] Refactor activemask macro to inline function
See also abandoned D66846, split into this diff and others.

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Reviewed By: jdoerfert, ABataev

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66851

llvm-svn: 370781
2019-09-03 16:31:30 +00:00
Jon Chesterfield 3294421926 Use target_impl functions to replace more inline asm
Summary:
Use target_impl functions to replace more inline asm
Follow on from D65836. Removes remaining asm shuffles and lanemask accessors
Also changes the types of target_impl bitwise functions to unsigned.

Reviewers: jdoerfert, ABataev, grokos, Hahnfeld, gregrodgers, ronlieb, hfinkel

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66809

llvm-svn: 370216
2019-08-28 15:04:06 +00:00
Jon Chesterfield 80f9a38a76 [libomptarget] Refactor syncthreads macro to inline function
Summary:
[libomptarget] Refactor syncthreads macro to inline function
See also abandoned D66846, split into this diff and others.

Rev 2 of D66855

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66861

llvm-svn: 370210
2019-08-28 14:22:35 +00:00
Jon Chesterfield be3d487313 [libomptarget] Refactor syncwarp macro to inline function
Summary:
[libomptarget] Refactor syncwarp macro to inline function
See also abandoned D66846, split into this diff and others.

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66857

llvm-svn: 370149
2019-08-28 02:02:53 +00:00
Jon Chesterfield e73e3013a6 Fix build break due to close brace lost in merge
llvm-svn: 370148
2019-08-28 01:56:26 +00:00
Jon Chesterfield 327aa81123 [libomptarget] Refactor shfl_down_sync macro to inline function
Summary:
[libomptarget] Refactor shfl_down_sync macro to inline function
See also abandoned D66846, split into this diff and others.

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66853

llvm-svn: 370146
2019-08-28 01:47:41 +00:00
Jon Chesterfield b9b712df82 [libomptarget] Refactor shfl_sync macro to inline function
Summary:
[libomptarget] Refactor shfl_sync macro to inline function
See also abandoned D66846, split into this diff and others.

Reviewers: jdoerfert, ABataev, grokos, ronlieb, gregrodgers

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66852

llvm-svn: 370144
2019-08-28 01:31:04 +00:00
Alexey Bataev da8b5cc9f1 [OPENMP][NVPTX]Add __kmpc_syncwarp(int32_t) function.
Summary:
Added function void __kmpc_syncwarp(int32_t) to expose it to the
compiler. It is required to fix the problem with the critical regions in
Cuda9.0+. We cannot use barrier in the critical region, but still need
to reconverge the threads in the warp after. This function allows to do
this.

Reviewers: grokos, jdoerfert

Subscribers: guansong, openmp-commits, kkwli0, caomhin

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D66672

llvm-svn: 369933
2019-08-26 17:32:45 +00:00
Alexey Bataev 0366168f3a [OPENMP][NVPTX]Use __syncwarp() to reconverge the threads.
Summary:
In Cuda 9.0 it is not guaranteed that threads in the warps are
convergent. We need to use __syncwarp() function to reconverge
the threads and to guarantee the memory ordering among threads in the
warps.
This is the first patch to fix the problem with the test
libomptarget/deviceRTLs/nvptx/src/sync.cu on Cuda9+.
This patch just replaces calls to __shfl_sync() function with the call
of __syncwarp() function where we need to reconverge the threads when we
try to modify the value of the parallel level counter.

Reviewers: grokos

Subscribers: guansong, jfb, jdoerfert, caomhin, kkwli0, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D65013

llvm-svn: 369796
2019-08-23 18:34:48 +00:00
Jonathan Peyton 57ae6b8e37 Force honoring nthreads-var and thread-limit-var inside teams construct on host
This patch fixes https://bugs.llvm.org/show_bug.cgi?id=42906, via adding
adjustment of number of threads on enter to the teams construct on host
according to user settings. This allows to pass checks and avoid assertions
at time of team of threads creation.

Patch by Andrey Churbanov

Differential Revision: https://reviews.llvm.org/D66351

llvm-svn: 369430
2019-08-20 19:39:17 +00:00
Jonas Hahnfeld d2ae0c4f44 [OpenMP] Enable warning about "implicit fallthrough"
Fix last warned location in ittnotify_static.cpp using the defined
macro KMP_FALLTHROUGH().

Differential Revision: https://reviews.llvm.org/D65871

llvm-svn: 369003
2019-08-15 13:26:55 +00:00
Jonas Hahnfeld 4d77e50e6e [OpenMP] Remove 'unnecessary parentheses'
The variables in kmp_lock.cpp are really arrays of function pointers
that return void or int, not pointers to functions that return void*
or int*. The other changes are only cosmetic.

Differential Revision: https://reviews.llvm.org/D65870

llvm-svn: 369002
2019-08-15 13:26:41 +00:00
Jonas Hahnfeld fb72a03f85 [OMPT] Resolve warnings because of ints in if conditions
The implementation status can only be one of
ompt_event_UNIMPLEMENTED = ompt_set_never = 1
ompt_event_MAY_ALWAYS = ompt_set_always = 5

In both cases, the condition was already true, so just remove
the check.

Differential Revision: https://reviews.llvm.org/D65869

llvm-svn: 369001
2019-08-15 13:26:29 +00:00
Jonas Hahnfeld dc23c832f4 [OpenMP] Turn on -Wall compiler warnings by default
Instead, maintain a list of disabled options to still build libomp and
libomptarget without warnings. This includes -Wno-error and -Wno-pedantic
to silence warnings that LLVM enables when building in-tree.

I tested the following compilers:
 * Clang 6.0, 7.0, 8.0
 * GCC 4.8.5 (CentOS 7), GCC 6, 7, 8, 9
 * Intel Compiler 16, 17, 18, 19

RFC thread on openmp-dev mailing list:
http://lists.llvm.org/pipermail/openmp-dev/2019-August/002668.html

Differential Revision: https://reviews.llvm.org/D65867

llvm-svn: 368999
2019-08-15 13:11:50 +00:00
Jon Chesterfield ed3324f6b6 Factor architecture dependent code out of loop.cu
Summary:
[libomptarget] Factor architecture dependent code out of loop.cu

Related to the patch series starting D64217. Added subscribers to said series as reviewers. This effort is smaller in scope.

This patch factors out just enough architecture dependent code from loop.cu to allow the same source to be used with amdgcn, given a different target_impl.h. Testing is that the same bitcode (modulo variable names) is generated for libomptarget before and after the refactor, for nvptx and the out of tree amdgcn.

Reviewers: jdoerfert, ABataev, bollu, jfb, tra, grokos, Hahnfeld, guansong, xtian, gregrodgers, ronlieb, hfinkel, gtbercea, guraypp, arpith-jacob

Reviewed By: jdoerfert, ABataev

Subscribers: dexonsmith, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D65836

llvm-svn: 368751
2019-08-13 21:41:47 +00:00
Andrey Churbanov 5eec1a9d32 Cleanup unused variable.
This patch fixes problem raised in post-review comments of the
https://reviews.llvm.org/D65285. Developers of ittnotify confirmed
that dll_path_ptr field of the __itt_global structure is never used
by ittnotify library, so it is safe to remove the dll_path array.

Differential Revision: https://reviews.llvm.org/D65885

llvm-svn: 368559
2019-08-12 12:37:30 +00:00
Gheorghe-Teodor Bercea 6c7b882e52 [OpenMP][libomptarget] Add support for close map modifier
Summary:
This patch adds support for the close map modifier.

The close map modifier will overwrite the unified shared memory requirement and create a device copy of the data.

Reviewers: ABataev, Hahnfeld, caomhin, grokos, jdoerfert, AlexEichenberger

Reviewed By: Hahnfeld, AlexEichenberger

Subscribers: guansong, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D65340

llvm-svn: 368488
2019-08-09 21:32:57 +00:00
Jonas Hahnfeld 7a0f2dc5a4 [libomptarget] Remove duplicate RTLRequiresFlags per device
We have one global RTLs.RequiresFlags, I don't see a need to make a
copy per device that the runtime manages. This was problematic anyway
because the copy happened during the first __tgt_register_lib(). This
made it impossible to call __tgt_register_requires() from normal user
funtions for testing.
Hence, this change also fixes unified_shared_memory/shared_update.c for
older versions of Clang that don't call __tgt_register_requires() before
__tgt_register_lib().

Differential Revision: https://reviews.llvm.org/D66019

llvm-svn: 368465
2019-08-09 19:20:39 +00:00
Gheorghe-Teodor Bercea a1d20506e7 [OpenMP][libomptarget] Add support for unified memory for regular maps
Summary:
This patch adds support for using unified memory in the case of regular maps that happen when a target region is offloaded to the device.

For cases where only a single version of the data is required then the host address can be used. When variables need to be privatized in any way or globalized, then the copy to the device is still required for correctness.

Reviewers: ABataev, jdoerfert, Hahnfeld, AlexEichenberger, caomhin, grokos

Reviewed By: Hahnfeld

Subscribers: mgorny, guansong, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D65001

llvm-svn: 368192
2019-08-07 17:29:45 +00:00
Jon Chesterfield ae0178bee7 Use forceinline. Necessary for nvcc to inline small functions within the bitcode library
Summary:
[libomptarget] Use forceinline. Necessary for nvcc to inline small functions within the bitcode library
Suggested in D65836

Reviewers: ABataev, jdoerfert, grokos, gregrodgers

Subscribers: openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D65876

llvm-svn: 368177
2019-08-07 15:24:12 +00:00
Alexey Bataev c10180ed8e [OPENMP][OFFLOADING]Fix the test, NFC.
llvm-svn: 368068
2019-08-06 18:13:39 +00:00
Jonathan Peyton 73d5abd809 [OpenMP] Add support for GOMP_*_nonmonotonic_* functions
Patch by Isuru Fernando

Differential Revision: https://reviews.llvm.org/D65714

llvm-svn: 367949
2019-08-05 23:23:52 +00:00
Hansang Bae dcdbe6515b [OpenMP] Fix broken build due to new OMPT tests
New OMPT tests with teams construct should be disabled for GCC as it
emits code with a GOMP entry not supported in the LLVM runtime.

Differential Revision: https://reviews.llvm.org/D65757

llvm-svn: 367939
2019-08-05 21:46:13 +00:00
Michael Kruse 78769ec403 [libomptarget] Harmonize emitting CUDA errors and general debug messages.
Ensures that CUDA fail reasons (such as "No CUDA-capable device detected")
are printed together with libomptarget's debug message
(e.g. "Error when setting CUDA context"). Previously, the former was
printed only in CMAKE_BUILD_TYPE=Debug builds while the latter was
enabled by LIBOMPTARGET_ENABLE_DEBUG.

With this change, also only call cuGetErrorString when the error will be
printed.

Suggested-by: Ye Luo <xw111luoye@gmail.com>

Differential Revision: https://reviews.llvm.org/D65687

llvm-svn: 367910
2019-08-05 19:12:10 +00:00
Michael Kruse 2c7a8eaf3d [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
Hansang Bae 67e93a1ae0 Add OMPT support for teams construct
This change adds OMPT support for events from teams construct.

Differential Revision: https://reviews.llvm.org/D64025

llvm-svn: 367746
2019-08-03 02:38:53 +00:00
Jonas Hahnfeld 52b87ac32f [OpenMP] Rename last file to cpp and remove LIBOMP_CFLAGS
All other files are already C++ and the build system has always
passed '-x c++' for C files, effectively compiling them as C++.

To stay warning free we need one fix in ittnotify_static.{c,cpp}:
The variable dll_path can be written to, so it must not be const.
GCC complained with -Wcast-qual and I think it's right.

Differential Revision: https://reviews.llvm.org/D65285

llvm-svn: 367343
2019-07-30 18:37:28 +00:00
Yi Kong 3d21a3af87 [openmp] Workaround bug in old Android pthread_attr_setstacksize
Round the stack size to a multiple of the page size. Older versions of
Android (until KitKat) would fail pthread_attr_setstacksize with
EINVAL if the stack size was not a multiple of the page size.

Patch by Dan Albert <danalbert@google.com>.

Test: Build, copied into the NDK, passed openmp test on ICS.
Bug: https://github.com/android-ndk/ndk/issues/9
llvm-svn: 367070
2019-07-25 22:29:55 +00:00
Jonas Hahnfeld baeab1fc44 [OpenMP] Fix build of stubs library, NFC.
Both Clang and GCC complained that they cannot initialize a return
object of type 'kmp_proc_bind_t' with an 'int'. While at it, also
fix a warning about missing parentheses thrown by Clang.

Differential Revision: https://reviews.llvm.org/D65284

llvm-svn: 367041
2019-07-25 17:51:24 +00:00
Alexey Bataev ca424d100c [OPENMP][NVPTX]Perform memory flush if number of threads to sync is 1 or less.
Summary:
According to the OpenMP standard, barrier operation must perform
implicit flush operation. Currently, if there is only one thread in the
team, barrier does not flush the memory. Patch fixes this problem.

Reviewers: grokos, gtbercea, kkwli0

Subscribers: guansong, jdoerfert, openmp-commits, caomhin

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D62398

llvm-svn: 367024
2019-07-25 15:02:28 +00:00
Jonas Hahnfeld 2488ae9df1 [OpenMP] RISCV64 port
This is a port of libomp for the RISC-V 64-bit Linux target.

We have tested this port on a HiFive Unleashed development board
using a downstream LLVM that has support for the missing bits in
upstream. As of now, all tests are passing, including OMPT.

Patch by Ferran Pallarès!

Differential Revision: https://reviews.llvm.org/D59880

llvm-svn: 367021
2019-07-25 14:36:20 +00:00
Jonas Hahnfeld 6e40ae8f3d [libomptarget] Handle offload policy in push_tripcount
If the first target region in a program calls the push_tripcount
function, libomptarget didn't handle the offload policy correctly.
This could lead to unexpected error messages as seen in
http://lists.llvm.org/pipermail/openmp-dev/2019-June/002561.html

To solve this, add a check calling IsOffloadDisabled() as all other
entry points already do. If this method returns false, libomptarget
is effectively disabled.

Differential Revision: https://reviews.llvm.org/D64626

llvm-svn: 366810
2019-07-23 14:20:48 +00:00
Jonas Hahnfeld a2748c74d6 [OMPT] Cleanup reset of exit_frame pointer
This is done at call-site and does not need to be handled in
__kmp_invoke_microtask. It was already absent from the x86
and x86_64 assembly, this patch removes it from the generic
implementation in z_Linux_util.cpp and adds documentation for
AArch64 and PPC64 that it's actually not needed. I can't test
on these architectures, so I don't want to change the code just
because it looks right :)

While at it, rename some variables for consistency and add a
check in test/ompt/parallel/normal.c that the pointer was reset
before entering the barrier.

Differential Revision: https://reviews.llvm.org/D64442

llvm-svn: 366721
2019-07-22 18:46:02 +00:00
Jonas Hahnfeld 4138b2f167 Delete empty file
This is a left-over from r356288 which was reviewed in D58989.

llvm-svn: 366716
2019-07-22 18:11:06 +00:00
Alexey Bataev da43861b4a [OpenMP][libomptarget] Suppress C++ 11 related warnings when building libomptarget-nvptx bitcode library, by Doru Bercea.
Summary: Pass -std=c++11 flag to compiler to suppress C++ 11 related warnings when building NVPTX bitcode library.

Reviewers: ABataev, caomhin, Hahnfeld

Reviewed By: ABataev, Hahnfeld

Subscribers: jdoerfert, Hahnfeld, jholewinski, mgorny, guansong, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D55772

llvm-svn: 366438
2019-07-18 13:54:01 +00:00
Ron Lieberman 59532488b1 [OPENMP] Resolve lost LoopTripCnt for subsequent loops in same thread.
Remove loopTripCnt from threaded device stack after consuming it.
Added a libomptarget DP message to aid in future debugging and to
validate the added testcase, which only runs in Debug build.

Differential Revision: https://reviews.llvm.org/D64808

llvm-svn: 366349
2019-07-17 17:07:52 +00:00
Jonathan Peyton aa5cdafa40 Remove REQUIRES OMP spec version within lit tests
This is a follow up patch to D64534 (r365963) which removed all OMP
spec versioning within the OpenMP runtime codebase.  This patch removes
REQUIRES: openmp-x.y lines from lit tests.

llvm-svn: 366341
2019-07-17 15:41:00 +00:00
Jonas Hahnfeld 1ff5535785 [OpenMP] Move header inclusion out of 'extern "C"'
This leads to problems when compiling C++ code with libc++ for Nvidia GPUs
because Clang now uses wrappers for math functions that might include
C++ templates not allowed in 'extern "C"'.

Differentiel Revision: https://reviews.llvm.org/D64625

llvm-svn: 366229
2019-07-16 17:16:43 +00:00
Alexey Bataev 85b9651edd [OPENMP][NVPTX]Fixed checks for cuda versions.
Summary:
We used CUDART_VERSION macro to check for the installed cuda version
but this macro is defined in cuda_runtime_api.h, which is not used by
project. Better to use CUDA_VERSION macro, which is defined in cuda.h.
Also, added the check if this macro is defined. If macro is undefined,
there is something wrong with the cuda configuration and we should not
continue the compilation.
This also fixes problems with runtime building in cuda 10+.

Reviewers: grokos

Subscribers: guansong, jdoerfert, caomhin, kkwli0, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D64648

llvm-svn: 366224
2019-07-16 16:07:10 +00:00
Alexey Bataev 42816107f7 [OPENMP]Fix threadid in __kmpc_omp_taskwait call for dependent target calls.
Summary:
We used to call __kmpc_omp_taskwait function with global threadid set to
0. It may crash the application at the runtime if the thread executing
 target region is not a master thread.

Reviewers: grokos, kkwli0

Subscribers: guansong, jdoerfert, caomhin, openmp-commits

Tags: #openmp

Differential Revision: https://reviews.llvm.org/D64571

llvm-svn: 366220
2019-07-16 15:51:32 +00:00
Jonathan Peyton e4b4f994d2 [OpenMP] Remove OMP spec versioning
Remove all older OMP spec versioning from the runtime and build system.

Patch by Terry Wilmarth

Differential Revision: https://reviews.llvm.org/D64534

llvm-svn: 365963
2019-07-12 21:45:36 +00:00
Jonas Hahnfeld aca476b296 [libomptarget] Fix typos and grammar in error messages, NFC.
llvm-svn: 365890
2019-07-12 10:21:55 +00:00
Jonas Hahnfeld 2dfc5179f6 [libomptarget-nvptx] Remove dead functions
These entry points are never called by Clang trunk nor clang-ykt. If
XL doesn't use them either, they can finally go away.

Differential Revision: https://reviews.llvm.org/D52700

llvm-svn: 365817
2019-07-11 20:12:51 +00:00