| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[mlir] Use free op create functions. NFC. (#157374) The builder create methods are deprecated: https://mlir.llvm.org/deprecation/. See https://discourse.llvm.org/t/psa-opty-create-now-with-100-more-tab-complete/87339. | 11 个月前 | |
[mlir] Add FP software implementation lowering pass: arith-to-apfloat (#167848) Reland pass and fix linker errors. --------- Co-authored-by: Maksim Levental <maksim.levental@gmail.com> | 9 个月前 | |
[mlir][LLVM] Resync memory effect attribute with LLVM IR (#168568) - Add missing locations, namely 'ErrnoMem', 'TargetMem0', and 'TargetMem1'. | 8 个月前 | |
Fix side effects for LLVM integer operations (udiv, sdiv) incorrectly marked as Pure (#166648) This MR modifies side effect traits of some integer arithmetic operations in the LLVM dialect. Prior to this MR, the LLVM dialect sdiv and udiv operations were marked as Pure through tblgen inheritance of the LLVM_ArithmeticOpBase class. The Pure trait allowed incorrect hoisting of sdiv/udiv operations by the loop-independent-code-motion pass. This MR modifies the sdiv and udiv LLVM operations to have traits and code motion behavior identical to their counterparts in the arith dialect, which were established by the commit/review below. https://github.com/llvm/llvm-project/commit/ed39825be48805b174d3177f1d8d41ed84784d18 https://reviews.llvm.org/D137814 | 9 个月前 | |
[MLIR][LLVM] Add bytecode support for several attributes (#162577) For a total of 20 attributes, 18 debug information related + 2 regular ones (loop and alias_scope). Quick background on how this work: if a given attribute isn't supported, by default its textual form is dumped into the bytecode. In order to get proper encoding, an attribute needs a tablegen description of it and its element. There's an additional rule here: if an attribute is only used by another attribute, it's user need also to have an encoding in order for it to be encoded. (e.g. DICompileUnitAttr only gets encoded while in DISubprogramAttr if the later also has an encoded form), otherwise text is used. For this reason, this PR does a bunch at the same time, otherwise there isn't really much to test (easy to break it down if needed though). The PR is tested against some of our internal apps, successfully round-tripping around 14Gb of llvm dialect text. Some interesting findings include a 800K mlir textual file that used to become 1.2G in bytecode format - now down to 100K due to proper encoding of debug info attributes. In the future we should find a way to merge this together in the attribute definitions (perhaps autogenerate the entries from LLVM attribute descriptions), seems like we can benefit from the boilerplate. It's not clear yet how to solve some of the tablegen issues; some fields require manual translation of flag values using LocalVar, others require custom getters, etc. Ideas on that front are welcome. A next natural step here is to add type support, LLVM structs can also lead to non-neglible disk footprint. | 10 个月前 | |
[MLIR][LLVM] Add bytecode support for several attributes (#162577) For a total of 20 attributes, 18 debug information related + 2 regular ones (loop and alias_scope). Quick background on how this work: if a given attribute isn't supported, by default its textual form is dumped into the bytecode. In order to get proper encoding, an attribute needs a tablegen description of it and its element. There's an additional rule here: if an attribute is only used by another attribute, it's user need also to have an encoding in order for it to be encoded. (e.g. DICompileUnitAttr only gets encoded while in DISubprogramAttr if the later also has an encoded form), otherwise text is used. For this reason, this PR does a bunch at the same time, otherwise there isn't really much to test (easy to break it down if needed though). The PR is tested against some of our internal apps, successfully round-tripping around 14Gb of llvm dialect text. Some interesting findings include a 800K mlir textual file that used to become 1.2G in bytecode format - now down to 100K due to proper encoding of debug info attributes. In the future we should find a way to merge this together in the attribute definitions (perhaps autogenerate the entries from LLVM attribute descriptions), seems like we can benefit from the boilerplate. It's not clear yet how to solve some of the tablegen issues; some fields require manual translation of flag values using LocalVar, others require custom getters, etc. Ideas on that front are welcome. A next natural step here is to add type support, LLVM structs can also lead to non-neglible disk footprint. | 10 个月前 | |
[llvm] Remove unused includes (NFC) (#148342) These are identified by misc-include-cleaner. I've filtered out those that break builds. Also, I'm staying away from llvm-config.h, config.h, and Compiler.h, which likely cause platform- or compiler-specific build failures. | 1 年前 | |
[mlir] Simplify Default cases in type switches. NFC. (#165767) Use default values instead of lambdas when possible. std::nullopt and nullptr can be used now because of https://github.com/llvm/llvm-project/pull/165724. | 9 个月前 | |
[mlir] Simplify unreachable type switch cases. NFC. (#162032) Use DefaultUnreachable from https://github.com/llvm/llvm-project/pull/161970. | 10 个月前 | |
[MLIR][LLVM] Support named barrier as a global variable type in llvm dialect (#169194) Enables amdgcn.named.barrier target extension type as a global variable type in MLIR. | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
[MLIR][NVVM] Add NVVMRequiresSM op traits (#126886) Motivation: Currently, the NVVMOps are not verified against the supported SM architectures. This can manifest as an ISel failure in the NVPTX LLVM backend during CodeGen to PTX ISA. This PR addresses this issue by adding verifier checks for Target-SM architectures in the NVVM Dialect itself, thereby catching the errors early on. Summary: * Parametric traits named NVVMRequiresSM and NVVMRequiresSMa are added to facilitate the version checks for typical and arch-accelerated versions respectively. * These traits can be attached to any NVVM Op to enable the checks for the particular Op. (example shown below) * An attribute interface called named TargetAttrVerifyInterface is added to the GPU dialect which any target attribute seeking to perform target-verification on the module can implement. * The checks are performed by the NVVMTargetAttr (implementing the TargetAttrVerifyInterface interface) when called from the GPU module verifier where it walks through the module and performs the checks for Ops with the NVVMRequiresSM traits. * A few Ops in NVVMOps.td have been updated to serve as examples. Example Usage: def NVVM_ReduxOp : NVVM_Op<"redux.sync"> {...} ----> def NVVM_ReduxOp : NVVM_Op<"redux.sync", [NVVMRequiresSM<80>]> {...} def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned"> {...} ----> def NVVM_WgmmaFenceAlignedOp : NVVM_Op<"wgmma.fence.aligned", [NVVMRequiresSMa<[90]>]> {...} --------- Co-authored-by: Guray Ozen <guray.ozen@gmail.com> | 1 年前 | |
[mlir][rocdl] Add rocdl inlining interface (#163058) All rocdl ops should be safe to inline. | 10 个月前 | |
[mlir] Migrate away from ArrayRef(std::nullopt) (NFC) (#144989) ArrayRef has a constructor that accepts std::nullopt. This constructor dates back to the days when we still had llvm::Optional. Since the use of std::nullopt outside the context of std::optional is kind of abuse and not intuitive to new comers, I would like to move away from the constructor and eventually remove it. This patch takes care of the mlir side of the migration, starting with straightforward places where I see ArrayRef or ValueRange nearby. Note that ValueRange has a constructor that forwards arguments to an ArrayRef constructor. | 1 年前 | |
[mlir][LLVM] Remove llvm deps from the LLVM dialect (#150692) This patch removes spurious includes of llvm/IR files, and unnecessary link components in the LLVM dialect. The only major dependencies still coming from LLVM are llvm::DataLayout, which is used by verifyDataLayoutString and some dwarf symbols in some attributes. Both of them should likely be removed in the future. Finally, I also removed one constructor from LLVM::AssumeOp that used [OperandBundleDefT](https://llvm.org/doxygen/classllvm_1_1OperandBundleDefT.html) without good reason and introduced a header unnecessarily. | 1 年前 | |
[MLIR][XeVM] blockload and blockstore ops should use scalar types (#161708) instead of single element vectors. XeVM type system does not support single element vectors. | 10 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 11 个月前 | ||
| 9 个月前 | ||
| 8 个月前 | ||
| 9 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 |