summaryrefslogtreecommitdiffstats
path: root/clang/test/OpenMP
Commit message (Collapse)AuthorAgeFilesLines
* [OpenMP 5.0] Parsing/sema support for from clause with mapper modifier.Michael Kruse2019-02-255-3/+49
| | | | | | | | | | | | | | | | | | | | This patch implements the parsing and sema support for the OpenMP 'from'-clause with potential user-defined mappers attached. User-defined mappers are a new feature in OpenMP 5.0. A 'from'-clause can have an explicit or implicit associated mapper, which instructs the compiler to generate and use customized mapping functions. An example is shown below: struct S { int len; int *d; }; #pragma omp declare mapper(id: struct S s) map(s, s.d[0:s.len]) struct S ss; #pragma omp target update from(mapper(id): ss) // use the mapper with name 'id' to map ss from device Contributed-by: Lingda Li <lildmh@gmail.com> Differential Revision: https://reviews.llvm.org/D58638 llvm-svn: 354817
* [OpenMP 5.0] Parsing/sema support for to clause with mapper modifier.Michael Kruse2019-02-225-8/+56
| | | | | | | | | | | | | | | | | | | This patch implements the parsing and sema support for OpenMP to clause with potential user-defined mappers attached. User defined mapper is a new feature in OpenMP 5.0. A to/from clause can have an explicit or implicit associated mapper, which instructs the compiler to generate and use customized mapping functions. An example is shown below: struct S { int len; int *d; }; #pragma omp declare mapper(id: struct S s) map(s, s.d[0:s.len]) struct S ss; #pragma omp target update to(mapper(id): ss) // use the mapper with name 'id' to map ss to device Contributed-by: <lildmh@gmail.com> Differential Revision: https://reviews.llvm.org/D58523 llvm-svn: 354698
* [OPENMP] Delayed diagnostics for VLA support.Alexey Bataev2019-02-221-1/+4
| | | | | | Generalized processing of the deferred diagnostics for OpenMP/CUDA code. llvm-svn: 354690
* Revert "[OPENMP] Delayed diagnostics for VLA support."Alexey Bataev2019-02-221-4/+1
| | | | | | | This reverts commit r354679 to fix the problem with the Windows buildbots llvm-svn: 354680
* [OPENMP] Delayed diagnostics for VLA support.Alexey Bataev2019-02-221-1/+4
| | | | | | Generalized processing of the deferred diagnostics for OpenMP/CUDA code. llvm-svn: 354679
* [OPENMP] Delay emission of the asm target-specific error messages.Alexey Bataev2019-02-201-0/+118
| | | | | | | | | | | | | | | | | Summary: Added the ability to emit target-specific builtin assembler error messages only in case if the function is really is going to be emitted for the device. Reviewers: rjmccall Subscribers: guansong, jdoerfert, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D58243 llvm-svn: 354486
* [OPENMP][NVPTX]Use faster teams reduction algorithm.Alexey Bataev2019-02-201-26/+940
| | | | | | | A faster way to reduce the values in teams reductions was found, the codegen is updated to use this faster algorithm and new runtime functions. llvm-svn: 354479
* [OpenMP 5.0] Parsing/sema support for map clause with mapper modifier.Michael Kruse2019-02-1915-29/+176
| | | | | | | | | | | | | | | | | | | This patch implements the parsing and sema support for OpenMP map clauses with potential user-defined mapper attached. User defined mapper is a new feature in OpenMP 5.0. A map clause can have an explicit or implicit associated mapper, which instructs the compiler to generate extra data mapping. An example is shown below: struct S { int len; int *d; }; #pragma omp declare mapper(id: struct S s) map(s, s.d[0:s.len]) struct S ss; #pragma omp target map(mapper(id) tofrom: ss) // use the mapper with name 'id' to map ss Contributed-by: Lingda Li <lildmh@gmail.com> Differential Revision: https://reviews.llvm.org/D58074 llvm-svn: 354347
* PR40642: Fix determination of whether the final statement of a statementRichard Smith2019-02-156-20/+26
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | expression is a discarded-value expression. Summary: We used to get this wrong in three ways: 1) During parsing, an expression-statement followed by the }) ending a statement expression was always treated as producing the value of the statement expression. That's wrong for ({ if (1) expr; }) 2) During template instantiation, various kinds of statement (most statements not appearing directly in a compound-statement) were not treated as discarded-value expressions, resulting in missing volatile loads (etc). 3) In all contexts, an expression-statement with attributes was not treated as producing the value of the statement expression, eg ({ [[attr]] expr; }). Also fix incorrect enforcement of OpenMP rule that directives can "only be placed in the program at a position where ignoring or deleting the directive would result in a program with correct syntax". In particular, a label (be it goto, case, or default) should not affect whether directives are permitted. Reviewers: aaron.ballman, rjmccall Subscribers: cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D57984 llvm-svn: 354090
* [OPENMP]Delay emission of the error messages for the exceptions.Alexey Bataev2019-02-081-0/+13
| | | | | | | | | | Fixed diagnostic emission for the exceptions support in case of the compilation of OpenMP code for the devices. From now on, it uses delayed diagnostics mechanism, previously used for CUDA only. It allow to diagnose not allowed used of exceptions only in functions that are going to be codegen'ed. llvm-svn: 353542
* Revert "[OPENMP]Initial support for the delayed diagnostics."Alexey Bataev2019-02-081-13/+0
| | | | | | | This reverts commit r353540. Erroneously committed, need to fix the message and description. llvm-svn: 353541
* [OPENMP]Initial support for the delayed diagnostics.Alexey Bataev2019-02-081-0/+13
| | | | | | | | | It is important to delay the emission of the diagnostic messages for the functions unless it is proved that the function is going to be used on the device side. It is required to support compilation with some of the target-specific system headers. llvm-svn: 353540
* [opaque pointer types] Cleanup CGBuilder's Create*GEP.James Y Knight2019-02-0815-58/+57
| | | | | | | | | | | | | | | | | | | | | | | Some of these functions take some extraneous arguments, e.g. EltSize, Offset, which are computable from the Type and DataLayout. Add some asserts to ensure that the computed values are consistent with the passed-in values, in preparation for eliminating the extraneous arguments. This also asserts that the Type is an Array for the calls named "Array" and a Struct for the calls named "Struct". Then, correct a couple of errors: 1. Using CreateStructGEP on an array type. (this causes the majority of the test differences, as struct GEPs are created with i32 indices, while array GEPs are created with i64 indices) 2. Passing the wrong Offset to CreateStructGEP in TargetInfo.cpp on x86-64 NACL (which uses 32-bit pointers). Differential Revision: https://reviews.llvm.org/D57766 llvm-svn: 353529
* [OPENMP] issue error messages for multiple teams contructs in a target constructKelvin Li2019-02-051-0/+14
| | | | | | | | | | | | | | | | The fix is to issue error messages if there are more than one teams construct inside a target constructs. #pragma omp target { #pragma omp teams { ... } #pragma omp teams { ... } } llvm-svn: 353186
* [OpenMP] Adding support to the mutexinoutset dep-typeSergi Mateo Bellido2019-02-0415-68/+107
| | | | | | | | | | | Summary: this commit adds support to a new dependence type introduced in OpenMP 5.0. The LLVM OpenMP RTL already supports this feature, so we only need to modify CLANG to take advantage of them. Differential Revision: https://reviews.llvm.org/D57576 llvm-svn: 353018
* [OpenMP 5.0] Parsing/sema support for "omp declare mapper" directive.Michael Kruse2019-02-014-0/+256
| | | | | | | | | | | | | | | | | This patch implements parsing and sema for "omp declare mapper" directive. User defined mapper, i.e., declare mapper directive, is a new feature in OpenMP 5.0. It is introduced to extend existing map clauses for the purpose of simplifying the copy of complex data structures between host and device (i.e., deep copy). An example is shown below: struct S { int len; int *d; }; #pragma omp declare mapper(struct S s) map(s, s.d[0:s.len]) // Memory region that d points to is also mapped using this mapper. Contributed-by: Lingda Li <lildmh@gmail.com> Differential Revision: https://reviews.llvm.org/D56326 llvm-svn: 352906
* [OPENMP]Fix PR40536: Do not emit __kmpc_push_target_tripcount if notAlexey Bataev2019-01-303-2/+5
| | | | | | | | | | required. Function __kmpc_push_target_tripcount should be emitted only if the offloading entry is going to be emitted (for use in tgt_target... functions). Otherwise, it should not be emitted. llvm-svn: 352669
* [OPENMP]Fix PR40513: lastprivate taskloop counter.Alexey Bataev2019-01-291-2/+15
| | | | | | | | We don't need to use the predetermined data-sharing attributes for the loop counters if the user explicitly specified correct data-sharing attributes for such variables. llvm-svn: 352543
* [OPENMP]Make the loop with unsigned counter countable.Alexey Bataev2019-01-293-5/+10
| | | | | | | | According to the report, better to keep the original strict compare operation as the loop condition with unsigned loop counters to make the loop countable. This allows further loop transformations. llvm-svn: 352526
* [OPENMP][NVPTX]Emit service debug variable for NVPTX.Alexey Bataev2019-01-281-0/+1
| | | | | | | In case of the empty module, the ptxas tool may emit error message about empty debug info sections. This patch fixes this bug. llvm-svn: 352421
* Emit !callback metadata and introduce the callback attributeJohannes Doerfert2019-01-191-3/+4
| | | | | | | | | | | | | | | | | | | | | | | | | | | With commit r351627, LLVM gained the ability to apply (existing) IPO optimizations on indirections through callbacks, or transitive calls. The general idea is that we use an abstraction to hide the middle man and represent the callback call in the context of the initial caller. It is described in more detail in the commit message of the LLVM patch r351627, the llvm::AbstractCallSite class description, and the language reference section on callback-metadata. This commit enables clang to emit !callback metadata that is understood by LLVM. It does so in three different cases: 1) For known broker functions declarations that are directly generated, e.g., __kmpc_fork_call for the OpenMP pragma parallel. 2) For known broker functions that are identified by their name and source location through the builtin detection, e.g., pthread_create from the POSIX thread API. 3) For user annotated functions that carry the "callback(callee, ...)" attribute. The attribute has to include the name, or index, of the callback callee and how the passed arguments can be identified (as many as the callback callee has). See the callback attribute documentation for detailed information. Differential Revision: https://reviews.llvm.org/D55483 llvm-svn: 351629
* [OpenMP] Avoid remainder operations for loop index values on a collapsed ↵Gheorghe-Teodor Bercea2019-01-094-36/+188
| | | | | | | | | | | | | | | | loop nest. Summary: Change the strategy for computing loop index variables after collapsing a loop nest via the collapse clause by replacing the expensive remainder operation with multiplications and additions. Reviewers: ABataev, caomhin Reviewed By: ABataev Subscribers: guansong, arphaman, cfe-commits Differential Revision: https://reviews.llvm.org/D56413 llvm-svn: 350759
* [OpenMP] Add flag for preventing the extension to 64 bits for the collapse ↵Gheorghe-Teodor Bercea2019-01-091-10/+15
| | | | | | | | | | | | | | | | loop counter Summary: Introduce a compiler flag for cases when the user knows that the collapsed loop counter can be safely represented using at most 32 bits. This will prevent the emission of expensive mathematical operations (such as the div operation) on the iteration variable using 64 bits where 32 bit operations are sufficient. Reviewers: ABataev, caomhin Reviewed By: ABataev Subscribers: hfinkel, kkwli0, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D55928 llvm-svn: 350758
* Incorrect implicit data-sharing for nested tasksAlexey Bataev2019-01-091-1/+19
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: There is a minor issue in how the implicit data-sharings for nested tasks are computed. For the following example: ``` int x; #pragma omp task shared(x) #pragma omp task x++; ``` We compute an implicit data-sharing of shared for `x` in the second task although I think that it should be firstprivate. Below you can find the part of the OpenMP spec that covers this example: - // 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 and that in the enclosing context is determined to be shared by all implicit tasks bound to the current team is shared.// - //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 firstprivate.// Since each implicit-task has its own copy of `x`, we shouldn't apply the first rule. Reviewers: ABataev Reviewed By: ABataev Subscribers: cfe-commits, rogfer01 Tags: #openmp Differential Revision: https://reviews.llvm.org/D56430 llvm-svn: 350734
* [OPENMP]Fix PR40191: Do not allow orphaned cancellation constructs.Alexey Bataev2019-01-082-40/+40
| | | | | | Prohibited use of the orphaned cancellation directives. llvm-svn: 350634
* [OPENMP]Add call to __kmpc_push_target_tripcount() function.Alexey Bataev2019-01-078-2/+37
| | | | | | | | | | Each we create the target regions with the teams distribute inner region, we can better estimate number of the teams required to execute the target region. Function __kmpc_push_target_tripcount() is used for purpose, which accepts device_id and the number of the iterations, performed by the associated loop. llvm-svn: 350571
* [OPENMP][NVPTX]Reduce number of barriers in reductions.Alexey Bataev2019-01-072-5/+0
| | | | | | | After the fix for the syncthreads we don't need to generate extra barriers for the parallel reductions. llvm-svn: 350530
* [OpenMP] Refactor const restriction for linearJoel E. Denny2019-01-0415-30/+30
| | | | | | | | | | | | | | As discussed in D56113, this patch refactors the implementation of the const restriction for linear to reuse a function introduced by D56113. A side effect is that, if a variable has mutable members, this diagnostic is now skipped, and the diagnostic for the variable not being an integer or pointer is reported instead. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D56299 llvm-svn: 350441
* [OpenMP] Refactor const restriction for reductionsJoel E. Denny2019-01-0432-448/+448
| | | | | | | | | | | | | As discussed in D56113, this patch refactors the implementation of the const restriction for reductions to reuse a function introduced by D56113. A side effect is that diagnostics sometimes now say "variable" instead of "list item" when a list item is a variable. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D56298 llvm-svn: 350440
* [OpenMP] Replace predetermined shared for const variableJoel E. Denny2019-01-0438-349/+357
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | The following appears in OpenMP 3.1 sec. 2.9.1.1 as a predetermined data-sharing attribute: > Variables with const-qualified type having no mutable member are > shared. It does not appear in OpenmP 4.0, 4.5, or 5.0. This patch removes the implementation of that attribute when the requested OpenMP version is greater than 3.1. One effect of that removal is that `default(none)` affects const variables without mutable members. Also, without this patch, if a const variable without mutable members was explicitly lastprivate or private, it was an error because it was predetermined shared. Now, clang instead complains that it's const without mutable fields, which is a more intelligible diagnostic. That should be fine for all of the above versions because they all have something like the following, which is quoted from OpenMP 5.0 sec. 2.19.3: > A variable that is privatized must not have a const-qualified type > unless it is of class type with a mutable member. This restriction does > not apply to the firstprivate clause. reduction and linear clauses already have separate checks for const variables. Future patches will merge the implementations. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D56113 llvm-svn: 350439
* [OPENMP][NVPTX]Use new functions from the runtime library.Alexey Bataev2019-01-0411-24/+24
| | | | | | Updated codegen to use the new functions from the runtime library. llvm-svn: 350415
* [OPENMP][NVPTX]Use __kmpc_barrier_simple_spmd(nullptr, 0) instead ofAlexey Bataev2019-01-036-51/+51
| | | | | | | | | | nvvm_barrier0. Use runtime functions instead of the direct call to the nvvm intrinsics. It allows to prevent some dangerous LLVM optimizations, that breaks the code for the NVPTX target. llvm-svn: 350328
* [OpenMP] Added support for explicit mapping of classes using 'this' pointer. ↵Patrick Lyster2019-01-023-2/+95
| | | | | | Differential revision: https://reviews.llvm.org/D55982 llvm-svn: 350252
* [OPENMP]Fix processing of the clauses on target combined directives.Alexey Bataev2018-12-281-0/+21
| | | | | | | | For constants with the predefined data-sharing clauses we may had troubles with the target combined directives. It may cause compiler crash in some corner cases. llvm-svn: 350127
* Pass a concrete triple for two OpenMP tests that depend on TLSNico Weber2018-12-262-26/+26
| | | | | | | | | | | | | | | Not all %itanium_abi_triple values support TLS. Makes OpenMP/declare_reduction_codegen.cpp, OpenMP/parallel_copyin_codegen.cpp for %itanium_abi_triples without TLS support. Alternatively we could pass -fnoopenmp-use-tls and tweak some of the CHECK lines, but this seems a bit simpler. Fixes PR40156. Differential Revision: https://reviews.llvm.org/D56086 llvm-svn: 350067
* [CodeGen] Generate llvm.loop.parallel_accesses instead of ↵Michael Kruse2018-12-2011-221/+228
| | | | | | | | | | | | | | 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
* [OPENMP]Mark the loop as started when initialized.Alexey Bataev2018-12-191-1/+12
| | | | | | | | | 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
* [OpenMP] Fix data sharing analysis in nested clauseJoel E. Denny2018-12-191-0/+10
| | | | | | | | | | | | | | | | | | | | | | 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
* [OPENMP] parsing and sema support for 'close' map-type-modifierKelvin Li2018-12-1812-24/+170
| | | | | | | | | | | | 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
* [OPENMP][NVPTX]Emit shared memory buffer for reduction as 128 bytesAlexey Bataev2018-12-1810-11/+11
| | | | | | | | | | | 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
* [OPENMP][NVPTX]Added extra sync point to the inter-warp copy function.Alexey Bataev2018-12-181-0/+4
| | | | | | | The parallel reduction operation requires an extra synchronization point in the inter-warp copy function to avoid divergence. llvm-svn: 349525
* [OPENMP][NVPTX]Improved interwarp copy function.Alexey Bataev2018-12-142-28/+21
| | | | | | | | | 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
* [OPENMP][NVPTX] Fix globalization of the mapped array sections.Alexey Bataev2018-12-061-19/+52
| | | | | | | | | 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
* [OPENMP][NVPTX]Fixed emission of the critical region.Alexey Bataev2018-12-041-1/+7
| | | | | | | | | | | | | 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
* [OPENMP][NVPTX]Mark __kmpc_barrier functions as convergent.Alexey Bataev2018-12-041-2/+8
| | | | | | | | __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
* Move AST tests into their own test directory; NFC.Aaron Ballman2018-11-301-84/+0
| | | | | | 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
* [OPENMP][NVPTX]Call get __kmpc_global_thread_num in worker afterAlexey Bataev2018-11-291-1/+1
| | | | | | | | | initialization. Function __kmpc_global_thread_num should be called only after initialization, not earlier. llvm-svn: 347919
* [OpenMP] Add a new version of the SPMD deinit kernel functionGheorghe-Teodor Bercea2018-11-298-20/+20
| | | | | | | | | | | | | | 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
* [OPENMP]Fix emission of the target regions in virtual functions.Alexey Bataev2018-11-281-0/+24
| | | | | | | | 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
* [OPENMP][NVPTX]Basic support for reductions across the teams.Alexey Bataev2018-11-271-948/+7
| | | | | | Added basic codegen support for the reductions across the teams. llvm-svn: 347715
OpenPOWER on IntegriCloud