Commit Graph

112 Commits

Author SHA1 Message Date
Rahul Joshi 549eac9d87 [MLIR] Remove unnecessary CHECK's from tests for which we do not run FileCheck.
Differential Revision: https://reviews.llvm.org/D90651
2020-11-02 15:21:33 -08:00
Alex Zinenko 03e6f40cdb [mlir] Do not print back 0 alignment in LLVM dialect 'alloca' op
The alignment attribute in the 'alloca' op treats the '0' value as 'unset'.
When parsing the custom form of the 'alloca' op, ignore the alignment attribute
with if its value is '0' instead of actually creating it and producing a
slightly different textually yet equivalent semantically form in the output.

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D90179
2020-10-26 23:19:20 +01:00
Mehdi Amini f05173d0bf Implement callee/caller type checking for llvm.call
This aligns the behavior with the standard call as well as the LLVM verifier.

Reviewed By: ftynse, dcaballe

Differential Revision: https://reviews.llvm.org/D88362
2020-10-04 20:15:06 +00:00
Alex Zinenko aec9e20a3e [mlir] introduce type constraints for operands of LLVM dialect operations
Historically, the operations in the MLIR's LLVM dialect only checked that the
operand are of LLVM dialect type without more detailed constraints. This was
due to LLVM dialect types wrapping LLVM IR types and having clunky verification
methods. With the new first-class modeling, it is possible to define type
constraints similarly to other dialects and use them to enforce some
correctness rules in verifiers instead of having LLVM assert during translation
to LLVM IR. This hardening discovered several issues where MLIR was producing
LLVM dialect operations that cannot exist in LLVM IR.

Depends On D85900

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D85901
2020-09-04 10:01:59 +02:00
River Riddle 431bb8b318 [mlir][ODS] Use c++ types for integer attributes of fixed width when possible.
Unsigned and Signless attributes use uintN_t and signed attributes use intN_t, where N is the fixed width. The 1-bit variants use bool.

Differential Revision: https://reviews.llvm.org/D86739
2020-09-01 13:43:32 -07:00
Kazuaki Ishizaki a23d055912 [mlir] NFC: fix trivial typo under test and tools
Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D86648
2020-08-27 15:37:42 +09:00
Alex Zinenko 9c4825ce28 [mlir] do not use llvm.cmpxchg with floats
According to the LLVM Language Reference, 'cmpxchg' accepts integer or pointer
types. Several MLIR tests were using it with floats as it appears possible to
programmatically construct and print such an instruction, but it cannot be
parsed back. Use integers instead.

Depends On D85899

Reviewed By: flaub, rriddle

Differential Revision: https://reviews.llvm.org/D85900
2020-08-17 15:44:23 +02:00
Alex Zinenko 168213f91c [mlir] Move data layout from LLVMDialect to module Op attributes
Legacy implementation of the LLVM dialect in MLIR contained an instance of
llvm::Module as it was required to parse LLVM IR types. The access to the data
layout of this module was exposed to the users for convenience, but in practice
this layout has always been the default one obtained by parsing an empty layout
description string. Current implementation of the dialect no longer relies on
wrapping LLVM IR types, but it kept an instance of DataLayout for
compatibility. This effectively forces a single data layout to be used across
all modules in a given MLIR context, which is not desirable. Remove DataLayout
from the LLVM dialect and attach it as a module attribute instead. Since MLIR
does not yet have support for data layouts, use the LLVM DataLayout in string
form with verification inside MLIR. Introduce the layout when converting a
module to the LLVM dialect and keep the default "" description for
compatibility.

This approach should be replaced with a proper MLIR-based data layout when it
becomes available, but provides an immediate solution to compiling modules with
different layouts, e.g. for GPUs.

This removes the need for LLVMDialectImpl, which is also removed.

Depends On D85650

Reviewed By: aartbik

