Commit Graph

194 Commits

Author SHA1 Message Date
Alexey Bataev 3015bcc62a [OPENMP] Generalize codegen for 'sections'-based directive.
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
2016-01-22 08:56:50 +00:00
Alexey Bataev 8524d15954 [OPENMP] Fix crash on reduction for complex variables.
reworked codegen for reduction operation for complex types to avoid crash

llvm-svn: 258394
2016-01-21 12:35:58 +00:00
Alexey Bataev 9619f04c0e [OPENMP 4.0] Fix for codegen of 'cancel' directive within 'sections' directive.
Allow to emit code for 'cancel' directive within 'sections' directive with single sub-section.

llvm-svn: 258307
2016-01-20 12:29:47 +00:00
Samuel Antao 7259076032 [OpenMP] Parsing + sema for "target exit data" directive.
Patch by Arpith Jacob. Thanks!

llvm-svn: 258177
2016-01-19 20:04:50 +00:00
Samuel Antao df67fc468e [OpenMP] Parsing + sema for "target enter data" directive.
Patch by Arpith Jacob. Thanks!

llvm-svn: 258165
2016-01-19 19:15:56 +00:00
Carlo Bertolli b4adf55e0f Add OpenMP dist_schedule clause to distribute directive and related regression tests.
llvm-svn: 257917
2016-01-15 18:50:31 +00:00
Samuel Antao ee8fb302f5 [OpenMP] Reapply rL256842: [OpenMP] Offloading descriptor registration and device codegen.
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
2016-01-06 13:42:12 +00:00
Samuel Antao 7d5de9a1ee [OpenMP] Revert rL256842: [OpenMP] Offloading descriptor registration and device codegen.
It was causing two regression, so I'm reverting until the cause is found.

llvm-svn: 256858
2016-01-05 19:16:13 +00:00
Samuel Antao 4d5f0bbea1 [OpenMP] Offloading descriptor registration and device codegen.
Summary:
In order to offloading work properly two things need to be in place:
- a descriptor with all the offloading information (device entry functions, and global variable) has to be created by the host and registered in the OpenMP offloading runtime library.
- all the device functions need to be emitted for the device and a convention has to be in place so that the runtime library can easily map the host ID of an entry point with the actual function in the device.

This patch adds support for these two things. However, only entry functions are being registered given that 'declare target' directive is not yet implemented.

About offloading descriptor:

The details of the descriptor are explained with more detail in http://goo.gl/L1rnKJ. Basically the descriptor will have fields that specify the number of devices, the pointers to where the device images begin and end (that will be defined by the linker), and also pointers to a the begin and end of table whose entries contain information about a specific entry point. Each entry has the type:
```
struct __tgt_offload_entry{
 void *addr;
 char *name;
 int64_t size;
};
```  
and will be implemented in a pre determined (ELF) section `.omp_offloading.entries` with 1-byte alignment, so that when all the objects are linked, the table is in that section with no padding in between entries (will be like a C array). The code generation ensures that all `__tgt_offload_entry` entries are emitted in the same order for both host and device so that the runtime can have the corresponding entries in both host and device in same index of the table, and efficiently implement the mapping.

The resulting descriptor is registered/unregistered with the runtime library using the calls `__tgt_register_lib` and `__tgt_unregister_lib`. The registration is implemented in a high priority global initializer so that the registration happens always before any initializer (that can potentially include target regions) is run.

The driver flag -omptargets= was created to specify a comma separated list of devices the user wants to support so that the new functionality can be exercised. Each device is specified with its triple.


About target codegen:

