Commit Graph

702 Commits

Author SHA1 Message Date
Vitaly Buka 36c9f5a58b [NFC][OpenMP] Simplify 2f9be69d84 2022-08-17 18:59:48 -07:00
Vitaly Buka 2f9be69d84 [OpenMP] Fix another after scope after D129608
https://lab.llvm.org/buildbot/#/builders/5/builds/26770
2022-08-13 12:13:54 -07:00
Vitaly Buka f385eaf48f [OpenMP] Fix use after scope after D129608
Broken builder https://lab.llvm.org/buildbot/#/builders/5/builds/26764
2022-08-13 09:40:51 -07:00
Jennifer Yu 2ca27206f9 [OpenMP] Fix segmentation fault when data field is used in is_device_pt
Currently, the field just emit map info for this pointer variable. It is
failed at run time. For the fields, the PartialStruct is created and it
needs call to emitCombinedEntry which create the base that covers all
the pieces.

The change is to generate map info as regular fields.

Differential Revision: https://reviews.llvm.org/D129608
2022-08-12 17:10:26 -07:00
Corentin Jabot 127bf44385 [Clang][C++20] Support capturing structured bindings in lambdas
This completes the implementation of P1091R3 and P1381R1.

This patch allow the capture of structured bindings
both for C++20+ and C++17, with extension/compat warning.

In addition, capturing an anonymous union member,
a bitfield, or a structured binding thereof now has a
better diagnostic.

We only support structured bindings - as opposed to other kinds
of structured statements/blocks. We still emit an error for those.

In addition, support for structured bindings capture is entirely disabled in
OpenMP mode as this needs more investigation - a specific diagnostic indicate the feature is not yet supported there.

Note that the rest of P1091R3 (static/thread_local structured bindings) was already implemented.

at the request of @shafik, i can confirm the correct behavior of lldb wit this change.

Fixes https://github.com/llvm/llvm-project/issues/54300
Fixes https://github.com/llvm/llvm-project/issues/54300
Fixes https://github.com/llvm/llvm-project/issues/52720

Reviewed By: aaron.ballman

Differential Revision: https://reviews.llvm.org/D122768
2022-08-04 10:12:53 +02:00
Corentin Jabot a274219600 Revert "[Clang][C++20] Support capturing structured bindings in lambdas"
This reverts commit 44f2baa380.

Breaks self builds and seems to have conformance issues.
2022-08-03 21:00:29 +02:00
Corentin Jabot 44f2baa380 [Clang][C++20] Support capturing structured bindings in lambdas
This completes the implementation of P1091R3 and P1381R1.

This patch allow the capture of structured bindings
both for C++20+ and C++17, with extension/compat warning.

In addition, capturing an anonymous union member,
a bitfield, or a structured binding thereof now has a
better diagnostic.

We only support structured bindings - as opposed to other kinds
of structured statements/blocks. We still emit an error for those.

In addition, support for structured bindings capture is entirely disabled in
OpenMP mode as this needs more investigation - a specific diagnostic indicate the feature is not yet supported there.

Note that the rest of P1091R3 (static/thread_local structured bindings) was already implemented.

at the request of @shafik, i can confirm the correct behavior of lldb wit this change.

Fixes https://github.com/llvm/llvm-project/issues/54300
Fixes https://github.com/llvm/llvm-project/issues/54300
Fixes https://github.com/llvm/llvm-project/issues/52720

Reviewed By: aaron.ballman

Differential Revision: https://reviews.llvm.org/D122768
2022-08-03 20:00:01 +02:00
Gabriel Ravier 5674a3c880 Fixed a number of typos
I went over the output of the following mess of a command:

(ulimit -m 2000000; ulimit -v 2000000; git ls-files -z |
 parallel --xargs -0 cat | aspell list --mode=none --ignore-case |
 grep -E '^[A-Za-z][a-z]*$' | sort | uniq -c | sort -n |
 grep -vE '.{25}' | aspell pipe -W3 | grep : | cut -d' ' -f2 | less)

