| Age | Commit message (Collapse) | Author | Files | Lines |
|---|
|
Recursive types can cause re-entrant mapper emission. The mapper
function is created by OpenMPIRBuilder before the callbacks run, so it
may already exist in the LLVM module even though it is not yet
registered in the ModuleTranslation mapping table. Reuse and register it
to break the recursion. Added offloading test.
|
|
Fix OpenMP mapper lowering by attaching user-defined/default mappers
only to the base parent entry, not combined/segment entries. This
prevents mapper calls with partial sizes. Added relevant tests.
|
|
implicit mappers (#177389)
Make implicit default mapper generation respect defaultmap categories so
unrelated defaultmap clauses no longer suppress mappers for derived
types.
Added related tests.
|
|
The compiler skips mapping of named constants (parameters) to OpenMP
target regions under the assumption that constants don't need to be
mapped. This assumption is not valid when array is accessed inside with
dynamic index. The problem can be seen with the following code:
```
module fir_lowering_check
implicit none
integer, parameter :: dp = selected_real_kind(15, 307)
real(dp), parameter :: arrays(2) = (/ 0.0, 0.0 /)
contains
subroutine test(hold)
integer, intent(in) :: hold
integer :: z
real(dp) :: temp
!$omp target teams distribute parallel do
do z = 1, 2
temp = arrays(hold)
end do
!$omp end target teams distribute parallel do
end subroutine test
end module fir_lowering_check
program main
use fir_lowering_check
implicit none
integer :: hold
hold = 1
call test(hold)
print *, "Finished"
end program main
```
It fails with the following error
`'hlfir.designate' op using value defined outside the region`
The fix is to allow mapping of constant arrays and map them as `to`.
|
|
The new tests added in https://github.com/llvm/llvm-project/pull/169623
are [failing](https://lab.llvm.org/buildbot/#/builders/225/builds/544)
on the Intel GPU runner, so XFAIL them as with the other tests.
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
|
|
PR #144635 enabled non-contiguous updates for both `update from` and
`update to` clauses, but tests for `update to` were missing. This PR
adds those missing tests to ensure coverage.
|
|
This test also depends on the line number. Following similar approach as
other with [[@LINE]] macro.
|
|
Summary:
This is only really meaningful for the NVPTX target. Not all build
environments support host LTO and these are redundant tests, just clean
this up and make it run faster.
|
|
being marked as parial maps (#175133)
The following test was triggering a runtime crash **on the host before
launching the kernel**:
```fortran
program test_omp_target_map_bug_v5
implicit none
type nested_type
real, allocatable :: alloc_field(:)
end type nested_type
type nesting_type
integer :: int_field
type(nested_type) :: derived_field
end type nesting_type
type(nesting_type) :: config
allocate(config%derived_field%alloc_field(1))
!$OMP TARGET ENTER DATA MAP(TO:config, config%derived_field%alloc_field)
!$OMP TARGET
config%derived_field%alloc_field(1) = 1.0
!$OMP END TARGET
deallocate(config%derived_field%alloc_field)
end program test_omp_target_map_bug_v5
```
In particular, the runtime was producing a segmentation fault when the
test is compiled with any optimization level > 0; if you compile with
-O0 the sample ran fine.
After debugging the runtime, it turned out the crash was happening at
the point where the runtime calls the default mapper emitted by the
compiler for `nesting_type; in particular at this point in the runtime:
https://github.com/llvm/llvm-project/blob/c62cd2877cc25a0d708ad22a70c2a57590449c4d/offload/libomptarget/omptarget.cpp#L307.
Bisecting the optimization pipeline using `-mllvm -opt-bisect-limit=N`,
the first pass that triggered the issue on `O1` was the `instcombine`
pass. Debugging this further, the issue narrows down to canonicalizing
`getelementptr` instructions from using struct types (in this case the
`nesting_type` in the sample above) to using addressing bytes (`i8`). In
particular, in `O0`, you would see something like this:
```llvm
define internal void @.omp_mapper._QQFnesting_type_omp_default_mapper(ptr noundef %0, ptr noundef %1, ptr noundef %2, i64 noundef %3, i64 noundef %4, ptr noundef %5) #6 {
entry:
%6 = udiv exact i64 %3, 56
%7 = getelementptr %_QFTnesting_type, ptr %2, i64 %6
....
}
```
```llvm
define internal void @.omp_mapper._QQFnesting_type_omp_default_mapper(ptr noundef %0, ptr noundef %1, ptr noundef %2, i64 noundef %3, i64 noundef %4, ptr noundef %5) #6 {
entry:
%6 = getelementptr i8, ptr %2, i64 %3
....
}
```
The `udiv exact` instruction emitted by the OMP IR Builder (see:
https://github.com/llvm/llvm-project/blob/c62cd2877cc25a0d708ad22a70c2a57590449c4d/llvm/lib/Frontend/OpenMP/OMPIRBuilder.cpp#L9154)
allows `instcombine` to assume that `%3` is divisible by the struct size
(here `56`) and, therefore, replaces the result of the division with
direct GEP on `i8` rather than the struct type.
However, the runtime was calling
`@.omp_mapper._QQFnesting_type_omp_default_mapper` not with `56` (the
proper struct size) but with `48`!
Debugging this further, I found that the size of `omp.map.info`
operation to which the default mapper is attached computes the value of
`48` because we set the map to partial (see:
https://github.com/llvm/llvm-project/blob/c62cd2877cc25a0d708ad22a70c2a57590449c4d/flang/lib/Optimizer/OpenMP/MapInfoFinalization.cpp#L1146
and
https://github.com/llvm/llvm-project/blob/c62cd2877cc25a0d708ad22a70c2a57590449c4d/mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp#L4501-L4512).
However, I think this is incorrect since the emitted mapper (and
user-defined mappers in general) are defined on the whole struct type
and should never be marked as partial. Hence, the fix in this PR.
|
|
The changes in line numbers caused a few CHECK macros to now fail. This
is fixed by this PR.
Build w/ breakages:
https://lab.llvm.org/staging/#/builders/105/builds/39748
|
|
We finally got our buildbot added (to staging, at least) so we want to
start running L0 tests in CI.
We need `check-offload` to pass though, so XFAIL everything failing.
There's a couple `UNSUPPORTED` as well, those are for sporadic fails.
Also make set the `gpu` and `intelgpu` LIT variables when testing the
`spirv64-intel` triple.
We have no DeviceRTL yet so basically everything fails, but we manage to
get
```
Total Discovered Tests: 432
Unsupported : 169 (39.12%)
Passed : 67 (15.51%)
Expectedly Failed: 196 (45.37%)
```
We still don't build the level zero plugin by default and these tests
don't run unless the plugin was built, so this has no effect on most
builds.
---------
Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
|
|
Typo exposed by recent `not` behavior change, we need to make sure we're
using the LLVM one.
Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
|
|
`not` behavior change in
https://github.com/llvm/llvm-project/pull/174298 requires `--crash`
passed now.
Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
|
|
While the main problem with the test is that it requires LLD, given that
it is unlikely to be testing anything meaningful for a CPU-only build,
just mark it as requiring GPU.
Fixes #100780
Signed-off-by: Michał Górny <mgorny@gentoo.org>
|
|
array sections (#157443)
### 1. ElementType deduction for pointer-based array sections
Problem: Pointer-based array sections were previously ignored during
`ElementType` deduction, leading to incorrect assumptions about array
item types.
This often resulted in out-of-bounds access, as seen in the assertion
failure:
```
Assertion `idx < size()' failed.
llvm-project/llvm/include/llvm/ADT/SmallVector.h:292:
reference llvm::SmallVectorTemplateCommon<llvm::Value *>::operatorsize_type
[T = llvm::Value *]
```
Fix: Added a check in clang/lib/CodeGen/CGOpenMPRuntime.cpp to ensure
`ElementType` is correctly detected for cases involving non-contiguous
updates with a base pointer.
Impact: Resolves failures in OpenMP_VV (formerly sollve_vv) and other
offload/clang-OpenMP tests:
All tests under:
https://github.com/OpenMP-Validation-and-Verification/OpenMP_VV/tree/master/tests/5.0/target_update
test_target_update_mapper_from_discontiguous.c
test_target_update_mapper_to_discontiguous.c
test_target_update_to_discontiguous.c
test_target_update_from_discontiguous.c
### 2. Zero-dimension propagation in struct member mappings
Problem: A zero-dimension entry for struct members introduced
inconsistencies in complex mapping logic within OMPIRBuilder.cpp.
Placeholder zeros propagated to emitNonContiguousDescriptor(), breaking
reverse indexing logic and corrupting IR:
Loops assume `Dims[I] >= 1`. When `Dims[I] == 0`:
Reverse indexing still stores pointers to uninitialized allocas or
mismatched slots. Runtime interprets `ArgSizes[I]` (derived from
`Dims[I])` as dimensionality, causing size/offset calculations to
collapse to zero → results in `size=0` async copy and plugin interface
errors.
Fix: Prepend a synthetic dimension of size 1 instead of appending a
zero, preserving correctness in `targetDataUpdate()` for non-contiguous
updates.
Impact: Added dedicated test cases that previously failed on main.
|
|
|
|
(#153683)
This adds support for using `ATTACH` map-type for proper
pointer-attachment when mapping list-items that have base-pointers.
For example, for the following:
```c
int *p;
#pragma omp target enter data map(p[1:10])
```
The following maps are now emitted by clang:
```
(A)
&p[0], &p[1], 10 * sizeof(p[1]), TO | FROM
&p, &p[1], sizeof(p), ATTACH
```
Previously, the two possible maps emitted by clang were:
```
(B)
&p[0], &p[1], 10 * sizeof(p[1]), TO | FROM
(C)
&p, &p[1], 10 * sizeof(p[1]), TO | FROM | PTR_AND_OBJ
````
(B) does not perform any pointer attachment, while (C) also maps the
pointer p, both of which are incorrect.
-----
With this change, we are using ATTACH-style maps, like `(A)`, for cases
where the expression has a base-pointer. For example:
```cpp
int *p, **pp;
S *ps, **pps;
... map(p[0])
... map(p[10:20])
... map(*p)
... map(([20])p)
... map(ps->a)
... map(pps->p->a)
... map(pp[0][0])
... map(*(pp + 10)[0])
```
#### Grouping of maps based on attach base-pointers
We also group mapping of clauses with the same base decl in the order of
the increasing complexity of their base-pointers, e.g. for something
like:
```
S **spp;
map(spp[0][0], spp[0][0].a), // attach-ptr: spp[0]
map(spp[0]), // attach-ptr: spp
map(spp), // attach-ptr: N/A
```
We first map `spp`, then `spp[0]` then `spp[0][0]` and `spp[0][0].a`.
This allows us to also group "struct" allocation based on their attach
pointers. This resolves the issues of us always mapping everything from
the beginning of the symbol `spp`. Each group is mapped independently,
and at the same level, like `spp[0][0]` and its member `spp[0][0].a`, we
still get map them together as part of the same contiguous struct
`spp[0][0]`. This resolves issue #141042.
#### use_device_ptr/addr fixes
The handling of `use_device_ptr/addr` was updated to use the attach-ptr
information, and works for many cases that were failing before. It has
to be done as part of this series because otherwise, the switch from
ptr_to_obj to attach-style mapping would have caused regressions in
existing use_device_ptr/addr tests.
#### Handling of attach-pointers that are members of implicitly mapped
structs:
* When a struct member-pointer, like `p` below, is a base-pointer in a
`map` clause on a target construct (like `map(p[0:1])`, and the base of
that struct is either the `this` pointer (implicitly or explicitly), or
a struct that is implicitly mapped on that construct, we add an implicit
`map(p)` so that we don't implicitly map the full struct.
```c
struct S { int *p;
void f1() {
#pragma omp target map(p[0:1]) // Implicitly map this->p, to ensure
// that the implicit map of `this[:]` does
// not map the full struct
printf("%p %p\n", &p, p);
}
```
#### Scope for improvement:
* We may be able to compute attach-ptr expr while collecting
component-lists in Sema.
* But we cache the computation results already, and `findAttachPtrExpr`
is fairly simple, and fast.
* There may be a better way to implement semantic expr comparison.
#### Needs future work:
* Attach-style maps not yet emitted for declare mappers.
* Mapping of class member references: We are still using PTR_AND_OBJ
maps for them. We will likely need to change that to handle
`ref_ptr/ref_ptee`, and `attach` map-type-modifier on them.
* Implicit capturing of "this" needs to map the full `this[0:1]` unless
there is an explicit map on one of the members, or a map with a member
as its base-pointer.
* Implicit map added for capturing a class member pointer needs to also
add a zero-length-array-section map.
* `use_device_addr` on array-sections-on-pointers need further
improvements (documented using FIXMEs)
#### Why a large PR
While it's unfortunate that this PR has gotten large and difficult to
review, the issue is that all the functional changes have to be made
together, to prevent regressions from partially implemented changes.
For example, the changes to capturing were previously done separately
(#145454), but they would still cause stability issues in absence of
full attach-mapping. And attach-mapping needs those changes to be able
to launch kernels.
We extracted the utilities and functions, like those for finding
attach-ptrs, or comparing exprs, out as a separate NFC PR that doesn't
call those functions, just adds them (#155625). Maybe the change that
adds a new error message for use_device_addr on array-sections with
non-var base-pointers could have been extracted out too (but that would
have had to be a follow-up change in that case, and we would get
comp-fails with this PR when the erroneous case was not
caught/diagnosed).
---------
Co-authored-by: Alex Duran <alejandro.duran@intel.com>
|
|
|
|
This commit makes the cuLaunchKernel call to pass the total arguments size without tail padding.
|
|
(#169331)" (#170851)
Add support for OpenMP is_device_ptr clause for target directives.
[MLIR][OpenMP] Add OpenMPToLLVMIRTranslation support for is_device_ptr
#169367 This PR adds support for the OpenMP is_device_ptr clause in the
MLIR to LLVM IR translation for target regions. The is_device_ptr clause
allows device pointers (allocated via OpenMP runtime APIs) to be used
directly in target regions without implicit mapping.
|
|
Reverts llvm/llvm-project#169331
|
|
Add support for OpenMP is_device_ptr clause for target directives.
[MLIR][OpenMP] Add OpenMPToLLVMIRTranslation support for is_device_ptr #169367
This PR adds support for the OpenMP is_device_ptr clause in the MLIR to LLVM IR translation for target regions. The is_device_ptr clause allows device pointers (allocated via OpenMP runtime APIs) to be used directly in target regions without implicit mapping.
|
|
These are C tests, not C++, so no function parameters means unspecified
number of parameters, not `void`.
These compile fine on the current tested offload targets because an
error is only
[thrown](https://github.com/llvm/llvm-project/blob/main/clang/lib/Sema/SemaDecl.cpp#L10695)
if the calling convention doesn't support variadic arguments, which they
happen to.
When compiling this test for other targets that do not support variadic
arguments, we get an error, which does not seem intentional.
Just add `void` to the parameter list.
---------
Signed-off-by: Nick Sarnie <nick.sarnie@intel.com>
|
|
This patch add support for lowering of custom reductions to MLIR. It
also enhances the capability of the pass to automatically mark functions
as "declare target" by traversing custom reduction initializers and
combiners.
|
|
(#119589)
While the infrastructure for declare target to/enter and link for
variables exists in the MLIR dialect and at the Flang level, the current
lowering from MLIR -> LLVM IR isn't in place, it's only in place for
variables that have the link clause applied.
This PR aims to extend that lowering to an initial implementation that
incorporates declare target to as well, which primarily requires changes
in the OpenMPToLLVMIRTranslation phase. However, a minor addition to the
OpenMP dialect was required to extend the declare target enumerator to
include a default None field as well.
This also requires a minor change to the Flang lowering's
MapInfoFinlization.cpp pass to alter the map type for descriptors to
deal with cases where a variable is marked declare to. Currently, when a
descriptor variable is mapped declare target to the descriptor component
can become attatched, and cannot be updated, this results in issues when
an unusual allocation range is specified (effectively an off-by X
error). The current solution is to map the descriptor always, as we
always require an up-to-date version of this data. However, this also
requires an interlinked PR that adds a more intricate type of mapping of
structures/record types that clang currently implements, to circumvent
the overwriting of the pointer in the descriptor.
3/3 required PRs to enable declare target to mapping, this PR should
pass all tests and provide an all green CI.
Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com
|
|
This PR introduces a new additional type of map lowering for record
types that Clang currently supports, in which a user can map a top-level
record type and then individual members with different mapping,
effectively creating a sort of "overlapping" mapping that we attempt to
cut around.
This is currently most predominantly used in Fortran, when mapping
descriptors and there data, we map the descriptor and its data with
separate map modifiers and "cut around" the pointer data, so that wedo
not overwrite it unless the runtime deems it a neccesary action based on
its reference counting mechanism. However, it is a mechanism that will
come in handy/trigger when a user explitily maps a record type (derived
type or structure) and then explicitly maps a member with a different
map type.
These additions were predominantly in the OpenMPToLLVMIRTranslation.cpp
file and phase, however, one Flang test that checks end-to-end IR
compilation (as far as we care for now at least) was altered.
2/3 required PRs to enable declare target to mapping, should look at PR
3/3 to check for full green passes (this one will fail a number due to
some dependencies).
Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com
|
|
(#140562)
This patch adds support to emit default declare mappers for implicit
mapping of derived types when not supplied by user. This especially
helps tackle mapping of allocatables of derived types.
|
|
After #163011 was merged, the tests in
[`offload/test/offloading/gpupgo`](https://github.com/llvm/llvm-project/compare/main...EthanLuisMcDonough:llvm-project:gpupgo-names-fix-pr?expand=1#diff-f769f6cebd25fa527bd1c1150cc64eb585c41cb8a8b325c2bc80c690e47506a1)
broke because the offload plugins were no longer able to find
`__llvm_prf_nm`. This pull request explicitly makes `__llvm_prf_nm`
visible to the host on GPU targets and reverses the changes made in
f7e9968a5ba99521e6e51161f789f0cc1745193f.
|
|
Summary:
I made that an unimplemented error, but forgot that it was used for this
environment variable.
|
|
Summary:
This was a lot of code that was only used for upstream LLVM builds of
AMDGPU offloading. We have a generic and fast `malloc` in `libc` now so
just use that. Simplifies code, can be added back if we start providing
alternate forms but I don't think there's a single use-case that would
justify it yet.
|
|
offsets (#165486)
Currently this is being calculated incorrectly and will result in
incorrect index offsets in more complicated array slices. This PR tries
to address it by refactoring and changing the calculation to be more
correct.
|
|
While people look into it, xfail the tests.
|
|
|
|
The test covers some of the identifier symbols in the interop runtime.
This test, for now, is to guard against complete breakage, which was the
result of the other `interop.c` test not being enabled on AMD and thus,
not caught by our buildbots.
|
|
(#154349)
This PR adds deferral of descriptor maps until they are necessary for
assumed dummy argument types. The intent is to avoid a problem where a
user can inadvertently map a temporary local descriptor to device
without their knowledge and proceed to never unmap it. This temporary
local descriptor remains lodged in OpenMP device memory and the next
time another variable or descriptor residing in the same stack address
is mapped we incur a runtime OpenMP map error as we try to remap the
same address.
This fix was discussed with the OpenMP committee and applies to OpenMP
5.2 and below, future versions of OpenMP can avoid this issue via the
attach semantics added to the specification.
|
|
types (#160766)
This PR adds support for nested derived types and their mappers to the
MapInfoFinalization pass.
- Generalize MapInfoFinalization to add child maps for arbitrarily
nested allocatables when a derived object is mapped via declare mapper.
- Traverse HLFIR designates rooted at the target block arg and build
full coordinate_of chains; append members with correct membersIndex.
This fixes #156461.
|
|
Summary:
Several of these tests have been failing for literal years. Ideally we
make efforts to fix this, but keeping these broken has had serious
consequences on our testing infrastructure where failures are the norm
so almost all test failures are disregarded. I made a tracking issue for
the ones that have been disabled.
https://github.com/llvm/llvm-project/issues/161265
|
|
Fortran no-loop test is supported only for GPU.
|
|
Enable the generation of no-loop kernels for Fortran OpenMP code. target
teams distribute parallel do pragmas can be promoted to no-loop kernels
if the user adds the -fopenmp-assume-teams-oversubscription and
-fopenmp-assume-threads-oversubscription flags.
If the OpenMP kernel contains reduction or num_teams clauses, it is not
promoted to no-loop mode.
The global OpenMP device RTL oversubscription flags no longer force
no-loop code generation for Fortran.
|
|
derived types" (#160759)
Reverts llvm/llvm-project#160116
|
|
types (#160116)
This PR adds support for nested derived types and their mappers to the
MapInfoFinalization pass.
- Generalize MapInfoFinalization to add child maps for arbitrarily
nested allocatables when a derived object is mapped via declare mapper.
- Traverse HLFIR designates rooted at the target block arg and build
full coordinate_of chains; append members with correct membersIndex.
This fixes #156461.
|
|
With declare mapper, the parent base entry was emitted as `TARGET_PARAM`
only. The mapper received a map-type without `to/from`, causing
components to degrade to `alloc`-only (no copies), breaking allocatable
payload mapping. This PR preserves the map-type bits from the parent.
This fixes #156466.
|
|
Adds end-to-end tests for `do concurrent` offloading to the device.
PR stack:
- https://github.com/llvm/llvm-project/pull/155754
- https://github.com/llvm/llvm-project/pull/155987
- https://github.com/llvm/llvm-project/pull/155992
- https://github.com/llvm/llvm-project/pull/155993 ◀️
- https://github.com/llvm/llvm-project/pull/157638
- https://github.com/llvm/llvm-project/pull/156610
- https://github.com/llvm/llvm-project/pull/156837
|
|
The refactoring lead to an additional data transfer. This changes the
assumed transfers in the check-strings to work with that changed
behavior.
|
|
Currently, there's a number of issues with mapping characters with LEN's
specified (strings effectively). They're represented as a char type in
FIR with a len parameter, and then later on they're expanded into an
array of characters when we're translating to the LLVM dialect. However,
we don't generate a bounds for these at lowering. The fix in this PR for
this is to generate a bounds from the LEN parameter and attatch it to
the map on lowering from FIR to the LLVM dialect when we encounter this
type.
|
|
(Re)enables a couple of tests that were disabled on AMDGPU for some
reason. Pass for me locally.
|
|
|
|
Add a new AutomapToTargetData pass. This gathers the declare target
enter variables which have the AUTOMAP modifier. And adds
omp.declare_target_enter/exit mapping directives for fir.alloca and
fir.free oeprations on the AUTOMAP enabled variables.
Automap Ref: OpenMP 6.0 section 7.9.7.
|
|
(#153048)"
This reverts commit 4e6d510eb3ec5b5e5ea234756ea1f0b283feee4a.
|
|
Add a new AutomapToTargetData pass. This gathers the declare target
enter variables which have the AUTOMAP modifier. And adds
omp.declare_target_enter/exit mapping directives for fir.alloca and
fir.free oeprations on the AUTOMAP enabled variables.
Automap Ref: OpenMP 6.0 section 7.9.7.
|