The target codegen is pretty much straightforward as it reuses completely the logic of the host version for the same target region. The tricky part is to identify the meaningful target regions in the device side. Unlike other programming models, like CUDA, there are no already outlined functions with attributes that mark what should be emitted or not. So, the information on what to emit is passed in the form of metadata in host bc file. This requires a new option to pass the host bc to the device frontend. Then everything is similar to what happens in CUDA: the global declarations emission is intercepted to check to see if it is an "interesting" declaration. The difference is that instead of checking an attribute, the metadata information in checked. Right now, there is only a form of metadata to pass information about the device entry points (target regions). A class `OffloadEntriesInfoManagerTy` was created to manage all the information and queries related with the metadata. The metadata looks like this:
```
!omp_offload.info = !{!0, !1, !2, !3, !4, !5, !6}

!0 = !{i32 0, i32 52, i32 77426347, !"_ZN2S12r1Ei", i32 479, i32 13, i32 4}
!1 = !{i32 0, i32 52, i32 77426347, !"_ZL7fstatici", i32 461, i32 11, i32 5}
!2 = !{i32 0, i32 52, i32 77426347, !"_Z9ftemplateIiET_i", i32 444, i32 11, i32 6}
!3 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 99, i32 11, i32 0}
!4 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 272, i32 11, i32 3}
!5 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 127, i32 11, i32 1}
!6 = !{i32 0, i32 52, i32 77426347, !"_Z3fooi", i32 159, i32 11, i32 2}
```
The fields in each metadata entry are (in sequence):
Entry 1) an ID of the type of metadata - right now only zero is used meaning "OpenMP target region".
Entry 2) a unique ID of the device where the input source file that contain the target region lives. 
Entry 3) a unique ID of the file where the input source file that contain the target region lives. 
Entry 4) a mangled name of the function that encloses the target region.
Entries 5) and 6) line and column number where the target region was found.
Entry 7) is the order the entry was emitted.

Entry 2) and 3) are required to distinguish files that have the same function name.
Entry 4) is required to distinguish different instances of the same declaration (usually templated ones)
Entries 5) and 6) are required to distinguish the particular target region in body of the function (it is possible that a given target region is not an entry point - if clause can evaluate always to zero - and therefore we need to identify the "interesting" target regions. )

This patch replaces http://reviews.llvm.org/D12306.

Reviewers: ABataev, hfinkel, tra, rjmccall, sfantao

Subscribers: FBrygidyn, piotr.rak, Hahnfeld, cfe-commits

Differential Revision: http://reviews.llvm.org/D12614

llvm-svn: 256842
2016-01-05 16:23:04 +00:00
Alexey Bataev a6f2a14b94 [OPENMP 4.5] Codegen for 'schedule' clause with monotonic/nonmonotonic modifiers.
OpenMP 4.5 adds support for monotonic/nonmonotonic modifiers in 'schedule' clause. Add codegen for these modifiers.

llvm-svn: 256666
2015-12-31 06:52:34 +00:00
Alexey Bataev 6f531ec0a2 [OPENMP] Remove explicit call for implicit barrier
#pragma omp parallel needs an implicit barrier that is currently done by an explicit call to __kmpc_barrier. However, the runtime already ensures a barrier in __kmpc_fork_call which currently leads to two barriers per region per thread.
Differential Revision: http://reviews.llvm.org/D15561

llvm-svn: 255992
2015-12-18 10:24:53 +00:00
Alexey Bataev 8ef3141127 [OPENMP] Fix for http://llvm.org/PR25878: Error compiling an OpenMP program
OpenMP codegen tried to emit the code for its constructs even if it was detected as a dead-code. Added checks to ensure that the code is emitted if the code is not dead.

llvm-svn: 255990
2015-12-18 07:58:25 +00:00
Alexey Bataev fc57d1601d [OPENMP 4.5] Codegen for 'hint' clause of 'critical' directive
OpenMP 4.5 defines 'hint' clause for 'critical' directive. Patch adds codegen for this clause.

llvm-svn: 255639
2015-12-15 10:55:09 +00:00
Alexey Bataev 28c75417b2 [OPENMP 4.5] Parsing/sema for 'hint' clause of 'critical' directive.
OpenMP 4.5 adds 'hint' clause to critical directive. Patch adds parsing/semantic analysis for this clause.

llvm-svn: 255625
2015-12-15 08:19:24 +00:00
Carlo Bertolli 6200a3d0f3 Add parse and sema of OpenMP distribute directive with all clauses except dist_schedule
llvm-svn: 255498
2015-12-14 14:51:25 +00:00
Alexey Bataev 33c56402d8 [OPENMP] Fix debug info for 'atomic' construct.
Debug info for statement under 'atomic' construct must point exactly to that statement, not the directive itself.