and proceeded to spend a few days looking at it to find probable typos
and fixed a few hundred of them in all of the llvm project (note, the
ones I found are not anywhere near all of them, but it seems like a
good start).

Differential Revision: https://reviews.llvm.org/D130827
2022-08-01 13:13:18 -04:00
Joseph Huber 5300263c70 [OpenMP] Add loop tripcount argument to kernel launch and remove push function
Previously we added the `push_target_tripcount` function to send the
loop tripcount to the device runtime so we knew how to configure the
teams / threads for execute the loop for a teams distribute construct.
This was implemented as a separate function mostly to avoid changing the
interface for backwards compatbility. Now that we've changed it anyway
and the new interface can take an arbitrary number of arguments via the
struct without changing the ABI, we can move this to the new interface.
This will simplify the runtime by removing unnecessary state between
calls.

Depends on D128550

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D128816
2022-07-08 14:44:16 -04:00
Joseph Huber 1fff116645 [OpenMP] Change OpenMP code generation for target region entries
This patch changes the code we generate to enter a target region on the
device. This is in-line with the new definition in the runtime that was
added previously. Additionally we implement this in the OpenMPIRBuilder
so that this code can be shared with Flang in the future.

Reviewed By: ABataev

Differential Revision: https://reviews.llvm.org/D128550
2022-07-08 14:44:11 -04:00
Ritanya B Bharadwaj 8322fe200d Adding support for target in_reduction
Implementing target in_reduction by wrapping target task with host task with in_reduction and if clause. This is in compliance with OpenMP 5.0 section: 2.19.5.6.
So, this

```
  for (int i=0; i<N; i++) {
    res = res+i
  }
```

will become

```

   #pragma omp task in_reduction(+:res) if(0)
   #pragma omp target map(res)
   for (int i=0; i<N; i++) {
     res = res+i
   }
```

Reviewed By: ABataev

Differential Revision: https://reviews.llvm.org/D125669
2022-06-27 10:36:46 -05:00
Kazu Hirata ca4af13e48 [clang] Don't use Optional::getValue (NFC) 2022-06-20 22:59:26 -07:00
Mike Rice 0a5cfbf7b2 [OpenMP] Use the align clause value from 'omp allocate' for globals
Refactor the code that handles the align clause of 'omp allocate' so
it can be used with globals as well as local variables.

Differential Revision: https://reviews.llvm.org/D126426
2022-05-26 09:51:48 -07:00
Mike Rice 239094cdee [OpenMP] Add codegen for 'omp_all_memory' reserved locator.
This creates an entry with address=nullptr and flag=0x80.
When an 'omp_all_memory' entry is specified any other 'out' or
'inout' entries are not needed and are not passed to the runtime.

Differential Revision: https://reviews.llvm.org/D126321
2022-05-24 15:26:23 -07:00
Mike Rice 9ba937112f [OpenMP] Add parsing/sema support for omp_all_memory reserved locator
Adds support for the reserved locator 'omp_all_memory' for use
in depend clauses with 'out' or 'inout' dependence-types.

Differential Revision: https://reviews.llvm.org/D125828
2022-05-24 10:28:59 -07:00
Mike Rice 772b0c44a4 [OpenMP] Fix mangling for linear parameters with negative stride
The 'n' character is used in place of '-' in the mangled name.

Differential Revision: https://reviews.llvm.org/D125406
2022-05-11 14:02:09 -07:00
Mike Rice 0dbaef61b5 [OpenMP] Fix mangling for linear modifiers with variable stride
This adds support for variable stride with the val, uval, and ref linear
modifiers.  Previously only the no modifer type ls<argno> was supported.

  val  -> Ls<argno>
  uval -> Us<argno>
  ref  -> Rs<argno>