Differential Revision: https://reviews.llvm.org/D85652
2020-08-17 15:12:36 +02:00
Alex Zinenko bae1517266 [mlir] Add verification to LLVM dialect types
Now that LLVM dialect types are implemented directly in the dialect, we can use
MLIR hooks for verifying type construction invariants. Implement the verifiers
and use them in the parser.

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D85663
2020-08-11 17:21:52 +02:00
Alex Zinenko 4e491570b5 [mlir] Remove LLVMTypeTestDialect
This dialect was introduced during the bring-up of the new LLVM dialect type
system for testing purposes. The main LLVM dialect now uses the new type system
and the test dialect is no longer necessary, so remove it.

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D85224
2020-08-05 14:39:36 +02:00
Alex Zinenko ec1f4e7c3b [mlir] switch the modeling of LLVM types to use the new mechanism
A new first-party modeling for LLVM IR types in the LLVM dialect has been
developed in parallel to the existing modeling based on wrapping LLVM `Type *`
instances. It resolves the long-standing problem of modeling identified
structure types, including recursive structures, and enables future removal of
LLVMContext and related locking mechanisms from LLVMDialect.

This commit only switches the modeling by (a) renaming LLVMTypeNew to LLVMType,
(b) removing the old implementaiton of LLVMType, and (c) updating the tests. It
is intentionally minimal. Separate commits will remove the infrastructure built
for the transition and update API uses where appropriate.

Depends On D85020

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D85021
2020-08-04 14:29:25 +02:00
Alex Zinenko 0c40af6b59 [mlir] First-party modeling of LLVM types
The current modeling of LLVM IR types in MLIR is based on the LLVMType class
that wraps a raw `llvm::Type *` and delegates uniquing, printing and parsing to
LLVM itself. This model makes thread-safe type manipulation hard and is being
progressively replaced with a cleaner MLIR model that replicates the type
system.  Introduce a set of classes reflecting the LLVM IR type system in MLIR
instead of wrapping the existing types. These are currently introduced as
separate classes without affecting the dialect flow, and are exercised through
a test dialect. Once feature parity is reached, the old implementation will be
gradually substituted with the new one.

Depends On D84171

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D84339
2020-08-03 15:45:29 +02:00
Alex Zinenko cba733edf5 [mlir] LLVM dialect: use addressof instead of constant to create function pointers
`llvm.mlir.constant` was originally introduced as an LLVM dialect counterpart
to `std.constant`. As such, it was supporting "function pointer" constants
derived from the symbol name. This is different from `std.constant` that allows
for creation of a "function" constant since MLIR, unlike LLVM IR, supports
this. Later, `llvm.mlir.addressof` was introduced as an Op that obtains a
constant pointer to a global in the LLVM dialect. It naturally extends to
functions (in LLVM IR, functions are globals) and should be used for defining
"function pointer" values instead.

Fixes PR46344.

Differential Revision: https://reviews.llvm.org/D82667
2020-06-29 12:21:33 +02:00
Alex Zinenko fbeceb9ced [mlir] Modernize LLVM dialect rountrip test
This test largely predates MLIR testing guidelines. Update it to match the
guidelines. In particular, avoid pattern-matching SSA value names, avoid
unnecessary CHECK-NEXT, relax assumptions about the form of SSA names.
Value-returning operations are still matched agaist _any_ name in order to
check that the operation indeed produces values.

Differential Revision: https://reviews.llvm.org/D82656
2020-06-29 09:47:36 +02:00
Alex Zinenko 6323065fd6 [mlir] support returning unranked memrefs
Initially, unranked memref descriptors in the LLVM dialect were designed only
to be passed into functions. An assertion was guarding against returning
unranked memrefs from functions in the standard-to-LLVM conversion. This is
insufficient for functions that wish to return an unranked memref such that the
caller does not know the rank in advance, and hence cannot allocate the
descriptor and pass it in as an argument.

Introduce a calling convention for returning unranked memref descriptors as
follows. An unranked memref descriptor always points to a ranked memref
descriptor stored on stack of the current function. When an unranked memref
descriptor is returned from a function, the ranked memref descriptor it points
to is copied to dynamically allocated memory, the ownership of which is
transferred to the caller. The caller is responsible for deallocating the
dynamically allocated memory and for copying the pointed-to ranked memref
descriptor onto its stack.

Provide default lowerings for std.return, std.call and std.indirect_call that
maintain the conversion defined above.

