aboutsummaryrefslogtreecommitdiff
Commit message (Collapse)AuthorAgeFilesLines
* [OpenMP] Pass mapping names to add components in a user defined mapperJoseph Huber2021-04-011-6/+8
| | | | | | | | | | | | | Summary: Currently the mapping names are not passed to the mapper components that set up the array region. This means array mappings will not have their names availible in the runtime. This patch fixes this by passing the argument name to the region correctly. This means that the mapped variable's name will be the declared mapper that placed it on the device. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D99681
* [OPENMP]Fix PR48740: OpenMP declare reduction in C does not require an ↵Alexey Bataev2021-03-301-3/+7
| | | | | | | | | | | | | initializer If no initializer-clause is specified, the private variables will be initialized following the rules for initialization of objects with static storage duration. Need to adjust the implementation to the current version of the standard. Differential Revision: https://reviews.llvm.org/D99539
* [OPENMP]Map data field with l-value reference types.Alexey Bataev2021-03-291-26/+80
| | | | | | | Added initial support dfor the mapping of the data members with l-value reference types. Differential Revision: https://reviews.llvm.org/D98812
* [OpenMP][InstrProfiling] Fix a missing instr profiling counterXun Li2021-03-251-1/+3
| | | | | | | | When emitting a function body there needs to be a instr profiling counter emitted. Otherwise instr profiling won't work for this function. Reviewed By: MaskRay Differential Revision: https://reviews.llvm.org/D98135
* [OPENMP]Fix PR48571: critical/master in outlined contexts cause crash.Alexey Bataev2021-03-241-12/+20
| | | | | | | | If emit inlined region for master/critical directives, no need to clear lambda/block context data, otherwise the variables cannot be found and it causes a crash at compile time. Differential Revision: https://reviews.llvm.org/D99280
* [OpenMP] Restore backwards compatibility for libomptargetJoseph Huber2021-03-111-1/+1
| | | | | | | | | | | | | Summary: The changes introduced in D87946 changed the API for libomptarget functions. `__kmpc_push_target_tripcount` was a function in Clang 11.x but was not given a backward-compatible interface. This change will require people using Clang 13.x or 12.x to recompile their offloading programs. Reviewed By: jdoerfert cchen Differential Revision: https://reviews.llvm.org/D98358
* [OPENMP]Fix PR48759: "fatal error" when compile with preprocessed file.Alexey Bataev2021-03-041-3/+7
| | | | | | | If the file in line directive does not exist on the system we need, to use the original file to get its file id. Differential Revision: https://reviews.llvm.org/D97945
* [OPENMP50]Mapping of the subcomponents with the 'default' mappers.Alexey Bataev2021-03-021-22/+47
| | | | | | | If the mapped structure has data members, which have 'default' mappers, need to map these members individually using their 'default' mappers. Differential Revision: https://reviews.llvm.org/D92195
* [OPENMP50]Allow overlapping mapping in target constructs.Alexey Bataev2021-02-161-295/+322
| | | | | | | | | | OpenMP 5.0 removed a lot of restriction for overlapped mapped items comparing to OpenMP 4.5. Patch restricts the checks for overlapped data mappings only for OpenMP 4.5 and less and reorders mapping of the arguments so, that present and alloc mappings are processed first and then all others. Differential Revision: https://reviews.llvm.org/D86119
* [OpenMP] Implement '#pragma omp tile', by Michael Kruse (@Meinersbur).Michael Kruse2021-02-161-0/+5
| | | | | | | | | | | | The tile directive is in OpenMP's Technical Report 8 and foreseeably will be part of the upcoming OpenMP 5.1 standard. This implementation is based on an AST transformation providing a de-sugared loop nest. This makes it simple to forward the de-sugared transformation to loop associated directives taking the tiled loops. In contrast to other loop associated directives, the OMPTileDirective does not use CapturedStmts. Letting loop associated directives consume loops from different capture context would be difficult. A significant amount of code generation logic is taking place in the Sema class. Eventually, I would prefer if these would move into the CodeGen component such that we could make use of the OpenMPIRBuilder, together with flang. Only expressions converting between the language's iteration variable and the logical iteration space need to take place in the semantic analyzer: Getting the of iterations (e.g. the overload resolution of `std::distance`) and converting the logical iteration number to the iteration variable (e.g. overload resolution of `iteration + .omp.iv`). In clang, only CXXForRangeStmt is also represented by its de-sugared components. However, OpenMP loop are not defined as syntatic sugar. Starting with an AST-based approach allows us to gradually move generated AST statements into CodeGen, instead all at once. I would also like to refactor `checkOpenMPLoop` into its functionalities in a follow-up. In this patch it is used twice. Once for checking proper nesting and emitting diagnostics, and additionally for deriving the logical iteration space per-loop (instead of for the loop nest). Differential Revision: https://reviews.llvm.org/D76342
* [Branch-Rename] Fix some linksxgupta2021-02-011-3/+3
| | | | | | | | According to the [[ https://foundation.llvm.org/docs/branch-rename/ | status of branch rename ]], the master branch of the LLVM repository is removed on 28 Jan 2021. Reviewed By: mehdi_amini Differential Revision: https://reviews.llvm.org/D95766
* [OpenMP][FIX] Enforce a function boundary for a new data environmentJohannes Doerfert2021-01-251-0/+8
| | | | | | | | | | | Whenever we enter a new OpenMP data environment we want to enter a function to simplify reasoning. Later we probably want to remove the entire specialization wrt. the if clause and pass the result to the runtime, for now this should fix PR48686. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D94315
* [OpenMP] Add support for mapping names in mapper APIJoseph Huber2021-01-211-6/+18
| | | | | | | | | Summary: The custom mapper API did not previously support the mapping names added previously. This means they were not present if a user requested debugging information while using the mapper functions. This adds basic support for passing the mapped names to the runtime library. Reviewers: jdoerfert Differential Revision: https://reviews.llvm.org/D94806
* [OPENMP]Do not use OMP_MAP_TARGET_PARAM for data movement directives.Alexey Bataev2021-01-191-20/+10
| | | | | | | | | | | | | OMP_MAP_TARGET_PARAM flag is used to mark the data that shoud be passed as arguments to the target kernels, nothing else. But the compiler still marks the data with OMP_MAP_TARGET_PARAM flags even if the data is passed to the data movement directives, like target data, target update etc. This flag is just ignored for this directives and the compiler does not need to emit it. Reviewed By: cchen Differential Revision: https://reviews.llvm.org/D91261
* [Clang][OpenMP] Fixed an issue that clang crashed when compiling OpenMP ↵Shilei Tian2021-01-191-10/+7
| | | | | | | | | | | | | | | | | | | | program in device only mode without host IR D94745 rewrites the `deviceRTLs` using OpenMP and compiles it by directly calling the device compilation. `clang` crashes because entry in `OffloadEntriesDeviceGlobalVar` is unintialized. Current design supposes the device compilation can only be invoked after host compilation with the host IR such that `clang` can initialize `OffloadEntriesDeviceGlobalVar` from host IR. This avoids us using device compilation directly, especially when we only have code wrapped into `declare target` which are all device code. The same issue also exists for `OffloadEntriesInfoManager`. In this patch, we simply initialized an entry if it is not in the maps. Not sure we need an option to tell the device compiler that it is invoked standalone. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D94871
* [OpenMP][AMDGPU] Use AMDGPU_KERNEL calling convention for entry functionPushpinder Singh2021-01-061-0/+2
| | | | | | | | | AMDGPU backend requires entry functions/kernels to have AMDGPU_KERNEL calling convention for proper linking. Reviewed By: JonChesterfield Differential Revision: https://reviews.llvm.org/D94060
* [OPENMP]Use the real pointer value as base, not indexed value.Alexey Bataev2020-11-201-2/+5
| | | | | | | | | | After fix for PR48174 the base pointer for pointer-based array-sections/array-subscripts will be emitted as `&ptr[idx]`, but actually it should be just `ptr`, i.e. the address stored in the ponter to point correctly to the beginning of the array. Currently it may lead to a crash in the runtime. Differential Revision: https://reviews.llvm.org/D91805
* [AMDGPU] Set the default globals address space to 1Alex Richardson2020-11-201-5/+6
| | | | | | | | | | | | This will ensure that passes that add new global variables will create them in address space 1 once the passes have been updated to no longer default to the implicit address space zero. This also changes AutoUpgrade.cpp to add -G1 to the DataLayout if it wasn't already to present to ensure bitcode backwards compatibility. Reviewed by: arsenm Differential Revision: https://reviews.llvm.org/D84345
* [OpenMP] Add Location Fields to Libomptarget Runtime for DebuggingJoseph Huber2020-11-191-14/+42
| | | | | | | | | Summary: Add support for passing source locations to libomptarget runtime functions using the ident_t struct present in the rest of the libomp API. This will allow the runtime system to give much more insightful error messages and debugging values. Reviewers: jdoerfert grokos Differential Revision: https://reviews.llvm.org/D87946
* [OpenMP] Add Passing in Original Declaration Names To Mapper APIJoseph Huber2020-11-181-53/+217
| | | | | | | | | Summary: This patch adds support for passing in the original delcaration name in the source file to the libomptarget runtime. This will allow the runtime to provide more intelligent debugging messages. This patch takes the original expression parsed from the OpenMP map / update clause and provides a textual representation if it was explicitly mapped, otherwise it takes the name of the variable declaration as a fallback. The information in passed to the runtime in a global array of strings that matches the existing ident_t source location strings using ";name;filename;column;row;;" Reviewers: jdoerfert Differential Revision: https://reviews.llvm.org/D89802
* [OPENMP]Fix PR48174: compile-time crash with target enter data on a global ↵Alexey Bataev2020-11-181-0/+4
| | | | | | | | | | struct. The compiler should treat array subscript with base pointer as a first pointer in complex data, it is used only for member expression with base pointer. Differential Revision: https://reviews.llvm.org/D91660
* [OPENMP]Fix PR48076: mapping of data member pointer.Alexey Bataev2020-11-171-3/+13
| | | | | | | | | | | If the data member pointer is mapped, the compiler tries to optimize the mapping of such data by discarding explicit mapping flags and trying to emit combined data instead. In some cases, this optimization is not quite correctly implemented and it leads to a program crash at the runtime. Instead, if the data member is mapped, just emit it as is and do not emit combined mapping flags for it. Differential Revision: https://reviews.llvm.org/D91552
* [OPENMP]Fix PR48076: Check map types array before accessing its front.Alexey Bataev2020-11-121-2/+3
| | | | | | | | Need to check if there are map types for the components before trying to access them when trying to modify type mappings for combined partial mappings. Differential Revision: https://reviews.llvm.org/D91370
* [OMPIRBuilder] Start 'Create' methods with lower case. NFC.Michael Kruse2020-11-091-4/+4
| | | | | | | | | | | | For consistency with the IRBuilder, OpenMPIRBuilder has method names starting with 'Create'. However, the LLVM coding style has methods names starting with lower case letters, as all other OpenMPIRBuilder already methods do. The clang-tidy configuration used by Phabricator also warns about the naming violation, adding noise to the reviews. This patch renames all `OpenMPIRBuilder::CreateXYZ` methods to `OpenMPIRBuilder::createXYZ`, and updates all in-tree callers. I tested check-llvm, check-clang, check-mlir and check-flang to ensure that I did not miss a caller. Reviewed By: mehdi_amini, fghanim Differential Revision: https://reviews.llvm.org/D91109
* [OpenMP] Fix -Wmisleading-indentation after D84192Fangrui Song2020-11-061-1/+2
|
* [OpenMP5.0] map item can be non-contiguous for target updatecchen2020-11-061-15/+318
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | In order not to modify the `tgt_target_data_update` information but still be able to pass the extra information for non-contiguous map item (offset, count, and stride for each dimension), this patch overload `arg` when the maptype is set as `OMP_MAP_DESCRIPTOR`. The origin `arg` is for passing the pointer information, however, the overloaded `arg` is an array of descriptor_dim: struct descriptor_dim { int64_t offset; int64_t count; int64_t stride }; and the array size is the same as dimension size. In addition, since we have count and stride information in descriptor_dim, we can replace/overload the `arg_size` parameter by using dimension size. For supporting `stride` in array section, we use a dummy dimension in descriptor to store the unit size. The formula for counting the stride in dimension D_n: `unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride`. Demonstrate how it works: ``` double arr[3][4][5]; D0: { offset = 0, count = 1, stride = 8 } // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 } // stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8 D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40 } // stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20 D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 } // stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200 // X here means we need to offload this data, therefore, runtime will transfer // data from offset 80, 96, 120, 136, 400, 416, 440, 456 // Runtime patch: https://reviews.llvm.org/D82245 // OOOOO OOOOO OOOOO // OOOOO OOOOO OOOOO // XOXOO OOOOO XOXOO // XOXOO OOOOO XOXOO ``` Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D84192
* [OpenMP] target nested `use_device_ptr() if()` and is_device_ptr trigger assertscchen2020-11-041-3/+11
| | | | | | | | | | | | | | | | | | | | | | | | | | Clang now asserts for the below case: ``` void clang::CodeGen::CGOpenMPRuntime::createOffloadEntriesAndInfoMetadata(): Assertion `std::get<0>(E) && "All ordered entries must exist!"' failed. ``` The reason why Clang hit the assert is because in `emitTargetDataCalls`, both `BeginThenGen` and `BeginElseGen` call `registerTargetRegionEntryInfo` and try to register the Entry in OffloadEntriesTargetRegion with same key. If changing the expression in if clause to any constant expression, then the assert disappear. (https://godbolt.org/z/TW7haj) The assert itself is to avoid user from accessing elements out of bound inside `OrderedEntries` in `createOffloadEntriesAndInfoMetadata`. In this patch, I add a check in `registerTargetRegionEntryInfo` to avoid register the target region more than once. A test case that triggers assert: https://godbolt.org/z/4cnGW8 Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D90704
* [Clang][OpenMP] Added the support for target data nowaitShilei Tian2020-10-281-2/+3
| | | | | | | | | | Previously we added support for target nowait, but target data nowait has not been supported yet. In this patch, target data nowait will also be wrapped into a task. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D90099
* [openmp] Use front() instead of *begin() to not hide bugs when CurTypes is ↵Benjamin Kramer2020-10-281-1/+1
| | | | empty.
* Revert "[OpenMP] Add Passing in Original Declaration Names To Mapper API"Benjamin Kramer2020-10-281-215/+52
| | | | | This reverts commit d981c7b7581efc3ef378709042100e75da0185a0 and a87d7b3d448a16e416d1980b9d6aea99e4c9900b. Test fails under msan.
* [OpenMP] Add Passing in Original Declaration Names To Mapper APIJoseph Huber2020-10-271-52/+215
| | | | | | | | | | | | | | | | | | Summary: This patch adds support for passing in the original delcaration name in the source file to the libomptarget runtime. This will allow the runtime to provide more intelligent debugging messages. This patch takes the original expression parsed from the OpenMP map / update clause and provides a textual representation if it was explicitly mapped, otherwise it takes the name of the variable declaration as a fallback. The information in passed to the runtime in a global array of strings that matches the existing ident_t source location strings using ";name;filename;column;row;;". See clang/test/OpenMP/target_map_names.cpp for an example of the generated output for a given map clause. Reviewers: jdoervert Differential Revision: https://reviews.llvm.org/D89802
* [Clang][OpenMP] Avoid unnecessary privatization of mapper array when there ↵Shilei Tian2020-10-271-16/+11
| | | | | | | | | | | | is no user defined mapper In current implementation, if it requires an outer task, the mapper array will be privatized no matter whether it has mapper. In fact, when there is no mapper, the mapper array only contains number of nullptr. In the libomptarget, the use of mapper array is `if (mappers_array && mappers_array[i])`, which means we can directly set mapper array to nullptr if there is no mapper. This can avoid unnecessary data copy. In this patch, the data privatization will not be emitted if the mapper array is nullptr. When it comes to the emit of task body, the nullptr will be used directly. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D90101
* [Clang][OpenMP] Fixed an issue of segment fault when using target nowaitShilei Tian2020-10-261-2/+1
| | | | | | | | | | | | The implementation of target nowait just wraps the target region into a task. The essential four parameters (base ptr, ptr, size, mapper) are taken as firstprivate such that they will be copied to the private location. When there is no user-defined mapper, the mapper variable will be nullptr. However, it will be still copied to the corresponding place. Therefore, a memcpy will be generated and the source pointer will be nullptr, causing a segmentation fault. The root cause is when calling `emitOffloadingArraysArgument`, the last argument `Options` has a field about whether it requires a task. It only takes depend clause into account. In this patch, the nowait clause is also included. There're two things that will be done in another patches: 1. target data nowait has not been supported yet. D90099 added the support. 2. When there is no mapper, the mapper array can be nullptr no matter whether it requires outer task or not. It can avoid an unnecessary data copy. This is an optimization that is covered in D90101. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D89844
* [Clang][OpenMP] Added support for nowait target in CodeGen via regular taskShilei Tian2020-09-251-1/+2
| | | | | | | | | | | | | | | | Previously for nowait target, CG emitted a function call to `__tgt_target_nowait`, etc. However, in OpenMP RTL, these functions just directly call the no-nowait version, which means nowait is not working as expected. OpenMP specification says a target is acutally a target task, which is an untied and detachable task. It is natural to go to the direction that generates a task for a nowait target. However, OpenMP task has a problem that it must be within to a parallel region; otherwise the task will be executed immediately. As a result, if we directly wrap to a regular task, the `target nowait` outside of a parallel region is still a synchronous version. In D77609, I added the support for unshackled task in OpenMP RTL. Basically, unshackled task is a task that is not bound to any parallel region. So all nowait target will be tranformed into an unshackled task. In order to distinguish from regular task, a new flag bit is set for unshackled task. This flag will be used by RTL for later process. Since all target tasks are allocated via `__kmpc_omp_target_task_alloc`, and in current `libomptarget`, `__kmpc_omp_target_task_alloc` just calls `__kmpc_omp_task_alloc`. Therefore, we can modify the flag in `__kmpc_omp_target_task_alloc` so that we don't need to modify the FE too much. If users choose to opt out the feature, they just need to use a RTL w/o support of unshackled threads. As a result, in this patch, the `target nowait` region is simply wrapped into a regular task. Later once we have RTL support for unshackled tasks, the wrapped tasks can be executed by unshackled threads w/o changes in the FE. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D78075
* [OpenMP 5.0] Fix user-defined mapper privatization in tasksAlexey Bataev2020-09-171-17/+37
| | | | | | This patch fixes the problem that user-defined mapper array is not correctly privatized inside a task. This problem causes openmp/libomptarget/test/offloading/target_depend_nowait.cpp fails. Differential Revision: https://reviews.llvm.org/D84470
* [OPENMP]Fix codegen for is_device_ptr component, captured by reference.Alexey Bataev2020-09-151-4/+6
| | | | | | | | | Need to map the component as TO instead of the literal, because need to pass a reference to a component if the pointer is overaligned. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D84887
* [OPENMP]Add support for allocate vars in untied tasks.Alexey Bataev2020-09-151-61/+107
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Local vars, marked with pragma allocate, mustbe allocate by the call of the runtime function and cannot be allocated as other local variables. Instead, we allocate a space for the pointer in private record and store the address, returned by kmpc_alloc call in this pointer. So, for untied tasks ``` #pragma omp task untied { S s; #pragma omp allocate(s) allocator(allocator) s = x; } ``` compiler generates something like this: ``` struct task_with_privates { S *ptr; }; void entry(task_with_privates *p) { S *s = p->s; switch(partid) { case 1: p->s = (S*)kmpc_alloc(); kmpc_omp_task(); br exit; case 2: *s = x; kmpc_omp_task(); br exit; case 2: ~S(s); kmpc_free((void*)s); br exit; } exit: } ``` Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D86558
* Assert we've found the size of each (non-overlapping) structure. NFCI.Simon Pilgrim2020-09-141-0/+1
| | | | Fixes clang static analyzer warning.
* [OPENMP] Fix PR47063: crash when trying to get captured statetment.Alexey Bataev2020-08-121-2/+1
| | | | | | Need to call getRawStmt() function instead, when trying to get inner associated statement for the executable directive. Not all directives use captured statements.
* [OPENMP]Fix PR37671: Privatize local(private) variables in untied tasks.Alexey Bataev2020-08-121-72/+120
| | | | | | | | | | | | | | In untied tasks, need to allocate the space for local variales, declared in task region, when the memory for task data is allocated. THe function can be interrupted and we can exit from the function in untied task switch. Need to keep the state of the local variables in this case. Also, the compiler should not call cleanup when exiting in untied task switch until the real exit out of the declaration scope is met during execution. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D84457
* [OPENMP]Do not add TGT_OMP_TARGET_PARAM flag to non-captured mapped arguments.Alexey Bataev2020-08-121-11/+19
| | | | | | | | | | If the arguments are mapped, but are actually not used in the target region, the compiler still adds attribute TGT_OMP_TARGET_PARAM for such arguments. It makes the libomptarget to add such parameters to the list of arguments, passed to the kernel at the runtime, and may lead to incorrect results/crashes during execution. Differential Revision: https://reviews.llvm.org/D85755
* Revert "[OPENMP]Fix PR37671: Privatize local(private) variables in untied ↵Alexey Bataev2020-08-121-111/+72
| | | | | | | tasks." This reverts commit ec9563c54ed25e9f9cbe60985399212d50bd801d to investigate compiler crash revelaed by the buildbots.
* [OPENMP]Fix PR37671: Privatize local(private) variables in untied tasks.Alexey Bataev2020-08-121-72/+111
| | | | | | | | | | | | | | | | | | | Summary: In untied tasks, need to allocate the space for local variales, declared in task region, when the memory for task data is allocated. THe function can be interrupted and we can exit from the function in untied task switch. Need to keep the state of the local variables in this case. Also, the compiler should not call cleanup when exiting in untied task switch until the real exit out of the declaration scope is met during execution. Reviewers: jdoerfert Subscribers: yaxunl, guansong, cfe-commits, sstefan1, caomhin Tags: #clang Differential Revision: https://reviews.llvm.org/D84457
* [OpenMP][NFC] Reuse OMPIRBuilder `struct ident_t` handling in ClangJohannes Doerfert2020-08-101-110/+17
| | | | | | | | | | | | | | | | | | Replace the `ident_t` handling in Clang with the methods offered by the OMPIRBuilder. This cuts down on the clang code as well as the differences between the two, making further transitions easier. Tests have changed but there should not be a real functional change. The most interesting difference is probably that we stop generating local ident_t allocations for now and just use globals. Given that this happens only with debug info, the location part of the `ident_t` is probably bigger than the test anyway. As the location part is already a global, we can avoid the allocation, memcpy, and store in favor of a constant global that is slightly bigger. This can be revisited if there are complications. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D80735
* [OpenMP] Fix `present` for exit from `omp target data`Joel E. Denny2020-08-051-5/+34
| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | Without this patch, the following example fails but shouldn't according to OpenMP TR8: ``` #pragma omp target enter data map(alloc:i) #pragma omp target data map(present, alloc: i) { #pragma omp target exit data map(delete:i) } // fails presence check here ``` OpenMP TR8 sec. 2.22.7.1 "map Clause", p. 321, L23-26 states: > If the map clause appears on a target, target data, target enter > data or target exit data construct with a present map-type-modifier > then on entry to the region if the corresponding list item does not > appear in the device data environment an error occurs and the > program terminates. There is no corresponding statement about the exit from a region. Thus, the `present` modifier should: 1. Check for presence upon entry into any region, including a `target exit data` region. This behavior is already implemented correctly. 2. Should not check for presence upon exit from any region, including a `target` or `target data` region. Without this patch, this behavior is not implemented correctly, breaking the above example. In the case of `target data`, this patch fixes the latter behavior by removing the `present` modifier from the map types Clang generates for the runtime call at the end of the region. In the case of `target`, we have not found a valid OpenMP program for which such a fix would matter. It appears that, if a program can guarantee that data is present at the beginning of a `target` region so that there's no error there, that data is also guaranteed to be present at the end. This patch adds a comment to the runtime to document this case. Reviewed By: grokos, RaviNarayanaswamy, ABataev Differential Revision: https://reviews.llvm.org/D84422
* [OpenMP][FIX] Consistently use OpenMPIRBuilder if requestedJohannes Doerfert2020-07-301-9/+36
| | | | | | | | | | | | | | | | | | | When we use the OpenMPIRBuilder for the parallel region we need to also use it to get the thread ID (among other things) in the body. This is because CGOpenMPRuntime::getThreadID() and CGOpenMPRuntime::emitUpdateLocation implicitly assumes that if they are called from within a parallel region there is a certain structure to the code and certain members of the OMPRegionInfo are initialized. It might make sense to initialize them even if we use the OpenMPIRBuilder but we would preferably get rid of such state instead. Bug reported by Anchu Rajendran Sudhakumari. Depends on D82470. Reviewed By: anchu-rajendran Differential Revision: https://reviews.llvm.org/D82822
* [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in ↵Alexey Bataev2020-07-301-27/+54
| | | | | | | | | | | | | | | | | target region. Need to map the base pointer for all directives, not only target data-based ones. The base pointer is mapped for array sections, array subscript, array shaping and other array-like constructs with the base pointer. Also, codegen for use_device_ptr clause was modified to correctly handle mapping combination of array like constructs + use_device_ptr clause. The data for use_device_ptr clause is emitted as the last records in the data mapping array. Reviewed By: ye-luo Differential Revision: https://reviews.llvm.org/D84767
* Revert "[OPENMP]Fix PR46824: Global declare target pointer cannot be ↵Alexey Bataev2020-07-301-52/+26
| | | | | | | accessed in target region." This reverts commit 142d0d3ed8e07aca2476bc4ecc1a12d15577a84a to investigate undefined behavior revealed by buildbots.
* [OPENMP]Fix PR46824: Global declare target pointer cannot be accessed in ↵Alexey Bataev2020-07-301-26/+52
| | | | | | | | | | | | | | | | target region. Need to map the base pointer for all directives, not only target data-based ones. The base pointer is mapped for array sections, array subscript, array shaping and other array-like constructs with the base pointer. Also, codegen for use_device_ptr clause was modified to correctly handle mapping combination of array like constructs + use_device_ptr clause. The data for use_device_ptr clause is emitted as the last records in the data mapping array. It applies only for global pointers. Differential Revision: https://reviews.llvm.org/D84767
* [OpenMP] Implement TR8 `present` motion modifier in Clang (1/2)Joel E. Denny2020-07-291-35/+50
| | | | | | | | | | This patch implements Clang front end support for the OpenMP TR8 `present` motion modifier for `omp target update` directives. The next patch in this series implements OpenMP runtime support. Reviewed By: ABataev Differential Revision: https://reviews.llvm.org/D84711