summaryrefslogtreecommitdiffstats
path: root/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
Commit message (Collapse)AuthorAgeFilesLines
* [OpenMP][NFCI] Use the libFrontend ProcBindKind in ClangJohannes Doerfert2019-12-261-1/+1
| | | | | | | | | | | | This removes the OpenMPProcBindClauseKind enum in favor of llvm::omp::ProcBindKind which lives in OpenMPConstants.h and was introduced in D70109. No change in behavior is expected. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D70289
* [OPENMP50]Support for declare variant directive for NVPTX target.Alexey Bataev2019-10-101-0/+12
| | | | | | | NVPTX does not support global aliases. Instead, we have to copy the full body of the variant function for the original function. llvm-svn: 374387
* [OpenMP] Add support for registering requires directives with the runtimeGheorghe-Teodor Bercea2019-05-211-1/+1
| | | | | | | | | | | | | | | | | | | | | | | Summary: This patch adds support for the registration of the requires directives with the runtime. Each requires directive clause will enable a particular flag to be set. The set of flags is passed to the runtime to be checked for compatibility with other such flags coming from other object files. The registration function is called whenever OpenMP is present even if a requires directive is not present. This helps detect cases in which requires directives are used inconsistently. Reviewers: ABataev, AlexEichenberger, caomhin Reviewed By: ABataev, AlexEichenberger Subscribers: jholewinski, guansong, jfb, jdoerfert, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D60568 llvm-svn: 361298
* [OPENMP]Codegen support for allocate directive on global variables.Alexey Bataev2019-03-211-0/+5
| | | | | | | | | For the global variables the allocate directive must specify only the predefined allocator. This allocator must be translated into the correct form of the address space for the targets that support different address spaces. llvm-svn: 356702
* [OPENMP]Remove unused parameter, NFC.Alexey Bataev2019-03-191-2/+1
| | | | | | | Parameter CodeGenModule &CGM is not required for CGOpenMPRuntime member functions, since class holds the reference to the CGM. llvm-svn: 356480
* [OPENMP]Target region: emit const firstprivates as globals with constantAlexey Bataev2019-03-051-0/+4
| | | | | | | | | | | | memory. If the variable with the constant non-scalar type is firstprivatized in the target region, the local copy is created with the data copying. Instead, we allocate the copy in the constant memory and avoid extra copying in the outlined target regions. This global copy is used in the target regions without loss of the performance. llvm-svn: 355418
* [OPENMP][NVPTX]Use faster teams reduction algorithm.Alexey Bataev2019-02-201-0/+6
| | | | | | | 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
* [opaque pointer types] Pass function types for runtime function calls.James Y Knight2019-02-051-7/+7
| | | | | | | | | | | | | Emit{Nounwind,}RuntimeCall{,OrInvoke} have been modified to take a FunctionCallee as an argument, and CreateRuntimeFunction has been modified to return a FunctionCallee. All callers have been updated. Additionally, CreateBuiltinFunction is removed, as it was redundant with CreateRuntimeFunction after some previous changes. Differential Revision: https://reviews.llvm.org/D57668 llvm-svn: 353184
* Cleanup: replace uses of CallSite with CallBase.James Y Knight2019-01-301-1/+0
| | | | llvm-svn: 352595
* Update the file headers across all of the LLVM projects in the monorepoChandler Carruth2019-01-191-4/+3
| | | | | | | | | | | | | | | | | to reflect the new license. We understand that people may be surprised that we're moving the header entirely to discuss the new license. We checked this carefully with the Foundation's lawyer and we believe this is the correct approach. Essentially, all code in the project is now made available by the LLVM project under our new license, so you will see that the license headers include that license only. Some of our contributors have contributed code under our old license, and accordingly, we have retained a copy of our old license notice in the top-level files in each project and repository. llvm-svn: 351636
* [OPENMP][NVPTX]Use __kmpc_barrier_simple_spmd(nullptr, 0) instead ofAlexey Bataev2019-01-031-0/+3
| | | | | | | | | | 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][NVPTX]Mark __kmpc_barrier functions as convergent.Alexey Bataev2018-12-041-0/+12
| | | | | | | | __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
* [OPENMP][NVPTX]Emit default locations with the correct Exec|RuntimeAlexey Bataev2018-11-261-0/+5
| | | | | | | | | | | 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
* [OPENMP][NVPTX]Emit default locations as constant with undefined mode.Alexey Bataev2018-11-211-0/+10
| | | | | | | | | | | | 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
* [OpenMP] Check target architecture supports unified shared memory for ↵Patrick Lyster2018-11-191-0/+5
| | | | | | requires directive. Differential Review: https://reviews.llvm.org/D54493 llvm-svn: 347214
* [OPENMP][NVPTX]Emit correct reduction code for teams/parallelAlexey Bataev2018-11-161-0/+4
| | | | | | | | | | | 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
* [OPENMP][NVPTX]Allow to use shared memory for theAlexey Bataev2018-11-091-3/+3
| | | | | | | | | | 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
* [OPENMP][NVPTX]Improve emission of the globalized variables forAlexey Bataev2018-11-021-0/+18
| | | | | | | | | | | | | | | | | | | 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
* [OPENMP] Support for mapping of the lambdas in target regions.Alexey Bataev2018-10-301-0/+5
| | | | | | | | | | 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
* [OpenMP][NVPTX] Use single loops when generating code for distribute ↵Gheorghe-Teodor Bercea2018-10-291-1/+1
| | | | | | | | | | | | | | | | 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
* [OPENMP][NVPTX]Reduce memory usage in target region.Alexey Bataev2018-10-121-0/+3
| | | | | | | Additional reduction of the global memory usage in the target regions without parallel regions. llvm-svn: 344413
* [OPENMP][NVPTX]Reduce memory usage in orphaned functions.Alexey Bataev2018-10-121-1/+3
| | | | | | | | | | | | 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
* [OPENMP][NVPTX] Support memory coalescing for globalized variables.Alexey Bataev2018-10-091-2/+14
| | | | | | | | | 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
* [OpenMP] Make default parallel for schedule in NVPTX target regions in SPMD ↵Gheorghe-Teodor Bercea2018-09-271-1/+6
| | | | | | | | | | | | | | | | mode achieve coalescing Summary: Set default schedule for parallel for loops to schedule(static, 1) when using SPMD mode on the NVPTX device offloading toolchain to ensure coalescing. Reviewers: ABataev, Hahnfeld, caomhin Reviewed By: ABataev Subscribers: jholewinski, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D52629 llvm-svn: 343260
* [OpenMP] Make default distribute schedule for NVPTX target regions in SPMD ↵Gheorghe-Teodor Bercea2018-09-271-0/+5
| | | | | | | | | | | | | | | | mode achieve coalescing Summary: For the OpenMP NVPTX toolchain choose a default distribute schedule that ensures coalescing on the GPU when in SPMD mode. This significantly increases the performance of offloaded target code and reduces the number of registers used on the GPU side. Reviewers: ABataev, caomhin, Hahnfeld Reviewed By: ABataev, Hahnfeld Subscribers: Hahnfeld, jholewinski, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D52434 llvm-svn: 343253
* [OPENMP] Fix PR38710: static functions are not emitted as implicitlyAlexey Bataev2018-08-301-2/+3
| | | | | | | | | 'declare target'. All the functions, referenced in implicit|explicit target regions must be emitted during code emission for the device. llvm-svn: 341093
* [OPENMP][NVPTX] Add support for lightweight runtime.Alexey Bataev2018-08-291-0/+1
| | | | | | | | If the target construct can be executed in SPMD mode + it is a loop based directive with static scheduling, we can use lightweight runtime support. llvm-svn: 340953
* [OPENMP, NVPTX] Fix globalization of the variables passed to orphanedAlexey Bataev2018-06-211-6/+6
| | | | | | | | | | parallel region. If the current construct requires sharing of the local variable in the inner parallel region, this variable must be globalized to avoid runtime crash. llvm-svn: 335285
* [OPENMP, NVPTX] Initial support for L2 parallelism in SPMD mode.Alexey Bataev2018-05-101-2/+12
| | | | | | | | Added initial support for L2 parallelism in SPMD mode. Note, though, that the orphaned parallel directives are not currently supported in SPMD mode. llvm-svn: 332016
* Remove \brief commands from doxygen comments.Adrian Prantl2018-05-091-20/+20
| | | | | | | | | | | | | | | | | | | This is similar to the LLVM change https://reviews.llvm.org/D46290. We've been running doxygen with the autobrief option for a couple of years now. This makes the \brief markers into our comments redundant. Since they are a visual distraction and we don't want to encourage more \brief markers in new code either, this patch removes them all. Patch produced by for i in $(git grep -l '\@brief'); do perl -pi -e 's/\@brief //g' $i & done for i in $(git grep -l '\\brief'); do perl -pi -e 's/\\brief //g' $i & done Differential Revision: https://reviews.llvm.org/D46320 llvm-svn: 331834
* [OPENMP, NVPTX] Codegen for critical construct.Alexey Bataev2018-05-071-0/+10
| | | | | | Added correct codegen for the critical construct on NVPTX devices. llvm-svn: 331652
* [OPENMP, NVPTX] Added support for L2 parallelism.Alexey Bataev2018-05-071-21/+27
| | | | | | | Added initial codegen for level 2, 3 etc. parallelism. Currently, all the second, the third etc. parallel regions will run sequentially. llvm-svn: 331642
* [OPENMP] General code improvements.Alexey Bataev2018-04-161-1/+1
| | | | llvm-svn: 330154
* Fix typos in clangAlexander Kornienko2018-04-061-1/+1
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Found via codespell -q 3 -I ../clang-whitelist.txt Where whitelist consists of: archtype cas classs checkk compres definit frome iff inteval ith lod methode nd optin ot pres statics te thru Patch by luzpaz! (This is a subset of D44188 that applies cleanly with a few files that have dubious fixes reverted.) Differential revision: https://reviews.llvm.org/D44188 llvm-svn: 329399
* [OPENMP] Added emission of offloading data sections for declare targetAlexey Bataev2018-03-301-1/+2
| | | | | | | | | | variables. Added emission of the offloading data sections for the variables within declare target regions + fixes emission of the declare target variables marked as declare target not within the declare target region. llvm-svn: 328888
* [OPENMP] Codegen for ctor|dtor of declare target variables.Alexey Bataev2018-03-281-1/+1
| | | | | | | | When the declare target variables are emitted for the device, constructors|destructors for these variables must emitted and registered by the runtime in the offloading sections. llvm-svn: 328705
* [OPENMP, NVPTX] Globalization of the private redeclarations.Alexey Bataev2018-03-201-1/+3
| | | | | | | | If the generic codegen is enabled and private copy of the original variable escapes the declaration context, this private copy should be globalized just like it was the original variable. llvm-svn: 327985
* [OPENMP, NVPTX] Improve globalization of the variables captured by value.Alexey Bataev2018-03-151-5/+16
| | | | | | | | | | | | | If the variable is captured by value and the corresponding parameter in the outlined function escapes its declaration context, this parameter must be globalized. To globalize it we need to get the address of the original parameter, load the value, store it to the global address and use this global address instead of the original. Patch improves globalization for parallel|teams regions + functions in declare target regions. llvm-svn: 327654
* [OpenMP] Add OpenMP data sharing infrastructure using global memoryGheorghe-Teodor Bercea2018-03-141-0/+33
| | | | | | | | | | | | | | | | | Summary: This patch handles the Clang code generation phase for the OpenMP data sharing infrastructure. TODO: add a more detailed description. Reviewers: ABataev, carlo.bertolli, caomhin, hfinkel, Hahnfeld Reviewed By: ABataev Subscribers: jholewinski, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D43660 llvm-svn: 327513
* [OpenMP] Remove implicit data sharing code gen that aims to use device ↵Gheorghe-Teodor Bercea2018-03-071-11/+0
| | | | | | | | | | | | | | | | shared memory Summary: Remove this scheme for now since it will be covered by another more generic scheme using global memory. This code will be worked into an optimization for the generic data sharing scheme. Removing this completely and then adding it via future patches will make all future data sharing patches cleaner. Reviewers: ABataev, carlo.bertolli, caomhin Reviewed By: ABataev Subscribers: jholewinski, guansong, cfe-commits Differential Revision: https://reviews.llvm.org/D43625 llvm-svn: 326948
* [OPENMP] Add debug info for generated functions.Alexey Bataev2018-01-041-1/+2
| | | | | | | Most of the generated functions for the OpenMP were generated with disabled debug info. Patch fixes this for better user experience. llvm-svn: 321816
* [OpenMP] Add implicit data sharing support when offloading to NVIDIA GPUs ↵Gheorghe-Teodor Bercea2017-11-211-0/+11
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | using OpenMP device offloading Summary: This patch is part of the development effort to add support in the current OpenMP GPU offloading implementation for implicitly sharing variables between a target region executed by the team master thread and the worker threads within that team. This patch is the first of three required for successfully performing the implicit sharing of master thread variables with the worker threads within a team. The remaining two patches are: - Patch D38978 to the LLVM NVPTX backend which ensures the lowering of shared variables to an device memory which allows the sharing of references; - Patch (coming soon) is a patch to libomptarget runtime library which ensures that a list of references to shared variables is properly maintained. A simple code snippet which illustrates an implicit data sharing situation is as follows: ``` #pragma omp target { // master thread only int v; #pragma omp parallel { // worker threads // use v } } ``` Variable v is implicitly shared from the team master thread which executes the code in between the target and parallel directives. The worker threads must operate on the latest version of v, including any updates performed by the master. The code generated in this patch relies on the LLVM NVPTX patch (mentioned above) which prevents v from being lowered in the thread local memory of the master thread thus making the reference to this variable un-shareable with the workers. This ensures that the code generated by this patch is correct. Since the parallel region is outlined the passing of arguments to the outlined regions must preserve the original order of arguments. The runtime therefore maintains a list of references to shared variables thus ensuring their passing in the correct order. The passing of arguments to the outlined parallel function is performed in a separate function which the data sharing infrastructure constructs in this patch. The function is inlined when optimizations are enabled. Reviewers: hfinkel, carlo.bertolli, arpith-jacob, Hahnfeld, ABataev, caomhin Reviewed By: ABataev Subscribers: cfe-commits, jholewinski Differential Revision: https://reviews.llvm.org/D38976 llvm-svn: 318773
* [OPENMP] Generalization of calls of the outlined functions.Alexey Bataev2017-08-141-1/+1
| | | | | | General improvement of the outlined functions calls. llvm-svn: 310840
* [OPENMP][DEBUG] Set proper address space info if required by target.Alexey Bataev2017-08-081-0/+20
| | | | | | | | | | | Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310387
* Revert "[OPENMP][DEBUG] Set proper address space info if required by target."Alexey Bataev2017-08-081-25/+0
| | | | | | This reverts commit r310377. llvm-svn: 310379
* [OPENMP][DEBUG] Set proper address space info if required by target.Alexey Bataev2017-08-081-0/+25
| | | | | | | | | | | Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310377
* Revert "[OPENMP][DEBUG] Set proper address space info if required by target."Alexey Bataev2017-08-081-25/+0
| | | | | | This reverts commit r310360. llvm-svn: 310364
* [OPENMP][DEBUG] Set proper address space info if required by target.Alexey Bataev2017-08-081-0/+25
| | | | | | | | | | | Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310360
* Revert "[OPENMP][DEBUG] Set proper address space info if required by target."Alexey Bataev2017-08-041-26/+0
| | | | | | This reverts commit r310104. llvm-svn: 310135
* [OPENMP][DEBUG] Set proper address space info if required by target.Alexey Bataev2017-08-041-0/+26
| | | | | | | | | | | Arguments, passed to the outlined function, must have correct address space info for proper Debug info support. Patch sets global address space for arguments that are mapped and passed by reference. Also, cuda-gdb does not handle reference types correctly, so reference arguments are represented as pointers. llvm-svn: 310104
OpenPOWER on IntegriCloud