Differential Revision: https://reviews.llvm.org/D125330
2022-05-10 14:12:44 -07:00
Mike Rice 1a02519bc5 [OpenMP] Add mangling support for linear modifiers (ref,uval,val)
Add mangling for linear parameters specified with ref, uval, and val
for 'omp declare simd' vector functions.

Add missing stride for linear this parameters.

Differential Revision: https://reviews.llvm.org/D125269
2022-05-10 09:56:55 -07:00
David Pagan 37471cf2c3 [clang][OpenMP] Local variable alignment incorrect with align clause
If alignment specified with align clause is less than natural alignment for
list item type, the alignment should be set to the natural alignment.

See OMP5.1 specification, page 185, lines 7-10

Differential Revision: https://reviews.llvm.org/D124676
2022-05-03 13:10:01 -07:00
Joseph Huber 643c9b22ef [OpenMP] Make generating offloading entries more generic
This patch moves the logic for generating the offloading entries to the
OpenMPIRBuilder. This makes it easier to re-use in other places, such as
for OpenMP support in Flang or using the same method for generating
offloading entires for other languages like Cuda.

Reviewed By: tianshilei1992

Differential Revision: https://reviews.llvm.org/D123460
2022-04-29 09:14:31 -04:00
Joseph Huber 9d3550c517 [OpenMP] Add AMDGPU calling convention to ctor / dtor functions
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
2022-03-25 22:44:20 -04:00
Joseph Huber 3c6d32ec6c [OpenMP] Make Ctor / Dtor functions have external visibility
The default construction of constructor functions by LLVM tends to make
them have internal linkage. When we call a ctor / dtor function in the
target region we are actually creating a kernel that is called at
registration. Because the ctor is a kernel we need to make sure it's
externally visible so we can actually call it. This prevented AMDGPU
from correctly using constructors while NVPTX could use them simply
because it ignored internal visibility.

Reviewed By: JonChesterfield

Differential Revision: https://reviews.llvm.org/D122504
2022-03-25 22:44:17 -04:00
Joseph Huber b9f67d44ba [OpenMP] Replace device kernel linkage with weak_odr
Currently the device kernels all have weak linkage to prevent linkage
errors on multiple defintions. However, this prevents some optimizations
from adequately analyzing them because of the nature of weak linkage.
This patch replaces the weak linkage with weak_odr linkage so we can
statically assert that multiple declarations of the same kernel will
have the same definition.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D122443
2022-03-25 11:29:15 -04:00
Jennifer Yu a6cdac48ff Eliminate extra set of simd variant function attribute.
Current clang generates extra set of simd variant function attribute
with extra 'v' encoding.
For example:
_ZGVbN2v__Z5add_1Pf vs _ZGVbN2vv__Z5add_1Pf
The problem is due to declaration of ParamAttrs following:
    llvm::SmallVector<ParamAttrTy, 8> ParamAttrs(ParamPositions.size());
where ParamPositions.size() is grown after following assignment:
    Pos = ParamPositions[PVD];
So the PVD is not find in ParamPositions.

The problem is ParamPositions need to set for each FD decl. To fix this

Move ParamPositions's init inside while loop for each FD.

Differential Revision: https://reviews.llvm.org/D122338
2022-03-24 13:27:28 -07:00
Dávid Bolvanský a683ba4ff5 [NFCI] Fix set-but-unused warning in CGOpenMPRuntime.cpp 2022-03-24 07:49:21 +01:00
Joseph Huber 0d16c23af1 [OpenMP] Do not create offloading entries for internal or hidden symbols
Currently we create offloading entries to register device variables with
the host. When we register a variable we will look up the symbol in the
device image and map the device address to the host address. This is a
problem when the symbol is declared with hidden visibility or internal
linkage. This means the symbol is not accessible externally and we
cannot get its address. We should still allow static variables to be
declared on the device, but ew should not create an offloading entry for
them so they exist independently on the host and device.