llvm-svn: 255487
2015-12-14 09:26:19 +00:00
NAKAMURA Takumi 2d5c6ddf74 Revert r255001, "Add parse and sema for OpenMP distribute directive and all its clauses excluding dist_schedule."
It causes memory leak. Some tests in test/OpenMP would fail.

llvm-svn: 255094
2015-12-09 04:35:57 +00:00
Alexey Bataev 382967a2e4 [OPENMP 4.5] Parsing/sema for 'num_tasks' clause.
OpenMP 4.5 adds directives 'taskloop' and 'taskloop simd'. These directives support clause 'num_tasks'. Patch adds parsing/semantic analysis for this clause.

llvm-svn: 255008
2015-12-08 12:06:20 +00:00
Carlo Bertolli b9bfa75b28 Add parse and sema for OpenMP distribute directive and all its clauses excluding dist_schedule.
llvm-svn: 255001
2015-12-08 04:21:03 +00:00
Alexey Bataev 1fd4aed26b [OPENMP 4.5] parsing/sema support for 'grainsize' clause.
OpenMP 4.5 adds 'taksloop' and 'taskloop simd' directives, which have 'grainsize' clause. Patch adds parsing/sema analysis of this clause.

llvm-svn: 254903
2015-12-07 12:52:51 +00:00
Alexey Bataev b825de17b7 [OPENMP 4.5] parsing/sema support for 'nogroup' clause.
OpenMP 4.5 adds 'taskloop' and 'taskloop simd' directives. These directives have new 'nogroup' clause. Patch adds basic parsing/sema support for this clause.

llvm-svn: 254899
2015-12-07 10:51:44 +00:00
Serge Pavlov 3a5614599a [PGO] Instrument only base constructors and destructors.
Constructors and destructors may be represented by several functions
in IR. Only base structors correspond to source code, others are
small pieces of code and eventually call the base variant. In this
case instrumentation of non-base structors has little sense, this
fix remove it. Now profile data of a declaration corresponds to
exactly one function in IR, it agrees with the current logic of the
profile data loading.

This change fixes PR24996.

Differential Revision: http://reviews.llvm.org/D15158

llvm-svn: 254876
2015-12-06 14:32:39 +00:00
Alexey Bataev 0a6ed84a0d [OPENMP 4.5] Parsing/sema support for 'omp taskloop simd' directive.
OpenMP 4.5 adds directive 'taskloop simd'. Patch adds parsing/sema analysis for 'taskloop simd' directive and its clauses.

llvm-svn: 254597
2015-12-03 09:40:15 +00:00
Samuel Antao 4af1b7b693 [OpenMP] Update target directive codegen to use 4.5 implicit data mappings.
Summary:
This patch implements the 4.5 specification for the implicit data maps. OpenMP 4.5 specification changes the default way data is captured into a target region. All the non-aggregate kinds are passed by value by default. This required activating the capturing by value during SEMA for the target region. All the non-aggregate values that can be encoded in the size of a pointer are properly casted and forwarded to the runtime library. On top of fixing the previous weird behavior for mapping pointers in nested data regions (an explicit map was always required), this also improves performance as the number of allocations/transactions to the device per non-aggregate map are reduced from two to only one - instead of passing a reference and the value, only the value passed.

Explicit maps will be added later on once firstprivate, private, and map clauses' SEMA and parsing are available.

Reviewers: hfinkel, rjmccall, ABataev

Subscribers: cfe-commits, carlo.bertolli

Differential Revision: http://reviews.llvm.org/D14940

llvm-svn: 254521
2015-12-02 17:44:43 +00:00
Alexey Bataev a056935a2f [OPENMP 4.5] Parsing/sema analysis for 'priority' clause.
OpenMP 4.5 defines new clause 'priority' for 'task', 'taskloop' and 'taskloop simd' directives. Added parsing and sema analysis for 'priority' clause in 'task' and 'taskloop' directives.

