summaryrefslogtreecommitdiffstats
path: root/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
Commit message (Collapse)AuthorAgeFilesLines
* Adjust License.txt file to use the LLVM licenseMehdi Amini2019-12-231-13/+4
| | | | PiperOrigin-RevId: 286906740
* NFC: Cleanup non-conforming usages of namespaces.River Riddle2019-12-181-1/+1
| | | | | | | * 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-181-9/+8
| | | | | | | | in `mlir` namespace. Aside from being cleaner, this also makes the codebase more consistent. PiperOrigin-RevId: 286206974
* Integrate from upstream at revision e4fce659a759.MLIR Team2019-12-171-0/+1
| | | | PiperOrigin-RevId: 285982330
* Remove LLVM dependency on mlir::Module and instead check Traits.Tres Popp2019-12-161-3/+4
| | | | 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
* 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
* Use llvm.func to define functions with wrapped LLVM IR function typeAlex Zinenko2019-10-101-3/+5
| | | | | | | | | | | | | | 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
* Change MLIR translation functions signatureLei Zhang2019-09-171-11/+2
| | | | | | | | | | 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 3 additional intrinsic ops to NVVM dialect, in preparation to implement ↵MLIR Team2019-08-271-3/+4
| | | | | | block-wide reduce. PiperOrigin-RevId: 265720077
* NFC: Move LLVMIR, SDBM, and StandardOps to the Dialect/ directory.River Riddle2019-08-191-2/+2
| | | | PiperOrigin-RevId: 264193915
* Translation to LLVM IR: use LogicalResult instead of boolAlex Zinenko2019-08-091-2/+2
| | | | | | | | 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
* 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
* NFC: Rename Module to ModuleOp.River Riddle2019-07-101-3/+4
| | | | | | Module is a legacy name that only exists as a typedef of ModuleOp. PiperOrigin-RevId: 257427248
* NFC: Rename Function to FuncOp.River Riddle2019-07-101-1/+1
| | | | PiperOrigin-RevId: 257293379
* NFC: Remove `Module::getFunctions` in favor of a general `getOps<T>`.River Riddle2019-07-081-1/+1
| | | | | | 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-041-4/+4
| | | | | | This interface was created before MLIR introduced LogicalResult. PiperOrigin-RevId: 256554557
* Replace the implementation of Function and Module with FuncOp and ModuleOp.River Riddle2019-07-031-1/+1
| | | | | | This is an important step in allowing for the top-level of the IR to be extensible. FuncOp and ModuleOp contain all of the necessary functionality, while using the existing operation infrastructure. As an interim step, many of the usages of Function and Module, including the name, will remain the same. In the future, many of these will be relaxed to allow for many different types of top-level operations to co-exist. PiperOrigin-RevId: 256427100
* NFC: Refactor Module to be value typed.River Riddle2019-07-021-5/+4
| | | | | | As with Functions, Module will soon become an operation, which are value-typed. This eases the transition from Module to ModuleOp. A new class, OwningModuleRef is provided to allow for owning a reference to a Module, and will auto-delete the held module on destruction. PiperOrigin-RevId: 256196193
* NFC: Refactor Function to be value typed.River Riddle2019-07-011-15/+16
| | | | | | Move the data members out of Function and into a new impl storage class 'FunctionStorage'. This allows for Function to become value typed, which will greatly simplify the transition of Function to FuncOp(given that FuncOp is also value typed). PiperOrigin-RevId: 255983022
* NVVM target: emit nvvm.annotations for kernel functionsAlex Zinenko2019-06-251-3/+27
| | | | | | | | | | 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
* Fix translation of NVVM special registers to intrinsics.Stephan Herhut2019-06-011-3/+3
| | | | | | | | | | The original implementation did not map the return value of the intrinsics call to the result value of the special register op. Uses of the result value hence hit a nullpointer. -- PiperOrigin-RevId: 250255436
* Remove unnecessary C++ specifier in CPP files. NFC.Jacques Pienaar2019-05-201-1/+1
| | | | | | | | These are only required in .h files to disambiguate between C and C++ header files. -- PiperOrigin-RevId: 248219135
* Add transformation of the NVVM dialect to an LLVM module. Only handlesStephan Herhut2019-05-061-0/+84
the generation of intrinsics out of NVVM index ops for now. -- PiperOrigin-RevId: 245933152
OpenPOWER on IntegriCloud