[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 ho

[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

show more ...

[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 runtim

[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

show more ...

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.

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

show more ...

[clang] Don't use Optional::getValue (NFC)

[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.

Diff

[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

show more ...

[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 a

[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

show more ...

[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.

Differ

[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

show more ...

[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

[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

[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

show more ...

[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 fo

[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

show more ...

[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 na

[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

show more ...

[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

[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

show more ...

[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 th

[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

show more ...

[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 i

[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

show more ...

[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
fr

[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

show more ...

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__Z

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

show more ...

[NFCI] Fix set-but-unused warning in CGOpenMPRuntime.cpp

[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 u

[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

show more ...

[CGOpenMPRuntime] Remove uses of deprecated Address constructor

And as these are the last remaining uses, also remove the
constructor itself.

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
a

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
47eb4f7dcd845878b16a53dadd765195b9c24b6e, where the used type is
important for determining alignment.

show more ...

[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

[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.

show more ...

[CGOpenMPRuntime] Reuse getDepobjElements() (NFC)

There were two more places repeating this code, reuse the helper.
This requires moving the static functions into the class.

Revert "[CodeGen] Avoid deprecated Address ctor in EmitLoadOfPointer()"

This reverts commit 767ec883e37510a247ea5695921876ef67cf5b3f.

This results in a some incorrect alignments which are not cover

Revert "[CodeGen] Avoid deprecated Address ctor in EmitLoadOfPointer()"

This reverts commit 767ec883e37510a247ea5695921876ef67cf5b3f.

This results in a some incorrect alignments which are not covered
by existing tests.

show more ...

[CGOpenMPRuntime] Remove some uses of deprecated Adddress ctor

[CGOpenMPRuntime] Remove some uses of deprecated Address ctor