summaryrefslogtreecommitdiffstats
path: root/openmp/libomptarget/deviceRTLs/nvptx/src
Commit message (Collapse)AuthorAgeFilesLines
...
* Factor architecture dependent code out of loop.cuJon Chesterfield2019-08-132-9/+49
| | | | | | | | | | | | | | | | | | | | | Summary: [libomptarget] Factor architecture dependent code out of loop.cu Related to the patch series starting D64217. Added subscribers to said series as reviewers. This effort is smaller in scope. This patch factors out just enough architecture dependent code from loop.cu to allow the same source to be used with amdgcn, given a different target_impl.h. Testing is that the same bitcode (modulo variable names) is generated for libomptarget before and after the refactor, for nvptx and the out of tree amdgcn. Reviewers: jdoerfert, ABataev, bollu, jfb, tra, grokos, Hahnfeld, guansong, xtian, gregrodgers, ronlieb, hfinkel, gtbercea, guraypp, arpith-jacob Reviewed By: jdoerfert, ABataev Subscribers: dexonsmith, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D65836 llvm-svn: 368751
* Use forceinline. Necessary for nvcc to inline small functions within the ↵Jon Chesterfield2019-08-071-1/+1
| | | | | | | | | | | | | | | | | | bitcode library Summary: [libomptarget] Use forceinline. Necessary for nvcc to inline small functions within the bitcode library Suggested in D65836 Reviewers: ABataev, jdoerfert, grokos, gregrodgers Subscribers: openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D65876 llvm-svn: 368177
* [OPENMP][NVPTX]Perform memory flush if number of threads to sync is 1 or less.Alexey Bataev2019-07-251-0/+3
| | | | | | | | | | | | | | | | | Summary: According to the OpenMP standard, barrier operation must perform implicit flush operation. Currently, if there is only one thread in the team, barrier does not flush the memory. Patch fixes this problem. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D62398 llvm-svn: 367024
* [OPENMP][NVPTX]Fixed checks for cuda versions.Alexey Bataev2019-07-161-2/+4
| | | | | | | | | | | | | | | | | | | | | Summary: We used CUDART_VERSION macro to check for the installed cuda version but this macro is defined in cuda_runtime_api.h, which is not used by project. Better to use CUDA_VERSION macro, which is defined in cuda.h. Also, added the check if this macro is defined. If macro is undefined, there is something wrong with the cuda configuration and we should not continue the compilation. This also fixes problems with runtime building in cuda 10+. Reviewers: grokos Subscribers: guansong, jdoerfert, caomhin, kkwli0, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D64648 llvm-svn: 366224
* [libomptarget-nvptx] Remove dead functionsJonas Hahnfeld2019-07-112-85/+1
| | | | | | | | | These entry points are never called by Clang trunk nor clang-ykt. If XL doesn't use them either, they can finally go away. Differential Revision: https://reviews.llvm.org/D52700 llvm-svn: 365817
* [OPENMP][NVPTX]Relax flush directive.Alexey Bataev2019-06-271-1/+1
| | | | | | | | | | | | | | | | | | | | | | | Summary: According to the OpenMP standard, flush makes a thread’s temporary view of memory consistent with memory and enforces an order on the memory operations of the variables explicitly specified or implied. According to the Cuda toolkit documentation (https://docs.nvidia.com/cuda/archive/8.0/cuda-c-programming-guide/index.html#memory-fence-functions), __threadfence() functions provides required functionality. __threadfence_system() also provides required functionality, but it also includes some extra functionality, like synchronization of page-locked host memory, synchronization for the host, etc. It is not required per the standard and we can use more relaxed version of memory fence operation. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jfb, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D62397 llvm-svn: 364572
* [OPENMP][CUDA]Use __syncthreads when compiled by nvcc and clang >= 9.0.Alexey Bataev2019-06-191-0/+5
| | | | | | | | | | | | | | | | Summary: The problems with __syncthreads() were fixed in clang >= 9.0 and the original __syncthreads() can be used instead of the ptx instruction. Reviewers: grokos Subscribers: guansong, jdoerfert, openmp-commits, kkwli0, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D63515 llvm-svn: 363807
* Revert "[OPENMP][NVPTX]Fix barriers and parallel level counters, NFC."Alexey Bataev2019-05-243-6/+3
| | | | | | This reverts commit r361421 to split the patch into 3 parts. llvm-svn: 361638
* [OPENMP][NVPTX]Fix barriers and parallel level counters, NFC.Alexey Bataev2019-05-223-3/+6
| | | | | | | | | | | | | | | | | | | | | Summary: Parallel level counter should be volatile to prevent some dangerous optimiations by the ptxas. Otherwise, ptxas optimizations lead to undefined behaviour in some cases. Also, use __threadfence() for #pragma omp flush and if the barrier should not be used (we have only one thread in the team), still perform flush operation since the standard requires implicit flush when executing barriers. Reviewers: gtbercea, kkwli0, grokos Subscribers: guansong, jfb, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D62199 llvm-svn: 361421
* [OPENMP][NVPTX]Simplify handling of thread limit, NFC.Alexey Bataev2019-05-136-26/+19
| | | | | | | | | | | | | | | | | Summary: Patch improves performance of the full runtime mode by moving threads limit counter to the shared memory. It also allows to save global memory. Reviewers: grokos, kkwli0, gtbercea Subscribers: guansong, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D61801 llvm-svn: 360584
* [OPENMP][NVPTX]Improve number of threads counter, NFC.Alexey Bataev2019-05-1011-83/+47
| | | | | | | | | | | | | | | | | Summary: Patch improves performance of the full runtime mode by moving number-of-threads counter to the shared memory. It also allows to save global memory. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jfb, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D61785 llvm-svn: 360457
* [OPENMP][NVPTX]Improve thread limit counter, NFC.Alexey Bataev2019-05-035-12/+6
| | | | | | | | | | | | | | | | | Summary: Patch improves performance of the full runtime mode by moving thread-limit counter to the shared memory. It also allows to save global memory. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jdoerfert, caomhin, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D61526 llvm-svn: 359922
* [OPENMP][NVPTX]Improved several standard OpenMP functions, NFC.Alexey Bataev2019-05-031-20/+19
| | | | | | | | | | | | | | | | | Summary: Used parallelLevel[] counter to simplify and improve implementation of the existing standard OpenMP functions. Functions are tested already in several tests, the patch is NFC. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jdoerfert, caomhin, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D61459 llvm-svn: 359892
* [OPENMP][NVPTX]Improve code by using parallel level counter.Alexey Bataev2019-05-029-190/+129
| | | | | | | | | | | | | | | | | | Summary: Previously for the different purposes we need to get the active/common parallel level and with full runtime we iterated over all the records to calculate this level. Instead, we can used the warp-based parallel level counters used in no-runtime mode. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jfb, jdoerfert, caomhin, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D61395 llvm-svn: 359822
* [OPENMP][NVPTX]Improve omp_get_max_threads() function.Alexey Bataev2019-05-021-5/+2
| | | | | | | | | | | | | | | | Summary: Function omp_get_max_threads() can always return 1 if current execution mode is SPMD. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jdoerfert, caomhin, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D61379 llvm-svn: 359792
* [OPENMP][NVPTX]Improved omp_get_thread_limit() function.Alexey Bataev2019-05-021-6/+3
| | | | | | | | | | | | | | | | Summary: Function omp_get_thread_limit() in SPMD mode can return the maximum available number of threads as a result. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D61378 llvm-svn: 359790
* [OPENMP][NVPTX]Correctly handle L2 parallelism in SPMD mode.Alexey Bataev2019-04-267-15/+28
| | | | | | | | | | | | | | | | | | | Summary: The parallelLevel counter must be on per-thread basis to fully support L2+ parallelism, otherwise we may end up with undefined behavior. Introduce the parallelLevel on per-warp basis using shared memory. It allows to avoid the problems with the synchronization and allows fully support L2+ parallelism in SPMD mode with no runtime. Reviewers: gtbercea, grokos Subscribers: guansong, jdoerfert, caomhin, kkwli0, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D60918 llvm-svn: 359341
* [OPENMP][NVPTX]Fix dynamic scheduling in L2+ SPMD parallel regions.Alexey Bataev2019-04-153-6/+18
| | | | | | | | | | | | | | | | | | | | Summary: If the kernel is executed in SPMD mode and the L2+ parallel for region with the dynamic scheduling is executed, dynamic scheduling functions are called. They expect full runtime support, but SPMD kernels may be executed without the full runtime. It leads to the runtime crash of the compiled program. Patch fixes this problem + fixes handling of the parallelism level in SPMD mode, which is required as part of this patch. Reviewers: gtbercea, kkwli0, grokos Subscribers: guansong, jdoerfert, openmp-commits, caomhin Tags: #openmp Differential Revision: https://reviews.llvm.org/D60578 llvm-svn: 358442
* [OpenMP][libomptarget] New reduction scheme for team reductionsGheorghe-Teodor Bercea2019-02-202-0/+148
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Summary: This patch adds a more sophisticated team reduction scheme to the OpenMP libomptarget-nvptx runtime. The scheme uses a fixed size global memory buffer whose length can be adjusted via compiler flag: ``` -fopenmp-cuda-teams-reduction-recs-num=1024 ``` The global buffer is a structure of arrays (with default size of 1024 each and controlled by the above flag), one array for each reduction variable. Values in the buffer are processed by the last team to finish executing the body of the target region. In addition to adding support for the new flag, the compiler also emits special functions used for the reduction of the intermediate reduction values. These changes will be added in a separate compiler patch following this one. Reviewers: ABataev, caomhin Reviewed By: ABataev Subscribers: guansong, jfb, jdoerfert, openmp-commits Tags: #openmp Differential Revision: https://reviews.llvm.org/D58409 llvm-svn: 354471
* Update more file headers across all of the LLVM projects in the monorepoChandler Carruth2019-01-1920-80/+60
| | | | | | | | | | | | | | | | | | to reflect the new license. These used slightly different spellings that defeated my regular expressions. 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: 351648
* [OpenMP][libomptarget] Use shared memory variable for tracking parallel levelGheorghe-Teodor Bercea2019-01-096-69/+21
| | | | | | | | | | | | | | Summary: Replace existing infrastructure for tracking parallel level using global memory with a per-team shared memory variable. This minimizes the impact of the overhead of tracking the parallel level for non-nested cases. Reviewers: ABataev, caomhin Reviewed By: ABataev Subscribers: guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D55773 llvm-svn: 350747
* [OPENMP][NVPTX]Fix dynamic scheduling.Alexey Bataev2019-01-072-22/+47
| | | | | | | | | | | | | | | Summary: Previous implementation may cause the runtime crash when the number of teams is > 1024. Patch fixes this problem + reduces number of the atomic operations by 32 times. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jfb, openmp-commits, caomhin Differential Revision: https://reviews.llvm.org/D56332 llvm-svn: 350524
* [OPENMP][NVPTX]General formatting/code improvement, NFC.Alexey Bataev2019-01-0410-81/+57
| | | | | | | | | | | | Summary: Formatting. Reviewers: gtbercea, grokos, kkwli0 Subscribers: guansong, openmp-commits, caomhin Differential Revision: https://reviews.llvm.org/D56290 llvm-svn: 350431
* [OPENMP][NVPTX]Improve performance + reduce number of used registers.Alexey Bataev2019-01-0413-83/+124
| | | | | | | | | | | | | | | | | Summary: Reduced number of the used register + improved performance propagating the information about current execution/data sharing mode directly from the compiler, where it is possible. In some cases, it requires new/reworked interfaces of the runtime external functions. Old functions are marked as deprecated. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, jfb, openmp-commits, caomhin Differential Revision: https://reviews.llvm.org/D56278 llvm-svn: 350405
* [OPENMP][NVPTX]Fix incompatibility of __syncthreads with LLVM, NFC.Alexey Bataev2019-01-034-6/+15
| | | | | | | | | | | | | | | | | | Summary: One of the LLVM optimizations, split critical edges, also clones tail instructions. This is a dangerous operation for __syncthreads() functions and this transformation leads to undefined behavior or incorrect results. Patch fixes this problem by replacing __syncthreads() function with the assembler instruction, which cost is too high and wich cannot be copied. Reviewers: grokos, gtbercea, kkwli0 Subscribers: guansong, openmp-commits, caomhin Differential Revision: https://reviews.llvm.org/D56274 llvm-svn: 350333
* [OPENMP][NVPTX]Added/fixed debugging messages, NFC.Alexey Bataev2018-12-283-3/+7
| | | | | | | | | | | | | | Summary: Added or fixed new/old debugging messages for the better diagnostics. Reviewers: gtbercea, kkwli0, grokos Reviewed By: grokos Subscribers: caomhin, guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D56102 llvm-svn: 350137
* [OPENMP][NVPTX]Fixed initialization of the data-sharing interface.Alexey Bataev2018-12-281-6/+9
| | | | | | | | | | | | | | | | Summary: Avoid using of the atomic loop to wait for the completion of the data-sharing interface initialization, use __shfl_sync instead for the communication within the warp to signal other threads in the warp about completion of the initialization. Reviewers: gtbercea, kkwli0, grokos Subscribers: guansong, jfb, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D56100 llvm-svn: 350129
* [OPENMP][NVPTX]Outline assert into noinline function, NFC.Alexey Bataev2018-12-281-8/+19
| | | | | | | | | | | | | | | | Summary: At high optimization level asserts lead to some unexpected results because of auto-inserted unreachable instructions. This outlining prevents some of such dangerous optimizations and leads to better stability. Reviewers: gtbercea, kkwli0, grokos Subscribers: guansong, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D56101 llvm-svn: 350128
* [OPENMP][NVPTX]Revert __kmpc_shuffle_int64 to its original form.Alexey Bataev2018-12-101-5/+0
| | | | | | | | | | | | | | Summary: Use the original shuffle implementation for __kmpc_shuffle_int64 since default implementation uses the same implementation. Reviewers: gtbercea Subscribers: guansong, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D55514 llvm-svn: 348772
* [OPENMP][NVPTX]Enable fast shuffles on 64bit values only if CUDA >= 9.Alexey Bataev2018-12-101-1/+11
| | | | | | | | | | | | | | Summary: Shuffle on 64bit data is allowed only for CUDA >= 9.0. Also, fixed the constant for the mask, need one extra L in the end. Reviewers: gtbercea, kkwli0 Subscribers: guansong, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D55440 llvm-svn: 348758
* [OPENMP][NVPTX]Save registers for optimized builds with enabled logging.Alexey Bataev2018-12-071-18/+14
| | | | | | | | | | | | | | | | | Summary: Introduced special noinline function log that allows to save some registers for optimized builds but with enabled logging. Also, it increases the stability of the optimized builds with inlined runtime. Reviewers: gtbercea, kkwli0 Reviewed By: gtbercea Subscribers: caomhin, guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D55436 llvm-svn: 348606
* [OPENMP][NVPTX]Correct type casting for printf args + simplified shfl64 ↵Alexey Bataev2018-12-0611-108/+140
| | | | | | | | | | | | | | | | function. Summary: Explicitly casted printf's args to the required types + simplified shfl64 function. Reviewers: gtbercea, kkwli0 Subscribers: guansong, jfb, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D55379 llvm-svn: 348521
* [OPENMP][NVPTX]Fix __kmpc_flush to flush the memory per system, not per block.Alexey Bataev2018-12-061-1/+1
| | | | | | | | | | | | | | | Summary: According to the standard, after memory flushing the changes in the memory must be visible to all the threads in all teams. Patch fixes this. Reviewers: gtbercea, kkwli0 Subscribers: guansong, jfb, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D55370 llvm-svn: 348491
* [OpenMP][libomptarget] Flush intermediate values during team reduction Gheorghe-Teodor Bercea2018-12-031-0/+1
| | | | | | | | | | | | | | Summary: Ensure intermediate values of a team reduction are flushed to memory. Reviewers: ABataev, caomhin Reviewed By: ABataev Subscribers: guansong, jfb, openmp-commits Differential Revision: https://reviews.llvm.org/D55219 llvm-svn: 348148
* [OPENMP][NVPTX]Make runtime compatible with the original runtime.Alexey Bataev2018-11-305-159/+87
| | | | | | | | | | | | | | | Summary: Reworked runtime to make it compatible with the requirements of the original runtime library. Also, simplified some code to reduce number of function calls. Reviewers: gtbercea, kkwli0 Subscribers: guansong, jfb, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D55130 llvm-svn: 348003
* [OpenMP][libomptarget] Add new version of SPMD deinit kernel function with ↵Gheorghe-Teodor Bercea2018-11-272-3/+8
| | | | | | | | | | | | | | | | argument Summary: To enable the compiler to optimize parts of the function that are not needed when runtime can be omitted, a new version of the SPMD deinit kernel function is needed. This function takes the runtime required flag as an argument. Reviewers: ABataev, kkwli0, caomhin Reviewed By: ABataev Subscribers: guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D54969 llvm-svn: 347714
* [OPENMP][NVPTX]Basic support for reductions across the teams.Alexey Bataev2018-11-272-0/+24
| | | | | | | | | | | | | | | Summary: Added functions __kmpc_nvptx_teams_reduce_nowait_simple and __kmpc_nvptx_teams_end_reduce_nowait_simple to implement basic support for reductions across the teams. Reviewers: gtbercea, kkwli0 Subscribers: guansong, jfb, caomhin, openmp-commits Differential Revision: https://reviews.llvm.org/D54967 llvm-svn: 347710
* [OpenMP][libomptarget] Refactor SPMD and runtime requirement checkingGheorghe-Teodor Bercea2018-11-279-171/+262
| | | | | | | | | | | | | | Summary: Refactor the checking for SPMD mode and whether the runtime is initialized or not. This uses constant flags which enables the runtime to optimize out unused sections of code that depend on these flags. Reviewers: ABataev, caomhin Reviewed By: ABataev Subscribers: guansong, jfb, openmp-commits Differential Revision: https://reviews.llvm.org/D54960 llvm-svn: 347698
* [OPENMP][NVPTX]Improved lock/critical constructs.Alexey Bataev2018-11-203-24/+11
| | | | | | | | | | | | Summary: Improved support for critical constructs + omp_..._lock... constructs. Reviewers: gtbercea, kkwli0, caomhin Subscribers: guansong, jfb, openmp-commits Differential Revision: https://reviews.llvm.org/D54766 llvm-svn: 347342
* [OPENMP][NVPTX]Fixed/improved support for globalization in team contexts.Alexey Bataev2018-11-026-71/+102
| | | | | | | | | | | | | | | | | | | Summary: Current globalization scheme works correctly only for SPMD+lightweight runtime mode and does not work for full runtime. Patch improves support for the globalization scheme + reduces global memory consumption in lightweight runtime mode. Patch adds runtime functions to work with the statically allocated global memory. It allows to improve performance and memory consumption. This global memory must be allocated by the compiler. Reviewers: grokos, kkwli0, gtbercea, caomhin Subscribers: guansong, jfb, openmp-commits Differential Revision: https://reviews.llvm.org/D53943 llvm-svn: 345976
* [OpenMP][libomptarget] Add runtime function for pushing coalesced global recordsGheorghe-Teodor Bercea2018-11-015-35/+45
| | | | | | | | | | | | | | Summary: In the case of coalesced global records, we need to push the exact data size passed in. This patch fixes this by outlining the common functionality of the previous push function and by adding a separate entry point for coalesced pushes. The pop function remains unchanged. Reviewers: ABataev, grokos, caomhin Reviewed By: ABataev, grokos Subscribers: jholewinski, cfe-commits, Hahnfeld, guansong, jfb, openmp-commits Differential Revision: https://reviews.llvm.org/D53141 llvm-svn: 345867
* [libomptarget-nvptx] reduction: Determine if runtime uninitializedJonas Hahnfeld2018-10-011-8/+10
| | | | | | | | | | Pass in the correct value of isRuntimeUninitialized() which solves parallel reductions as reported on the mailing list. For reference: r333285 did the same for loop scheduling. Differential Revision: https://reviews.llvm.org/D52725 llvm-svn: 343476
* [libomptarget-nvptx] Align data sharing stackJonas Hahnfeld2018-09-301-0/+7
| | | | | | | | | | | | NVPTX requires addresses of pointer locations to be 8-byte aligned or there will be an exception during runtime. This could happen without this patch as shown in the added test: getId() requires 4 byte of stack and putValueInParallel() uses 16 bytes to store the addresses of the captured variables. Differential Revision: https://reviews.llvm.org/D52655 llvm-svn: 343402
* [libomptarget-nvptx] Fix ancestor_thread_num and team_size (non-SPMD)Jonas Hahnfeld2018-09-301-4/+8
| | | | | | | | | | | | | | According to OpenMP 4.5, p250:12-14: If the requested nest level is outside the range of 0 and the nest level of the current thread, as returned by the omp_get_level routine, the routine returns -1. The SPMD code path will need a similar fix. Differential Revision: https://reviews.llvm.org/D51787 llvm-svn: 343401
* [libomptarget-nvptx] Ignore calls to dynamic APIJonas Hahnfeld2018-09-292-37/+6
| | | | | | | | | | | | | | | There is no support and according to the OpenMP 4.5, p238:7-9: For implementations that do not support dynamic adjustment of the number of threads this routine has no effect: the value of dyn-var remains false. Add a test that cancellation and nested parallelism aren't supported either. Differential Revision: https://reviews.llvm.org/D51785 llvm-svn: 343381
* [libomptarget-nvptx] Fix number of threads in parallelJonas Hahnfeld2018-09-292-84/+45
| | | | | | | | | | | | If there is no num_threads() clause we must consider the nthreads-var ICV. Its value is set by omp_set_num_threads() and can be queried using omp_get_max_num_threads(). The rewritten code now closely resembles the algorithm given in the OpenMP standard. Differential Revision: https://reviews.llvm.org/D51783 llvm-svn: 343380
* [OpenMP][libomptarget] Set the frame pointer then test empty slot conditionGheorghe-Teodor Bercea2018-09-251-3/+3
| | | | | | | | | | | | | | Summary: NFC - just fixing a bug: the empty slot test was before the re-setting of the Stack pointer. Reviewers: ABataev, caomhin, Hahnfeld Reviewed By: ABataev Subscribers: guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D52122 llvm-svn: 343006
* [OpenMP][libomptarget] Simplify warp master selection for data sharingGheorghe-Teodor Bercea2018-09-251-2/+2
| | | | | | | | | | | | | | | | | | Summary: There is currently no supported situation where the warp master is not the first thread in the warp. This also avoids the device execution from hanging on Volta GPUs when ballot_sync is called by a number of threads that is less that the size of a warp. Reviewers: ABataev, caomhin, grokos Reviewed By: grokos Subscribers: guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D50188 llvm-svn: 342972
* [OPENMP][NVPTX] Add support for lastprivates/reductions handling in SPMD ↵Alexey Bataev2018-09-216-1/+74
| | | | | | | | | | | | | | | | | | | constructs with lightweight runtime. Summary: We need the support for per-team shared variables to support codegen for lastprivates/reductions. Patch adds this support by using shared memory if the total size of the reductions/lastprivates is <= 128 bytes, then pre-allocated buffer in global memory if size is <= 4K bytes,or uses malloc/free, otherwise. Reviewers: gtbercea, kkwli0, grokos Subscribers: guansong, openmp-commits Differential Revision: https://reviews.llvm.org/D51875 llvm-svn: 342737
* [libomptarget-nvptx] Remove last mentions of __kmpc_print_*Jonas Hahnfeld2018-09-081-12/+0
| | | | | | | Their implementation was removed during review, delete their prototype declarations. llvm-svn: 341748
OpenPOWER on IntegriCloud