| Commit message (Collapse) | Author | Age | Files | Lines |
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
llvm.mem.parallel_loop_access metadata.
Instead of generating llvm.mem.parallel_loop_access metadata, generate
llvm.access.group on instructions and llvm.loop.parallel_accesses on
loops. There is one access group per generated loop.
This is clang part of D52116/r349725.
Differential Revision: https://reviews.llvm.org/D52117
llvm-svn: 349823
|
|
|
|
|
|
|
|
|
| |
Need to mark the loop as started when the initialization statement is
found. It is required to prevent possible incorrect loop iteraton
variable detection during template instantiation and fix the compiler
crash during the codegen.
llvm-svn: 349657
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Without this patch, clang doesn't complain that X needs explicit data
sharing attributes in the following:
```
#pragma omp target teams default(none)
{
#pragma omp parallel num_threads(X)
;
}
```
However, clang does produce that complaint after the braces are
removed. With this patch, clang complains in both cases.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D55861
llvm-svn: 349635
|
|
|
|
|
|
|
|
|
|
|
|
| |
A map clause with the close map-type-modifier is a hint to
prefer that the variables are mapped using a copy into faster
memory.
Patch by Ahsan Saghir (saghir)
Differential Revision: https://reviews.llvm.org/D55719
llvm-svn: 349551
|
|
|
|
|
|
|
|
|
|
|
| |
buffer.
Seems to me, nvlink has a bug with the proper support of the weakly
linked symbols. It does not allow to define several shared memory buffer
with the different sizes even with the weak linkage. Instead we always
use 128 bytes buffer to prevent nvlink from the error message emission.
llvm-svn: 349540
|
|
|
|
|
|
|
| |
The parallel reduction operation requires an extra synchronization point
in the inter-warp copy function to avoid divergence.
llvm-svn: 349525
|
|
|
|
|
|
|
|
|
| |
Inlined runtime with the current implementation of the interwarp copy
function leads to the undefined behavior because of the not quite
correct implementation of the barriers. Start using generic
__kmpc_barier function instead of the custom made barriers.
llvm-svn: 349192
|
|
|
|
|
|
|
|
|
| |
If the array section is based on pointer and this sections is mapped in
target region + then it is used in the inner parallel region, it also
must be globalized as the pointer itself is passed by value, not by
reference.
llvm-svn: 348492
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Critical regions in NVPTX are the constructs, which, generally speaking,
are not supported by the NVPTX target. Instead we're using special
technique to handle the critical regions. Currently they are supported
only within the loop and all the threads in the loop must execute the
same critical region.
Inside of this special regions the regions still must be emitted as
critical, to avoid possible data races between the teams +
synchronization must use __kmpc_barrier functions.
llvm-svn: 348272
|
|
|
|
|
|
|
|
| |
__kmpc_barrier runtime functions must be marked as convergent to prevent
some dangerous optimizations. Also, for NVPTX target all barriers must
be emitted as simple barriers.
llvm-svn: 348271
|
|
|
|
|
|
| |
This moves everything primarily testing the functionality of -ast-dump and -ast-print into their own directory, rather than leaving the tests spread around the testing directory.
llvm-svn: 348017
|
|
|
|
|
|
|
|
|
| |
initialization.
Function __kmpc_global_thread_num should be called only after
initialization, not earlier.
llvm-svn: 347919
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Summary: This patch adds a new runtime for the SPMD deinit kernel function which replaces the previous function. The new function takes as argument the flag which signals whether the runtime is required or not. This enables the compiler to optimize out the part of the deinit function which are not needed.
Reviewers: ABataev, caomhin
Reviewed By: ABataev
Subscribers: jholewinski, guansong, cfe-commits
Differential Revision: https://reviews.llvm.org/D54970
llvm-svn: 347915
|
|
|
|
|
|
|
|
| |
Fixed emission of the target regions found in the virtual functions.
Previously we may end up with the situation when those regions could be
skipped.
llvm-svn: 347793
|
|
|
|
|
|
| |
Added basic codegen support for the reductions across the teams.
llvm-svn: 347715
|
|
|
|
|
|
|
|
|
|
|
| |
modes.
If the region is inside target|teams|distribute region, we can emit the
locations with the correct info for execution mode and runtime mode.
Patch adds this ability to the NVPTX codegen to help the optimizer to
produce better code.
llvm-svn: 347583
|
|
|
|
|
|
|
|
|
|
|
|
| |
For the NVPTX target default locations should be emitted as constants +
additional info must be emitted in the reserved_2 field of the ident_t
structure. The 1st bit controls the execution mode and the 2nd bit
controls use of the lightweight runtime. The combination of the bits for
Non-SPMD mode + lightweight runtime represents special undefined mode,
used outside of the target regions for orphaned directives or functions.
Should allow and additional optimization inside of the target regions.
llvm-svn: 347425
|
|
|
|
|
|
|
|
|
| |
Loop-control variables with the default data-sharing attributes should
not be captured in the OpenMP region as they are private by default.
Also, default attributes should be emitted for such variables in the
inner OpenMP regions for the correct data sharing during codegen.
llvm-svn: 347409
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
forms of random access iterator
In OpenMP 4.5, only 4 relational operators are supported: <, <=, >,
and >=. This work is to enable support for relational operator
!= (not-equal) as one of the canonical forms.
Patch by Anh Tuyen Tran
Differential Revision: https://reviews.llvm.org/D54441
llvm-svn: 347405
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
This patch adjusts a test not to depend on deprecated FileCheck
behavior that permits overlapping matches within a block of CHECK-DAG
directives. Thus, this patch also removes uses of FileCheck's
-allow-deprecated-dag-overlap command-line option.
There were two issues in this test:
1. There were sets of patterns for store instructions in which a
pattern X could match a superset of a pattern Y. While X appeared
before Y, Y's intended match appeared before X's intended match. The
result was that X matched Y's intended match. Under the old
overlapping behavior, Y also matched Y's intended match. Under the
new non-overlapping behavior, Y had nothing left to match. This patch
fixes this by gathering these sets in one place and putting the most
specific patterns (Y) before the more general patterns (X).
2. The CHECK-DAG patterns involving the variables CBPADDR3 and
CBPADDR4 were the same, but there was only one match in the text, so
CBPADDR4 patterns had nothing to match under the new non-overlapping
behavior. Moreover, a preceding related series of directives had
variables (SADDR0, BPADDR0, etc.) numbered only 0 through 4, but this
series had variables numbered 0 through 5. Assuming CBPADDR4's
directives were not intended, this patch removes them.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D54765
llvm-svn: 347351
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
This patch adjusts a test not to depend on deprecated FileCheck
behavior that permits overlapping matches within a block of CHECK-DAG
directives. Thus, this patch also removes uses of FileCheck's
-allow-deprecated-dag-overlap command-line option.
Specifically, the FileCheck variables DBG_LOC_START, DBG_LOC_END, and
DBG_LOC_CANCEL were all set to the same value. As a result, three
TERM_DEBUG-DAG patterns, one for each variable, all matched the same
text under the old overlapping behavior. Under the new
non-overlapping behavior, that's not permitted. This patch's solution
is to replace these variables with one variable and replace these
patterns with one pattern.
Reviewed By: ABataev
Differential Revision: https://reviews.llvm.org/D54764
llvm-svn: 347350
|
|
|
|
|
|
| |
requires directive. Differential Review: https://reviews.llvm.org/D54493
llvm-svn: 347214
|
|
|
|
|
|
|
|
|
| |
If lambda is used inside of the OpenMP region and captures `this`, we
should recapture it in the OpenMP region also. But we should do this
only if the OpenMP region is used in the context of the same class, just
like the lambda.
llvm-svn: 347096
|
|
|
|
|
|
|
|
|
|
|
| |
reductions.
Fixed previously committed code for the reduction support in
teams/parallel constructs taking into account new design of the NVPTX
support in the compiler. Teams reduction are not fully functional yet,
it is going to be fixed in the following patches.
llvm-svn: 347081
|
|
|
|
|
|
|
|
|
| |
If the statements between target|teams|distribute directives does not
require execution in master thread, like constant expressions, null
statements, simple declarations, etc., such construct can be xecuted in
SPMD mode.
llvm-svn: 346551
|
|
|
|
|
|
|
|
|
|
| |
target|teams|distribute variables.
If the total size of the variables, declared in target|teams|distribute
regions, is less than the maximal size of shared memory available, the
buffer is allocated in the shared memory.
llvm-svn: 346507
|
|
|
|
|
|
|
|
| |
The base pointer for the lambda mapping must point to the lambda capture
placement and pointer must point to the captured variable itself. Patch
fixes this problem.
llvm-svn: 346408
|
|
|
|
|
|
|
|
| |
Fixed lookup for the target regions in unused virtual functions + fixed
processing of the global variables not marked as declare target but
emitted during debug info emission.
llvm-svn: 346343
|
|
|
|
|
|
|
|
| |
Coalesced memory access requires use of the new function
`__kmpc_data_sharing_coalesced_push_stack` instead of the
`__kmpc_data_sharing_push_stack`.
llvm-svn: 345991
|
|
|
|
|
|
| |
The previously used combination `PTR_AND_OBJ | PRIVATE` could be used for mapping of some data in Fortran. Changed it to `PTR_AND_OBJ | LITERAL`.
llvm-svn: 345982
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
target/teams/distribute regions.
Target/teams/distribute regions exist for all the time the kernel is
executed. Thus, if the variable is declared in their context and then
escape it, we can allocate global memory statically instead of
allocating it dynamically.
Patch captures all the globalized variables in target/teams/distribute
contexts, merges them into the records, one per each target region.
Those records are then joined into the union, one per compilation unit
(to save the global memory). Those units are organized into
2 x dimensional arrays, where the first dimension is
the number of blocks per SM and the second one is the number of SMs.
Runtime functions manage this global memory space between the executing
teams.
llvm-svn: 345978
|
|
|
|
|
|
| |
Also renamed test files relating to 'requires'. Differntial review: https://reviews.llvm.org/D53513
llvm-svn: 345967
|
|
|
|
|
|
|
|
|
|
| |
Added support for mapping of lambdas in the target regions. It scans all
the captures by reference in the lambda, implicitly maps those variables
in the target region and then later reinstate the addresses of
references in lambda to the correct addresses of the captured|privatized
variables.
llvm-svn: 345609
|
|
|
|
|
|
|
|
|
|
|
| |
being shared.
According to the standard, the variables with unspecified data-sharing
attributes in presence of `default(none)` clause must be reported to
users. Compiler did not generate error reports for the variables used in
other OpenMP regions. Patch fixes this.
llvm-svn: 345533
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Summary: Iteration variable must be strictly less than the number of iterations. This fixes a bug introduced by previous patch D53448.
Reviewers: ABataev, caomhin
Reviewed By: ABataev
Subscribers: guansong, cfe-commits
Differential Revision: https://reviews.llvm.org/D53827
llvm-svn: 345527
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
parallel for
Summary: This patch adds a new code generation path for bound sharing directives containing distribute parallel for. The new code generation scheme applies to chunked schedules on distribute and parallel for directives. The scheme simplifies the code that is being generated by eliminating the need for an outer for loop over chunks for both distribute and parallel for directives. In the case of distribute it applies to any sized chunk while in the parallel for case it only applies when chunk size is 1.
Reviewers: ABataev, caomhin
Reviewed By: ABataev
Subscribers: jholewinski, guansong, cfe-commits
Differential Revision: https://reviews.llvm.org/D53448
llvm-svn: 345509
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Summary: This patch enables the choosing of the default schedule for parallel for loops even in non-SPMD cases.
Reviewers: ABataev, caomhin
Reviewed By: ABataev
Subscribers: jholewinski, guansong, cfe-commits
Differential Revision: https://reviews.llvm.org/D53443
llvm-svn: 345507
|
|
|
|
|
|
|
|
| |
If the loop counter is not declared in the context of the loop and it is
private, such loop counters should not be captured in the outlined
regions.
llvm-svn: 345505
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
Summary:
This is a simple test of the parallel for code generation. It will be used to showcase the change introduced by patch D53443.
Reviewers: ABataev, caomhin
Reviewed By: ABataev
Subscribers: guansong, cfe-commits
Differential Revision: https://reviews.llvm.org/D53772
llvm-svn: 345417
|
|
|
|
|
|
|
|
|
|
|
| |
According to the OpenMP standard, In a task generating construct, if no
default clause is present, a variable for which the data-sharing
attribute is not determined by the rules above is firstprivatized.
Compiler tries to implement this, but if the variable is not directly
used in the task context, this variable may not be firstprivatized.
Patch fixes this problem.
llvm-svn: 345277
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| |
function
Summary:
For the following code:
```
int i;
#pragma omp taskloop
for (i = 0; i < 100; ++i)
{}
#pragma omp taskloop nogroup
for (i = 0; i < 100; ++i)
{}
```
Clang emits the following LLVM IR:
```
...
call void @__kmpc_taskgroup(%struct.ident_t* @0, i32 %0)
%2 = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @0, i32 %0, i32 1, i64 80, i64 8, i32 (i32, i8*)* bitcast (i32 (i32, %struct.kmp_task_t_with_privates*)* @.omp_task_entry. to i32 (i32, i8*)*))
...
call void @__kmpc_taskloop(%struct.ident_t* @0, i32 %0, i8* %2, i32 1, i64* %8, i64* %9, i64 %13, i32 0, i32 0, i64 0, i8* null)
call void @__kmpc_end_taskgroup(%struct.ident_t* @0, i32 %0)
...
%15 = call i8* @__kmpc_omp_task_alloc(%struct.ident_t* @0, i32 %0, i32 1, i64 80, i64 8, i32 (i32, i8*)* bitcast (i32 (i32, %struct.kmp_task_t_with_privates.1*)* @.omp_task_entry..2 to i32 (i32, i8*)*))
...
call void @__kmpc_taskloop(%struct.ident_t* @0, i32 %0, i8* %15, i32 1, i64* %21, i64* %22, i64 %26, i32 0, i32 0, i64 0, i8* null)
```
The first set of instructions corresponds to the first taskloop construct. It is important to note that the implicit taskgroup region associated with the taskloop construct has been materialized in our IR: the `__kmpc_taskloop` occurs inside a taskgroup region. Note also that this taskgroup region does not exist in our second taskloop because we are using the `nogroup` clause.
The issue here is the 4th argument of the kmpc_taskloop call, starting from the end, is always a zero. Checking the LLVM OpenMP RT implementation, we see that this argument corresponds to the nogroup parameter:
```
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);
```
So basically we always tell to the RT to do another taskgroup region. For the first taskloop, this means that we create two taskgroup regions. For the second example, it means that despite the fact we had a nogroup clause we are going to have a taskgroup region, so we unnecessary wait until all descendant tasks have been executed.
Reviewers: ABataev
Reviewed By: ABataev
Subscribers: rogfer01, cfe-commits
Differential Revision: https://reviews.llvm.org/D53636
llvm-svn: 345180
|
|
|
|
|
|
|
|
|
| |
The compiler is crashing if we trying to post-capture the fields
implicitly captured inside of the task constructs. Seems, this kind of
processing is not supported and such fields should not be
firstprivatized.
llvm-svn: 345177
|
|
|
|
|
|
|
| |
This reverts commit https://reviews.llvm.org/rL344150 which causes
MachineOutliner related failures on the ppc64le multistage buildbot.
llvm-svn: 344526
|
|
|
|
|
|
|
| |
Additional reduction of the global memory usage in the target regions
without parallel regions.
llvm-svn: 344413
|
|
|
|
|
|
|
|
|
|
|
|
| |
if the function has globalized variables and called in context of
target/teams/distribute regions, it does not need to globalize 32
copies of the same variables for memory coalescing, it is enough to
have just one copy, because there is parallel region.
Patch does this by adding call for `__kmpc_parallel_level` function and
checking its return value. If the code sees that the parallel level is
0, then only one variable is allocated, not 32.
llvm-svn: 344356
|
|
|
|
|
|
|
|
|
|
|
| |
target/teams/distribute regions.
Previously introduced globalization scheme that uses memory coalescing
scheme may increase memory usage fr the variables that are devlared in
target/teams/distribute contexts. We don't need 32 copies of such
variables, just 1. Patch reduces memory use in this case.
llvm-svn: 344273
|
|
|
|
|
|
| |
Differential Revision: https://reviews.llvm.org/D53079
llvm-svn: 344249
|
|
|
|
|
|
|
|
|
| |
This is currently a clang extension and a resolution
of the defect report in the C++ Standard.
Differential Revision: https://reviews.llvm.org/D46441
llvm-svn: 344150
|
|
|
|
|
|
|
|
|
| |
Added support for memory coalescing for better performance for
globalized variables. From now on all the globalized variables are
represented as arrays of 32 elements and each thread accesses these
elements using `tid & 31` as index.
llvm-svn: 344049
|
|
|
|
|
|
|
|
|
| |
mode.
__kmpc_global_thread_num() should be called before initialization of the
runtime.
llvm-svn: 343857
|