summaryrefslogtreecommitdiffstats
path: root/mlir/test/Target
Commit message (Collapse)AuthorAgeFilesLines
* [mlir] Fix translation of splat constants to LLVM IRAlex Zinenko2020-01-141-0/+28
| | | | | | | | | | | | | | | | 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] Floating constants for import-llvmAlex Zinenko2019-12-271-0/+29
| | | | | | | | | | | | | | | | | | | | | 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
* Add integer bit-shift operations to the standard dialect.Manuel Freiberger2019-12-221-0/+6
| | | | | | | | | | | | | | | | | | | Rename the 'shlis' operation in the standard dialect to 'shift_left'. Add tests for this operation (these have been missing so far) and add a lowering to the 'shl' operation in the LLVM dialect. Add also 'shift_right_signed' (lowered to LLVM's 'ashr') and 'shift_right_unsigned' (lowered to 'lshr'). The original plan was to name these operations 'shift.left', 'shift.right.signed' and 'shift.right.unsigned'. This works if the operations are prefixed with 'std.' in MLIR assembly. Unfortunately during import the short form is ambigous with operations from a hypothetical 'shift' dialect. The best solution seems to omit dots in standard operations for now. Closes tensorflow/mlir#226 PiperOrigin-RevId: 286803388
* Added LLVM ops and lowering phases from standard dialect for FAbs, FCeil, ↵Marcel Koester2019-12-181-0/+42
| | | | | | | | | | | Cos, FNeg, CopySign. Added test cases for the newly added LLVM operations and lowering features. Closes tensorflow/mlir#300 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/300 from dfki-jugr:std_to_llvm da6168bbc1a369ae2e99ad3881fdddd82f075dd4 PiperOrigin-RevId: 286231169
* Introduce prefetch op: affine -> std -> llvm intrinsicUday Bondhugula2019-12-181-2/+8
| | | | | | | | | | | | | | | | | | | Introduce affine.prefetch: op to prefetch using a multi-dimensional subscript on a memref; similar to affine.load but has no effect on semantics, but only on performance. Provide lowering through std.prefetch, llvm.prefetch and map to llvm's prefetch instrinsic. All attributes reflected through the lowering - locality hint, rw, and instr/data cache. affine.prefetch %0[%i, %j + 5], false, 3, true : memref<400x400xi32> Signed-off-by: Uday Bondhugula <uday@polymagelabs.com> Closes tensorflow/mlir#225 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/225 from bondhugula:prefetch 4c3b4e93bc64d9a5719504e6d6e1657818a2ead0 PiperOrigin-RevId: 286212997
* Automated rollback of commit f68ac464d818629e0fe10c23b44ac782d64a12d2Christian Sigg2019-12-121-8/+8
| | | | PiperOrigin-RevId: 285162061
* Switch from shfl.bfly to shfl.down.Christian Sigg2019-12-121-8/+8
| | | | | | | 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
* Add std.log* and llvm.intr.log* that correspond to the LLVMIR intrinsicsNicolas Vasilache2019-12-111-0/+33
| | | | PiperOrigin-RevId: 285073483
* Introduce Linkage attribute to the LLVM dialectAlex Zinenko2019-12-022-15/+70
| | | | | | | | | | | 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
* Allow LLVM::ExtractElementOp to have non-i32 indices.MLIR Team2019-11-251-4/+13
| | | | | | Also change the text format a bit, so that indices are braced by squares. PiperOrigin-RevId: 282437095
* mlir-translate: support -verify-diagnosticsAlex Zinenko2019-11-071-0/+6
| | | | | | | | | | | 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
* Add support for the LLVM FNeg instructionEric Schweitz2019-11-061-0/+3
| | | | | | | Closes tensorflow/mlir#216 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/216 from schweitzpgi:llvmir-fneg-op f9b5f185845d671b745ab6fc213d5d9aff044b34 PiperOrigin-RevId: 278795325
* [llvm] Allow GlobalOp to take a region for complex initializersJames Molloy2019-11-052-0/+17
| | | | | | | | | | | | | | | 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-051-0/+146
| | | | | | | | | | | | | | | | 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
* [NVVM] Add mma.sync operation.MLIR Team2019-11-041-0/+9
| | | | PiperOrigin-RevId: 278440547
* [mlir][llvm] Add missing cast opsJames Molloy2019-11-011-0/+19
| | | | | | Also adds a builder method for fcmp, identical to that for icmp. PiperOrigin-RevId: 277923158
* Get active source lane predicate from shuffle instruction.Christian Sigg2019-10-191-0/+10
| | | | | | | | 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
* LLVM Dialect: introduce llvm.mlir.null operationAlex Zinenko2019-10-111-0/+8
| | | | | | | | 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
* Use llvm.func to define functions with wrapped LLVM IR function typeAlex Zinenko2019-10-104-56/+57
| | | | | | | | | | | | | | 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
* Add exp operation to LLVMOPs.td.Alexander Belyaev2019-10-091-3/+13
| | | | PiperOrigin-RevId: 273718958
* [ROCm] Fix the return type for the device function calls from i32 to i64.Deven Desai2019-10-081-12/+12
| | | | | | | | | 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
* Add integer shift ops to LLVM dialect.Christian Sigg2019-09-301-0/+6
| | | | PiperOrigin-RevId: 272140049
* [ROCm] Adding ROCDL Dialect.Deven Desai2019-09-271-0/+35
| | | | | | | | | | | | | 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
* Add integer sign- and zero-extension and truncation to standard.Manuel Freiberger2019-09-211-0/+10
| | | | | | | | | | | | This adds sign- and zero-extension and truncation of integer types to the standard dialects. This allows to perform integer type conversions without having to go to the LLVM dialect and introduce custom type casts (between standard and LLVM integer types). Closes tensorflow/mlir#134 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/134 from ombre5733:sext-zext-trunc-in-std c7657bc84c0ca66b304e53ec03797e09152e4d31 PiperOrigin-RevId: 270479722
* Make GlobalOp's value attribute optional.Christian Sigg2019-09-211-0/+3
| | | | | | 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-0/+3
| | | | PiperOrigin-RevId: 270012505
* Add warpsize and laneid intrinsics to NVVM dialect.MLIR Team2019-09-091-0/+4
| | | | PiperOrigin-RevId: 268041263
* Add support for array-typed constants.MLIR Team2019-09-041-2/+8
| | | | PiperOrigin-RevId: 267121729
* LLVM dialect: prefix auxiliary operations with "mlir."Alex Zinenko2019-09-031-85/+85
| | | | | | | | | | 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-12/+34
| | | | | | block-wide reduce. PiperOrigin-RevId: 265720077
* Add FPToSI/FPExt/FPTrunc cast ops to the LLVM dialect.Eric Schweitz2019-08-261-0/+10
| | | | | | Closes tensorflow/mlir#99 PiperOrigin-RevId: 265538731
* Add iterator support to ElementsAttr and SparseElementsAttr.River Riddle2019-08-221-1/+8
| | | | | | 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-8/+1
| | | | PiperOrigin-RevId: 264672975
* Add iterator support to ElementsAttr and SparseElementsAttr.River Riddle2019-08-211-1/+8
| | | | | | 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
* LLVM dialect: prefix operations that correspond to intrinsics with "intr."Alex Zinenko2019-08-201-2/+2
| | | | | | | | | | LLVM intrinsics have an open name space and their names can potentially overlap with names of LLVM instructions (LLVM intrinsics are functions, not instructions). In MLIR, LLVM intrinsics are modeled as operations, so it needs to make sure their names cannot clash with the instructions. Use the "intr." prefix for intrinsics in the LLVM dialect. PiperOrigin-RevId: 264372173
* Add alignment support for llvm.allocaNicolas Vasilache2019-08-181-0/+10
| | | | | | Extend the LLVM dialect AllocaOp with an alignment attribute. PiperOrigin-RevId: 264068306
* LLVM dialect: introduce fmuladd intrinsic as operationAlex Zinenko2019-08-131-0/+15
| | | | | | | | | | | This operation is important to achieve decent performance in computational kernels. In LLVM, it is implemented as an intrinsic (through function declaration and function call). Thanks to MLIR's extendable set of operations, it does not have to differentiate between built-ins and intrinsics, so fmuladd is introduced as a general type-polymorphic operation. Custom printing and parsing will be added later. PiperOrigin-RevId: 263106305
* LLVM dialect: introduce llvm.addressof to access globalsAlex Zinenko2019-08-121-0/+17
| | | | | | | | | | | 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
* LLVM dialect and translation: support global stringsAlex Zinenko2019-08-091-0/+3
| | | | | | | | | | | 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/+9
| | | | | | | Add support for translating recently introduced llvm.global operations to global variables in the LLVM IR proper. PiperOrigin-RevId: 262564700
* Add support for vector ops in the LLVM dialectNicolas Vasilache2019-08-091-0/+11
| | | | | | | | | | | | This CL is step 1/n towards building a simple, programmable and portable vector abstraction in MLIR that can go all the way down to generating assembly vector code via LLVM's opt and llc tools. This CL adds the 3 instructions `llvm.extractelement`, `llvm.insertelement` and `llvm.shufflevector` as documented in the LLVM LangRef "Vector Instructions" section. The "Experimental Vector Reduction Intrinsics" are left out for now and can be added in the future on a per-need basis. Appropriate roundtrip and LLVM Target tests are added. PiperOrigin-RevId: 262542095
* Add support for floating-point comparison 'fcmp' to the LLVM dialect.Nagy Mostafa2019-08-081-0/+33
| | | | | | | | 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
* Add the LLVM IR unreachable instruction to the LLVMIR dialect.Eric Schweitz2019-08-081-0/+4
| | | | | | | | http://llvm.org/docs/LangRef.html#unreachable-instruction Closes tensorflow/mlir#64 PiperOrigin-RevId: 262301557
* [mlir-translate] Fix test suite.Denis Khalikov2019-08-051-10/+10
| | | | | | | | | | llvm ir printer was changed at LLVM r367755. Prints value numbers for unnamed functions argument. Closes tensorflow/mlir#67 COPYBARA_INTEGRATE_REVIEW=https://github.com/tensorflow/mlir/pull/67 from denis0x0D:sandbox/fix_mlir_translate ae46844e66f34a02e0cf86782ddadc5bce58b30d PiperOrigin-RevId: 261640048
* Change the attribute dictionary syntax to separate name and value with '='.River Riddle2019-06-251-2/+2
| | | | | | | | | | | The current syntax separates the name and value with ':', but ':' is already overloaded by several other things(e.g. trailing types). This makes the syntax difficult to parse in some situtations: Old: "foo: 10 : i32" New: "foo = 10 : i32" PiperOrigin-RevId: 255097928
* Modify the syntax of the the ElementsAttrs to print the type as a colon type.River Riddle2019-06-251-1/+1
| | | | | | | | | This is the standard syntax for types on operations, and is also already used by IntegerAttr and FloatAttr. Example: dense<5> : tensor<i32> dense<[3]> : tensor<1xi32> PiperOrigin-RevId: 255069157
* NVVM target: emit nvvm.annotations for kernel functionsAlex Zinenko2019-06-251-1/+11
| | | | | | | | | | PTX backend in LLVM expects additional module-level metadata `!nvvm.annotations` that lists functions that can be used as GPU kernels. Generate this metadata based on the `gpu.kernel` attribute attached to functions. This attribute is added automatically by the kernel outlining pass in the GPU dialect lowering flow. PiperOrigin-RevId: 254957345
* Refactor SplatElementsAttr to inherit from DenseElementsAttr as opposed to ↵River Riddle2019-06-191-1/+1
| | | | | | being a separate Attribute type. DenseElementsAttr provides a better internal representation for splat values as well as better API for accessing elements. PiperOrigin-RevId: 253138287
* Add support for llvm.constant with StringAttr as value.Stephan Herhut2019-06-011-0/+6
| | | | | | | | | These are translated to an llvm::ConstantDataArray on translation to llvm IR proper. -- PiperOrigin-RevId: 249813111
* Add LLVM::IntToPtrOp and LLVM::PtrToIntOp to LLVM dialect.Stephan Herhut2019-06-011-0/+9
| | | | | | -- PiperOrigin-RevId: 249604199
OpenPOWER on IntegriCloud