llvm-svn: 254398
2015-12-01 10:17:31 +00:00
Alexey Bataev 49f6e78d71 [OPENMP 4.5] Parsing/sema analysis for 'taskloop' directive.
Adds initial parsing and semantic analysis for 'taskloop' directive.

llvm-svn: 254367
2015-12-01 04:18:41 +00:00
Kelvin Li a15fb1a110 [OpenMP] Parsing and sema support for thread_limit clause.
http://reviews.llvm.org/D15029

llvm-svn: 254207
2015-11-27 18:47:36 +00:00
Kelvin Li 099bb8c65d [OpenMP] Parsing and sema support for num_teams clause
http://reviews.llvm.org/D14802

llvm-svn: 254019
2015-11-24 20:50:12 +00:00
Kelvin Li 0bff7afab5 [OpenMP] Parsing and sema support for map clause
http://reviews.llvm.org/D14134 

llvm-svn: 253849
2015-11-23 05:32:03 +00:00
NAKAMURA Takumi 811a09ec5b CGStmtOpenMP.cpp: Prune redundant \param. [-Wdocumentation]
llvm-svn: 249698
2015-10-08 16:41:42 +00:00
Alexey Bataev f24e7b1f60 [OPENMP 4.1] Codegen for array sections/subscripts in 'reduction' clause.
OpenMP 4.1 adds support for array sections/subscripts in 'reduction' clause. Patch adds codegen for this feature.

llvm-svn: 249672
2015-10-08 09:10:53 +00:00
Samuel Antao bed3c46632 [OpenMP] Target directive host codegen.
This patch implements the outlining for offloading functions for code 
annotated with the OpenMP target directive. It uses a temporary naming 
of the outlined functions that will have to be updated later on once 
target side codegen and registration of offloading libraries is 
implemented - the naming needs to be made unique in the produced 
library.

llvm-svn: 249148
2015-10-02 16:14:20 +00:00
Alexey Bataev 5f600d6a49 [OPENMP 4.1] Codegen for ‘simd’ clause in ‘ordered’ directive.
Description.
If the simd clause is specified, the ordered regions encountered by any thread will use only a single SIMD lane to execute the ordered regions in the order of the loop iterations.
Restrictions.
An ordered construct with the simd clause is the only OpenMP construct that can appear in the simd region.

An ordered directive with ‘simd’ clause is generated as an outlined function and corresponding function call to prevent this part of code from vectorization later in backend.

llvm-svn: 248772
2015-09-29 03:48:57 +00:00
Alexey Bataev d14d1e6f25 [OPENMP 4.1] Add 'simd' clause for 'ordered' directive.
Parsing and sema analysis for 'simd' clause in 'ordered' directive.
Description
If the simd clause is specified, the ordered regions encountered by any thread will use only a single SIMD lane to execute the ordered
regions in the order of the loop iterations.
Restrictions
An ordered construct with the simd clause is the only OpenMP construct that can appear in the simd region

llvm-svn: 248696
2015-09-28 06:39:35 +00:00
Alexey Bataev 346265e3bc [OPENMP 4.1] Add 'threads' clause for '#pragma omp ordered'.
OpenMP 4.1 extends format of '#pragma omp ordered'. It adds 3 additional clauses: 'threads', 'simd' and 'depend'.
If no clause is specified, the ordered construct behaves as if the threads clause had been specified. If the threads clause is specified, the threads in the team executing the loop region execute ordered regions sequentially in the order of the loop iterations.
The loop region to which an ordered region without any clause or with a threads clause binds must have an ordered clause without the parameter specified on the corresponding loop directive.

llvm-svn: 248569
2015-09-25 10:37:12 +00:00
Alexey Bataev 87933c7ced [OPENMP 4.0] Add 'if' clause for 'cancel' directive.
Add parsing, sema analysis and codegen for 'if' clause in 'cancel' directive.

