Summary: Test for checking if the mapping is performed correctly. This is a test initially included in Patch https://reviews.llvm.org/D29905
Reviewers: Hahnfeld, carlo.bertolli, caomhin
Reviewed By: Hahnfeld
Subscribers: tra, cfe-commits
Differential Revision: https://reviews.llvm.org/D38040
llvm-svn: 314210
directives.
If the variable is used in the target-based region but is not found in
any private|mapping clause, then generate implicit firstprivate|map
clauses for these implicitly mapped variables.
llvm-svn: 314205
If the captured variable has re-declaration we may end up with the
situation where the captured variable is the re-declaration while the
referenced variable is the canonical declaration (or vice versa). In
this case we may generate wrong code. Patch fixes this situation.
llvm-svn: 313995
This is to fix PR31620. MaxAtomicInlineWidth is set to 128 for x86_64. However
for target without cx16 support, 128 atomic operation will generate __sync_*
libcalls. The patch set MaxAtomicInlineWidth to 64 if the target doesn't support
cx16.
Differential Revision: https://reviews.llvm.org/D38046
llvm-svn: 313992
If the captured variable has some redeclarations we may run into the
situation where the redeclaration is used instead of the canonical
declaration and we may consider this variable as one not captured
before.
llvm-svn: 313880
When the value specified for n in ordered(n) is larger than the number of loops a segmentation fault can occur in one of two ways when attempting to print out a diagnostic for an associated depend(sink : vec):
1) The iteration vector vec contains less than n items
2) The iteration vector vec contains a variable that is not a loop control variable
This patch addresses both of these issues.
Differential Revision: https://reviews.llvm.org/D38049
llvm-svn: 313675
__kmpc_for_static_fini().
Added special flags for calls of __kmpc_for_static_fini(), like previous
ly for __kmpc_for_static_init(). Added flag OMP_IDENT_WORK_DISTRIBUTE
for distribute cnstruct, OMP_IDENT_WORK_SECTIONS for sections-based
constructs and OMP_IDENT_WORK_LOOP for loop-based constructs in
location flags.
llvm-svn: 312642
move constructor.
Previously user-defined reduction initializer was considered as an
assignment expression, not as initializer. Fixed this by treating the
initializer expression as an initializer.
llvm-svn: 312638
step>1.
If the loop is a loot with random access iterators and the iteration
construct is represented it += n, then the compiler crashed because of
reusing of the same MaterializedTemporaryExpr around N. Patch fixes it
by using the expression as written, without any special kind of
wrappings.
llvm-svn: 312292
Capturing of the global variables occurs only in target regions. Patch
fixes it and allows capturing of globals in all target executable
directives.
llvm-svn: 312024
Summary:
Most DIExpressions are empty or very simple. When they are complex, they
tend to be unique, so checking them inline is reasonable.
This also avoids the need for CodeGen passes to append to the
llvm.dbg.mir named md node.
See also PR22780, for making DIExpression not be an MDNode.
Reviewers: aprantl, dexonsmith, dblaikie
Subscribers: qcolombet, javed.absar, eraman, hiraditya, llvm-commits
Differential Revision: https://reviews.llvm.org/D37075
llvm-svn: 311594
of class fails to map class static variable.
If the global variable is captured and it has several redeclarations,
sometimes it may lead to a compiler crash. Patch fixes this by working
only with canonical declarations.
llvm-svn: 311479
If worksharing construct has at least one linear item, an implicit
synchronization point must be emitted to avoid possible conflict with
the loading/storing values to the original variables. Added implicit
barrier if the linear item is found before actual start of the
worksharing construct.
llvm-svn: 311013
If exceptions are enabled, there may be a problem with the codegen of
the finalization functions from OpenMP runtime. It happens because of
the problem with the getting of thread identifier value. Patch tries to
fix it by using the result of the call of function
__kmpc_global_thread_num() rather than loading of value of outlined
function parameter.
llvm-svn: 311007
When translating arguments for NVPTX target it is not taken into account
that function may have variable number of arguments. Patch fixes this
problem.
llvm-svn: 310920
__kmpc_for_static_init().
OpenMP 5.0 will include OpenMP Tools interface that requires distinguishing different worksharing constructs.
Since the same entry point (__kmp_for_static_init(ident_t *loc,
kmp_int32 global_tid,........)) is called in case static
loop/sections/distribute it is suggested using 'flags' field of the
ident_t structure to pass the type of the construct.
llvm-svn: 310865
name.
If the host code is compiled with the debug info, while the target
without, there is a problem that the compiler is unable to find the
debug wrapper. Patch fixes this problem by emitting special name for the
debug version of the code.
llvm-svn: 310511
Converting a _Complex type to a real one simply discards the imaginary part.
This can easily lead to loss of information so for safety (and GCC
compatibility) this patch disallows that when the conversion would be implicit.
The one exception is bool, which actually compares both real and imaginary
parts and so is safe.
llvm-svn: 310427
Arguments, passed to the outlined function, must have correct address
space info for proper Debug info support. Patch sets global address
space for arguments that are mapped and passed by reference.
Also, cuda-gdb does not handle reference types correctly, so reference
arguments are represented as pointers.
llvm-svn: 310387
Arguments, passed to the outlined function, must have correct address
space info for proper Debug info support. Patch sets global address
space for arguments that are mapped and passed by reference.
Also, cuda-gdb does not handle reference types correctly, so reference
arguments are represented as pointers.
llvm-svn: 310377
Arguments, passed to the outlined function, must have correct address
space info for proper Debug info support. Patch sets global address
space for arguments that are mapped and passed by reference.
Also, cuda-gdb does not handle reference types correctly, so reference
arguments are represented as pointers.
llvm-svn: 310360
Arguments, passed to the outlined function, must have correct address
space info for proper Debug info support. Patch sets global address
space for arguments that are mapped and passed by reference.
Also, cuda-gdb does not handle reference types correctly, so reference
arguments are represented as pointers.
llvm-svn: 310104
'#pragma pack (pop)' and suspicious uses of '#pragma pack' in included files
The second recommit (r309106) was reverted because the "non-default #pragma
pack value chages the alignment of struct or union members in the included file"
warning proved to be too aggressive for external projects like Chromium
(https://bugs.chromium.org/p/chromium/issues/detail?id=749197). This recommit
makes the problematic warning a non-default one, and gives it the
-Wpragma-pack-suspicious-include warning option.
The first recommit (r308441) caused a "non-default #pragma pack value might
change the alignment of struct or union members in the included file" warning
in LLVM itself. This recommit tweaks the added warning to avoid warnings for
#includes that don't have any records that are affected by the non-default
alignment. This tweak avoids the previously emitted warning in LLVM.
Original message:
This commit adds a new -Wpragma-pack warning. It warns in the following cases:
- When a translation unit is missing terminating #pragma pack (pop) directives.
- When entering an included file if the current alignment value as determined
by '#pragma pack' directives is different from the default alignment value.
- When leaving an included file that changed the state of the current alignment
value.
rdar://10184173
Differential Revision: https://reviews.llvm.org/D35484
llvm-svn: 309386
When an omp for loop is canceled the constructed objects are being destructed
twice.
It looks like the desired code is:
{
Obj o;
If (cancelled) branch-through-cleanups to cancel.exit.
}
[cleanups]
cancel.exit:
__kmpc_for_static_fini
br cancel.cont (*)
cancel.cont:
__kmpc_barrier
return
The problem seems to be the branch to cancel.cont is currently also going
through the cleanups calling them again. This change just does a direct branch
instead.
Patch By: michael.p.rice@intel.com
Differential Revision: https://reviews.llvm.org/D35854
llvm-svn: 309288
The warning fires on non-suspicious code in Chromium. Reverting until a
solution is figured out.
> Recommit r308327 2nd time: Add a warning for missing
> '#pragma pack (pop)' and suspicious uses of '#pragma pack' in included files
>
> The first recommit (r308441) caused a "non-default #pragma pack value might
> change the alignment of struct or union members in the included file" warning
> in LLVM itself. This recommit tweaks the added warning to avoid warnings for
> #includes that don't have any records that are affected by the non-default
> alignment. This tweak avoids the previously emitted warning in LLVM.
>
> Original message:
>
> This commit adds a new -Wpragma-pack warning. It warns in the following cases:
>
> - When a translation unit is missing terminating #pragma pack (pop) directives.
> - When entering an included file if the current alignment value as determined
> by '#pragma pack' directives is different from the default alignment value.
> - When leaving an included file that changed the state of the current alignment
> value.
>
> rdar://10184173
>
> Differential Revision: https://reviews.llvm.org/D35484
llvm-svn: 309186
'#pragma pack (pop)' and suspicious uses of '#pragma pack' in included files
The first recommit (r308441) caused a "non-default #pragma pack value might
change the alignment of struct or union members in the included file" warning
in LLVM itself. This recommit tweaks the added warning to avoid warnings for
#includes that don't have any records that are affected by the non-default
alignment. This tweak avoids the previously emitted warning in LLVM.
Original message:
This commit adds a new -Wpragma-pack warning. It warns in the following cases:
- When a translation unit is missing terminating #pragma pack (pop) directives.
- When entering an included file if the current alignment value as determined
by '#pragma pack' directives is different from the default alignment value.
- When leaving an included file that changed the state of the current alignment
value.
rdar://10184173
Differential Revision: https://reviews.llvm.org/D35484
llvm-svn: 309106
If the member declaration is captured in the OMPCapturedExprDecl, we may
loose data-sharing attribute info for this declaration. Patch fixes this
bug.
llvm-svn: 308629
This seems to have broken the sanitizer-x86_64-linux buildbot. Reverting until
it's fixed, especially since this landed just before the 5.0 branch.
> This commit adds a new -Wpragma-pack warning. It warns in the following cases:
>
> - When a translation unit is missing terminating #pragma pack (pop) directives.
> - When entering an included file if the current alignment value as determined
> by '#pragma pack' directives is different from the default alignment value.
> - When leaving an included file that changed the state of the current alignment
> value.
>
> rdar://10184173
>
> Differential Revision: https://reviews.llvm.org/D35484
llvm-svn: 308455
and suspicious uses of '#pragma pack' in included files
This commit adds a new -Wpragma-pack warning. It warns in the following cases:
- When a translation unit is missing terminating #pragma pack (pop) directives.
- When entering an included file if the current alignment value as determined
by '#pragma pack' directives is different from the default alignment value.
- When leaving an included file that changed the state of the current alignment
value.
rdar://10184173
Differential Revision: https://reviews.llvm.org/D35484
llvm-svn: 308441
of '#pragma pack' in included files
This commit adds a new -Wpragma-pack warning. It warns in the following cases:
- When a translation unit is missing terminating #pragma pack (pop) directives.
- When entering an included file if the current alignment value as determined
by '#pragma pack' directives is different from the default alignment value.
- When leaving an included file that changed the state of the current alignment
value.
rdar://10184173
Differential Revision: https://reviews.llvm.org/D35484
llvm-svn: 308327
If taskloop directive has no associated nogroup clause, it must emitted
inside implicit taskgroup block. Runtime supports it, but we need to
generate implicit taskgroup block explicitly to support future
reductions codegen.
llvm-svn: 307822
Added checks for the reduction clauses in the taskloop directives:
1. Only addressable items must be used in reduction clauses.
2. Reduction clauses cannot be used with nogroup clauses.
llvm-svn: 307693
Combined directives like 'target parallel' have two captured statements.
Sema has to check the right one from the right direction.
Previously, Sema::IsOpenMPCapturedByRef would return false for mapped
scalars on combined directives. This results in a wrong signature of
the outlined function which triggers an assertion:
void llvm::CallInst::init(llvm::FunctionType *, llvm::Value *, ArrayRef<llvm::Value *>, ArrayRef<OperandBundleDef>, const llvm::Twine &): Assertion `(i >= FTy->getNumParams() || FTy->getParamType(i) == Args[i]->getType()) && "Calling a function with a bad signature!"' failed.
Fixes PR30975 (and PR31985). New function was taken from clang-ykt.
Differential Revision: https://reviews.llvm.org/D34888
llvm-svn: 306956
Currently, if the some of the parameters are captured by value, this
argument is converted to uintptr_t type and thus we loosing the debug
info about real type of the argument (captured variable):
```
void @.outlined_function.(uintptr %par);
...
%a = alloca i32
%a.casted = alloca uintptr
%cast = bitcast uintptr* %a.casted to i32*
%a.val = load i32, i32 *%a
store i32 %a.val, i32 *%cast
%a.casted.val = load uintptr, uintptr* %a.casted
call void @.outlined_function.(uintptr %a.casted.val)
...
```
To resolve this problem, in debug mode a speciall external wrapper
function is generated, that calls the outlined function with the correct
parameters types:
```
void @.wrapper.(uintptr %par) {
%a = alloca i32
%cast = bitcast i32* %a to uintptr*
store uintptr %par, uintptr *%cast
%a.val = load i32, i32* %a
call void @.outlined_function.(i32 %a)
ret void
}
void @.outlined_function.(i32 %par);
...
%a = alloca i32
%a.casted = alloca uintptr
%cast = bitcast uintptr* %a.casted to i32*
%a.val = load i32, i32 *%a
store i32 %a.val, i32 *%cast
%a.casted.val = load uintptr, uintptr* %a.casted
call void @.wrapper.(uintptr %a.casted.val)
...
```
llvm-svn: 306697
According to OpenMP 5.0 at least one 'map' or 'use_device_ptr' clause
must be specified for 'target data' construct. Patch adds support for
this feature.
llvm-svn: 304216
If the variable is marked as TLS variable and target device does not
support TLS, the error is emitted for the variable even if it is not
used in target regions. Patch fixes this and allows to use the values of
the TLS variables in target regions.
llvm-svn: 303768
Currently clang checks for default data sharing attributes only for
variables captured in OpenMP regions by reference. Patch adds checks for
variables captured by value.
llvm-svn: 303077
Summary:
First, getCurFunction looks through blocks and lambdas, which is wrong.
Inside a lambda, va_start should refer to the lambda call operator
prototype. This fixes PR32737.
Second, we shouldn't use any of the getCur* methods, because they look
through contexts that we don't want to look through (EnumDecl,
CapturedStmtDecl). We can use CurContext directly as the calling
context.
Finally, this code assumed that CallExprs would never appear outside of
code contexts (block, function, obj-c method), which is wrong. Struct
member initializers are an easy way to create and parse exprs in a
non-code context.
Reviewers: rsmith
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D32761
llvm-svn: 302188
https://reviews.llvm.org/D32807
This patch allows the map modifier 'always' to be separated by the map type (to, from, tofrom) only by a whitespace, rather than strictly by a comma as in current trunk.
llvm-svn: 302031
CheckForIntOverflow used to implement a whitelist of top-level expressions to
send to the constant expression evaluator, which handled many more expressions
than the CheckForIntOverflow whitelist did.
llvm-svn: 301742
If some function template is instantiated during handling of OpenMP
code, currently it may cause crash of compiler because of trying of
capturing variables in non-capturing function scopes. Patch fixes this
bug.
llvm-svn: 301416
[OpenMP] Initial implementation of code generation for pragma 'distribute parallel for' on host
https://reviews.llvm.org/D29508
This patch makes the following additions:
It abstracts away loop bound generation code from procedures associated with pragma 'for' and loops in general, in such a way that the same procedures can be used for 'distribute parallel for' without the need for a full re-implementation.
It implements code generation for 'distribute parallel for' and adds regression tests. It includes tests for clauses.
It is important to notice that most of the clauses are implemented as part of existing procedures. For instance, firstprivate is already implemented for 'distribute' and 'for' as separate pragmas. As the implementation of 'distribute parallel for' is based on the same procedures, then we automatically obtain implementation for such clauses without the need to add new code. However, this requires regression tests that verify correctness of produced code.
llvm-svn: 301340
https://reviews.llvm.org/D29508
This patch makes the following additions:
1. It abstracts away loop bound generation code from procedures associated with pragma 'for' and loops in general, in such a way that the same procedures can be used for 'distribute parallel for' without the need for a full re-implementation.
2. It implements code generation for 'distribute parallel for' and adds regression tests. It includes tests for clauses.
It is important to notice that most of the clauses are implemented as part of existing procedures. For instance, firstprivate is already implemented for 'distribute' and 'for' as separate pragmas. As the implementation of 'distribute parallel for' is based on the same procedures, then we automatically obtain implementation for such clauses without the need to add new code. However, this requires regression tests that verify correctness of produced code.
Looking forward to comments.
llvm-svn: 301223
https://reviews.llvm.org/D32237
This patch prepares sema with additional fields to support all those composite and combined constructs of OpenMP that include pragma 'distribute' and 'for', such as 'distribute parallel for'. It also extends the regression tests for 'distribute parallel for' and adds a new one.
llvm-svn: 300802
If the type of the captured variable is a pointer(s) to variably
modified type, this type was not processed correctly. Need to drill into
the type, find the innermost variably modified array type and convert it
to canonical parameter type.
llvm-svn: 299868
checkNestingOfRegions uses CancelRegion to determine whether cancel and
cancellation point are valid in the given nesting. This leads to unuseful
diagnostics if CancelRegion is invalid. The given test case has produced:
region cannot be closely nested inside 'parallel' region
As a solution, introduce checkCancelRegion and call it first to get the
expected error:
one of 'for', 'parallel', 'sections' or 'taskgroup' is expected
Differential Revision: https://reviews.llvm.org/D30135
llvm-svn: 295808
With tasks, the cancel may happen in another task. This has a different
region info which means that we can't find it here.
Differential Revision: https://reviews.llvm.org/D30091
llvm-svn: 295474
This resolves a deadlock with the cancel directive when there is no explicit
cancellation point. In that case, the implicit barrier acts as cancellation
point. After removing the barrier after cancel, the now unmatched barrier for
the explicit cancellation point has to go as well.
This has probably worked before rL255992: With the calls for the explicit
barrier, it was sure that all threads passed a barrier before exiting.
Reported by Simon Convent and Joachim Protze!
Differential Revision: https://reviews.llvm.org/D30088
llvm-svn: 295473
This patch implements codegen for the reduction clause on
any teams construct for elementary data types. It builds
on parallel reductions on the GPU. Subsequently,
the team master writes to a unique location in a global
memory scratchpad. The last team to do so loads and
reduces this array to calculate the final result.
This patch emits two helper functions that are used by
the OpenMP runtime on the GPU to perform reductions across
teams.
Patch by Tian Jin in collaboration with Arpith Jacob
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29879
llvm-svn: 295335
This patch implements codegen for the reduction clause on
any parallel construct for elementary data types. An efficient
implementation requires hierarchical reduction within a
warp and a threadblock. It is complicated by the fact that
variables declared in the stack of a CUDA thread cannot be
shared with other threads.
The patch creates a struct to hold reduction variables and
a number of helper functions. The OpenMP runtime on the GPU
implements reduction algorithms that uses these helper
functions to perform reductions within a team. Variables are
shared between CUDA threads using shuffle intrinsics.
An implementation of reductions on the NVPTX device is
substantially different to that of CPUs. However, this patch
is written so that there are minimal changes to the rest of
OpenMP codegen.
The implemented design allows the compiler and runtime to be
decoupled, i.e., the runtime does not need to know of the
reduction operation(s), the type of the reduction variable(s),
or the number of reductions. The design also allows reuse of
host codegen, with appropriate specialization for the NVPTX
device.
While the patch does introduce a number of abstractions, the
expected use case calls for inlining of the GPU OpenMP runtime.
After inlining and optimizations in LLVM, these abstractions
are unwound and performance of OpenMP reductions is comparable
to CUDA-canonical code.
Patch by Tian Jin in collaboration with Arpith Jacob
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29758
llvm-svn: 295333
This patch implements codegen for the reduction clause on
any parallel construct for elementary data types. An efficient
implementation requires hierarchical reduction within a
warp and a threadblock. It is complicated by the fact that
variables declared in the stack of a CUDA thread cannot be
shared with other threads.
The patch creates a struct to hold reduction variables and
a number of helper functions. The OpenMP runtime on the GPU
implements reduction algorithms that uses these helper
functions to perform reductions within a team. Variables are
shared between CUDA threads using shuffle intrinsics.
An implementation of reductions on the NVPTX device is
substantially different to that of CPUs. However, this patch
is written so that there are minimal changes to the rest of
OpenMP codegen.
The implemented design allows the compiler and runtime to be
decoupled, i.e., the runtime does not need to know of the
reduction operation(s), the type of the reduction variable(s),
or the number of reductions. The design also allows reuse of
host codegen, with appropriate specialization for the NVPTX
device.
While the patch does introduce a number of abstractions, the
expected use case calls for inlining of the GPU OpenMP runtime.
After inlining and optimizations in LLVM, these abstractions
are unwound and performance of OpenMP reductions is comparable
to CUDA-canonical code.
Patch by Tian Jin in collaboration with Arpith Jacob
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29758
llvm-svn: 295319
https://reviews.llvm.org/D29501
It looks like I forgot to remove a FIXME comment with the associated statement. The test does not need it and it gives the wrong impression of being an incomplete test.
llvm-svn: 294195
Support for CUDA printf is exploited to support printf for
an NVPTX OpenMP device.
To reflect the support of both programming models, the file
CGCUDABuiltin.cpp has been renamed to CGGPUBuiltin.cpp, and
the call EmitCUDADevicePrintfCallExpr has been renamed to
EmitGPUDevicePrintfCallExpr.
Reviewers: jlebar
Differential Revision: https://reviews.llvm.org/D17890
llvm-svn: 293444
This is a simple patch to teach OpenMP codegen to emit the construct
in Generic mode.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29143
llvm-svn: 293183
This patch adds support for the proc_bind clause on the Spmd construct
'target parallel' on the NVPTX device. Since the parallel region is created
upon kernel launch, this clause can be safely ignored on the NVPTX device at
codegen time for level 0 parallelism.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29128
llvm-svn: 293069
The thread_limit-clause on the combined directive applies to the
'teams' region of this construct. We modify the ThreadLimitClause
class to capture the clause expression within the 'target' region.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29087
llvm-svn: 293049
The num_teams-clause on the combined directive applies to the
'teams' region of this construct. We modify the NumTeamsClause
class to capture the clause expression within the 'target' region.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29085
llvm-svn: 293048
This patch adds support for codegen of 'target teams' on the host.
This combined directive has two captured statements, one for the
'teams' region, and the other for the 'parallel'.
This target teams region is offloaded using the __tgt_target_teams()
call. The patch sets the number of teams as an argument to
this call.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29084
llvm-svn: 293005
This patch adds support for codegen of 'target teams' on the host.
This combined directive has two captured statements, one for the
'teams' region, and the other for the 'parallel'.
This target teams region is offloaded using the __tgt_target_teams()
call. The patch sets the number of teams as an argument to
this call.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29084
llvm-svn: 293001
This patch adds support for the Spmd construct 'target parallel' on the
NVPTX device. This involves ignoring the num_threads clause on the device
since the number of threads in this combined construct is already set on
the host through the call to __tgt_target_teams().
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29083
llvm-svn: 292999
The num_threads-clause on the combined directive applies to the
'parallel' region of this construct. We modify the NumThreadsClause
class to capture the clause expression within the 'target' region.
The offload runtime call for 'target parallel' is changed to
__tgt_target_teams() with 1 team and the number of threads set by
this clause or a default if none.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29082
llvm-svn: 292997
The DSAChecker code in SemaOpenMP looks at the captured statement
associated with an OpenMP directive. A combined directive such as
'target parallel' has nested capture statements, which have to be
fully traversed before executing the DSAChecker. This is a patch
to perform the traversal for such combined directives.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D29026
llvm-svn: 292794
with SEH and openmp
In some cituations (during codegen for Windows SEH constructs)
CodeGenFunction instance may have CurFn equal to nullptr. OpenMP related
code does not expect such situation during cleanup.
llvm-svn: 292590
The if-clause on the combined directive potentially applies to both the
'target' and the 'parallel' regions. Codegen'ing the if-clause on the
combined directive requires additional support because the expression in
the clause must be captured by the 'target' capture statement but not
the 'parallel' capture statement. Note that this situation arises for
other clauses such as num_threads.
The OMPIfClause class inherits OMPClauseWithPreInit to support capturing
of expressions in the clause. A member CaptureRegion is added to
OMPClauseWithPreInit to indicate which captured statement (in this case
'target' but not 'parallel') captures these expressions.
To ensure correct codegen of captured expressions in the presence of
combined 'target' directives, OMPParallelScope was added to 'parallel'
codegen.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28781
llvm-svn: 292437
This patch adds codegen for the 'target parallel' directive on the NVPTX
device. We term offload OpenMP directives such as 'target parallel' and
'target teams distribute parallel for' as SPMD constructs. SPMD constructs,
in contrast to Generic ones like the plain 'target', can never contain
a serial region.
SPMD constructs can be handled more efficiently on the GPU and do not
require the Warp Loop of the Generic codegen scheme. This patch adds
SPMD codegen support for 'target parallel' on the NVPTX device and can
be reused for other SPMD constructs.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28755
llvm-svn: 292428
This patch adds support for codegen of 'target parallel' on the host.
It is also the first combined directive that requires two or more
captured statements. Support for this functionality is included in
the patch.
A combined directive such as 'target parallel' has two captured
statements, one for the 'target' and the other for the 'parallel'
region. Two captured statements are required because each has
different implicit parameters (see SemaOpenMP.cpp). For example,
the 'parallel' has 'global_tid' and 'bound_tid' while the 'target'
does not. The patch adds support for handling multiple captured
statements based on the combined directive.
When codegen'ing the 'target parallel' directive, the 'target'
outlined function is created using the outer captured statement
and the 'parallel' outlined function is created using the inner
captured statement.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28753
llvm-svn: 292419
This patch adds support for codegen of 'target parallel' on the host.
It is also the first combined directive that requires two or more
captured statements. Support for this functionality is included in
the patch.
A combined directive such as 'target parallel' has two captured
statements, one for the 'target' and the other for the 'parallel'
region. Two captured statements are required because each has
different implicit parameters (see SemaOpenMP.cpp). For example,
the 'parallel' has 'global_tid' and 'bound_tid' while the 'target'
does not. The patch adds support for handling multiple captured
statements based on the combined directive.
When codegen'ing the 'target parallel' directive, the 'target'
outlined function is created using the outer captured statement
and the 'parallel' outlined function is created using the inner
captured statement.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28753
llvm-svn: 292374
* Do not initialize these variables when initializing the rest of the
thread_locals in the TU; they have unordered initialization so they can be
initialized by themselves.
This fixes a rejects-valid bug: we would make the per-variable initializer
function internal, but put it in a comdat keyed off the variable, resulting
in link errors when the comdat is selected from a different TU (as the per
TU TLS init function tries to call an init function that does not exist).
* On Darwin, when we decide that we're not going to emit a thread wrapper
function at all, demote its linkage to External. Fixes a verifier failure
on explicit instantiation of a thread_local variable on Darwin.
llvm-svn: 291865
This patch is to implement sema and parsing for 'target teams distribute simd’ pragma.
Differential Revision: https://reviews.llvm.org/D28252
llvm-svn: 291579
Summary:
This patch introduces support for the execution of parallel constructs in a target
region on the NVPTX device. Parallel regions must be in the lexical scope of the
target directive.
The master thread in the master warp signals parallel work for worker threads in worker
warps on encountering a parallel region.
Note: The patch does not yet support capture of arguments in a parallel region so
the test cases are simple.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28145
llvm-svn: 291565
This patch is to add support of the 'is_device_ptr' clause with the 'target parallel for' pragma.
Differential Revision: https://reviews.llvm.org/D28255
llvm-svn: 291540
This patch is to add support of the 'is_device_ptr' clause with the 'target parallel for simd' pragma.
Differential Revision: https://reviews.llvm.org/D28402
llvm-svn: 291537
In C++11, a destructor's implicit exception-spec is nothrow.
The IR for the destructor's invocation changed from invoke to call.
Differential Revision: https://reviews.llvm.org/D28425
llvm-svn: 291458
Summary:
This patch adds two fields to the offload entry descriptor. One field is meant to signal Ctors/Dtors and `link` global variables, and the other is reserved for runtime library use.
Currently, these fields are only filled with zeros in the current code generation, but that will change when `declare target` is added.
The reason, we are adding these fields now is to make the code generation consistent with the runtime library proposal under review in https://reviews.llvm.org/D14031.
Reviewers: ABataev, hfinkel, carlo.bertolli, kkwli0, arpith-jacob, Hahnfeld
Subscribers: cfe-commits, caomhin, jholewinski
Differential Revision: https://reviews.llvm.org/D28298
llvm-svn: 291124
This patch includes updates for codegen of the target region for the NVPTX
device. It moves initializers from the compiler to the runtime and updates
the worker loop to assume parallel work is retrieved from the runtime. A
subsequent patch will update the codegen to retrieve the parallel work using
calls to the runtime. It includes the removal of the inline attribute
for the worker loop and disabling debug info in it.
This allows codegen for a target directive and serial execution on the
NVPTX device.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28125
llvm-svn: 291121
This patch includes updates for codegen of the target region for the NVPTX
device. It moves initializers from the compiler to the runtime and updates
the worker loop to assume parallel work is retrieved from the runtime. A
subsequent patch will update the codegen to retrieve the parallel work using
calls to the runtime. It includes the removal of the inline attribute
for the worker loop and disabling debug info in it.
This allows codegen for a target directive and serial execution on the
NVPTX device.
Reviewers: ABataev
Differential Revision: https://reviews.llvm.org/D28125
llvm-svn: 290983
https://reviews.llvm.org/D17840
This patch enables private, firstprivate, and lastprivate clauses for the OpenMP distribute directive.
Regression tests differ from the similar case of the same clauses on the for directive, by removing a reference to two global variables g and g1. This is necessary because: 1. a distribute pragma is only allowed inside a target region; 2. referring a global variable (e.g. g and g1) in a target region requires the program to enclose the variable in a "declare target" region; 3. declare target pragmas, which are used to define a declare target region, are currently unavailable in clang (patch being prepared).
For this reason, I moved the global declarations into local variables.
llvm-svn: 290898
This patch is to implement sema and parsing for 'target teams distribute parallel for simd’ pragma.
Differential Revision: https://reviews.llvm.org/D28202
llvm-svn: 290862
This patch is to implement sema and parsing for 'target teams distribute parallel for’ pragma.
Differential Revision: https://reviews.llvm.org/D28160
llvm-svn: 290725
This patch is to implement sema and parsing for 'target teams distribute' pragma.
Differential Revision: https://reviews.llvm.org/D28015
llvm-svn: 290508
Much to my surprise, '-disable-llvm-optzns' which I thought was the
magical flag I wanted to get at the raw LLVM IR coming out of Clang
deosn't do that. It still runs some passes over the IR. I don't want
that, I really want the *raw* IR coming out of Clang and I strongly
suspect everyone else using it is in the same camp.
There is actually a flag that does what I want that I didn't know about
called '-disable-llvm-passes'. I suspect many others don't know about it
either. It both does what I want and is much simpler.
This removes the confusing version and makes that spelling of the flag
an alias for '-disable-llvm-passes'. I've also moved everything in Clang
to use the 'passes' spelling as it seems both more accurate (*all* LLVM
passes are disabled, not just optimizations) and much easier to remember
and spell correctly.
This is part of simplifying how Clang drives LLVM to make it cleaner to
wire up to the new pass manager.
Differential Revision: https://reviews.llvm.org/D28047
llvm-svn: 290392
program
Offload related code is not quite ready yet, but some simple examples
must not crash the compiler. Patch fixes the problem in offloading code
with exceptions.
llvm-svn: 290364
based coverage compilation
Added source location info to captured expression declaration + fixed
source location info for loop based directives.
llvm-svn: 290181
Directive name modifiers in 'if' clause are allowed only for OpenMP 4.5
and higher + in OpenMP 4.5 parsing procedure emits error message if ':'
is not found after directive name modifier.
llvm-svn: 290175
This patch is to add support of the 'is_device_ptr' clause in the 'target parallel' pragma.
Differential Revision: https://reviews.llvm.org/D27821
llvm-svn: 289989
This patch is to implement sema and parsing for 'teams distribute parallel for' pragma.
Differential Revision: https://reviews.llvm.org/D27345
llvm-svn: 289179
This patch is to implement sema and parsing for 'teams distribute parallel for simd' pragma.
Differential Revision: https://reviews.llvm.org/D27084
llvm-svn: 288294
If member expression is used in the task region and the base expression
is a DeclRefExp and the variable used in this ref expression is private,
it should be marked as implicitly firstprivate inside this region. Patch
fixes this issue.
llvm-svn: 288039
If 'omp cancel' construct is used in a worksharing construct it may
cause hanging of the software in case if reduction clause is used. Patch fixes this problem by avoiding extra reduction processing for branches that were canceled.
llvm-svn: 287227
Summary:
r286944 introduced bugs detected by ASAN as use-after-return.
r287025 have not fixed them completely.
This reverts commit r286944 and r287025.
Reviewers: ABataev
Subscribers: cfe-commits
Differential Revision: https://reviews.llvm.org/D26720
llvm-svn: 287069
Reworked fix after comments from Richard Smith. We must skip all
CapturedDecl-based contexts when trying to get correct mangling number
context.
llvm-svn: 286953
If 'omp cancel' construct is used in a worksharing construct it may cause
hanging of the software in case if reduction clause is used. Patch fixes
this problem by avoiding extra reduction processing for branches that
were canceled.
llvm-svn: 286944
Output generated by option -ast-print looks like C/C++ code, and it
really is for plain C. For C++ the produced output was not valid C++
code, but the differences were small. With this change the output
is fixed and can be compiled. Tests are changed so that output produced
by -ast-print is compiled again with the same flags and both outputs are
compared.
Option -ast-print is extensively used in clang tests but it itself
was tested poorly, existing tests only checked that compiler did not
crash. There are unit tests in file DeclPrinterTest.cpp, but they test
only terse output mode.
Differential Revision: https://reviews.llvm.org/D26452
llvm-svn: 286439
constructs.
For __real/__imag unary expressions clang emits lvalue with the
associated type from the original complex expression, but not the
underlying builtin integer or float type. This causes crash in codegen
for atomic constructs, if __real/__imag expression are used in atomic
constructs.
llvm-svn: 286129
After some changes in codegen capturing of VLA variables in OpenMP regions was broken, causing compiler crash. Patch fixes this issue.
llvm-svn: 286103
After some changes in codegen capturing of VLA variables in OpenMP
regions was broken, causing compiler crash. Patch fixes this issue.
llvm-svn: 286098
This has two significant effects:
1) Direct relational comparisons between null pointer constants (0 and nullopt)
and pointers are now ill-formed. This was always the case for C, and it
appears that C++ only ever permitted by accident. For instance, cases like
nullptr < &a
are now rejected.
2) Comparisons and conditional operators between differently-cv-qualified
pointer types now work, and produce a composite type that both source
pointer types can convert to (when possible). For instance, comparison
between 'int **' and 'const int **' is now valid, and uses an intermediate
type of 'const int *const *'.
Clang previously supported #2 as an extension.
We do not accept the cases in #1 as an extension. I've tested a fair amount of
code to check that this doesn't break it, but if it turns out that someone is
relying on this, we can easily add it back as an extension.
This is a re-commit of r284800.
llvm-svn: 284890
This has two significant effects:
1) Direct relational comparisons between null pointer constants (0 and nullopt)
and pointers are now ill-formed. This was always the case for C, and it
appears that C++ only ever permitted by accident. For instance, cases like
nullptr < &a
are now rejected.
2) Comparisons and conditional operators between differently-cv-qualified
pointer types now work, and produce a composite type that both source
pointer types can convert to (when possible). For instance, comparison
between 'int **' and 'const int **' is now valid, and uses an intermediate
type of 'const int *const *'.
Clang previously supported #2 as an extension.
We do not accept the cases in #1 as an extension. I've tested a fair amount of
code to check that this doesn't break it, but if it turns out that someone is
relying on this, we can easily add it back as an extension.
llvm-svn: 284800
Tests fall into one of the following categories:
- The requirement was unnecessary
- Additional quoting was required for backslashes in paths (see "sed -e
's/\\/\\\\/g'") in the sanitizer tests.
- OpenMP used 'REQUIRES: shell' as a proxy for the test failing on
Windows. Those tests fail there reliably, so use XFAIL instead.
I tried not to remove shell requirements that were added to suppress
flaky test failures, but if I screwed up, we can add it back as needed.
llvm-svn: 284793
There was a bug in the implementation of captured statements. If it has
a lambda expression in it and the same lambda expression is used outside
the captured region, clang produced an error:
```
error: definition with same mangled name as another definition
```
Here is an example:
```
struct A {
template <typename L>
void g(const L&) { }
};
template<typename T>
void f() {
{
A().g([](){});
}
A().g([](){});
}
int main() {
f<void>();
}
```
Error report:
```
main.cpp:3:10: error: definition with same mangled name as another
definition
void g(const L&) { }
^
main.cpp:3:10: note: previous definition is here
```
Patch fixes this bug.
llvm-svn: 284229
access, by Erich Keane
OpenMP creates a variable array type with a a null size-expr. The Debug
generation failed to due to this. This patch corrects the openmp
implementation, updates the tests, and adds a new one for this
condition.
Differential Revision: https://reviews.llvm.org/D25373
llvm-svn: 284110
This reverts commit r279003 as it breaks some of our buildbots (e.g.
clang-cmake-aarch64-quick, clang-x86_64-linux-selfhost-modules).
The error is in OpenMP/teams_distribute_simd_ast_print.cpp:
clang: /home/buildslave/buildslave/clang-cmake-aarch64-quick/llvm/include/llvm/ADT/DenseMap.h:527:
bool llvm::DenseMapBase<DerivedT, KeyT, ValueT, KeyInfoT, BucketT>::LookupBucketFor(const LookupKeyT&, const BucketT*&) const
[with LookupKeyT = clang::Stmt*; DerivedT = llvm::DenseMap<clang::Stmt*, long unsigned int>;
KeyT = clang::Stmt*; ValueT = long unsigned int;
KeyInfoT = llvm::DenseMapInfo<clang::Stmt*>;
BucketT = llvm::detail::DenseMapPair<clang::Stmt*, long unsigned int>]:
Assertion `!KeyInfoT::isEqual(Val, EmptyKey) && !KeyInfoT::isEqual(Val, TombstoneKey) &&
"Empty/Tombstone value shouldn't be inserted into map!"' failed.
llvm-svn: 279045
This patch is to implement sema and parsing for 'teams distribute simd’ pragma.
This patch is originated by Carlo Bertolli.
Differential Revision: https://reviews.llvm.org/D23528
llvm-svn: 279003
Summary: This patch adds support for the is_device_ptr clause. It expands SEMA to use the mappable expression logic that can only be tested with code generation in place and check conflicts with other data sharing related clauses using the mappable expressions infrastructure.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: https://reviews.llvm.org/D22788
llvm-svn: 276978
Summary: This patch adds support for the use_device_ptr clause. It includes changes in SEMA that could not be tested without codegen, namely, the use of the first private logic and mappable expressions support.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: https://reviews.llvm.org/D22691
llvm-svn: 276977
Summary: This patch add support to map pointers through references in class members. Although a reference does not have storage that a user can access, it still has to be mapped in order to get the deep copy right and the dereferencing code work properly.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: https://reviews.llvm.org/D22787
llvm-svn: 276934
Summary:
This patch fixes a bug in the map of array sections whose base is a reference to a pointer. The existing mapping support was not prepared to deal with it, causing the compiler to crash.
Mapping a reference to a pointer enjoys the same characteristics of a regular pointer, i.e., it is passed by value. Therefore, the reference has to be materialized in the target region.
Reviewers: hfinkel, carlo.bertolli, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: https://reviews.llvm.org/D22690
llvm-svn: 276933
This patch introduces a new cmake variable: CLANG_DEFAULT_RTLIB, thru
which we can specify a default value for -rtlib (libgcc or
compiler-rt) at build time, just like how we set the default C++
stdlib thru CLANG_DEFAULT_CXX_STDLIB.
With these two options, we can configure clang to build binaries on
Linux that have no runtime dependence on any gcc libs (libstdc++ or
libgcc_s).
Patch by Lei Zhang!
Differential Revision: https://reviews.llvm.org/D22663
llvm-svn: 276848
The OpenMP spec mandates that 'a teams construct must be contained within a
target construct'. Currently, this scenario is not diagnosed. This patch is
to add check for orphaned teams construct and issue an error message.
Differential Revision: https://reviews.llvm.org/D22785
llvm-svn: 276726
OpenMP 4.5 removed the restriction that array section lower bound must be non negative.
This change is to allow negative values for array section based on pointers.
For array section based on array type there is still a restriction: "The array section must be a subset of the original array."
Patch by David S.
Differential Revision: https://reviews.llvm.org/D22481
llvm-svn: 276177
Clang misdiagnoses atomic captures cases that contains parens.
i.e.
int v, int *p;
#pragma omp atomic capture
{ v = (*p); (*p)++; }
Patch by David S.
Differential Revision: https://reviews.llvm.org/D22487
llvm-svn: 276167
Having the following code pattern will result in incorrect diagnostic
int main() {
int arr[10];
#pragma omp target data map(arr[:])
#pragma omp target map(arr)
{}
}
t.cpp:4:24: error: original storage of expression in data environment is shared
but data environment do not fully contain mapped expression storage
#pragma omp target map(arr)
^~~
t.cpp:3:29: note: used here
#pragma omp target data map(arr[:])
^~~~~~
1 error generated.
Patch by David S.
Differential Revision: https://reviews.llvm.org/D22075
llvm-svn: 275926
target_parallel_for_simd_collapse_messages.cpp and target_parallel_for_simd_ordered_messages.cpp give different diagnostic messages in compiling with -std=c++11. The test cases are updated to make it compatible.
Differential Revision: https://reviews.llvm.org/D22417
llvm-svn: 275805
This patch adds the check for specifying both simdlen and safelen clauses on the 'distribute simd' or 'distribute parallel for simd' constructs.
Differential Revision: https://reviews.llvm.org/D22384
llvm-svn: 275529
This patch is to implement sema and parsing for 'target parallel for simd' pragma.
Differential Revision: http://reviews.llvm.org/D22096
llvm-svn: 275365
This patch is to add two additional tests for testing 'distribute parallel for simd' pragma with disallowed clauses and loops.
Differential Revision: http://reviews.llvm.org/D22169
llvm-svn: 275315
This patch is to add two additional tests for testing 'distribute simd' pragma with disallowed clauses and loops.
Differential Revision: http://reviews.llvm.org/D22176
llvm-svn: 275306
http://reviews.llvm.org/D21904
This patch is similar to the implementation of 'private' clause: it adds a list of private pointers to be used within the target data region to store the device pointers returned by the runtime.
Please refer to the following document for a full description of what the runtime witll return in this case (page 10 and 11):
https://github.com/clang-omp/OffloadingDesign
I am happy to answer any question related to the runtime interface to help reviewing this patch.
llvm-svn: 275271
Summary: This patch is an implementation of sema and parsing for the OpenMP composite pragma 'distribute simd'.
Differential Revision: http://reviews.llvm.org/D22007
llvm-svn: 274604
Summary: This patch is an implementation of sema and parsing for the OpenMP composite pragma 'distribute parallel for simd'.
Differential Revision: http://reviews.llvm.org/D21977
llvm-svn: 274530
Summary: This patch changes the options used by offloading to start with -fopenmp instead of -fomp. This makes the option naming more consistent and materializes a suggestion by Richard Smith in http://reviews.llvm.org/D9888.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, ABataev
Subscribers: kkwli0, cfe-commits, caomhin
Differential Revision: http://reviews.llvm.org/D21841
llvm-svn: 274283
Alexey Bataev