offloading
- This patch adds support for thread_limit clause on target directive according to OpenMP 51 [2.14.5]
- The idea is to create an outer task for target region, when there is a thread_limit clause, and manipulate the thread_limit of task instead. This way, thread_limit will be applied to all the relevant constructs enclosed by the target region.
Differential Revision: https://reviews.llvm.org/D152054
OMPTransformDirectiveScopeRAII doesn't have user-written copy
constructor/assignment operator but it frees memory in the destructor.
Delete these members since doesn't seem that OMPTransformDirectiveScopeRAII
objects are intended for copy.
Reviewed By: tahonermann, ABataev
Differential Revision: https://reviews.llvm.org/D155849
This patch removes the constraint of the `interop` directive where only a single
`depend` clause was allowed.
Differential Revision: https://reviews.llvm.org/D155692
This patch migrates the UseDevicePtr and UseDeviceAddr clause related code for handling privatisation from Clang codegen to the OMPIRBuilder
Depends on D150860
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D152554
This patch renames the `OpenMPIRBuilderConfig` flags to reduce confusion over
their meaning. `IsTargetCodegen` becomes `IsGPU`, whereas `IsEmbedded` becomes
`IsTargetDevice`. The `-fopenmp-is-device` compiler option is also renamed to
`-fopenmp-is-target-device` and the `omp.is_device` MLIR attribute is renamed
to `omp.is_target_device`. Getters and setters of all these renamed properties
are also updated accordingly. Many unit tests have been updated to use the new
names, but an alias for the `-fopenmp-is-device` option is created so that
external programs do not stop working after the name change.
`IsGPU` is set when the target triple is AMDGCN or NVIDIA PTX, and it is only
valid if `IsTargetDevice` is specified as well. `IsTargetDevice` is set by the
`-fopenmp-is-target-device` compiler frontend option, which is only added to
the OpenMP device invocation for offloading-enabled programs.
Differential Revision: https://reviews.llvm.org/D154591
The loop directive is a descriptive construct which allows the compiler
flexibility in how it generates code for the directive's associated
loop(s). See OpenMP specification 5.2 [257:8-9].
Codegen added in this patch for the combined 'loop' directives are:
'target teams loop' -> 'target teams distribute parallel for'
'teams loop' -> 'teams distribute parallel for'
'target parallel loop' -> 'target parallel for'
'parallel loop' -> 'parallel for'
NOTE: The implementation of the 'loop' directive itself is unchanged.
Differential Revision: https://reviews.llvm.org/D145823
Partial progress towards replacing uses of CreateElementBitCast, as it
no longer does what its name suggests.
Reviewed By: barannikov88
Differential Revision: https://reviews.llvm.org/D154229
* Add `Address::withElementType()` as a replacement for
`CGBuilderTy::CreateElementBitCast`.
* Partial progress towards replacing `CreateElementBitCast`, as it no
longer does what its name suggests. Either replace its uses with
`Address::withElementType()`, or remove them if no longer needed.
* Remove unused parameter 'Name' of `CreateElementBitCast`
Reviewed By: barannikov88, nikic
Differential Revision: https://reviews.llvm.org/D153196
This patch prefixes omp outlined helpers and reduction funcs
with the original function's name.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D140722
This patch attempts to prefix omp outlined helpers and reduction funcs
with the original function's name.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D140722
Code generation support for 'parallel masked' directive.
The `EmitOMPParallelMaskedDirective` was implemented.
In addition, the appropiate device functions were added.
Fix#59939.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D143527
If use_device_ptr/use_device_addr clauses are used on target data
directive and no device was specified during the compilation, only host
part should be emitted. But it still required to emit captured decls for
partially mapped data fields.
Differential Revision: https://reviews.llvm.org/D144993
There was an assertion triggering when invoking a captured member whose
initializer was in a blase class. This patch fixes it by allowing the
assertion on implicit casts to the base class rather than only the base
class itself.
Fixes https://github.com/llvm/llvm-project/issues/61027
Reviewed By: tianshilei1992
Differential Revision: https://reviews.llvm.org/D144873
This looked like a plausibly correct out of tree patch.
The changed testcases with the pragmas stripped out only use
inbounds GEPs so I assume this is correct.
Support for taskwait nowait clause with placeholder for runtime changes.
Reviewed By: cchen, ABataev
Differential Revision: https://reviews.llvm.org/D131830
Added codegen for `omp error` directive.
This is to generate IR to call:
void __kmpc_error(ident_t *loc, int severity, const char *message);
Differential Revision: https://reviews.llvm.org/D139166
This patch mechanically replaces None with std::nullopt where the
compiler would warn if None were deprecated. The intent is to reduce
the amount of manual work required in migrating from Optional to
std::optional.
This is part of an effort to migrate from llvm::Optional to
std::optional:
https://discourse.llvm.org/t/deprecating-llvm-optional-x-hasvalue-getvalue-getvalueor/63716
Error directive is allowed in both declared and executable contexts.
The function ActOnOpenMPAtClause is called in both places during the
parsers.
Adding a param "bool InExContext" to identify context which is used to
emit error massage.
Differential Revision: https://reviews.llvm.org/D137851
Currently generation of align assumptions for OpenMP simd construct is done
outside OMPIRBuilder for C code and it is not supported for Fortran.
According to OpenMP 5.0 standard (2.9.3) only pointers and arrays can be
aligned for C code.
If given aligned variable is pointer, then Clang generates the following set
of the LLVM IR isntructions to support simd align clause:
; memory allocation for pointer address:
%A.addr = alloca ptr, align 8
; some LLVM IR code
; Alignment instructions (alignment is equal to 32):
%0 = load ptr, ptr %A.addr, align 8
call void @llvm.assume(i1 true) [ "align"(ptr %0, i64 32) ]
If given aligned variable is array, then Clang generates the following set
of the LLVM IR isntructions to support simd align clause:
; memory allocation for array:
%B = alloca [10 x i32], align 16
; some LLVM IR code
; Alignment instructions (alignment is equal to 32):
%arraydecay = getelementptr inbounds [10 x i32], ptr %B, i64 0, i64 0
call void @llvm.assume(i1 true) [ "align"(ptr %arraydecay, i64 32) ]
OMPIRBuilder was modified to generate aligned assumptions. It generates only
llvm.assume calls. Frontend is responsible for generation of aligned pointer
and getting the default alignment value if user does not specify it in aligned
clause.
Unit and regression tests were added to check if aligned clause was handled correctly.
Differential Revision: https://reviews.llvm.org/D133578
Reviewed By: jdoerfert
If 'order(concurrent)' clause is specified, then the iterations of SIMD loop
can be executed concurrently.
This patch adds support for LLVM IR codegen via OMPIRBuilder for SIMD loop
with 'order(concurrent)' clause. The functionality added to OMPIRBuilder is
similar to the functionality implemented in 'CodeGenFunction::EmitOMPSimdInit'.
Reviewed By: jdoerfert
Differential Revision: https://reviews.llvm.org/D134046
Signed-off-by: Dominik Adamski <dominik.adamski@amd.com>
Parallel regions are outlined as functions with capture variables explicitly generated as distinct parameters in the function's argument list. That complicates the fork_call interface in the OpenMP runtime: (1) the fork_call is variadic since there is a variable number of arguments to forward to the outlined function, (2) wrapping/unwrapping arguments happens in the OpenMP runtime, which is sub-optimal, has been a source of ABI bugs, and has a hardcoded limit (16) in the number of arguments, (3) forwarded arguments must cast to pointer types, which complicates debugging. This patch avoids those issues by aggregating captured arguments in a struct to pass to the fork_call.
Reviewed By: jdoerfert, jhuber6, ABataev
Differential Revision: https://reviews.llvm.org/D102107
This patch adds OMPIRBuilder support for the safelen clause for the
simd directive.
Reviewed By: shraiysh, Meinersbur
Differential Revision: https://reviews.llvm.org/D131526
Currently if an OpenMP program uses `linear` clause, and is compiled with
optimization, `llvm.lifetime.end` for variables listed in `linear` clause are
emitted too early such that there could still be uses after that. Let's take the
following code as example:
```
// loop.c
int j;
int *u;
void loop(int n) {
int i;
for (i = 0; i < n; ++i) {
++j;
u = &j;
}
}
```
We compile using the command:
```
clang -cc1 -fopenmp-simd -O3 -x c -triple x86_64-apple-darwin10 -emit-llvm loop.c -o loop.ll
```
The following IR (simplified) will be generated:
```
@j = local_unnamed_addr global i32 0, align 4
@u = local_unnamed_addr global ptr null, align 8
define void @loop(i32 noundef %n) local_unnamed_addr {
entry:
%j = alloca i32, align 4
%cmp = icmp sgt i32 %n, 0
br i1 %cmp, label %simd.if.then, label %simd.if.end
simd.if.then: ; preds = %entry
call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %j)
store ptr %j, ptr @u, align 8
call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %j)
%0 = load i32, ptr %j, align 4
store i32 %0, ptr @j, align 4
br label %simd.if.end
simd.if.end: ; preds = %simd.if.then, %entry
ret void
}
```
The most important part is:
```
call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %j)
%0 = load i32, ptr %j, align 4
store i32 %0, ptr @j, align 4
```
`%j` is still loaded after `@llvm.lifetime.end.p0(i64 4, ptr nonnull %j)`. This
could cause the backend incorrectly optimizes the code and further generates
incorrect code. The root cause is, when we emit a construct that could have
`linear` clause, it usually has the following pattern:
```
EmitOMPLinearClauseInit(S)
{
OMPPrivateScope LoopScope(*this);
...
EmitOMPLinearClause(S, LoopScope);
...
(void)LoopScope.Privatize();
...
}
EmitOMPLinearClauseFinal(S, [](CodeGenFunction &) { return nullptr; });
```
Variables that need to be privatized are added into `LoopScope`, which also
serves as a RAII object. When `LoopScope` is destructed and if optimization is
enabled, a `@llvm.lifetime.end` is also emitted for each privatized variable.
However, the writing back to original variables in `linear` clause happens after
the scope in `EmitOMPLinearClauseFinal`, causing the issue we see above.
A quick "fix" seems to be, moving `EmitOMPLinearClauseFinal` inside the scope.
However, it doesn't work. That's because the local variable map has been updated
by `LoopScope` such that a variable declaration is mapped to the privatized
variable, instead of the actual one. In that way, the following code will be
generated:
```
%0 = load i32, ptr %j, align 4
store i32 %0, ptr %j, align 4
call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %j)
```
Well, now the life time is correct, but apparently the writing back is broken.
In this patch, a new function `OMPPrivateScope::restoreMap` is added and called
before calling `EmitOMPLinearClauseFinal`. This can make sure that
`EmitOMPLinearClauseFinal` can find the orignal varaibls to write back.
Fixes#56913.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D131272
Scope of changes:
1) Added new function to generate loop versioning
2) Added support for if clause to applySimd function
2) Added tests which confirm that lowering is successful
If ifCond is specified, then collapsed loop is duplicated and if branch
is added. Duplicated loop is executed if simd ifCond is evaluated to false.
Reviewed By: Meinersbur
Differential Revision: https://reviews.llvm.org/D129368
Signed-off-by: Dominik Adamski <dominik.adamski@amd.com>
This patch makes use of OMPIRBuilder support for codegen of taskgroup
construct in clang.
Depends on D128203
Reviewed By: Meinersbur
Differential Revision: https://reviews.llvm.org/D129992
This patch adds OMPIRBuilder support for the simdlen clause for the
simd directive. It uses the simdlen support in OpenMPIRBuilder when
it is enabled in Clang. Simdlen is lowered by OpenMPIRBuilder by
generating the loop.vectorize.width metadata.
Reviewed By: jdoerfert, Meinersbur
Differential Revision: https://reviews.llvm.org/D129149
D127041 introduced the support for `fmax` and `fmin` such that we can also reprent
`atomic compare` and `atomic compare capture` with `atomicrmw` instruction. This
patch simply lifts the limitation we set before.
Depend on D127041.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D127042
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
This patch adds the codegen support for `atomic compare capture` in clang.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D120290
This patch removes all `IgnoreImpCasts` in Sema, and only uses it if necessary. If the expression is not of the same type as the pointer value, a cast is inserted.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D126602
Without this patch, arguments to the
`llvm::OpenMPIRBuilder::AtomicOpValue` initializer are reversed.
Reviewed By: ABataev, tianshilei1992
Differential Revision: https://reviews.llvm.org/D126619