llvm-svn: 247976
2015-09-18 08:07:34 +00:00
Alexey Bataev 25e5b44654 [OPENMP] Emit __kmpc_cancel_barrier() and code for 'cancellation point' only if 'cancel' is found.
Patch improves codegen for OpenMP constructs. If the OpenMP region does not have internal 'cancel' construct, a call to 'void __kmpc_barrier()' runtime function is generated for all implicit/explicit barriers. If the region has inner 'cancel' directive, then
```
if (__kmpc_cancel_barrier())
  exit from outer construct;
```
code is generated.
Also, the code for 'canellation point' directive is not generated if parent directive does not have 'cancel' directive.

llvm-svn: 247681
2015-09-15 12:52:43 +00:00
Alexey Bataev c71a4099cf [OPENMP] Preserve alignment of the original variables for the captured references.
Patch makes codegen to preserve alignment of the shared variables captured and used in OpenMP regions.

llvm-svn: 247401
2015-09-11 10:29:41 +00:00
Alexey Bataev 2377fe95c6 [OPENMP] Outlined function for parallel and other regions with list of captured variables.
Currently all variables used in OpenMP regions are captured into a record and passed to outlined functions in this record. It may result in some poor performance because of too complex analysis later in optimization passes. Patch makes to emit outlined functions for parallel-based regions with a list of captured variables. It reduces code for 2*n GEPs, stores and loads at least.
Codegen for task-based regions remains unchanged because runtime requires that all captured variables are passed in captured record.

llvm-svn: 247251
2015-09-10 08:12:02 +00:00
John McCall 7f416cc426 Compute and preserve alignment more faithfully in IR-generation.
Introduce an Address type to bundle a pointer value with an
alignment.  Introduce APIs on CGBuilderTy to work with Address
values.  Change core APIs on CGF/CGM to traffic in Address where
appropriate.  Require alignments to be non-zero.  Update a ton
of code to compute and propagate alignment information.

As part of this, I've promoted CGBuiltin's EmitPointerWithAlignment
helper function to CGF and made use of it in a number of places in
the expression emitter.

The end result is that we should now be significantly more correct
when performing operations on objects that are locally known to
be under-aligned.  Since alignment is not reliably tracked in the
type system, there are inherent limits to this, but at least we
are no longer confused by standard operations like derived-to-base
conversions and array-to-pointer decay.  I've also fixed a large
number of bugs where we were applying the complete-object alignment
to a pointer instead of the non-virtual alignment, although most of
these were hidden by the very conservative approach we took with
member alignment.

Also, because IRGen now reliably asserts on zero alignments, we
should no longer be subject to an absurd but frustrating recurring
bug where an incomplete type would report a zero alignment and then
we'd naively do a alignmentAtOffset on it and emit code using an
alignment equal to the largest power-of-two factor of the offset.

We should also now be emitting much more aggressive alignment
attributes in the presence of over-alignment.  In particular,
field access now uses alignmentAtOffset instead of min.

Several times in this patch, I had to change the existing
code-generation pattern in order to more effectively use
the Address APIs.  For the most part, this seems to be a strict
improvement, like doing pointer arithmetic with GEPs instead of
ptrtoint.  That said, I've tried very hard to not change semantics,
but it is likely that I've failed in a few places, for which I
apologize.

ABIArgInfo now always carries the assumed alignment of indirect and
indirect byval arguments.  In order to cut down on what was already
a dauntingly large patch, I changed the code to never set align
attributes in the IR on non-byval indirect arguments.  That is,
we still generate code which assumes that indirect arguments have
the given alignment, but we don't express this information to the
backend except where it's semantically required (i.e. on byvals).
This is likely a minor regression for those targets that did provide
this information, but it'll be trivial to add it back in a later
patch.

I partially punted on applying this work to CGBuiltin.  Please
do not add more uses of the CreateDefaultAligned{Load,Store}
APIs; they will be going away eventually.

llvm-svn: 246985
2015-09-08 08:05:57 +00:00
Alexey Bataev caacd53dde [OPENMP] Fix for http://llvm.org/PR24674: assertion failed and and abort trap
Fix processing of shared variables with reference types in OpenMP constructs. Previously, if the variable was not marked in one of the private clauses, the reference to this variable was emitted incorrectly and caused an assertion later.