Fixes #54309

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D122352
2022-03-23 18:27:16 -04:00
Nikita Popov c070d5ceff [CGOpenMPRuntime] Remove uses of deprecated Address constructor
And as these are the last remaining uses, also remove the
constructor itself.
2022-03-23 12:40:44 +01:00
Nikita Popov 8b62dd3cd6 Reapply [CodeGen] Avoid deprecated Address ctor in EmitLoadOfPointer()
This requires some adjustment in caller code, because there was
a confusion regarding the meaning of the PtrTy argument: This
argument is the type of the pointer being loaded, not the addresses
being loaded from.

Reapply after fixing the specified pointer type for one call in
47eb4f7dcd, where the used type is
important for determining alignment.
2022-03-23 12:06:11 +01:00
Nikita Popov 47eb4f7dcd [CGOpenMPRuntime] Specify correct type in EmitLoadOfPointerLValue()
Perform a bitcast first, so we can specify the correct pointer type
inf EmitLoadOfPointerLValue(), rather than using a dummy void pointer.
2022-03-23 11:51:14 +01:00
Nikita Popov ba2be802b0 [CGOpenMPRuntime] Reuse getDepobjElements() (NFC)
There were two more places repeating this code, reuse the helper.
This requires moving the static functions into the class.
2022-03-23 11:31:49 +01:00
Nikita Popov 27f6cee12d Revert "[CodeGen] Avoid deprecated Address ctor in EmitLoadOfPointer()"
This reverts commit 767ec883e3.

This results in a some incorrect alignments which are not covered
by existing tests.
2022-03-23 10:24:39 +01:00
Nikita Popov cd6d9ae263 [CGOpenMPRuntime] Remove some uses of deprecated Adddress ctor 2022-03-22 16:29:35 +01:00
Nikita Popov 4f5640cad3 [CGOpenMPRuntime] Remove some uses of deprecated Address ctor 2022-03-22 15:35:45 +01:00
Nikita Popov 767ec883e3 [CodeGen] Avoid deprecated Address ctor in EmitLoadOfPointer()
This requires some adjustment in caller code, because there was
a confusion regarding the meaning of the PtrTy argument: This
argument is the type of the pointer being loaded, not the addresses
being loaded from.
2022-03-22 09:42:31 +01:00
Nikita Popov b6f85d8539 [CodeGen][OpenMP] Use correct type in EmitLoadOfPointer()
Rather than using a dummy void pointer type, we should specify the
correct private type and perform the bitcast beforehand rather than
afterwards. This way, the Address will have correct alignment
information.
2022-03-21 12:08:05 +01:00
Nikita Popov 52cc65d474 [OpenMPRuntime] Specify correct pointer type
Rather than specifying a dummy type in EmitLoadOfPointer() and
then casting it to the correct one, we should instead specify the
correct type and cast beforehand. Otherwise the computed alignment
will be incorrect.
2022-03-18 14:25:51 +01:00
Johannes Doerfert f02550bdd9 Reapply "[OpenMP][FIX] Allow device constructors for AMD GPU"
This reverts commit a597d6a780 and
reapplies 07b1766461.

In AMD GPU device code the globals are in AS(1). Before, we crashed if
the global was a structure. Now we simply cast away the AS before we
generate the code to initialize the global.

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

Fixes: https://github.com/llvm/llvm-project/issues/54421
2022-03-17 12:53:47 -05:00
Nikita Popov 6c0af92612 [CodeGen] Avoid some pointer element type accesses 2022-03-17 16:36:14 +01:00
Johannes Doerfert a597d6a780 Revert "[OpenMP][FIX] Allow device constructors for AMD GPU"
This reverts commit 07b1766461 as it broke
the buildbots:
    https://lab.llvm.org/buildbot#builders/193/builds/8594
2022-03-16 17:35:54 -05:00
Johannes Doerfert 07b1766461 [OpenMP][FIX] Allow device constructors for AMD GPU
In AMD GPU device code the globals are in AS(1). Before, we crashed if
the global was a structure. Now we simply cast away the AS before we
generate the code to initialize the global.