This convention is additionally exercised by a runtime test to guard against
memory errors.

Differential Revision: https://reviews.llvm.org/D82647
2020-06-26 15:37:37 +02:00
George Mitenkov a2edbd8170 [MLIR][LLVMDialect] Added bitreverse and ctpop intrinsics
Introduced `llvm.intr.bitreverse` and `llvm.intr.ctpop` LLVM bit
intrinsics to LLVM dialect. These intrinsics help with SPIR-V to
LLVM conversion, allowing a direct mapping from `spv.BitReverse`
and `spv.BitCount` respectively. Tests are added to `roundtrip.mlir`
and `llvm-intrinsics.mlir`.

Differential Revision: https://reviews.llvm.org/D82285
2020-06-23 14:25:35 -04:00
Stephan Herhut 2416e28c25 [mlir] Add support for alignment annotations to the LLVM dialect to LLVM translation.
Summary:
With this change, a function argument attribute of the form
"llvm.align" = <int> will be translated to the corresponding align
attribute in LLVM by the ModuleConversion.

Differential Revision: https://reviews.llvm.org/D82161
2020-06-19 16:36:06 +02:00
Jacques Pienaar b0921f68e1 [mlir] Add verify method to adaptor
This allows verifying op-indepent attributes (e.g., attributes that do not require the op to have been created) before constructing an operation. These include checking whether required attributes are defined or constraints on attributes (such as I32 attribute). This is not perfect (e.g., if one had a disjunctive constraint where one part relied on the op and the other doesn't, then this would not try and extract the op independent from the op dependent).

The next step is to move these out to a trait that could be verified earlier than in the generated method. The first use case is for inferring the return type while constructing the op. At that point you don't have an Operation yet and that ends up in one having to duplicate the same checks, e.g., verify that attribute A is defined before querying A in shape function which requires that duplication. Instead this allows one to invoke a method to verify all the traits and, if this is checked first during verification, then all other traits could use attributes knowing they have been verified.

It is a little bit funny to have these on the adaptor, but I see the adaptor as a place to collect information about the op before the op is constructed (e.g., avoiding stringly typed accessors, verifying what is possible to verify before the op is constructed) while being cheap to use even with constructed op (so layer of indirection between the op constructed/being constructed). And from that point of view it made sense to me.

Differential Revision: https://reviews.llvm.org/D80842
2020-06-05 09:47:37 -07:00
George Mitenkov 7293dd5b40 Added pow intrinsic to LLVMIR dialect
Added pow intrinsic to LLVMIR dialect. Added a roundrip test for it.

Differential Revision: https://reviews.llvm.org/D80248
2020-05-25 07:57:33 -04:00
jerryyin 9c53ac08de [mlir][rocdl] Exposing buffer load/store intrinsic
Summary:
* Updated ROCDLOps tablegen
* Added parsing and printing function for new intrinsic
* Added unit tests

Reviewers: ftynse

Subscribers: mehdi_amini, rriddle, jpienaar, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, stephenneuendorffer, Joonsoo, grosul1, frgossen, Kayjukh, jurahul, llvm-commits

Tags: #llvm

Differential Revision: https://reviews.llvm.org/D80233
2020-05-21 14:14:35 +00:00
MaheshRavishankar 49e6c19100 [mlir][StandardToLLVM] Add SinOp to LLVM dialect and lowering of std.sin to this op.
Differential Revision: https://reviews.llvm.org/D79505
2020-05-12 23:15:25 -07:00
rtayl 58cb88733f [mlir][rocdl] Add xdlops intrinsics to rocdl dialect
Summary: This adds xdlops (mfma) to the rocdl dialect and also tests the translation to llvm ir.

Reviewers: ftynse

Subscribers: mehdi_amini, rriddle, jpienaar, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, stephenneuendorffer, Joonsoo, grosul1, frgossen, Kayjukh, llvm-commits

Tags: #llvm #mlir

Differential Revision: https://reviews.llvm.org/D79642
2020-05-11 10:08:58 -04:00
Wen-Heng (Jack) Chung bc23c1d85e [mlir][rocdl] add rocdl.barier op.
- Add rocdl.barrier op.
- Lower gpu.barier to rocdl.barrier in -convert-gpu-to-rocdl.

Differential Revision: https://reviews.llvm.org/D79126
2020-05-04 10:35:01 +02:00
Alex Zinenko 129cf84e69 [mlir] LLVM dialect: support globals without linkage keyword, assuming 'external'
Similarly to actual LLVM IR, and to `llvm.mlir.func`, allow the custom syntax
of `llvm.mlir.global` to omit the linkage keyword. If omitted, the linkage is
assumed to be external. This makes the modeling of globals in the LLVM dialect
more consistent, both within the dialect and with LLVM IR.

Differential Revision: https://reviews.llvm.org/D78096
2020-04-15 10:58:32 +02:00
River Riddle 0359b86d8b [mlir][ODS] Add support for variadic regions.
Summary: This revision adds support for marking the last region as variadic in the ODS region list with the VariadicRegion directive.

Differential Revision: https://reviews.llvm.org/D77455
2020-04-05 01:03:38 -07:00
Mehdi Amini bab5bcf8fd Add a flag on the context to protect against creation of operations in unregistered dialects
Differential Revision: https://reviews.llvm.org/D76903
2020-03-30 19:37:31 +00:00
Shraiysh Vaishay ff77397fcf [mlir] Added llvm.resume and personality functions in LLVM IR Dialect
`llvm.resume` is similar to `llvm.return` except that has to be exactly
one operand and that should be derived from a `llvm.landingpad`
instruction.  Any function having `llvm.landingpad` instruction must
have a personality attribute.

Example:
LLVM IR
```
define dso_local i32 @main() personality i32 (...)* @__gxx_personality_v0 {
  invoke void @foo(i32 42)
          to label %3 unwind label %1

1:                                                ; preds = %0
  %2 = landingpad i8*
          catch i8** @_ZTIi
          catch i8* bitcast (i8** @_ZTIi to i8*)
  resume i8* %2

3:                                                ; preds = %0
  ret i32 1
}
```

MLIR - LLVM IR Dialect

```
llvm.func @main() -> !llvm.i32 attributes {personality = @__gxx_personality_v0} {
    %0 = llvm.mlir.constant(1 : i32) : !llvm.i32
    %1 = llvm.mlir.addressof @_ZTIi : !llvm<"i8**">
    %2 = llvm.bitcast %1 : !llvm<"i8**"> to !llvm<"i8*">
    %3 = llvm.mlir.addressof @_ZTIi : !llvm<"i8**">
    %4 = llvm.mlir.constant(42 : i32) : !llvm.i32
    llvm.invoke @foo(%4) to ^bb2 unwind ^bb1 : (!llvm.i32) -> ()
  ^bb1:	// pred: ^bb0
    %5 = llvm.landingpad (catch %3 : !llvm<"i8**">) (catch %2 : !llvm<"i8*">) : !llvm<"i8*">
    llvm.resume %5 : !llvm<"i8*">
  ^bb2:	// pred: ^bb0
    llvm.return %0 : !llvm.i32
  }
```

Differential Revision: https://reviews.llvm.org/D71888
2020-03-19 13:14:25 +01:00
Sagar Jain 76cf14035b [MLIR] Added llvm.fence
This patch adds llvm.fence. I tried not to change the syntax much.

syntax:

LLVM IR
`fence [syncscope("<target-scope>")] <ordering>`

MLIR LLVM Dialect

`llvm.fence [syncscope("<target-scope>")] <ordering>`

example:
LLVM IR: `fence syncscope("agent") seq_cst`
MLIR: `llvm.fence syncscope("agent") seq_cst`

Differential Revision: https://reviews.llvm.org/D75645
2020-03-17 17:53:37 +01:00
Alex Zinenko e119980f3f [mlir] LLVM dialect: move ensureDistinctSuccessors out of std->LLVM conversion
MLIR supports terminators that have the same successor block with different
block operands, which cannot be expressed in the LLVM's phi-notation as the
block identifier is used to tell apart the predecessors. This limitation can be
worked around by branching to a new block instead, with this new block
unconditionally branching to the original successor and forwarding the
argument. Until now, this transformation was performed during the conversion
from the Standard to the LLVM dialect. This does not scale well to multiple
dialects targeting the LLVM dialect as all of them would have to be aware of
this limitation and perform the preparatory transformation. Instead, do it as a
separate pass and run it immediately before the translation.

Differential Revision: https://reviews.llvm.org/D75619
2020-03-17 15:22:14 +01:00
Lei Zhang 5b2cc6c3d0 [mlir][ods] Improve integer signedness modelling
A previous commit added support for integer signedness in C++
IntegerType. This change introduces ODS definitions for
integer types and integer (element) attributes w.r.t. signedness.

This commit also updates various existing definitions' descriptions
to mention signless where suitable to make it more clear.

Positive and non-negative integer attributes are removed to avoid
the explosion of subclasses. Instead, one should use more atmoic
constraints together with Confined to model that. For example,
`Confined<..., [IntPositive]>`.

Differential Revision: https://reviews.llvm.org/D75610
2020-03-04 15:05:42 -05:00
Sagar Jain d85821dfa6 [MLIR] Added llvm.freeze
This patch adds llvm.freeze & processes undef constants from LLVM IR.

Syntax:
LLVM IR
`<result> = freeze ty <val>`

MLIR LLVM Dialect:
`llvm.freeze val attr-dict : type`

Example:
LLVM IR: `%3 = freeze i32 5`
MLIR: `%6 = llvm.freeze %5 : !llvm.i32`

Differential Revision: https://reviews.llvm.org/D75329
2020-03-02 10:24:01 +01:00
Tim Shen b762bbd4c8 [MLIR] change NVVM.mma.sync to the most useful variant.
Summary:
the .row.col variant turns out to be the popular one, contrary to what I
thought as .row.row. Since .row.col is so prevailing (as I inspect
cuDNN's behavior), I'm going to remove the .row.row support here, which
makes the patch a little bit easier.

Reviewers: ftynse

Subscribers: jholewinski, bixia, sanjoy.google, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, liufengdb, Joonsoo, llvm-commits

Tags: #llvm

Differential Revision: https://reviews.llvm.org/D74655
2020-02-18 17:57:04 -08:00
River Riddle 528adb2e48 [mlir][NFC] Use declarative format for several operations in LLVM and Linalg dialects
Differential Revision: https://reviews.llvm.org/D73503
2020-01-30 11:43:41 -08:00
River Riddle 82170d5619 [mlir] Update various operations to declaratively specify their assembly format.
Summary:
This revision switches over many operations to use the declarative methods for defining the assembly specification. This updates operations in the NVVM, ROCDL, Standard, and VectorOps dialects.

Differential Revision: https://reviews.llvm.org/D73407
2020-01-30 11:43:40 -08:00
Shraiysh Vaishay d242aa245c [MLIR] Added llvm.invoke and llvm.landingpad
Summary:
I have tried to implement `llvm.invoke` and `llvm.landingpad`.

  # `llvm.invoke` is similar to `llvm.call` with two successors added, the first one is the normal label and the second one is unwind label.
  # `llvm.launchpad` takes a variable number of args with either `catch` or `filter` associated with them. Catch clauses are not array types and filter clauses are array types. This is same as the criteria used by LLVM (4f82af81a0/llvm/include/llvm/IR/Instructions.h (L2866))

Examples:
LLVM IR
```
define i32 @caller(i32 %a) personality i8* bitcast (i32 (...)* @__gxx_personality_v0 to i8*) {
    invoke i32 @foo(i32 2) to label %success unwind label %fail

  success:
    ret i32 2

  fail:
    landingpad {i8*, i32} catch i8** @_ZTIi catch i8** null catch i8* bitcast (i8** @_ZTIi to i8*) filter [1 x i8] [ i8 1 ]
    ret i32 3
}
```
MLIR LLVM Dialect
```
llvm.func @caller(%arg0: !llvm.i32) -> !llvm.i32 {
  %0 = llvm.mlir.constant(3 : i32) : !llvm.i32
  %1 = llvm.mlir.constant("\01") : !llvm<"[1 x i8]">
  %2 = llvm.mlir.addressof @_ZTIi : !llvm<"i8**">
  %3 = llvm.bitcast %2 : !llvm<"i8**"> to !llvm<"i8*">
  %4 = llvm.mlir.null : !llvm<"i8**">
  %5 = llvm.mlir.addressof @_ZTIi : !llvm<"i8**">
  %6 = llvm.mlir.constant(2 : i32) : !llvm.i32
  %7 = llvm.invoke @foo(%6) to ^bb1 unwind ^bb2 : (!llvm.i32) -> !llvm.i32
^bb1:	// pred: ^bb0
  llvm.return %6 : !llvm.i32
^bb2:	// pred: ^bb0
  %8 = llvm.landingpad (catch %5 : !llvm<"i8**">) (catch %4 : !llvm<"i8**">) (catch %3 : !llvm<"i8*">) (filter %1 : !llvm<"[1 x i8]">) : !llvm<"{ i8*, i32 }">
  llvm.return %0 : !llvm.i32
}
```

Signed-off-by: Shraiysh Vaishay <cs17btech11050@iith.ac.in>

Differential Revision: https://reviews.llvm.org/D72006
2020-01-30 12:55:28 +01:00
Frank Laub fffea2842d [MLIR] LLVM Dialect: add llvm.cmpxchg and improve llvm.atomicrmw custom parser
Summary:
Add a `llvm.cmpxchg` op as a counterpart to LLVM IR's `cmpxchg` instruction.
Note that the `weak`, `volatile`, and `syncscope` attributes are not yet supported.

This will be useful for upcoming parallel versions of affine.for and generally
for reduction-like semantics (especially for reductions that can't make use
of `atomicrmw`, e.g. `fmax`).

Reviewers: ftynse, nicolasvasilache

Reviewed By: ftynse

Subscribers: merge_guards_bot, jfb, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, llvm-commits

Tags: #llvm

Differential Revision: https://reviews.llvm.org/D72995
2020-01-21 01:09:42 -08:00
Mehdi Amini fdb9cc7dc5 Fix printer for llvm.addressof symbol name that need escaping
Differential Revision: https://reviews.llvm.org/D73065
2020-01-20 22:09:18 +00:00
Frank Laub ee2de95507 [MLIR] LLVM dialect: modernize and cleanups
Summary:
Modernize some of the existing custom parsing code in the LLVM dialect.
While this reduces some boilerplate code, it also reduces the precision
of the diagnostic error messges.

Reviewers: ftynse, nicolasvasilache, rriddle

Reviewed By: rriddle

Subscribers: merge_guards_bot, mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, arpith-jacob, mgester, lucyrfox, liufengdb, llvm-commits

Tags: #llvm

Differential Revision: https://reviews.llvm.org/D72967
2020-01-17 17:11:50 -08:00
Frank Laub 60a0c612df [MLIR] LLVM dialect: Add llvm.atomicrmw
Summary:
This op is the counterpart to LLVM's atomicrmw instruction. Note that
volatile and syncscope attributes are not yet supported.

This will be useful for upcoming parallel versions of `affine.for` and generally
for reduction-like semantics.

Differential Revision: https://reviews.llvm.org/D72741
2020-01-17 21:17:14 +01:00
Eric Schweitz 37e2560d3d [Flang][mlir] add a band-aid to support the creation of mutually recursive types when lowering to LLVM IR
Summary:
This is a temporary implementation to support Flang.  The LLVM-IR parser
will need to be extended in some way to support recursive types.  The
exact approach here is still a work-in-progress.

Unfortunately, this won't pass roundtrip testing yet. Adding a comment
to the test file as a reminder.

Differential Revision: https://reviews.llvm.org/D72542
2020-01-17 21:17:06 +01:00
Christian Sigg 9b85582682 Automated rollback of commit f68ac464d8
PiperOrigin-RevId: 285162061
2019-12-12 03:48:38 -08:00
Christian Sigg f68ac464d8 Switch from shfl.bfly to shfl.down.
Both work for the current use case, but the latter allows implementing
prefix sums and is a little easier to understand for partial warps.

PiperOrigin-RevId: 285145287
2019-12-12 01:28:01 -08:00
Alex Zinenko 58adf99ed1 LLVM::AddressOfOp: properly take into account the address space
The AddressOf operation in the LLVM dialect return a pointer to a global
variable. The latter may be in a non-default address space as indicated by the
"addr_space" attribute. Check that the address space of the pointer returned by
AddressOfOp matches that of the referenced GlobalOp. Update the AddressOfOp
builder to respect this constraint.

PiperOrigin-RevId: 284138860
2019-12-06 01:09:13 -08:00
Alex Zinenko fdbb99cd62 Add linkage support to LLVMFuncOp
A recent commit introduced the Linkage attribute to the LLVM dialect and used
it in the Global Op. Also use it in LLVMFuncOp. As per LLVM Language Reference,
if the linkage attribute is omitted, the function is assumed to have external
linkage.

PiperOrigin-RevId: 283493299
2019-12-03 00:26:44 -08:00
Alex Zinenko d5e627f84b Introduce Linkage attribute to the LLVM dialect
LLVM IR supports linkage on global objects such as global variables and
functions. Introduce the Linkage attribute into the LLVM dialect, backed by an
integer storage. Use this attribute on LLVM::GlobalOp and make it mandatory.
Implement parsing/printing of the attribute and conversion to LLVM IR.

See tensorflow/mlir#277.

PiperOrigin-RevId: 283309328
2019-12-02 03:28:10 -08:00
MLIR Team 1012c492f0 Allow LLVM::ExtractElementOp to have non-i32 indices.
Also change the text format a bit, so that indices are braced by squares.

PiperOrigin-RevId: 282437095
2019-11-25 14:44:52 -08:00
Alex Zinenko b5af3784a6 Don't force newline before function attributes
Due to legacy reasons, a newline character followed by two spaces was always
inserted before the attributes of the function Op in pretty form. This breaks
formatting when functions are nested in some other operations. Don't print the
newline and just put the attributes on the same line, which is also more
consistent with module Op. Line breaking aware of indentation can be introduced
separately into the parser if deemed useful.

PiperOrigin-RevId: 281721793
2019-11-21 05:08:19 -08:00
Eric Schweitz 0d545921ea Add support for the LLVM FNeg instruction
Closes tensorflow/mlir#216

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/216 from schweitzpgi:llvmir-fneg-op f9b5f185845d671b745ab6fc213d5d9aff044b34
PiperOrigin-RevId: 278795325
2019-11-06 00:02:10 -08:00
James Molloy 250a11ae0f [llvm] Allow GlobalOp to take a region for complex initializers
This allows GlobalOp to either take a value attribute (for simple constants) or a region that can
contain IR instructions (that must be constant-foldable) to create a ConstantExpr initializer.

Example:
  // A complex initializer is constructed with an initializer region.
  llvm.mlir.global constant @int_gep() : !llvm<"i32*"> {
    %0 = llvm.mlir.addressof @g2 : !llvm<"i32*">
    %1 = llvm.mlir.constant(2 : i32) : !llvm.i32
    %2 = llvm.getelementptr %0[%1] : (!llvm<"i32*">, !llvm.i32) -> !llvm<"i32*">
    llvm.return %2 : !llvm<"i32*">
  }
PiperOrigin-RevId: 278717836
2019-11-05 15:11:01 -08:00
MLIR Team 1f43d0d000 [NVVM] Add mma.sync operation.
PiperOrigin-RevId: 278440547
2019-11-04 12:36:37 -08:00
James Molloy 96531e2f87 [mlir][llvm] Add missing cast ops
Also adds a builder method for fcmp, identical to that for icmp.

PiperOrigin-RevId: 277923158
2019-11-01 07:32:09 -07:00
River Riddle 2b61b7979e Convert the Canonicalize and CSE passes to generic Operation Passes.
This allows for them to be used on other non-function, or even other function-like, operations. The algorithms are already generic, so this is simply changing the derived pass type. The majority of this change is just ensuring that the nesting of these passes remains the same, as the pass manager won't auto-nest them anymore.

PiperOrigin-RevId: 276573038
2019-10-24 15:01:09 -07:00
River Riddle 9ac459e871 Add a Symbol trait to simplify defining operations that represent symbols.
This trait provides accessors for the name, symbol use list methods, verification, with more to be added.

PiperOrigin-RevId: 275864554
2019-10-21 09:58:59 -07:00
Kazuaki Ishizaki f28c5aca17 Fix minor spelling tweaks (NFC)
Closes tensorflow/mlir#175

PiperOrigin-RevId: 275726876
2019-10-20 09:44:36 -07:00
Christian Sigg c3e56cd12c Get active source lane predicate from shuffle instruction.
nvvm.shfl.sync.bfly optionally returns a predicate whether source lane was active. Support for this was added to clang in https://reviews.llvm.org/D68892.

Add an optional 'pred' unit attribute to the instruction to return this predicate. Specify this attribute in the partial warp reduction so we don't need to manually compute the predicate.

PiperOrigin-RevId: 275616564
2019-10-19 01:53:25 -07:00
Lei Zhang 057dc41bf6 Allow '_' when pretty printing dialect symbols
'_' is used frequently enough as the separator of words in symbols.
We should allow it in dialect symbols when considering pretty printing.

Also updated LangRef.md regarding pretty form.

PiperOrigin-RevId: 275312494
2019-10-17 12:24:18 -07:00
Eric Schweitz a3d084848d Add LLVM IR dialect hooks for FP128 and X86_FP80 types
Closes tensorflow/mlir#184

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/184 from schweitzpgi:more-float-types ca27d00510a86ffc9c79c65fb3a0193b5ea097a0
PiperOrigin-RevId: 274288813
2019-10-11 18:35:33 -07:00
Alex Zinenko 71b82bcbf6 LLVM Dialect: introduce llvm.mlir.null operation
Similarly to `llvm.mlir.undef`, this auxiliary operation creates an SSA value
that corresponds to `null` in LLVM IR.  This operation is necessary to model
sizeof(<...>) behavior when allocating memory.

PiperOrigin-RevId: 274158760
2019-10-11 06:32:24 -07:00
Deven Desai fee40fef5c [ROCm] Adding ROCDL Dialect.
This commit introduces the ROCDL Dialect (i.e. the ROCDL ops + the code to lower those ROCDL ops to LLWM intrinsics/functions). Think of ROCDL Dialect as analogous to the NVVM Dialect, but for AMD GPUs. This patch contains just the essentials needed to get a simple example up and running. We expect to make further additions to the ROCDL Dialect.

This is the first of 3 commits, the follow-up will be:
 * add a pass that lowers GPU Dialect to ROCDL Dialect
 * add a "mlir-rocm-runner" utility

Closes tensorflow/mlir#146

COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/146 from deven-amd:deven-rocdl-dialect e78e8005c75a78912631116c78dc844fcc4b0de9
PiperOrigin-RevId: 271511259
2019-09-27 00:22:32 -07:00
Christian Sigg 33a3a91ba2 Make GlobalOp's value attribute optional.
Make GlobalOp's value attribute an OptionalAttr. Change code that uses the value to handle 'nullopt'. Translate an unitialized value attribute to llvm::UndefValue.

PiperOrigin-RevId: 270423646
2019-09-21 01:20:28 -07:00
MLIR Team e79bfefb89 Add address space attribute to LLVMIR's GlobalOp.
PiperOrigin-RevId: 270012505
2019-09-19 04:50:46 -07:00
Alex Zinenko c6f8adad8e Move LLVMIR dialect tests from test/LLVMIR to test/Dialect and test/Conversion
This follows up on the recent restructuring that moved the dialects under
lib/Dialect and inter-dialect conversions to lib/Conversion. Originally, the
tests for both the LLVMIR dialect itself and the conversion from Standard to
LLVMIR dialect lived under test/LLVMIR.  This no longer reflects the code
structure.  Move the tests to either test/Dialect/LLVMIR or
test/Conversion/StandardToLLVM depending on the features they exercise.

PiperOrigin-RevId: 267159219
2019-09-04 08:38:18 -07:00