summaryrefslogtreecommitdiffstats
path: root/mlir/lib/Target
Commit message (Collapse)AuthorAgeFilesLines
* [mlir] Fix translation of splat constants to LLVM IRAlex Zinenko2020-01-141-1/+8
| | | | | | | | | | | | | | | | Summary: When converting splat constants for nested sequential LLVM IR types wrapped in MLIR, the constant conversion was erroneously assuming it was always possible to recursively construct a constant of a sequential type given only one value. Instead, wait until all sequential types are unpacked recursively before constructing a scalar constant and wrapping it into the surrounding sequential type. Subscribers: mehdi_amini, rriddle, jpienaar, burmako, shauheen, antiagainst, nicolasvasilache, arpith-jacob, mgester, lucyrfox, aartbik, liufengdb, llvm-commits Tags: #llvm Differential Revision: https://reviews.llvm.org/D72688
* [mlir] NFC: Remove Value::operator* and Value::operator-> now that Value is ↵River Riddle2020-01-112-3/+3
| | | | | | | | | | properly value-typed. Summary: These were temporary methods used to simplify the transition. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D72548
* [mlir] Fix -Wrange-loo-analysis warningsFangrui Song2020-01-011-1/+1
| | | | | | | | | | | | for (const auto &x : llvm::zip(..., ...)) -> for (auto x : llvm::zip(..., ...)) The return type of zip() is a wrapper that wraps a tuple of references. > warning: loop variable 'p' is always a copy because the range of type 'detail::zippy<detail::zip_shortest, ArrayRef<long> &, ArrayRef<long> &>' does not return a reference [-Wrange-loop-analysis]
* [mlir] Floating constants for import-llvmAlex Zinenko2019-12-271-0/+6
| | | | | | | | | | | | | | | | | | | | | Summary: `mlir-translate -import-llvm test.ll` was going into segmentation fault if `test.ll` had `float` or `double` constants. For example, ``` %3 = fadd double 3.030000e+01, %0 ``` Now, it is handled in `Importer::getConstantAsAttr` (similar behaviour as normal integers) Added tests for FP arithmetic Reviewers: ftynse, mehdi_amini Reviewed By: ftynse, mehdi_amini Subscribers: shauheen, mehdi_amini, rriddle, jpienaar, burmako, llvm-commits Tags: #llvm Differential Revision: https://reviews.llvm.org/D71912
* NFC: Replace ValuePtr with Value and remove it now that Value is value-typed.River Riddle2019-12-232-23/+23
| | | | | | ValuePtr was a temporary typedef during the transition to a value-typed Value. PiperOrigin-RevId: 286945714
* Adjust License.txt file to use the LLVM licenseMehdi Amini2019-12-235-65/+20
| | | | PiperOrigin-RevId: 286906740
* NFC: Introduce new ValuePtr/ValueRef typedefs to simplify the transition to ↵River Riddle2019-12-222-24/+24
| | | | | | | | | | Value being value-typed. This is an initial step to refactoring the representation of OpResult as proposed in: https://groups.google.com/a/tensorflow.org/g/mlir/c/XXzzKhqqF_0/m/v6bKb08WCgAJ This change will make it much simpler to incrementally transition all of the existing code to use value-typed semantics. PiperOrigin-RevId: 286844725
* Detemplatize ModuleTranslation::lookupValuesAlex Zinenko2019-12-191-0/+10
| | | | | | | | | This function template has been introduced in the early days of MLIR to work around the absence of common type for ranges of values (operands, block argumeents, vectors, etc). Core IR now provides ValueRange for exactly this purpose. Use it instead of the template parameter. PiperOrigin-RevId: 286431338
* NFC: Remove forbidden include of <iostream>River Riddle2019-12-181-2/+0
| | | | | See: https://llvm.org/docs/CodingStandards.html#include-iostream-is-forbidden PiperOrigin-RevId: 286226467
* NFC: Cleanup non-conforming usages of namespaces.River Riddle2019-12-183-25/+22
| | | | | | | * Fixes use of anonymous namespace for static methods. * Uses explicit qualifiers(mlir::) instead of wrapping the definition with the namespace. PiperOrigin-RevId: 286222654
* NFC: Remove unnecessary 'llvm::' prefix from uses of llvm symbols declared ↵River Riddle2019-12-184-23/+21
| | | | | | | | in `mlir` namespace. Aside from being cleaner, this also makes the codebase more consistent. PiperOrigin-RevId: 286206974
* Move function template definition to the header file. NFCAlex Zinenko2019-12-181-11/+0
| | | | | | | | | The definition of the function template LLVM::ModuleTranslation::lookupValues has been located in a source file. As long as it has been the only file that actually called into the function, this did not cause any problem. However, it creates linking issues if the function is used from other translation units. PiperOrigin-RevId: 286203078
* Integrate from upstream at revision e4fce659a759.MLIR Team2019-12-172-0/+2
| | | | PiperOrigin-RevId: 285982330
* Remove LLVM dependency on mlir::Module and instead check Traits.Tres Popp2019-12-163-14/+17
| | | | PiperOrigin-RevId: 285724678
* Automated rollback of commit f68ac464d818629e0fe10c23b44ac782d64a12d2Christian Sigg2019-12-121-5/+5
| | | | PiperOrigin-RevId: 285162061
* Switch from shfl.bfly to shfl.down.Christian Sigg2019-12-121-5/+5
| | | | | | | 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
* Introduce Linkage attribute to the LLVM dialectAlex Zinenko2019-12-022-6/+71
| | | | | | | | | | | 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
* Add support for nested symbol references.River Riddle2019-11-111-2/+2
| | | | | | | | | | | | | | | | | | This change allows for adding additional nested references to a SymbolRefAttr to allow for further resolving a symbol if that symbol also defines a SymbolTable. If a referenced symbol also defines a symbol table, a nested reference can be used to refer to a symbol within that table. Nested references are printed after the main reference in the following form: symbol-ref-attribute ::= symbol-ref-id (`::` symbol-ref-id)* Example: module @reference { func @nested_reference() } my_reference_op @reference::@nested_reference Given that SymbolRefAttr is now more general, the existing functionality centered around a single reference is moved to a derived class FlatSymbolRefAttr. Followup commits will add support to lookups, rauw, etc. for scoped references. PiperOrigin-RevId: 279860501
* mlir-translate: support -verify-diagnosticsAlex Zinenko2019-11-071-8/+8
| | | | | | | | | | | MLIR translation tools can emit diagnostics and we want to be able to check if it is indeed the case in tests. Reuse the source manager error handlers provided for mlir-opt to support the verification in mlir-translate. This requires us to change the signature of the functions that are registered to translate sources to MLIR: it now takes a source manager instead of a memory buffer. PiperOrigin-RevId: 279132972
* [llvm] Allow GlobalOp to take a region for complex initializersJames Molloy2019-11-052-19/+37
| | | | | | | | | | | | | | | 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
* [llvm] Add initial import of LLVM modules to mlir-translateJames Molloy2019-11-052-1/+597
| | | | | | | | | | | | | | | | This adds an importer from LLVM IR or bitcode to the LLVM dialect. The importer is registered with mlir-translate. Known issues exposed by this patch but not yet fixed: * Globals' initializers are attributes, which makes it impossible to represent a ConstantExpr. This will be fixed in a followup. * icmp returns i32 rather than i1. * select and a couple of other instructions aren't implemented. * llvm.cond_br takes its successors in a weird order. The testing here is known to be non-exhaustive. I'd appreciate feedback on where this functionality should live. It looks like the translator *from MLIR to LLVM* lives in Target/, but the SPIR-V deserializer lives in Dialect/ which is why I've put this here too. PiperOrigin-RevId: 278711683
* Get active source lane predicate from shuffle instruction.Christian Sigg2019-10-191-0/+11
| | | | | | | | 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
* Translation to LLVM: check the validity of module-level OpsAlex Zinenko2019-10-101-0/+8
| | | | | | | | | | | Translation to LLVM expects the entry module to have only specific types of ops that correspond to LLVM IR entities allowed in a module. Currently those are restricted to functions and globals. Introduce an additional check at the module level. Inside individual functions, the check for supported Ops is already performed, but it accepts all LLVM dialect Ops and wouldn't be immediately applicable at the module level. PiperOrigin-RevId: 274058651
* Use llvm.func to define functions with wrapped LLVM IR function typeAlex Zinenko2019-10-104-51/+17
| | | | | | | | | | | | | | This function-like operation allows one to define functions that have wrapped LLVM IR function type, in particular variadic functions. The operation was added in parallel to the existing lowering flow, this commit only switches the flow to use it. Using a custom function type makes the LLVM IR dialect type system more consistent and avoids complex conversion rules for functions that previously had to use the built-in function type instead of a wrapped LLVM IR dialect type and perform conversions during the analysis. PiperOrigin-RevId: 273910855
* [ROCm] Fix the return type for the device function calls from i32 to i64.Deven Desai2019-10-081-1/+1
| | | | | | | | | This is matching what the runtime library is expecting. Closes tensorflow/mlir#171 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/171 from deven-amd:deven-rocdl-device-func-i64 80762629a8c34e844ebdc542b34dd783990db9db PiperOrigin-RevId: 273640767
* Fix typos, NFC.Christian Sigg2019-10-041-1/+1
| | | | PiperOrigin-RevId: 272851237
* [ROCm] Adding ROCDL Dialect.Deven Desai2019-09-272-0/+133
| | | | | | | | | | | | | 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
* Make GlobalOp's value attribute optional.Christian Sigg2019-09-211-2/+4
| | | | | | 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
* Add address space attribute to LLVMIR's GlobalOp.MLIR Team2019-09-191-3/+6
| | | | PiperOrigin-RevId: 270012505
* Change MLIR translation functions signatureLei Zhang2019-09-172-22/+4
| | | | | | | | | | This CL changes translation functions to take MemoryBuffer as input and raw_ostream as output. It is generally better to avoid handling files directly in a library (unless the library is specifically for file manipulation) and we can unify all file handling to the mlir-translate binary itself. PiperOrigin-RevId: 269625911
* Add support for array-typed constants.MLIR Team2019-09-041-9/+21
| | | | PiperOrigin-RevId: 267121729
* LLVM dialect: prefix auxiliary operations with "mlir."Alex Zinenko2019-09-031-1/+2
| | | | | | | | | | Some of the operations in the LLVM dialect are required to model the LLVM IR in MLIR, for example "constant" operations are needed to declare a constant value since MLIR, unlike LLVM, does not support immediate values as operands. To avoid confusion with actual LLVM operations, we prefix such axuiliary operations with "mlir.". PiperOrigin-RevId: 266942838
* Add 3 additional intrinsic ops to NVVM dialect, in preparation to implement ↵MLIR Team2019-08-271-3/+4
| | | | | | block-wide reduce. PiperOrigin-RevId: 265720077
* Add iterator support to ElementsAttr and SparseElementsAttr.River Riddle2019-08-221-2/+2
| | | | | | This will allow iterating the values of a non-opaque ElementsAttr, with all of the types currently supported by DenseElementsAttr. This should help reduce the amount of specialization on DenseElementsAttr. PiperOrigin-RevId: 264968151
* Automated rollback of commit b9dc2e481818315f2f0d87455349f497f6118a4cRiver Riddle2019-08-211-2/+2
| | | | PiperOrigin-RevId: 264672975
* Add iterator support to ElementsAttr and SparseElementsAttr.River Riddle2019-08-211-2/+2
| | | | | | This will allow iterating the values of a non-opaque ElementsAttr, with all of the types currently supported by DenseElementsAttr. This should help reduce the amount of specialization on DenseElementsAttr. PiperOrigin-RevId: 264637293
* NFC: Move LLVMIR, SDBM, and StandardOps to the Dialect/ directory.River Riddle2019-08-192-4/+4
| | | | PiperOrigin-RevId: 264193915
* Use unreachable post switch rather than default case.Jacques Pienaar2019-08-121-4/+2
| | | | | | | | Prefer to enumerate all cases in the switch instead of using default to allow compiler to flag missing cases. This also avoids -Wcovered-switch-default warning. PiperOrigin-RevId: 262935972
* LLVM dialect: introduce llvm.addressof to access globalsAlex Zinenko2019-08-121-9/+23
| | | | | | | | | | | This instruction is a local counterpart of llvm.global that takes a symbol reference to a global and produces an SSA value containing the pointer to it. Used in combination, these two operations allow one to use globals with other operations expecting SSA values. At a cost of IR indirection, we make sure the functions don't implicitly capture the surrounding SSA values and remain suitable for parallel processing. PiperOrigin-RevId: 262908622
* Translation to LLVM IR: use LogicalResult instead of boolAlex Zinenko2019-08-092-32/+28
| | | | | | | | The translation code predates the introduction of LogicalResult and was relying on the obsolete LLVM convention of returning false on success. Change it to use MLIR's LogicalResult abstraction instead. NFC. PiperOrigin-RevId: 262589432
* LLVM dialect and translation: support global stringsAlex Zinenko2019-08-091-0/+11
| | | | | | | | | | | Unlike regular constant values, strings must be placed in some memory and referred to through a pointer to that memory. Until now, they were not supported in function-local constant declarations with `llvm.constant`. Introduce support for global strings using `llvm.global`, which would translate them into global arrays in LLVM IR and thus make sure they have some memory allocated for storage. PiperOrigin-RevId: 262569316
* Translation to LLVM: support llvm.globalAlex Zinenko2019-08-091-0/+10
| | | | | | | Add support for translating recently introduced llvm.global operations to global variables in the LLVM IR proper. PiperOrigin-RevId: 262564700
* Add support for floating-point comparison 'fcmp' to the LLVM dialect.Nagy Mostafa2019-08-081-0/+39
| | | | | | | | This adds support for fcmp to the LLVM dialect and adds any necessary lowerings, as well as support for EDSCs. Closes tensorflow/mlir#69 PiperOrigin-RevId: 262475255
* Move GPU dialect to {lib,include/mlir}/DialectAlex Zinenko2019-07-251-1/+1
| | | | | | | Per tacit agreement, individual dialects should now live in lib/Dialect/Name with headers in include/mlir/Dialect/Name and tests in test/Dialect/Name. PiperOrigin-RevId: 259896851
* Decouple LLVM dialect from Standard dialectAlex Zinenko2019-07-161-12/+11
| | | | | | | | | | | | | | | Due to the absence of ODS support for enum attributes, the implementation of the LLVM dialect `icmp` operation was reusing the comparison predicate from the Standard dialect, creating an avoidable library dependency. With ODS support and ICmpPredicate attribute recently introduced, the dependency is no longer justified. Update the Standard to LLVM convresion to also convert the CmpIPredicate into LLVM::ICmpPredicate and remove the unnecessary includes. Note that the MLIRLLVMIR library did not explicitly depend on MLIRStandardOps, requiring dependees of MLIRLLVMIR to also depend on MLIRStandardOps, which should no longer be the case. PiperOrigin-RevId: 258148456
* Rename FunctionAttr to SymbolRefAttr.River Riddle2019-07-121-2/+2
| | | | | | This allows for the attribute to hold symbolic references to other operations than FuncOp. This also allows for removing the dependence on FuncOp from the base Builder. PiperOrigin-RevId: 257650017
* NFC: Rename Module to ModuleOp.River Riddle2019-07-103-6/+7
| | | | | | Module is a legacy name that only exists as a typedef of ModuleOp. PiperOrigin-RevId: 257427248
* NFC: Rename Function to FuncOp.River Riddle2019-07-102-6/+6
| | | | PiperOrigin-RevId: 257293379
* NFC: Remove `Module::getFunctions` in favor of a general `getOps<T>`.River Riddle2019-07-082-3/+3
| | | | | | Modules can now contain more than just Functions, this just updates the iteration API to reflect that. The 'begin'/'end' methods have also been updated to iterate over opaque Operations. PiperOrigin-RevId: 257099084
* Make TranslateFromMLIRFunction type return LogicalResult instead of boolAlex Zinenko2019-07-042-8/+8
| | | | | | This interface was created before MLIR introduced LogicalResult. PiperOrigin-RevId: 256554557
OpenPOWER on IntegriCloud