llvm-svn: 246846
2015-09-04 11:26:21 +00:00
Alexey Bataev 7371aa365c [OPENMP 4.1] Codegen for extended format of 'if' clause.
Fixed codegen for extended format of 'if' clauses with special 'directive-name-modifier' + ast-print tests for extended format of 'if' clause.

llvm-svn: 246748
2015-09-03 08:45:56 +00:00
Benjamin Kramer fc600dc2ec [OpenMP] Make the filetered clause iterator a real iterator and type safe.
This replaces the filtered generic iterator with a type-specfic one based
on dyn_cast instead of comparing the kind enum. This allows us to use
range-based for loops and eliminates casts. No functionality change
intended.

llvm-svn: 246384
2015-08-30 15:12:28 +00:00
Alexey Bataev 45bfad51d8 [OPENMP 4.1] Add codegen for 'simdlen' clause.
Add emission of metadata for simd loops in presence of 'simdlen' clause.
If 'simdlen' clause is provided without 'safelen' clause, the vectorizer width for the loop is set to value of 'simdlen' clause + all read/write ops in loop are marked with '!llvm.mem.parallel_loop_access' metadata.
If 'simdlen' clause is provided along with 'safelen' clause, the vectorizer width for the loop is set to value of 'simdlen' clause + all read/write ops in loop are not marked with '!llvm.mem.parallel_loop_access' metadata.
If 'safelen' clause is provided without 'simdlen' clause, the vectorizer width for the loop is set to value of 'safelen' clause + all read/write ops in loop are not marked with '!llvm.mem.parallel_loop_access' metadata.

llvm-svn: 245697
2015-08-21 12:19:04 +00:00
Alexey Bataev 66b15b505f [OPENMP 4.1] Initial support for 'simdlen' clause.
Add parsing/sema analysis for 'simdlen' clause in simd directives. Also add check that if both 'safelen' and 'simdlen' clauses are specified, the value of 'simdlen' parameter is less than the value of 'safelen' parameter.

llvm-svn: 245692
2015-08-21 11:14:16 +00:00
Alexey Bataev bd9fec1eaa [OPENMP 4.1] Allow variables with reference types in private clauses.
OpenMP 4.1 allows to use variables with reference types in all private clauses (private, firstprivate, lastprivate, linear etc.). Patch allows to use such variables and fixes codegen for linear variables with reference types.

llvm-svn: 245268
2015-08-18 06:47:21 +00:00
Alexey Bataev b08f89ffc1 [OPENMP] Fix for http://llvm.org/PR24371: Assert failure compiling blender 2.75.
blender uses statements expression in condition of the loop under control of the '#pragma omp parallel for'. This condition is used several times in different expressions required for codegen of the loop directive. If there are some variables defined in statement expression, it fires an assert during codegen because of redefinition of the same variables.
We have to rebuild several expression to be sure that all variables are unique.

llvm-svn: 245041
2015-08-14 12:25:37 +00:00
Michael Wong b5c1698994 This patch fixes the assert in emitting captured code in the target data construct.
This is on behalf of Kelvin Li.
http://reviews.llvm.org/D11475

llvm-svn: 244569
2015-08-11 04:52:01 +00:00
Filipe Cabecinhas 7af183d841 Propagate SourceLocations through to get a Loc on float_cast_overflow
Summary:
float_cast_overflow is the only UBSan check without a source location attached.
This patch propagates SourceLocations where necessary to get them to the
EmitCheck() call.

Reviewers: rsmith, ABataev, rjmccall

Subscribers: cfe-commits

Differential Revision: http://reviews.llvm.org/D11757

llvm-svn: 244568
2015-08-11 04:19:28 +00:00
Michael Wong e710d5459e This patch commits OpenMP 4 target device clauses
This is committed on behalf of Kelvin Li
http://reviews.llvm.org/D11469?id=31227

llvm-svn: 244325
2015-08-07 16:16:36 +00:00