Differential Revision: https://reviews.llvm.org/D121837
2022-03-16 17:04:28 -05:00
David Blaikie c0a6433f2b Simplify OpenMP Lambda use
* Use default ref capture for non-escaping lambdas (this makes
  maintenance easier by allowing new uses, removing uses, having
  conditional uses (such as in assertions) not require updates to an
  explicit capture list)
* Simplify addPrivate API not to take a lambda, since it calls it
  unconditionally/immediately anyway - most callers are simply passing
  in a named value or short expression anyway and the lambda syntax just
  adds noise/overhead

Reviewed By: ABataev

Differential Revision: https://reviews.llvm.org/D121077
2022-03-07 18:23:20 +00:00
Akira Hatanaka d112cc2756 [NFC][Clang][OpaquePtr] Remove the call to Address::deprecated in
CreatePointerBitCastOrAddrSpaceCast

Differential Revision: https://reviews.llvm.org/D120757
2022-03-02 08:58:00 -08:00
Alexey Bataev d04d9220e1 [OPENMP]Fix PR50347: Mapping of global scope deep object fails.
Changed the we handle llvm::Constants in sizes arrays. ConstExprs and
GlobalValues cannot be used as initializers, need to put them at the
runtime, otherwise there wight be the compilation errors.

Differential Revision: https://reviews.llvm.org/D105297
2022-02-25 10:54:24 -08:00
Alexey Bataev ca6fa71b7e Revert "[OPENMP]Fix PR50347: Mapping of global scope deep object fails."
This reverts commit 638938117a. Need to
fix reported fail https://lab.llvm.org/buildbot/#/builders/193/builds/7496
2022-02-24 12:04:39 -08:00
Alexey Bataev 638938117a [OPENMP]Fix PR50347: Mapping of global scope deep object fails.
Changed the we handle llvm::Constants in sizes arrays. ConstExprs and
GlobalValues cannot be used as initializers, need to put them at the
runtime, otherwise there wight be the compilation errors.

Differential Revision: https://reviews.llvm.org/D105297
2022-02-24 11:49:14 -08:00
Joseph Huber 119d71cb73 [OpenMP][NFC] Address warnings and lint messages in CGOpenMPRuntime
Summary:
This patch addressed the warnings and linting messages for the
CGOpenMPRuntime.cpp file. This was causing some -Werror builds to fail.
2022-02-23 18:07:25 -05:00
Reid Kleckner 1d1b089c5d Fix more unused lambda capture warnings, NFC 2022-02-23 14:07:04 -08:00
Reid Kleckner cd37594c03 Fix unused lambda capture warning, NFC 2022-02-23 14:01:01 -08:00
Joseph Huber 2b97b16f29 [OpenMP] Add option to make offloading mandatory
Currently when we generate OpenMP offloading code we always make
fallback code for the CPU. This is necessary for implementing features
like conditional offloading and ensuring that unhandled pragmas don't
result in missing symbols. However, this is problematic for a few cases.
For offloading tests we can silently fail to the host without realizing
that offloading failed. Additionally, this makes it impossible to
provide interoperabiility to other offloading schemes like HIP or CUDA
because those methods do not provide any such host fallback guaruntee.
this patch adds the `-fopenmp-offload-mandatory` flag to prevent
generating the fallback symbol on the CPU and instead replaces the
function with a dummy global and the failed branch with 'unreachable'.

Reviewed By: ABataev

Differential Revision: https://reviews.llvm.org/D120353
2022-02-23 16:45:36 -05:00
Nikita Popov 5065076698 [CodeGen] Rename deprecated Address constructor
To make uses of the deprecated constructor easier to spot, and to
ensure that no new uses are introduced, rename it to
Address::deprecated().

While doing the rename, I've filled in element types in cases
where it was relatively obvious, but we're still left with 135
calls to the deprecated constructor.
2022-02-17 11:26:42 +01:00