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
[OpenMP] Initial implementation of parse and sema for composite pragma 'distribute parallel for'
This patch is an initial implementation for #distribute parallel for.
The main differences that affect other pragmas are:
The implementation of 'distribute parallel for' requires blocking of the associated loop, where blocks are "distributed" to different teams and iterations within each block are scheduled to parallel threads within each team. To implement blocking, sema creates two additional worksharing directive fields that are used to pass the team assigned block lower and upper bounds through the outlined function resulting from 'parallel'. In this way, scheduling for 'for' to threads can use those bounds.
As a consequence of blocking, the stride of 'distribute' is not 1 but it is equal to the blocking size. This is returned by the runtime and sema prepares a DistIncrExpr variable to hold that value.
As a consequence of blocking, the global upper bound (EnsureUpperBound) expression of the 'for' is not the original loop upper bound (e.g. in for(i = 0 ; i < N; i++) this is 'N') but it is the team-assigned block upper bound. Sema creates a new expression holding the calculation of the actual upper bound for 'for' as UB = min(UB, PrevUB), where UB is the loop upper bound, and PrevUB is the team-assigned block upper bound.
llvm-svn: 273884
http://reviews.llvm.org/D21564
This patch is an initial implementation for #distribute parallel for.
The main differences that affect other pragmas are:
The implementation of 'distribute parallel for' requires blocking of the associated loop, where blocks are "distributed" to different teams and iterations within each block are scheduled to parallel threads within each team. To implement blocking, sema creates two additional worksharing directive fields that are used to pass the team assigned block lower and upper bounds through the outlined function resulting from 'parallel'. In this way, scheduling for 'for' to threads can use those bounds.
As a consequence of blocking, the stride of 'distribute' is not 1 but it is equal to the blocking size. This is returned by the runtime and sema prepares a DistIncrExpr variable to hold that value.
As a consequence of blocking, the global upper bound (EnsureUpperBound) expression of the 'for' is not the original loop upper bound (e.g. in for(i = 0 ; i < N; i++) this is 'N') but it is the team-assigned block upper bound. Sema creates a new expression holding the calculation of the actual upper bound for 'for' as UB = min(UB, PrevUB), where UB is the loop upper bound, and PrevUB is the team-assigned block upper bound.
llvm-svn: 273705
Summary:
This patch fixes an issue detected when firstprivate variables are passed to an OpenMP outlined function vararg list. Currently they are not compatible with what the runtime library expects causing malfunction in some targets.
This patch fixes the issue by moving the casting logic already in place for offloading to the common code that creates the outline function and arguments and updates the regression tests accordingly.
Reviewers: hfinkel, arpith-jacob, carlo.bertolli, kkwli0, ABataev
Subscribers: cfe-commits, caomhin
Differential Revision: http://reviews.llvm.org/D21150
llvm-svn: 272900
instantiation.
Added checks for non-dependent context when trygin to capture
non-constant schedule chunk expression for proper codegen of outlined
functions.
llvm-svn: 272775
'simd' modifier.
Runtime library defines new schedule constant kmp_sch_static_balanced_chunked = 45 for static loop-based directives static with chunk adjustment (e.g., simd). Added codegen for this kind of schedule.
llvm-svn: 271204
directives.
'kmp_task_t' record type added a new field for 'priority' clause and
changed the representation of pointer to destructors for privates used
within loop-based directives.
Old representation:
typedef struct kmp_task { /* GEH: Shouldn't this be
aligned somehow? */
void *shareds; /**< pointer to block of
pointers to shared vars */
kmp_routine_entry_t routine; /**< pointer to routine
to call for executing task */
kmp_int32 part_id; /**< part id for the
task */
kmp_routine_entry_t destructors; /* pointer to function to
invoke deconstructors of firstprivate C++ objects */
/* private vars */
} kmp_task_t;
New representation:
typedef struct kmp_task { /* GEH: Shouldn't this be
aligned somehow? */
void *shareds; /**< pointer to block of
pointers to shared vars */
kmp_routine_entry_t routine; /**< pointer to routine
to call for executing task */
kmp_int32 part_id; /**< part id for the
task */
kmp_cmplrdata_t data1; /* Two known
optional additions: destructors and priority */
kmp_cmplrdata_t data2; /* Process
destructors first, priority second */
/* future data */
/* private vars */
} kmp_task_t;
Also excessive initialization of 'destructors' fields to 'null' was
removed from codegen if it is known that no destructors shal be used.
Currently a special bit is used in 'kmp_tasking_flags_t' bitfields
('destructors_thunk' bitfield).
llvm-svn: 271201
Summary:
Target regions require globals to be captured. This patch fixes a bug exposed when that happens in a template function.
Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev
Subscribers: guansong, ABataev, cfe-commits, caomhin, fraggamuffin
Differential Revision: http://reviews.llvm.org/D18110
llvm-svn: 271001
Summary: This patch implements the code generation for the `target update` directive. The implemntation relies on the logic already in place for target data standalone directives, i.e. target enter/exit data.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D20650
llvm-svn: 270886
Summary:
The patch contains the parsing and sema support for the `from` clause.
Patch based on the original post by Kelvin Li.
Reviewers: hfinkel, carlo.bertolli, kkwli0, arpith-jacob, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D18488
llvm-svn: 270882
Summary:
The patch contains the parsing and sema support for the `to` clause.
Patch based on the original post by Kelvin Li.
Reviewers: carlo.bertolli, hfinkel, kkwli0, arpith-jacob, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D18597
llvm-svn: 270880
Summary:
This patch is to add parsing and sema support for `target update` directive. Support for the `to` and `from` clauses will be added by a different patch. This patch also adds support for other clauses that are already implemented upstream and apply to `target update`, e.g. `device` and `if`.
This patch is based on the original post by Kelvin Li.
Reviewers: hfinkel, carlo.bertolli, kkwli0, arpith-jacob, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D15944
llvm-svn: 270878
Summary: If a variable is implicitly mapped (doesn't show in a map clause), the runtime library has to be informed if the corresponding capture shows up in first-private clause, so that the storage previously allocated in the device is used. This patch adds the support for that.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D20112
llvm-svn: 270870
Summary: This patch changes the bits used to specify the map types according to the latest version of the libomptarget document and add the support for zero size array section when pointers are being implicitly mapped. This completes the missing new 4.5 map semantics.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: caomhin, cfe-commits
Differential Revision: http://reviews.llvm.org/D20111
llvm-svn: 270868
OpenMP version.
If '-fopenmp' option is provided '-fopenmp-version=' allows to control,
which version of OpenMP must be supported. Currently it affects only the
value of _OPENMP define.
llvm-svn: 270838
OpenMP 4.5.
According to OpenMP 4.5 the _OPENMP macro name is defined to have the decimal value yyyymm where yyyy and mm are the year and month designations of the version of the OpenMP API that the implementation supports. Clang supports OpenMP 4.5 so updated value of _OPENMP macro to 201511.
llvm-svn: 270822
directives.
If firstprivate variable is is captured by value in outlined region and then used as firstprivate variable in inner worksharing directive, the copy for this firstprivate variable was not created. Fixed this bug.
llvm-svn: 270536
For better performance and to unify code with offloading part we pass
scalar firstprivate values by value, instead of by reference. It will
remove some extra copying operations.
llvm-svn: 269751
directives.
OpenMP 4.5 supports clause 'priority' in task-based directives. Patch
adds initial codegen support for this clause in codegen.
llvm-svn: 269050
If private variables require destructors call at the deletion of the
task, additional flag in task flags must be set. Patch fixes this
problem.
llvm-svn: 269039
schedule modifiers.
Runtime library expects some additional data in schedule argument for
loop-based directives, that have additional schedule modifiers
'monotonic|nonmonotonic'.
llvm-svn: 269035
Support OpenMP version 4.5 syntax for #pragma omp declare target.
Syntax:
#pragma omp declare target (extended-list) new-line
or
#pragma omp declare target clause[ [,] clause ... ] new-line
Where clause is one of the following:
to(extended-list)
link(list)
Differential Revision: http://reviews.llvm.org/D20011
llvm-svn: 268925
Summary: `getAssociatedStmt()` contains an assertion that assumes the statement always exists. In device code scanning, we need to look into the associated statement therefore we check its existence. This patch replaces `getAssociatedStmt` by `hasAssociatedStmt` so that we do not trigger the assertion for directives that happen not to have an associated statement (e.g target enter/exit data).
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, caomhin
Differential Revision: http://reviews.llvm.org/D19812
llvm-svn: 268870
OpenMP 4.5 adds taskloop/taskloop simd directives. These directives
allow to use lastprivate clause. Patch adds codegen for this clause.
llvm-svn: 268618
http://reviews.llvm.org/D18474
This patch fixes a bug in code generation of the correct OpenMP runtime library call in presence of target and teams, when target is separated by teams with multiple curly brackets. The current implementation will not be able to see the teams directive inside target and issue a call to tgt_target instead of the correct one tgt_target_teams.
llvm-svn: 267972
OpenMP 4.5 defines 'taskloop simd' directive, which is combined
directive for 'taskloop' and 'simd' directives. Patch adds initial
codegen support for this directive and its 2 basic clauses 'safelen' and
'simdlen'.
llvm-svn: 267872
directive.
OpenMP 4.5 defines 'taskloop' directive and 2 additional clauses
'grainsize' and 'num_tasks' for this directive. Patch adds codegen for
these clauses.
These clauses are generated as arguments of the '__kmpc_taskloop'
libcall and are encoded the following way:
void __kmpc_taskloop(ident_t *loc, int gtid, kmp_task_t *task, int if_val, kmp_uint64 *lb, kmp_uint64 *ub, kmp_int64 st, int nogroup, int sched, kmp_uint64 grainsize, void *task_dup);
If 'grainsize' is specified, 'sched' argument must be set to '1' and
'grainsize' argument must be set to the value of the 'grainsize' clause.
If 'num_tasks' is specified, 'sched' argument must be set to '2' and
'grainsize' argument must be set to the value of the 'num_tasks' clause.
It is possible because these 2 clauses are mutually exclusive and can't
be used at the same time on the same directive.
If none of these clauses is specified, 'sched' argument must be set to
'0'.
llvm-svn: 267862
Summary:
This patch adds support for the target exit data directive code generation.
Given that, apart from the employed runtime call, target exit data requires the same code generation pattern as target enter data, the OpenMP codegen entry point was renamed and reused for both.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17369
llvm-svn: 267814
Summary: This patch adds support for the target enter data directive code generation.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17368
llvm-svn: 267812
Summary:
This patch adds support for the target data directive code generation.
Part of the already existent functionality related with data maps is moved to a new function so that it could be reused.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17367
llvm-svn: 267811
Summary:
Implement codegen for the map clause. All the new list items in 4.5 specification are supported.
Fix bug in the generation of array sections that was exposed by some of the map clause tests: for pointer types the offsets have to be calculated from the pointee not the pointer.
Reviewers: hfinkel, kkwli0, carlo.bertolli, arpith-jacob, ABataev
Subscribers: ABataev, cfe-commits, caomhin, fraggamuffin
Differential Revision: http://reviews.llvm.org/D16749
llvm-svn: 267808
declare reductions.
If reduction clause is applied to instance of class with user-defined
reduction operation without initialization clause, it may cause a crash.
Patch fixes this issue.
llvm-svn: 267695
Currently there is a problem with codegen of inlined directives inside
lambdas, it may cause a crash during codegen because of incorrect
capturing of variables. Patch fixes this problem.
llvm-svn: 267677
The taskloop construct specifies that the iterations of one or more associated loops will be executed in parallel using OpenMP tasks. The iterations are distributed across tasks created by the construct and scheduled to be executed.
The next code will be generated for the taskloop directive:
#pragma omp taskloop num_tasks(N) lastprivate(j)
for( i=0; i<N*GRAIN*STRIDE-1; i+=STRIDE ) {
int th = omp_get_thread_num();
#pragma omp atomic
counter++;
#pragma omp atomic
th_counter[th]++;
j = i;
}
Generated code:
task = __kmpc_omp_task_alloc(NULL,gtid,1,sizeof(struct
task),sizeof(struct shar),&task_entry);
psh = task->shareds;
psh->pth_counter = &th_counter;
psh->pcounter = &counter;
psh->pj = &j;
task->lb = 0;
task->ub = N*GRAIN*STRIDE-2;
task->st = STRIDE;
__kmpc_taskloop(
NULL, // location
gtid, // gtid
task, // task structure
1, // if clause value
&task->lb, // lower bound
&task->ub, // upper bound
STRIDE, // loop increment
0, // 1 if nogroup specified
2, // schedule type: 0-none, 1-grainsize, 2-num_tasks
N, // schedule value (ignored for type 0)
(void*)&__task_dup_entry // tasks duplication routine
);
llvm-svn: 267395
LLVM really wants a debug location on every inlinable call in a function
with debug info, because it otherwise cannot set up inlining debug info.
This change applies an artificial line 0 debug location (which is how
DWARF marks automatically generated code that has no corresponding
source code) to the .__kmpc_global_dtor_. functions to avoid the
LLVM Verifier complaining.
llvm-svn: 267369
causes code generation failure.
The codegen part of firstprivate clause for member decls used type of
original variable without skipping reference type from
OMPCapturedExprDecl. Patch fixes this problem.
llvm-svn: 267125
If loop control variable for simd-based directives is explicitly marked
as linear/lastprivate in clauses, codegen for such construct would
crash. Patch fixes this problem.
llvm-svn: 267101
If the untied clause is present on a task construct, any thread in the
team can resume the task region after a suspension. Patch adds proper
codegen for untied tasks.
llvm-svn: 266853
If the untied clause is present on a task construct, any thread in the
team can resume the task region after a suspension. Patch adds proper
codegen for untied tasks.
llvm-svn: 266754
If the untied clause is present on a task construct, any thread in the team can resume the task region after a suspension. Patch adds proper codegen for untied tasks.
llvm-svn: 266722
With this patch compiler emits warning if it tries to make implicit instantiation
of a template but cannot find the template definition. The warning can be suppressed
by explicit instantiation declaration or by command line options
-Wundefined-var-template and -Wundefined-func-template. The implementation follows
the discussion of http://reviews.llvm.org/D12326.
Differential Revision: http://reviews.llvm.org/D16396
llvm-svn: 266719
The linear clause declares one or more list items to be private to a SIMD lane and to have a linear relationship with respect to the iteration space of a loop.
'linear' '(' <linear-list> [ ':' <linear-step> ] ')'
When a linear-step expression is specified in a linear clause it must be
either a constant integer expression or an integer-typed parameter that is specified in a uniform clause on the directive.
The special this pointer can be used as if was one of the arguments to the function in any of the linear, aligned, or uniform clauses.
llvm-svn: 266056
The aligned clause declares that the object to which each list item points is aligned to the number of bytes expressed in the optional parameter of the aligned clause.
'aligned' '(' <argument-list> [ ':' <alignment> ] ')'
The optional parameter of the aligned clause, alignment, must be a constant positive integer expression. If no optional parameter is specified, implementation-defined default alignments for SIMD instructions on the target platforms are assumed.
The special this pointer can be used as if was one of the arguments to the function in any of the linear, aligned, or uniform clauses.
llvm-svn: 266052
OpenMP 4.0 defines clause 'uniform' in 'declare simd' directive:
'uniform' '(' <argument-list> ')'
The uniform clause declares one or more arguments to have an invariant value for all concurrent invocations of the function in the execution of a single SIMD loop.
The special this pointer can be used as if was one of the arguments to the function in any of the linear, aligned, or uniform clauses.
llvm-svn: 266041
construct.
OpenMP 4.0 defines '#pragma omp declare simd' construct that may have
associated 'simdlen' clause with constant positive expression as an
argument:
simdlen(<const_expr>)
Patch adds parsin and semantic analysis for simdlen clause.
llvm-svn: 265668
Add parsing, sema analysis for 'declare target' construct for OpenMP 4.0
(4.5 support will be added in separate patch).
The declare target directive specifies that variables, functions (C, C++
and Fortran), and subroutines (Fortran) are mapped to a device. The declare
target directive is a declarative directive. In Clang declare target is
implemented as implicit attribute for the declaration.
The syntax of the declare target directive is as follows:
#pragma omp declare target
declarations-definition-seq
#pragma omp end declare target
Based on patch from Michael Wong http://reviews.llvm.org/D15321
llvm-svn: 265530
This patch implements the teams directive for the NVPTX backend. It is different from the host code generation path as it:
Does not call kmpc_fork_teams. All necessary teams and threads are started upon touching the target region, when launching a CUDA kernel, and their execution is coordinated through sequential and parallel regions within the target region.
Does not call kmpc_push_num_teams even if a num_teams of thread_limit clause is present. Setting the number of teams and the thread limit is implemented by the nvptx-related runtime.
Please note that I am now passing a Clang Expr * to emitPushNumTeams instead of the originally chosen llvm::Value * type. The reason for that is that I want to avoid emitting expressions for num_teams and thread_limit if they are not needed in the target region.
http://reviews.llvm.org/D17963
llvm-svn: 265304
directives.
OpenMP 4.5 allows privatization of non-static data members in non-static
member functions. Patch allows to use and implicit privatization of data
members used as counters in loop-based directives.
llvm-svn: 265121
Initial parsing/sema/serialization/deserialization support for '#pragma
omp declare simd' directive.
The 'declare simd' construct can be applied to a function to enable the
creation of one or more versions that can process multiple arguments
using SIMD instructions from a single invocation from a SIMD loop.
If the function has any declarations, then the declare simd construct
for any declaration that has one must be equivalent to the one specified
for the definition. Otherwise, the result is unspecified.
This pragma can be applied many times to the same declaration.
Internally this pragma is represented as an attribute. But we need special processing for this pragma because it must be used before function declaration, this directive is applied to.
Differential Revision: http://reviews.llvm.org/D10599
llvm-svn: 264853
For better support of some specific GNU extensions some extra
transformation of AST nodes were introduced. These transformations are
very hard to handle. The code is improved in handling of these
extensions by using captured expressions construct.
llvm-svn: 264709
Solution unifies interface of RegionCodeGenTy type to allow insert
runtime-specific code before/after main codegen action defined in
CGStmtOpenMP.cpp file. Runtime should not define its own RegionCodeGenTy
for general OpenMP directives, but must be allowed to insert its own
(required) code to support target specific codegen.
llvm-svn: 264700
Solution unifies interface of RegionCodeGenTy type to allow insert
runtime-specific code before/after main codegen action defined in
CGStmtOpenMP.cpp file. Runtime should not define its own RegionCodeGenTy
for general OpenMP directives, but must be allowed to insert its own
(required) code to support target specific codegen.
llvm-svn: 264576
Solution unifies interface of RegionCodeGenTy type to allow insert
runtime-specific code before/after main codegen action defined in
CGStmtOpenMP.cpp file. Runtime should not define its own RegionCodeGenTy
for general OpenMP directives, but must be allowed to insert its own
(required) code to support target specific codegen.
llvm-svn: 264569
Summary:
This patch adds base support for codegen of the target directive on the NVPTX device.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D17877
Reworked test case after buildbot failure on windows.
Updated patch to integrate r263837 and test case nvptx_target_firstprivate_codegen.cpp.
llvm-svn: 264018
This patch implements the following aspects:
It extends sema to check that a variable is not reference in both a map clause and firstprivate or private. This is needed to ensure correct functioning at codegen level, apart from being useful for the user.
It implements firstprivate for target in codegen. The implementation applies to both host and nvptx devices.
It adds regression tests for codegen of firstprivate, host and device side when using the host as device, and nvptx side.
Please note that the regression test for nvptx codegen is missing VLAs. This is because VLAs currently require saving and restoring the stack which appears not to be a supported operation by nvptx backend.
It adds a check in sema regression tests for target map, firstprivate, and private clauses.
http://reviews.llvm.org/D18203
llvm-svn: 263837
Summary:
Reworked test case after buildbot failure on windows.
This patch adds base support for codegen of the target directive on the NVPTX device.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D17877
llvm-svn: 263783
OpenMP 4.5 allows privatization of non-static data members in non-static
member functions. This patch adds support of private data members in
'copyprivate' clauses.
llvm-svn: 263706
OpenMP 4.0 allows to define custom reduction operations using '#pragma
omp declare reduction' construct. Patch allows to use this custom
defined reduction operations in 'reduction' clauses.
llvm-svn: 263701
This patch adds support for codegen of private clause of target and a regression test for host code generation, when the host is used as target device. I believe that code generation for nvptx backend would not require anything additional or different to what is done for the host.
http://reviews.llvm.org/D18105
llvm-svn: 263654
Summary:
This patch adds base support for codegen of the target directive on the NVPTX device.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D17877
llvm-svn: 263587
Summary:
This patch adds base support for codegen of the target directive on the NVPTX device.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D17877
llvm-svn: 263552
Summary:
The current offloading implementation is using -omptargets and -omp-host-ir-file-path options in the frontend. This causes the user a lot of trouble due to to the conflicts with the -o option. E.g. if the user misspells omptargets he will end up with a file with a weird name.
This patches replaces these two options with -fomptargets and -fomp-host-ir-file-path to avoid these issues, and it is also more consistent with the other options like -fopenmp.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, caomhin, fraggamuffin
Differential Revision: http://reviews.llvm.org/D18112
llvm-svn: 263442
Summary: In some cases it can be proved statically that multidimensional array section refer to contiguous storage and can therefore be allowed in a map clause. This patch adds support for those cases in SEMA.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17547
llvm-svn: 263019
OpenMP 4.5 allows privatization of non-static data members in OpenMP
constructs. Patch adds proper codegen support for data members in
'linear' clause
llvm-svn: 263003
This patch provide basic implementation of codegen for teams directive, excluding all clauses except dist_schedule. It also fixes parts of AST reader/writer to enable correct pre-compiled header handling.
http://reviews.llvm.org/D17170
llvm-svn: 262832
This patch provide basic implementation of codegen for teams directive, excluding all clauses except dist_schedule. It also fixes parts of AST reader/writer to enable correct pre-compiled header handling.
http://reviews.llvm.org/D17170
llvm-svn: 262741
Add code generation support for firstprivate and private clauses of teams on the host. Add extensive regression tests including lambda functions and vla testing.
http://reviews.llvm.org/D17582
llvm-svn: 262663
Summary:
This patch implements the launching of a target region in the presence of a nested teams region, i.e calls tgt_target_teams with the required arguments gathered from the enclosed teams directive.
The actual codegen of the region enclosed by the teams construct will be contributed in a separate patch.
Reviewers: hfinkel, arpith-jacob, kkwli0, carlo.bertolli, ABataev
Subscribers: cfe-commits, caomhin, fraggamuffin
Differential Revision: http://reviews.llvm.org/D17019
llvm-svn: 262625
Add parsing, sema analysis and serialization/deserialization for 'declare reduction' construct.
User-defined reductions are defined as
#pragma omp declare reduction( reduction-identifier : typename-list : combiner ) [initializer ( initializer-expr )]
These custom reductions may be used in 'reduction' clauses of OpenMP constructs. The combiner specifies how partial results can be combined into a single value. The
combiner can use the special variable identifiers omp_in and omp_out that are of the type of the variables being reduced with this reduction-identifier. Each of them will
denote one of the values to be combined before executing the combiner. It is assumed that the special omp_out identifier will refer to the storage that holds the resulting
combined value after executing the combiner.
As the initializer-expr value of a user-defined reduction is not known a priori the initializer-clause can be used to specify one. Then the contents of the initializer-clause
will be used as the initializer for private copies of reduction list items where the omp_priv identifier will refer to the storage to be initialized. The special identifier
omp_orig can also appear in the initializer-clause and it will refer to the storage of the original variable to be reduced.
Differential Revision: http://reviews.llvm.org/D11182
llvm-svn: 262582
OpenMP 4.5 allows to privatize data members of current class in member
functions. Patch adds initial support for privatization of data members
in 'linear' clause, no codegen support.
llvm-svn: 262578
OpenMP 4.5 allows to privatize non-static data members of current class
in non-static member functions. Patch supports codegen for non-static
data members in 'reduction' clauses.
llvm-svn: 262460
OpenMP 4.5 allows to privatize non-static data members of current class
in non-static member functions. Patch adds initial parsing/semantic
analysis for data members support in 'reduction' clauses.
llvm-svn: 262199
Summary: The map modifier 'delete' is parser in c++ mode as a delete keyword, which requires special handling in the map clause parsing.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17629
llvm-svn: 262094
OpenMP 4.5 allows to privatize non-static member decls in non-static
member functions. Patch captures such decls by reference in general (for
bitfields, by value) and then operates with this capture. For bitfields,
at the end of codegen for lastprivates original bitfield is updated with the value of captured copy.
llvm-svn: 261824
OpenMP 4.5 allows to use non-static data members of current class in
non-static member functions in different kind of clauses. Patch adds
support for data members in 'shared' clause.
llvm-svn: 261820
OpenMP 4.5 allows to privatize non-static data members of current class
in non-static member functions. Patch adds initial support for data
members.
llvm-svn: 261412
Patch fixes bug with codegen for lastprivate loop counters. Also it may
improve performance for lastprivates calculations in some cases.
llvm-svn: 261209
Loop-based directives allow to use iterators as loop counters. Iterators are allowed to define their own operators. This patch allows to use compound assignment operators for iterators.
llvm-svn: 260957
Expressions inside 'schedule'|'dist_schedule' clause must be captured in
combined directives to avoid possible crash during codegen. Patch
improves handling of such constructs
llvm-svn: 260954
Sync barrier will be emitted after generation of firstprivate variables
only if one of the firstprivate vars is used in lastprivate clause.
llvm-svn: 260877
Summary:
Unlike other outlined regions in OpenMP, offloading entry points have to have be visible (external linkage) for the device side. Using dots in the names of the entries can be therefore problematic for some toolchains, e.g. NVPTX.
Also the patch drops the column information in the unique name of the entry points. The parsing of directives ignore unknown tokens, preventing several target regions to be implemented in the same line. Therefore, the line information is sufficient for the name to be unique. Also, the preprocessor printer does not preserve the column information, causing offloading-entry detection issues if the host uses an integrated preprocessor and the target doesn't (or vice versa).
Reviewers: hfinkel, arpith-jacob, carlo.bertolli, kkwli0, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D17179
llvm-svn: 260837
OpenMP 4.5 introduces privatization of non-static data members of current class in non-static member functions.
To correctly handle such kind of privatization a new (pseudo)declaration VarDecl-based node is added. It allows to reuse an existing code for capturing variables in Lambdas/Block/Captured blocks of code for correct privatization and codegen.
llvm-svn: 260077
Summary:
Different devices may in some cases require different code generation schemes in order to implement OpenMP. This is required not only for performance reasons, but also because it may not be possible to have the current (default) implementation working for these devices. E.g. GPU's cannot implement the same scheme a target such as powerpc or x86b would use, in the sense that it does not have the ability to fork threads, instead all the threads are always executing and need to be managed by the implementation.
This patch proposes a reorganization of the code in the OpenMP code generation to pave the way to have specialized implementation of OpenMP support. More than a "real" patch this is more a request for comments in order to understand if what is proposed is acceptable or if there are better/easier ways to do it.
In this patch part of the common OpenMP codegen infrastructure is moved to a new file under a new namespace (CGOpenMPCommon) so it can be shared between the default implementation and the specialized one. When CGOpenMPRuntime is created, an attempt to select a specialized implementation is done.
In the patch a specialization for nvptx targets is done which currently checks if the target is an OpenMP device and trap if it is not.
Let me know comments suggestions you may have.
Reviewers: hfinkel, carlo.bertolli, arpith-jacob, kkwli0, ABataev
Subscribers: Hahnfeld, cfe-commits, fraggamuffin, caomhin, jholewinski
Differential Revision: http://reviews.llvm.org/D16784
llvm-svn: 259977
Codegen for array sections/array subscripts worked only for expressions with arrays as base. Patch fixes codegen for bases with pointer/reference types.
llvm-svn: 259776
Summary:
This patch adds parsing + sema for the target parallel for directive along with testcases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16759
llvm-svn: 259654
Summary:
This patch enhances Sema to check for the following restriction:
OpenMP 4.5 [2.17 Nesting of Regions]
If a target, target update, target data, target enter data, or
target exit data construct is encountered during execution of a
target region, the behavior is unspecified.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16758
llvm-svn: 259464
Summary:
This patch enhances Sema to check for the following restriction:
OpenMP 4.5 [2.17 Nesting of Regions]
If a target, target update, target data, target enter data, or
target exit data construct is encountered during execution of a
target region, the behavior is unspecified.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16758
llvm-svn: 259366
Summary:
This patch adds parsing + sema for the target parallel directive and its clauses along with testcases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16553
Rebased to current trunk and updated test cases.
llvm-svn: 258832
Summary:
This patch adds parsing + sema for the defaultmap clause associated with the target directive (among others).
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16527
llvm-svn: 258817
Summary:
Extend support in the map clause SEMA for the expressions supported in the OpenMP 4.5 specification, namely member expressions.
Fix some bugs in the previous implementation of SEMA related with expressions that do not consist of single variable references.
Fix bug in parsing when the expression in the map clause do not start with an identifier: accept any expression in the map clause and check for validity in SEMA instead of just ignoring it.
Reviewers: hfinkel, kkwli0, arpith-jacob, carlo.bertolli, ABataev
Subscribers: cfe-commits, fraggamuffin, caomhin
Differential Revision: http://reviews.llvm.org/D16385
llvm-svn: 258543
Summary:
Accept depend clause on target exit data directive in sema and add test cases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16401
llvm-svn: 258502
If 'sections' directive has only one sub-section, the code for 'single'-based directive was emitted. Removed this codegen, because it causes crashes in different cases.
llvm-svn: 258495
Summary:
Accept depend clause on target enter data directive in sema and add test cases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16400
llvm-svn: 258466
Summary:
Accept depend clause on target directive in sema and add test cases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16375
llvm-svn: 258460
Summary:
Accept nowait clause on target exit data directive in sema and add test cases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16362
llvm-svn: 258459
Summary:
Accept nowait clause on target enter data directive in sema and add test cases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16361
llvm-svn: 258457
Summary:
Allow nowait clause on target directive in sema and add test cases.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16358
llvm-svn: 258441
Summary:
Adds the following restriction in the OpenMP specifications.
OpenMP [2.10.1, Restrictions, p. 97]
At least one map clause must appear on the directive.
Reviewers: ABataev
Differential Revision: http://reviews.llvm.org/D16341
llvm-svn: 258425
OpenMP 4.5 allows to use non-static members of current class in non-static member functions in 'private' clause. Patch adds initial support for privatizing data members.
llvm-svn: 258299
Support for the following OpenMP 4.5 restriction on 'target enter data' and 'target exit data':
- A map-type must be specified in all map clauses.
I have to save 'IsMapTypeImplicit' when parsing a map clause to support this constraint and for more informative error messages. This helps me support the following case:
#pragma omp target enter data map(r) // expected-error {{map type must be specified for '#pragma omp target enter data'}}
and distinguish it from:
#pragma omp target enter data map(tofrom: r) // expected-error {{map type 'tofrom' is not allowed for '#pragma omp target enter data'}}
Patch by Arpith Jacob. Thanks!
llvm-svn: 258179
This commit is a follow-up to r251734, r251476, and r249735, which fixes
a bug where function attributes were not attached to thread local
wrapper functions.
rdar://problem/20828324
llvm-svn: 257865
Fixes processing of declarative directives and standalone executable directives. Declarative directives should not be allowed as an immediate statements and standalone executable directives are allowed to be used in case-stmt constructs.
llvm-svn: 257586
- Allow device ID to be signed.
- Add missing semicolon to some of the CHECK directives.
Thanks to Amjad Aboud for detecting the issue.
llvm-svn: 257065
This patch attempts to fix the regressions identified when the patch was committed initially.
Thanks to Michael Liao for identifying the fix in the offloading metadata generation
related with side effects in evaluation of function arguments.
llvm-svn: 256933
Arpith Chacko Jacob