Previously, GEPOp relies on `findKnownStructIndices` to check if a GEP
index should be static. The truth is, `findKnownStructIndices` can only
tell you a GEP index _might_ be indexing into a struct (which should use
a static GEP index). But GEPOp::build and GEPOp::verify are falsely
taking this information as a certain answer, which creates many false
alarms like the one depicted in
`test/Target/LLVMIR/Import/dynamic-gep-index.ll`.
The solution presented here adopts a new verification scheme: When we're
recursively checking the child element types of a struct type, instead
of checking every child types, we only check the one dictated by the
(static) GEP index value. We also combine "refinement" logics --
refine/promote struct index mlir::Value into constants -- into the very
verification process since they have lots of logics in common. The
resulting code is more concise and less brittle.
We also hide GEPOp::findKnownStructIndices since most of the
aforementioned logics are already encapsulated within GEPOp::build and
GEPOp::verify, we found little reason for findKnownStructIndices (or the
new findStructIndices) to be public.
Differential Revision: https://reviews.llvm.org/D124935
Complex nested in other types is perfectly fine, just nested structs
aren't supported. Instead of checking whether there's nesting just check
whether the struct we're dealing with is a complex number.
Differential Revision: https://reviews.llvm.org/D125381
By analogy with the NVGPU dialect, introduce an AMDGPU dialect for
AMD-specific intrinsic wrappers.
The dialect initially includes wrappers around the raw buffer intrinsics.
On AMD GPUs, a memref can be converted to a "buffer descriptor" that
allows more precise control of memory access, such as by allowing for
out of bounds loads/stores to be replaced by 0/ignored without adding
additional conditional logic, which is important for performance.
The repository currently contains a limited conversion from
transfer_read/transfer_write to Mubuf intrinsics, which are an older,
deprecated intrinsic for the same functionality.
The new amdgpu.raw_buffer_* ops allow these operations to be used
explicitly and for including metadata such as whether the target
chipset is an RDNA chip or not (which impacts the interpretation of
some bits in the buffer descriptor), while still maintaining an
MLIR-like interface.
(This change also exposes the floating-point atomic add intrinsic.)
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D122765
Add attribute to be able to generate the intrinsic version of async copy
generating a copy with l1 bypass. This correspond to
cp.async.cg.shared.global in ptx.
Differential Revision: https://reviews.llvm.org/D125241
A large DenseElementsAttr of i1could trigger a bug in printer/parser roundtrip.
Ex. A DenseElementsAttr of i1 with 200 elements will print as Hex format of length 400 before the fix. However, when parsing the printed text, an error will be triggered. After fix, the printed length will be 50.
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D122925
The NVVM dialect test coverage for all possible type/shape combinations
in the `nvvm.mma.sync` op is mostly complete. However, there were tests
missing for TF32 datatype support. This change adds tests for the one
relevant shape/type combination. This uncovered a small bug in the op
verifier, which this change also fixes.
Differential Revision: https://reviews.llvm.org/D124975
Inside processInstruction, we assign the translated mlir::Value to a
reference previously taken from the corresponding entry in instMap.
However, instMap (a DenseMap) might resize after the entry reference was
taken, rendering the assignment useless since it's assigning to a
dangling reference. Here is a (pseudo) snippet that shows the concept:
```
// inst has type llvm::Instruction *
Value &v = instMap[inst];
...
// op is one of the operands of inst, has type llvm::Value *
processValue(op);
// instMap resizes inside processValue
...
translatedValue = b.createOp<Foo>(...);
// v is already a dangling reference at this point!
// The following assignment is bogus.
v = translatedValue;
```
Nevertheless, after we stop caching llvm::Constant into instMap, there
is only one case that can cause processValue to resize instMap: If the
operand is a llvm::ConstantExpr. In which case we will insert the
derived llvm::Instruction into instMap.
To trigger instMap to resize, which is a DenseMap, the threshold depends
on the ratio between # of map entries and # of (hash) buckets. More specifically,
it resizes if (# of map entries / # of buckets) >= 0.75.
In this case # of map entries is equal to # of LLVM instructions, and # of
buckets is the power-of-two upperbound of # of map entries. Thus, eventually
in the attaching test case (test/Target/LLVMIR/Import/incorrect-instmap-assignment.ll),
we picked 96 and 128 for the # of map entries and # of buckets, respectively.
(We can't pick numbers that are too small since DenseMap used inlined
storage for small number of entries). Therefore, the ConstantExpr in the
said test case (i.e. a GEP) is the 96-th llvm::Value cached into the
instMap, triggering the issue we're discussing here on its enclosing
instruction (i.e. a load).
This patch fixes this issue by calling `operator[]` everytime we need to
update an entry.
Differential Revision: https://reviews.llvm.org/D124627
This patch add supports for translating FCmp and more kinds of FP
constants in addition to 32 & 64-bit ones. However, we can't express
ppc_fp128 constants right now because the semantics for its underlying
APFloat is `S_PPCDoubleDouble` but mlir::FloatType doesn't support such
semantics right now.
Differential Revision: https://reviews.llvm.org/D124630
This adds a cast operation that allows to perform an explicit type
conversion. The cast op is emitted as a C-style cast. It can be applied
to integer, float, index and EmitC types.
Reviewed By: jpienaar
Differential Revision: https://reviews.llvm.org/D123514
Constants in MLIR are not globally unique, unlike that in LLVM IR.
Therefore, reusing previous-translated constants might cause the user
operations not being dominated by the constant (because the
previous-translated ones can be placed in arbitrary place)
This indeed misses some opportunities where we actually can reuse a
previous-translated constants, but verbosity is not our first priority
here.
Differential Revision: https://reviews.llvm.org/D124404
More specifically, the llvm::Instruction generated by
llvm::ConstantExpr::getAsInstruction. Such Instruction will be deleted
right away, but it's possible that when getAsInstruction is called
again, it will create a new Instruction that has the same address with
the one we just deleted. Thus, we shouldn't keep it in the `instMap` to
avoid a conflicting index that triggers an assertion in
processInstruction.
Differential Revision: https://reviews.llvm.org/D124402
And move importer test files from `test/Target/LLVMIR` into
`test/Target/LLVMIR/Import`.
We simply translate struct-type ConstantAggregate(Zero) into a
serious of `llvm.insertvalue` operations against a `llvm.undef` root.
Note that this doesn't affect the original logics on translating
vector/array-type ConstantAggregate values.
Differential Revision: https://reviews.llvm.org/D124399
The callback is expected to create a branch to the ContinuationBB (sometimes called FiniBB in some lambdas) argument when finishing. This creates problems:
1. The InsertPoint used for CodeGenIP does not need to be the end of a block. If it is not, a naive callback will insert a branch instruction into the middle of the block.
2. The BasicBlock the CodeGenIP is pointing to may or may not have a terminator. There is an conflict where to branch to if the block already has a terminator.
3. Some API functions work only with block having a terminator. Some workarounds have been used to insert a temporary terminator that is removed again.
4. Some callbacks are sensitive to whether the BasicBlock has a terminator or not. This creates a callback ordering problem where different callback may have different behaviour depending on whether a previous callback created a terminator or not. The problem also exists for FinalizeCallbackTy where some callbacks do create branch to another "continue" block, but unlike BodyGenCallbackTy does not receive the target as argument. This is not addressed in this patch.
With this patch, the callback receives an CodeGenIP into a BasicBlock where to insert instructions. If it has to insert control flow, it can split the block at that position as needed but otherwise no separate ContinuationBB is needed. In particular, a callback can be empty without breaking the emitted IR. If the caller needs the control flow to branch to a specific target, it can insert the branch instruction itself and pass an InsertPoint before the terminator to the callback.
Certain frontends such as Clang may expect the current IRBuilder position to be at the end of a basic block. In this case its callbacks must split the block at CodeGenIP before setting the IRBuilder position such that the instructions after CodeGenIP are moved to another basic block and before returning create a new branch instruction to the split block.
Some utility functions such as `splitBB` are supporting correct splitting of BasicBlocks, independent of whether they have a terminator or not, returning/setting the InsertPoint of an IRBuilder to the end of split predecessor block, and optionally omitting creating a branch to the split successor block to be added later.
Reviewed By: kiranchandramohan
Differential Revision: https://reviews.llvm.org/D118409
This patch handles empty hint value for critical and atomic constructs.
This also adds checks and tests for hint clause on atomic constructs.
Reviewed By: peixin, kiranchandramohan, NimishMishra
Differential Revision: https://reviews.llvm.org/D123186
This reverts commit af0285122f.
The test "libomp::loop_dispatch.c" on builder
openmp-gcc-x86_64-linux-debian fails from time-to-time.
See #54969. This patch is unrelated.
The OMPScheduleType enum stores the constants from libomp's internal sched_type in kmp.h and are used by several kmp API functions. The enum values have an internal structure, namely each scheduling algorithm (e.g.) exists in four variants: unordered, orderend, normerge unordered, and nomerge ordered.
This patch (basically a followup to D114940) splits the "ordered" and "nomerge" bits into separate flags, as was already done for the "monotonic" and "nonmonotonic", so we can apply bit flags operations on them. It also now contains all possible combinations according to kmp's sched_type. Deriving of the OMPScheduleType enum from clause parameters has been moved form MLIR's OpenMPToLLVMIRTranslation.cpp to OpenMPIRBuilder to make available for clang as well. Since the primary purpose of the flag is the binary interface to libomp, it has been made more private to LLVMFrontend. The primary interface for generating worksharing-loop using OpenMPIRBuilder code becomes `applyWorkshareLoop` which derives the OMPScheduleType automatically and calls the appropriate emitter function.
While this is mostly a NFC refactor, it still applies the following functional changes:
* The logic from OpenMPToLLVMIRTranslation to derive the OMPScheduleType also applies to clang. Most notably, it now applies the nonmonotonic flag for non-static schedules by default.
* In OpenMPToLLVMIRTranslation, the nonmonotonic default flag was previously not applied if the simd modifier was used. I assume this was a bug, since the effect was due to `loop.schedule_modifier()` returning `mlir::omp::ScheduleModifier::none` instead of `llvm::Optional::None`.
* In OpenMPToLLVMIRTranslation, the nonmonotonic default flag was set even if ordered was specified, in breach to what the comment before citing the OpenMP specification says. I assume this was an oversight.
The ordered flag with parameter was not considered in this patch. Changes will need to be made (e.g. adding/modifying function parameters) when support for it is added. The lengthy names of the enum values can be discussed, for the moment this is avoiding reusing previously existing enum value names such as `StaticChunked` to avoid confusion.
Reviewed By: peixin
Differential Revision: https://reviews.llvm.org/D123403
LLVM IR is moving towards adoption of opaque pointer types. These require extra
information to be passed when constructing some operations, in particular GEP
and Alloca. Adapt the builders of said operations and modify the translation
code to handle both opaque and non-opaque pointers.
This incidentally adds the translation for Alloca alignment and fixes the translation
of struct-related GEP indices that must be constant.
Reviewed By: wsmoses
Differential Revision: https://reviews.llvm.org/D123792
LLVM IR has introduced and is moving forward with the concept of opaque
pointers, i.e. pointer types that are not carrying around the pointee type.
Instead, memory-related operations indicate the type of the data being accessed
through the opaque pointer. Introduce the initial support for opaque pointers
in the LLVM dialect:
- `LLVMPointerType` to support omitting the element type;
- alloca/load/store/gep to support opaque pointers in their operands and
results; this requires alloca and gep to store the element type as an
attribute;
- memory-related intrinsics to support opaque pointers in their operands;
- translation to LLVM IR for the ops above is no longer using methods
deprecated in LLVM API due to the introduction of opaque pointers.
Unlike LLVM IR, MLIR can afford to support both opaque and non-opaque pointers
at the same time and simplify the transition. Translation to LLVM IR of MLIR
that involves opaque pointers requires the LLVMContext to be configured to
always use opaque pointers.
Reviewed By: wsmoses
Differential Revision: https://reviews.llvm.org/D123310
This patch adds thread_local to llvm.mlir.global and adds translation for dso_local and addr_space to and from LLVM IR.
Reviewed By: Mogball
Differential Revision: https://reviews.llvm.org/D123412
This supports the threadprivate directive in OpenMP dialect following
the OpenMP 5.1 [2.21.2] standard. Also lowering to LLVM IR using OpenMP
IRBduiler.
Reviewed By: kiranchandramohan, shraiysh, arnamoy10
Differential Revision: https://reviews.llvm.org/D123350
This patch supports ordered clause specified without parameter in
worksharing-loop directive in the OpenMPIRBuilder and lowering MLIR to
LLVM IR.
Reviewed By: Meinersbur
Differential Revision: https://reviews.llvm.org/D114940
This patch adds MLIR NVVM support for the various NVPTX `mma.sync`
operations. There are a number of possible data type, shape,
and other attribute combinations supported by the operation, so a
custom assebmly format is added and attributes are inferred where
possible.
Reviewed By: ThomasRaoux
Differential Revision: https://reviews.llvm.org/D122410
This patch adds the nowait parameter to `createSingle` in
OpenMPIRBuilder and handling for IR generation from OpenMP Dialect.
Also added tests for the same.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D122371
This patch adds translation from omp.single to LLVM IR.
Depends on D122288
Reviewed By: ftynse, kiranchandramohan
Differential Revision: https://reviews.llvm.org/D122297
This patch
- adds assembly format for `omp.wsloop` operation
- removes the `parseClauses` clauses as it is not required anymore
This is expected to be the final patch in a series of patches for replacing
parsers for clauses with `oilist`.
Reviewed By: Mogball
Differential Revision: https://reviews.llvm.org/D121367
This patch adds translation from `omp.atomic.capture` to LLVM IR. Also
added tests for the same.
Depends on D121546
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D121554
This patch fixes the condition for emitting atomic update using
`atomicrmw` instruction or compare-exchange loop.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D121546
This commit moves FuncOp out of the builtin dialect, and into the Func
dialect. This move has been planned in some capacity from the moment
we made FuncOp an operation (years ago). This commit handles the
functional aspects of the move, but various aspects are left untouched
to ease migration: func::FuncOp is re-exported into mlir to reduce
the actual API churn, the assembly format still accepts the unqualified
`func`. These temporary measures will remain for a little while to
simplify migration before being removed.
Differential Revision: https://reviews.llvm.org/D121266
During MLIR translation to LLVMIR if an inlineable call has an UnkownLoc we get this error message:
```
inlinable function call in a function with debug info must have a !dbg location
call void @callee()
```
There is code that checks for this case and strips debug information to avoid this situation. I'm expanding this code to handle the case where an debug location points at a UnknownLoc. For example, a NamedLoc whose child location is an UnknownLoc.
Reviewed By: ftynse
Differential Revision: https://reviews.llvm.org/D121633
Patch adds a new operation for the SIMD construct. The op is designed to be very similar to the existing `wsloop` operation, so that the `CanonicalLoopInfo` of `OpenMPIRBuilder` can be used.
Reviewed By: shraiysh
Differential Revision: https://reviews.llvm.org/D118065
This patch moves the testcases from
`mlir/test/Target/LLVMIR/openmp-llvm-bad-schedule-modifier.mlir` to
`mlir/test/Dialect/OpenMP/invalid.mlir` as they test the verifier
(not the translation to LLVM IR).
Reviewed By: NimishMishra
Differential Revision: https://reviews.llvm.org/D120877
This patch adds lowering from omp.atomic.update to LLVM IR. Whenever a
special LLVM IR instruction is available for the operation, `atomicrmw`
instruction is emitted, otherwise a compare-exchange loop based update
is emitted.
Depends on D119522
Reviewed By: ftynse, peixin
Differential Revision: https://reviews.llvm.org/D119657
These passes generally don't rely on any special aspects of FuncOp, and moving allows
for these passes to be used in many more situations. The passes that obviously weren't
relying on invariants guaranteed by a "function" were updated to be generic pass, the
rest were updated to be FunctionOpinterface InterfacePasses.
The test updates are NFC switching from implicit nesting (-pass -pass2) form to
the -pass-pipeline form (generic passes do not implicitly nest as op-specific passes do).
Differential Revision: https://reviews.llvm.org/D121190
The current StandardToLLVM conversion patterns only really handle
the Func dialect. The pass itself adds patterns for Arithmetic/CFToLLVM, but
those should be/will be split out in a followup. This commit focuses solely
on being an NFC rename.
Aside from the directory change, the pattern and pass creation API have been renamed:
* populateStdToLLVMFuncOpConversionPattern -> populateFuncToLLVMFuncOpConversionPattern
* populateStdToLLVMConversionPatterns -> populateFuncToLLVMConversionPatterns
* createLowerToLLVMPass -> createConvertFuncToLLVMPass
Differential Revision: https://reviews.llvm.org/D120778
Min-Yih Hsu