| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[mlir][nvgpu] Mark TMA descriptor as MemWriteAt in tma.async.store (#79427) | 2 年前 | |
[mlir][nvgpu] Add tma last dim bytes check (#153451) Add the check the number of bytes in the last dimension of Tma must be a multiple of 16. | 1 年前 | |
[mlir][NFC] Remove trailing whitespaces from *.td and *.mlir files. This is generated by running sed --in-place 's/[[:space:]]\+$//' mlir/**/*.td sed --in-place 's/[[:space:]]\+$//' mlir/**/*.mlir Reviewed By: rriddle, dcaballe Differential Revision: https://reviews.llvm.org/D138866 | 3 年前 | |
[mlir][NFC] Remove trailing whitespaces from *.td and *.mlir files. This is generated by running sed --in-place 's/[[:space:]]\+$//' mlir/**/*.td sed --in-place 's/[[:space:]]\+$//' mlir/**/*.mlir Reviewed By: rriddle, dcaballe Differential Revision: https://reviews.llvm.org/D138866 | 3 年前 | |
| 1 年前 | ||
[mlir][NVGPU] Add support for structured sparsity MMA variants This change adds a new NVGPU operation that targets the PTX mma.sp.sync instruction variants. A lowering to NVVM is provided using inline assembly. Reviewed By: ThomasRaoux, manishucsd Differential Revision: https://reviews.llvm.org/D137202 | 3 年前 | |
[mlir][nvgpu] Make phaseParity of mbarrier.try_wait i1 (#81460) Currently, phaseParity argument of nvgpu.mbarrier.try_wait.parity is index. This can cause a problem if it's passed any value different than 0 or 1. Because the PTX instruction only accepts even or odd phase. This PR makes phaseParity argument i1 to avoid misuse. Here is the information from PTX doc: The .parity variant of the instructions test for the completion of the phase indicated by the operand phaseParity, which is the integer parity of either the current phase or the immediately preceding phase of the mbarrier object. An even phase has integer parity 0 and an odd phase has integer parity of 1. So the valid values of phaseParity operand are 0 and 1. See for more information: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-mbarrier-test-wait-mbarrier-try-wait | 2 年前 | |
[mlir][Transforms] GreedyPatternRewriteDriver: Do not CSE constants during iterations (#75897) The GreedyPatternRewriteDriver tries to iteratively fold ops and apply rewrite patterns to ops. It has special handling for constants: they are CSE'd and sometimes moved to parent regions to allow for additional CSE'ing. This happens in OperationFolder. To allow for efficient CSE'ing, OperationFolder maintains an internal lookup data structure to find the existing constant ops with the same value for each IsolatedFromAbove region: c++ /// A mapping between an insertion region and the constants that have been /// created within it. DenseMap<Region *, ConstantMap> foldScopes; Rewrite patterns are allowed to modify operations. In particular, they may move operations (including constants) from one region to another one. Such an IR rewrite can make the above lookup data structure inconsistent. We encountered such a bug in a downstream project. This bug materialized in the form of an op that uses the result of a constant op from a different IsolatedFromAbove region (that is not accessible). This commit changes the behavior of the GreedyPatternRewriteDriver such that OperationFolder is used to CSE constants at the beginning of each iteration (as the worklist is populated), but no longer during an iteration. OperationFolder is no longer used after populating the worklist, so we do not have to care about inconsistent state in the OperationFolder due to IR rewrites. The GreedyPatternRewriteDriver now performs the op folding by itself instead of calling OperationFolder::tryToFold. This change changes the order of constant ops in test cases, but not the region in which they appear. All broken test cases were fixed by turning CHECK into CHECK-DAG. Alternatives considered: The state of OperationFolder could be partially invalidated with every notifyOperationModified notification. That is more fragile than the solution in this commit because incorrect rewriter API usage can lead to missing notifications and hard-to-debug IsolatedFromAbove violations. (It did not fix the above mention bug in a downstream project, which could be due to incorrect rewriter API usage or due to another conceptual problem that I missed.) Moreover, ops are frequently getting modified during a greedy pattern rewrite, so we would likely keep invalidating large parts of the state of OperationFolder over and over. Migration guide: Turn CHECK into CHECK-DAG in test cases. Constant ops are no longer folded during a greedy pattern rewrite. If you rely on folding (and rematerialization) of constant ops during a greedy pattern rewrite, turn the folder into a pattern. | 2 年前 | |
| 1 年前 | ||
[MLIR][SCF] Add support for pipelining dynamic loops (#74350) Support loops without static boundaries. Since the number of iteration is not known we need to predicate prologue and epilogue in case the number of iterations is smaller than the number of stages. This patch includes work from @chengjunlu | 2 年前 |