| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[mlir][amdgpu] Add lowerings for ScaledExtPacked816 (#168123) * Adds lowerings for amdgpy.scaled_ext_packed816 * updates verifiers | 8 个月前 | |
[mlir][Affine] Add nsw to lowering of AffineMulExpr. (#121535) Since index operations have no set bitwidth, it is ill-defined to use signed/unsigned wrapping behavior. The corollary to which is that it is always safe to add nsw/nuw to lowering of affine ops. Also add a folder to fold div(s|u)i (mul (a, v), v) -> a Signed-off-by: MaheshRavishankar <mravisha@amd.com> | 1 年前 | |
[mlir][ArithToAMDGPU] limit scaling truncf/extf support to gfx950 (#155431) The current chip guard fails to prevent scaling_extf/truncf patterns from being applied on gfx1100 which does not have scaling support. --------- Signed-off-by: Muzammiluddin Syed <muzasyed@amd.com> | 10 个月前 | |
[mlir][arith] Add support for sitofp, uitofp to ArithToAPFloat (#169284) Add support for arith.sitofp and arith.uitofp. | 8 个月前 | |
[mlir][ArmSME] Rename slice move operations to insert/extract_tile_slice (#106755) This renames: - arm_sme.move_tile_slice_to_vector to arm_sme.extract_tile_slice - arm_sme.move_vector_to_tile_slice to arm_sme.insert_tile_slice The new names are more consistent with the rest of MLIR and should be easier to understand. The current names (to me personally) are hard to parse and easy to mix up when skimming through code. Additionally, the syntax for insert_tile_slice has changed from: mlir %4 = arm_sme.insert_tile_slice %0, %1, %2 : vector<[16]xi8> into vector<[16]x[16]xi8> To: mlir %4 = arm_sme.insert_tile_slice %0, %1[%2] : vector<[16]xi8> into vector<[16]x[16]xi8> This is for consistency with extract_tile_slice, but also helps with readability as it makes it clear which operand is the index. | 1 年前 | |
[mlir][emitc] Isolate expressions from above (#155641) The expression op is currently not isolated from above. This served its original usage as an optional, translation-oriented op, but is becoming less convenient now that expressions appear earlier in the emitc compilation flow and are gaining use as components of other emitc ops. This patch therefore adds the isolated-from-above trait to expressions. Syntactically, the only change is in the expression's signature which now includes the values being used in the expression as arguments and their types. The region's argument's names shadow the used values to keep the def-use relations clear. | 10 个月前 | |
Add 'exact' flag to arith.shrui/shrsi/divsi/divui operations (#165923) This MR adds support for the exact flag to the arith.shrui/shrsi/divsi/divui operations. The semantics are identical to those of the LLVM dialect and the LLVM language reference. This MR also modifies the mechanism for converting arith dialect **attributes** to corresponding **properties** in the LLVM dialect. (As a specific example, the integer overflow flags nsw/nuw are **properties** in the LLVM dialect, as opposed to attributes.) Previously, attribute converter classes were required to have a specific method to support integer overflow flags: C++ template <typename SourceOp, typename TargetOp> class AttrConvertPassThrough { public: ... LLVM::IntegerOverflowFlags getOverflowFlags() const { return LLVM::IntegerOverflowFlags::none; } }; This method was required, even for arith source operations that did not use integer overflow flags (e.g. AttrConvertFastMathToLLVM). This MR modifies the interface required by arith dialect attribute converters to instead provide a (possibly NULL) properties attribute: C++ template <typename SourceOp, typename TargetOp> class AttrConvertPassThrough { public: ... Attribute getPropAttr() const { return {}; } }; For arith operations with attributes that map to LLVM dialect **properties**, the attribute converter can create a DictionaryAttr containing target properties and return that attribute from the attribute converter's getPropAttr() method. The arith attribute conversion framework will set the propertiesAttr of an OperationState, and the target operation's setPropertiesFromAttr() method will be invoked to set the properties when the target operation is created. The AttrConvertOverflowToLLVM class in this MR uses the new approach. | 8 个月前 | |
[mlir][spirv] Add pattern matching for arith.index_cast i1 to index for ArithToSPIRV (#155729) Currently, arith.index_cast gets converted to OpSConvert: https://github.com/llvm/llvm-project/blob/9bf5bf3baf3c7aec82cdd235c6a2fd57b4dd55ab/mlir/lib/Conversion/ArithToSPIRV/ArithToSPIRV.cpp#L1331 [OpSConvert requires its operands to be of integer type](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSConvert), which poses an issue for i1 since SPIRV distinguishes between booleans and integers. As a result, the following example doesn't get converted, leaving behind illegal ops: %0 = arith.index_cast %arg0 : i1 to index This PR adds additional logic to convert arith.index_casts to SPIRV dialect when casting from i1 to index. Converting index_casts from index to i1 is a part of https://github.com/llvm/llvm-project/pull/156031. | 10 个月前 | |
[AArch64][SME] Improve codegen for aarch64.sme.cnts* when not in streaming mode (#154761) Builtins for reading the streaming vector length are canonicalised to use the aarch64.sme.cntsd intrinisic and a multiply, i.e. - cntsb -> cntsd * 8 - cntsh -> cntsd * 4 - cntsw -> cntsd * 2 This patch also removes the LLVM intrinsics for cnts[b,h,w], and adds patterns to improve codegen when cntsd is multiplied by a constant. | 10 个月前 | |
[mlir][sme] Use signed comparison in ArmSMEToSCF (#156834) This change switches from ult to slt in the comparison generated by TileLoadOpWithMaskAndPadNonZeroConversion. From the updated test: %[[ROW_IS_ACTIVE:.*]] = arith.cmpi slt, %[[TILE_SLICE_INDEX]], %[[NUM_ROWS]] : index Here: - %[[TILE_SLICE_INDEX]] is always non-negative. - %[[NUM_ROWS]] represents the number of remaining rows. %[[NUM_ROWS]] is computed as: mlir %c-4 = arith.constant -4 : index %c-4_vscale = arith.muli %c-4, %vscale_11 : index %num_rows_remaining = arith.addi %num_rows_init, %c-4_vscale : index (inserted by the "arm-sme-vector-legalization" pass, see VectorLegalization.cpp). Because of the subtraction, %num_rows_remaining can be negative. Therefore, the comparison must be signed (slt) rather than unsigned (ult). Fixes: https://github.com/iree-org/iree/issues/21714 | 10 个月前 | |
[MLIR][AsyncToLLVM] Remove typed pointer support (#70731) This commit removes the support for lowering Async to LLVM dialect with typed pointers. Typed pointers have been deprecated for a while now and it's planned to soon remove them from the LLVM dialect. Related PSA: https://discourse.llvm.org/t/psa-removal-of-typed-pointers-from-the-llvm-dialect/74502 | 2 年前 | |
[mlir][bufferization] Unranked memref support for clone (#94757) bufferization.clone does not currently support lowering to memref for unranked memrefs. This interferes with bufferizing unranked tensors at boundaries where a clone operation is needed. func.func @foo(%input: memref<*xf32>, %shape: memref<?xindex>) -> memref<*xf32> { %reshape = memref.reshape %input(%shape) : (memref<*xf32>, memref<?xindex>) -> memref<*xf32> %copy = bufferization.clone %reshape : memref<*xf32> to memref<*xf32> return %copy : memref<*xf32> } Patterns such as that are possibly when bufferizing functions with input and output unranked tensors. The clone operation currently fails to legalize during the bufferization-to-memref conversion with unranked memrefs. This change modifies the conversion of bufferization.clone to memref to generate the runtime calculations and allocation to allow for cloning an unranked memref. | 2 年前 | |
[mlir][Transforms] Dialect Conversion Driver without Rollback (#151865) This commit improves the allowPatternRollback flag handling in the dialect conversion driver. Previously, this flag was used to merely detect cases that are incompatible with the new One-Shot Dialect Conversion driver. This commit implements the driver itself: when the flag is set to "false", all IR changes are materialized immediately, bypassing the IRRewrite and ConversionValueMapping infrastructure. A few selected test cases now run with both the old and the new driver. RFC: https://discourse.llvm.org/t/rfc-a-new-one-shot-dialect-conversion-driver/79083 | 11 个月前 | |
[mlir][complex] Convert complex.tan to libm ctan call (#78250) We can convert complex.tan op to [ctan/ctanf](https://sourceware.org/newlib/libm.html#ctan) function in libm in the complex to libm conversion. | 2 年前 | |
[MLIR] Add new complex.powi op (#158722) This PR adds a new complex.powi operation to MLIR's complex dialect for computing complex numbers raised to integer powers. Key changes include: - Addition of the new PowiOp operation definition in the Complex dialect - Integration with algebraic simplification passes for optimization - Support for conversion to ROCDL library calls - Updates to Flang frontend to generate the new operation This depends on #158642. | 10 个月前 | |
[mlir][spirv][complex] Support convert complex.constant Reviewed By: kuhar Differential Revision: https://reviews.llvm.org/D151622 | 3 年前 | |
| 9 个月前 | ||
[mlir][Transforms] Dialect Conversion Driver without Rollback (#151865) This commit improves the allowPatternRollback flag handling in the dialect conversion driver. Previously, this flag was used to merely detect cases that are incompatible with the new One-Shot Dialect Conversion driver. This commit implements the driver itself: when the flag is set to "false", all IR changes are materialized immediately, bypassing the IRRewrite and ConversionValueMapping infrastructure. A few selected test cases now run with both the old and the new driver. RFC: https://discourse.llvm.org/t/rfc-a-new-one-shot-dialect-conversion-driver/79083 | 11 个月前 | |
[mlir][cfg-to-scf] Fix invalid transformation when value is used in a subregion (#67544) The current loop-reduce-form transformation incorrectly assumes that any value that is used in a block that isn't in the set of loop blocks is a block outside the loop. This is correct for a pure CFG but is incorrect if operations with subregions are present. In that case, a use may be in a subregion of an operation part of the loop and incorrectly deemed outside the loop. This would later lead to transformations with code that does not verify. This PR fixes that issue by checking the transitive parent block that is in the same region as the loop rather than the immediate parent block. | 2 年前 | |
| 2 年前 | ||
[mlir][bufferize] Make buffer-results-to-out-params support only functions that are neither public nor extern (#162441) The callers of public or extern functions are unknown, so their function signatures cannot be changed. | 9 个月前 | |
[MLIR][Affine] Extend getVectorReductionOp to support xor/maxnumf/minnumf (#163310) This PR extends the getVectorReductionOp function, which is used by the affine vectorizer, to also recognize and support xor/maxnumf/minnumf reduction operations. | 8 个月前 | |
[mlir][EmitC] Add pass that combines all available emitc conversions (#117549) | 1 年前 | |
[mlir][LLVM] Resync memory effect attribute with LLVM IR (#168568) - Add missing locations, namely 'ErrnoMem', 'TargetMem0', and 'TargetMem1'. | 8 个月前 | |
[mlir][spirv] Add 8-bit float type emulation (#148811) 8-bit floats are not supported in SPIR-V. They are emulated as 8-bit integer during conversion. | 11 个月前 | |
[mlir][Vector] Move vector.mask canonicalization to folder (#140324) This MR moves the canonicalization that elides empty vector.mask ops to folders. | 1 年前 | |
[MLIR][GPU] Generalize gpu.printf op lowering to LLVM call pattern. (#164297) Existing pattern for lowering gpu.printf op to LLVM call uses fixed function name and calling convention. Those two should be exposed as pass option to allow supporting Intel Compute Runtime for GPU. Also adds gpu.printf op pattern to GPU to LLVMSPV pass. It may appear out of place, but integration test is added to XeVM integration test as that is the current best folder for testing with Intel Compute Runtime. Test should be moved in the future if a better test folder is added. | 9 个月前 | |
| 8 个月前 | ||
[mlir][GPU] Generalize gpu.printf to not need gpu.module (#161266) In order to make the gpu.printf => [various LLVM calls] passes less order-dependent and to allow downstreams that don't use gpu.module to use gpu.printf, allow the flowerings for such prints to target the nearest SymbolTable instead. | 10 个月前 | |
[mlir][gpu][spirv] Add conversion for gpu.subgroup_mma_elementwise mulf (#158832) gpu.subgroup_mma_elementwise supports mulf op type. Add conversion for it. | 10 个月前 | |
[mlir][Transforms] Dialect Conversion Driver without Rollback (#151865) This commit improves the allowPatternRollback flag handling in the dialect conversion driver. Previously, this flag was used to merely detect cases that are incompatible with the new One-Shot Dialect Conversion driver. This commit implements the driver itself: when the flag is set to "false", all IR changes are materialized immediately, bypassing the IRRewrite and ConversionValueMapping infrastructure. A few selected test cases now run with both the old and the new driver. RFC: https://discourse.llvm.org/t/rfc-a-new-one-shot-dialect-conversion-driver/79083 | 11 个月前 | |
[mlir][Transforms][NFC] Dialect Conversion: Resolve insertion point TODO (#95653) Remove a TODO in the dialect conversion code base when materializing unresolved conversions: // FIXME: Determine a suitable insertion location when there are multiple // inputs. The implementation used to select an insertion point as follows: - If the cast has exactly one operand: right after the definition of the SSA value. - Otherwise: right before the cast op. However, it is not necessary to change the insertion point. Unresolved materializations (UnrealizedConversionCastOp) are built during buildUnresolvedArgumentMaterialization or buildUnresolvedTargetMaterialization. In the former case, the op is inserted at the beginning of the block. In the latter case, only one operand is supported in the dialect conversion, and the op is inserted right after the definition of the SSA value. I.e., the UnrealizedConversionCastOp is already inserted at the right place and it is not necessary to change the insertion point for the resolved materialization op. Note: The IR change changes slightly because the unrealized_conversion_cast ops at the beginning of a block are no longer doubly-inverted (by setting the insertion to the beginning of the block when inserting the unrealized_conversion_cast and again when inserting the resolved conversion op). All affected test cases were fixed by using CHECK-DAG instead of CHECK. Also improve the quality of multiple test cases that did not check for the correct operands. Note: This commit is in preparation of decoupling the argument/source/target materialization logic of the type converter from the dialect conversion (to reduce its complexity and make that functionality usable from a new dialect conversion driver). | 2 年前 | |
[mlir] fix LLVM type converter for structs (#73231) Existing implementation of the LLVM type converter for LLVM structs containing incompatible types was attempting to change identifiers of the struct in case of name clash post-conversion (all identified structs have different names post-conversion since one cannot change the body of the struct once initialized). Beyond a trivial error of not updating the counter in renaming, this approach was broken for recursive structs that can't be made aware of the renaming and would use the pre-existing struct with clashing name instead. For example, given !llvm.struct<"_Converted.foo", (struct<"_Converted.foo">, f32)> the following type !llvm.struct<"foo", (struct<"foo", index>)> would incorrectly convert to !llvm.struct<"_Converted_1.foo", (struct<"_Converted.foo", (struct<"_Converted.foo">, f32)>)> Remove this incorrect renaming and simply refuse to convert types if it would lead to identifier clashes for structs with different bodies. Document the expectation that such generated names are reserved and must not be present in the input IR of the converter. If we ever actually need to use handle such cases, this can be achieved by temporarily renaming structs with reserved identifiers to an unreserved name and back in a pre/post-processing pass that does _not_ use the type conversion infra. | 2 年前 | |
[mlir][mpi] fixing in-place and 0d mpi.all_reduce (#134225) * inplace allreduce needs special MPI token MPI_IN_PLACE as send buffer * 0d tensors have no sizes/strides in LLVM memref struct | 1 年前 | |
[mlir][EmitC] Add MathToEmitC pass for math function lowering to EmitC (#113799) This commit introduces a new MathToEmitC conversion pass that lowers selected math operations from the Math dialect to the emitc.call_opaque operation in the EmitC dialect. **Supported Math Operations:** The following operations are converted: - math.floor -> emitc.call_opaque<"floor"> - math.round -> emitc.call_opaque<"round"> - math.exp -> emitc.call_opaque<"exp"> - math.cos -> emitc.call_opaque<"cos"> - math.sin -> emitc.call_opaque<"sin"> - math.acos -> emitc.call_opaque<"acos"> - math.asin -> emitc.call_opaque<"asin"> - math.atan2 -> emitc.call_opaque<"atan2"> - math.ceil -> emitc.call_opaque<"ceil"> - math.absf -> emitc.call_opaque<"fabs"> - math.powf -> emitc.call_opaque<"pow"> **Target Language Standards:** The pass supports targeting different language standards: - C99: Generates calls with suffixes (e.g., floorf, fabsf) for single-precision floats. - CPP11: Prepends std:: to functions (e.g., std::floor, std::fabs). **Design Decisions:** The pass uses emitc.call_opaque instead of emitc.call to better emulate C-style function overloading. emitc.call_opaque does not require a unique type signature, making it more suitable for operations like <math.h> functions that may be overloaded for different types. This design choice ensures compatibility with C/C++ conventions. | 1 年前 | |
| 1 年前 | ||
[MLIR] Add sincos op to math dialect (#160772) Now that sincos is a supported intrinsic in the LLVM dialect (#160561) we are able to add the corresponding operation in the math dialect and add conversion patterns for LLVM and NVVM. We have several benchmarks that use sine and cosine in hot-loops, and saving some calculations by performing them together can benefit performance. We would like to have a way to represent sincos in the math dialect. | 9 个月前 | |
[mlir][Vector] add vector.insert canonicalization pattern to convert a chain of insertions to vector.from_elements (#142944) ## Description This change introduces a new canonicalization pattern for the MLIR Vector dialect that optimizes chains of insertions. The optimization identifies when a vector is **completely** initialized through a series of vector.insert operations and replaces the entire chain with a single vector.from_elements operation. Please be aware that the new pattern **doesn't** work for poison vectors where only **some** elements are set, as MLIR doesn't support partial poison vectors for now. **New Pattern: InsertChainFullyInitialized** * Detects chains of vector.insert operations. * Validates that all insertions are at static positions, and all intermediate insertions have only one use. * Ensures the entire vector is **completely** initialized. * Replaces the entire chain with a single vector.from_elementts operation. **Refactored Helper Function** * Extracted calculateInsertPosition from foldDenseElementsAttrDestInsertOp to avoid code duplication. ## Example // Before: %v1 = vector.insert %c10, %v0[0] : i64 into vector<2xi64> %v2 = vector.insert %c20, %v1[1] : i64 into vector<2xi64> // After: %v2 = vector.from_elements %c10, %c20 : vector<2xi64> It also works for multidimensional vectors. // Before: %v1 = vector.insert %cv0, %v0[0] : vector<3xi64> into vector<2x3xi64> %v2 = vector.insert %cv1, %v1[1] : vector<3xi64> into vector<2x3xi64> // After: %0:3 = vector.to_elements %arg1 : vector<3xi64> %1:3 = vector.to_elements %arg2 : vector<3xi64> %v2 = vector.from_elements %0#0, %0#1, %0#2, %1#0, %1#1, %1#2 : vector<2x3xi64> --------- Co-authored-by: Yang Bai <yangb@nvidia.com> Co-authored-by: Andrzej Warzyński <andrzej.warzynski@gmail.com> | 11 个月前 | |
[MLIR][ROCDL] Add math.clampf -> rocdl.fmed3 conversion (#163520) Added Pattern for lowering Math::ClampFOp to ROCDL::FMED3. Also added chipet option to MathToRocdl pass to check for arch support ISA instructions Solves [#15072](https://github.com/llvm/llvm-project/issues/157052) Reapplies https://github.com/llvm/llvm-project/pull/160100 Un-reverts the merged https://github.com/llvm/llvm-project/pull/163259, and fixes the error. --------- Signed-off-by: Keshav Vinayak Jha <keshavvinayakjha@gmail.com> | 9 个月前 | |
[MLIR][SPIRV] Add spirv.IsFinite and lower math.{isfinite,isinf,isnan} to spirv. (#151552) This patch adds support for lowering several float classification ops from the Math dialect to the SPIR-V dialect. ### Highlights: - Introduced a new spirv.IsFinite operation corresponding to the SPIR-V OpIsFinite instruction. - Lowered math.isfinite, math.isinf, and math.isnan to SPIR-V using CheckedElementwiseOpPattern. - Added corresponding tests for op definition and conversion lowering. This addresses the discussion in: https://github.com/llvm/llvm-project/issues/150778 --- Let me know if any additional adjustments are needed! --------- Co-authored-by: Jakub Kuderski <kubakuderski@gmail.com> | 11 个月前 | |
| 1 年前 | ||
[MLIR][MathToXeVM] Remove requirement for ModuleOp op type for MathToXeVM (#163619) This dependency is not needed. This PR removes it and checks that MathToXeVM works and respects pass pipelines specified by the user. | 9 个月前 | |
[mlir][emitc] Fix ineffective tests (#168197) These tests were only checking the specialized prefix, leaving common code unchecked (and incorrect). Checked code was also not using patterns for SSA values. | 8 个月前 | |
[MLIR] Fix FileCheck annotations and add 32-bit coverage (#167149) Addresses #93154 by ensuring the test RUN/FILECHECK prefixes across mlir/test/Conversion/FuncToLLVM/ and related files match their expected outputs, and adds an index-bitwidth=32 variant of expand-then-convert-to-llvm.mlir that exercises @collapse_shape_dynamic_with_non_identity_layout through the same MemRef-to-LLVM pipeline so the 32-bit layout descriptors are verified. | 8 个月前 | |
[mlir][memref-to-spirv]: Remap Image Load Coordinates (#160495) When converting a memref.load from the image address space to a spirv.ImageFetch ensure that we correctly map the load indices to width, height and depth. The lowering currently assumes a linear image tiling, that is row-major memory layout. This allows us to support any memref layout that is a permutation of the dimensions, more complex layouts are not currently supported. Because the ordering of the dimensions in the vector passed to image fetch is the opposite to that in the memref directions a final reversal of the mapped dimensions is always required. --------- Signed-off-by: Jack Frankland <jack.frankland@arm.com> | 10 个月前 | |
[MLIR][NVVM] Update mbarrier Ops to use AnyTypeOf[] (3/3) (#167567) This is a follow-up of PR #165558 and #165993. This patch updates the remaining two Ops to use the AnyTypeOf[] construct, completing the migration for the mbarrier family of Ops. mbarrier.arrive.expect_tx mbarrier.try_wait.parity Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 8 个月前 | |
[MLIR][NVVM] Update mbarrier Ops to use AnyTypeOf[] (3/3) (#167567) This is a follow-up of PR #165558 and #165993. This patch updates the remaining two Ops to use the AnyTypeOf[] construct, completing the migration for the mbarrier family of Ops. mbarrier.arrive.expect_tx mbarrier.try_wait.parity Signed-off-by: Durgadoss R <durgadossr@nvidia.com> | 8 个月前 | |
[flang][acc] Update assembly formats to include asyncOnly, async, and wait (#140122) The async implementation is inconsistent in terms of the assembly format. While renaming UpdateOp's async to asyncOnly, this PR handles asyncOnly along with async operands in every operation. Regarding EnterDataOp and ExitDataOp, they do not accept device types; thus, the async and the wait clauses without values lead to the async and the wait attributes (not asyncOnly nor waitOnly). This PR also processes them with async and wait operands all together. | 1 年前 | |
[Flang][OpenMP][Dialect] Swap to using MLIR dialect enum to encode map flags (#164043) This PR shifts from using the LLVM OpenMP enumerator bit flags to an OpenMP dialect specific enumerator. This allows us to better represent map types that wouldn't be of interest to the LLVM backend and runtime in the dialect. Primarily things like ref_ptr/ref_ptee/ref_ptr_ptee/atach_none/attach_always/attach_auto which are of interest to the compiler for certrain transformations (primarily in the FIR transformation passes dealing with mapping), but the runtime has no need to know about them. It also means if another OpenMP implementation comes along they won't need to stick to the same bit flag system LLVM chose/do leg work to address it. | 9 个月前 | |
[mlir] Fix test expectation. | 2 年前 | |
[mlir][ptr] Extend ptr_add operation to support shaped operands (#156374) This patch extends ptr_add to work with shaped types with value semantics, both for the offsets and base. Concretely this patch makes the following changes: - Supports scalar-to-scalar, scalar-to-shaped, shaped-to-scalar, and shaped-to-shaped combinations - Adds InferTypeOpInterface for automatic result type deduction - Adds tests for LLVM IR translation with vector operands Example: mlir func.func @ptr_add_tensor_2d(%ptrs: tensor<4x8x!ptr.ptr<#ptr.generic_space>>, %offsets: tensor<4x8xindex>) -> tensor<4x8x!ptr.ptr<#ptr.generic_space>> { %res = ptr.ptr_add %ptrs, %offsets : tensor<4x8x!ptr.ptr<#ptr.generic_space>>, tensor<4x8xindex> %res1 = ptr.ptr_add nuw %ptrs, %offsets : tensor<4x8x!ptr.ptr<#ptr.generic_space>>, tensor<4x8xindex> return %res : tensor<4x8x!ptr.ptr<#ptr.generic_space>> } The motivation behind this patch is to lay the groundwork for enabling triton styled loads and stores, and their variants. --------- Co-authored-by: Mehdi Amini <joker.eph@gmail.com> | 10 个月前 | |
[mlir][Transforms] Fix crash in reconcile-unrealized-casts (#158298) The reconcile-unrealized-casts pass used to crash when the input contains circular chains of unrealized_conversion_cast ops. Furthermore, the reconcileUnrealizedCasts helper functions used to erase ops that were not passed via the castOps operand. Such ops are now preserved. That's why some integration tests had to be changed. Also avoid copying the set of all unresolved materializations in convertOperations. This commit is in preparation of turning RewriterBase::replaceOp into a non-virtual function. This is a re-upload of #158067, which was reverted due to CI failures. Note for LLVM integration: If you are seeing tests that are failing with error: LLVM Translation failed for operation: builtin.unrealized_conversion_cast, you may have to add the -reconcile-unrealized-casts pass to your pass pipeline. (Or switch to the -convert-to-llvm pass instead of combining the various -convert-*-to-llvm passes.) --------- Co-authored-by: Mehdi Amini <joker.eph@gmail.com> | 10 个月前 | |
[mlir][SCF] Add pass option to deactivate pattern rollback (#168481) Add a pass option to convert-scf-to-cf to deactivate pattern rollback for better performance. The lowering patterns from SCF->CF to benefit a lot from this feature because splitBlock is expensive in the rollback driver. | 8 个月前 | |
[mlir][emitc] Add emitc.do op to the dialect (#143008) This patch adds: - Emission of the corresponding ops in the CppEmitter - Conversion from the SCF dialect to the EmitC dialect for the ops - Corresponding tests | 9 个月前 | |
[MLIR][SCFToGPU] Guard operands before AffineApplyOp::create to avoid crash (#167959) This fixes a crash in SCF→GPU when building the per‑dim index for mapped scf.parallel. **Change**: - Map step/lb through cloningMap, then run ensureLaunchIndependent. - If either is still unavailable at launch scope, emit a match‑failure; otherwise build the affine.apply. **Why this is correct:** - Matches how the pass already handles launch bounds; avoids creating an op with invalid operands and replaces a segfault with a clear diagnostic. **Tests**: - Added two small regressions that lower to gpu.launch and exercise the affine.apply path. Fixes : #167654 Signed-off-by: Shashi Shankar <shashishankar1687@gmail.com> | 8 个月前 | |
[flang][OpenMP] Enable tiling (#143715) This patch enables tiling in flang. In MLIR tiling is handled by changing the the omp.loop_nest op to be able to represent both collapse and tiling, so the flang front-end will combine the nested constructs into a single MLIR op. The MLIR->LLVM-IR lowering of the LoopNestOp is enhanced to first do the tiling if present, then collapse. | 10 个月前 | |
[mlir][spirv] Rework type capability queries (#160113) * Fix infinite recursion with nested structs. * Drop ::getCapbilities function from derived types, so that there's only one entry point that queries type extensions. * Move all capability logic to a new helper class -- this way the ::getCapabilities functions can't diverge across concrete types and 'convenience types' like CompositeType. Fixes: #159963 | 10 个月前 | |
| 8 个月前 | ||
[mlir] Use arith max or min ops instead of cmp + select (#82178) I believe the semantics should be the same, but this saves 1 op and simplifies the code. For example, the following two instructions: %2 = cmp sgt %0, %1 %3 = select %2, %0, %1 Are equivalent to: %2 = maxsi %0 %1 | 2 年前 | |
[NFC][mlir][mesh,shard] Fixing misnomers in mesh dialect, renaming 'mesh' dialect to 'shard' (#150177) Dialect to 'shard' (discourse 87053) - dialect name mesh -> shard - (device) mesh -> (device) grid - spmdize -> partition A lot of diffs, but simple renames only. @tkarna @yaochengji | 1 年前 | |
| 1 年前 | ||
| 1 年前 | ||
[mlir][tosa] Avoid requirement to specify enum name for enum attributes (#155695) This commit updates the printer/parser of operations that use enum type attributes such that they aren't required to specify the full enum name e.g. rounding_mode = #tosa.rounding_mode<SINGLE_ROUND> becomes rounding_mode = SINGLE_ROUND. For compatibility, text with #tosa.rounding_mode<...> will still be accepted by the parser. --------- Signed-off-by: Tai Ly <tai.ly@arm.com> Co-authored-by: Annu Singh <annu4444.as@gmail.com> | 11 个月前 | |
[mlir][tosa] Support RescaleOp with dynamic extension in TosaToLinalg (#155967) The shift, multiplier, inputZp, and outputZp can be either constant or non-constant, depending on whether dynamic extension is enabled. When these values are non-constant, they are added as inputs to linalg::GenericOp, and corresponding affine maps are appended to the indexingMaps. The commit help to pass following Tosa conformance tests. rescale_22x20_i32_outi8_sc0_rmS_pc0_iu0_ou0_dyn rescale_31x18_i8_outi8_sc0_rmS_pc0_iu1_ou0_dyn rescale_20x19_i16_outi8_sc0_rmS_pc0_iu1_ou0_dyn | 9 个月前 | |
[tosa] Change VariableOp to align with spec (#142240) This fixes Tosa VariableOp to align with spec 1.0 - add var_shape attribute to store shape of variable type - change type attribute to store element type of variable type - add a builder so previous construction calls still work - fix up level check of rank to be on variable type instead of initial value which is optional - add level check of size for variable type - add lit tests for variable op's without initial values - add lit test for variable op with fixed rank but unknown dimension - add invalid lit test for variable op with unranked type Signed-off-by: Tai Ly <tai.ly@arm.com> | 1 年前 | |
[MLIR][TOSA] Update IfOp print/parse to support ranked condition tens… (#149791) …or and optional block arguments This change extends the TOSA cond_if operation's print and parse logic to handle the following: - The condition operand may now have any rank, as long as the total number of elements sums to 1. %1 = tosa.cond_if %0 : tensor<1x1x1xi1> -> tensor<4xf32> - The then and else regions can now include optional block arguments. The updated IR syntax reflects this: %1 = tosa.cond_if %0 (%arg2 = %arg0, %arg3 = %arg1) : tensor<i1> (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32> - Removed parentheses around single result types in the printed representation, aligning with the AsmPrinter conventions. Co-authored-by: Luke Hutton <luke.hutton@arm.com> | 1 年前 | |
[mlir][tosa] Fix indexing in TosaToTensor (#140906) Changed the indexing used in the extractOp from one that is intended for 0d tensors to one that is intended for 1d tensors. --------- Co-authored-by: Shay Kleiman <shay.kleiman@mobileye.com> | 1 年前 | |
[mlir][Transforms] Dialect Conversion Driver without Rollback (#151865) This commit improves the allowPatternRollback flag handling in the dialect conversion driver. Previously, this flag was used to merely detect cases that are incompatible with the new One-Shot Dialect Conversion driver. This commit implements the driver itself: when the flag is set to "false", all IR changes are materialized immediately, bypassing the IRRewrite and ConversionValueMapping infrastructure. A few selected test cases now run with both the old and the new driver. RFC: https://discourse.llvm.org/t/rfc-a-new-one-shot-dialect-conversion-driver/79083 | 11 个月前 | |
[mlir][spirv] Fix some issues related to converting ub.poison to SPIR-V (#125905) This is a follow-up to 5df62bdc9be9c258c5ac45c8093b71e23777fa0e. That commit should not have needed to make the vector.insert and vector.extract conversions to SPIR-V directly handle the static poison index case, as there is a fold from those to ub.poison, and a conversion pattern from ub.poison to spirv.Undef, however: - The ub.poison fold result could not be materialized by the vector dialect (fixed as of d13940ee263ff50b7a71e21424913cc0266bf9d4). - The conversion pattern wasn't being populated in VectorToSPIRVPass, which is used by the tests. This commit changes this. - The ub.poison to spirv.Undef pattern rejected non-scalar types, which prevented its use for vector results. It is unclear why this restriction existed; a remark in D156163 said this was to avoid converting "user types", but it is not obvious why these shouldn't be permitted (the SPIR-V specification allows OpUndef for all types except OpTypeVoid). This commit removes this restriction. With these fixed, this commit removes the redundant static poison index handling, and updates the tests. | 1 年前 | |
[mlir][amx] Direct AMX data transfers (#154114) Extends Vector to AMX conversion to attempt populating AMX tiles directly from memory. When possible, contraction producers and consumers are replaced by AMX tile data transfer operations. This shortens data path by skipping intermediate register loads and stores. | 11 个月前 | |
[MLIR][Vector] Remove vector.splat (#162167) vector.splat has been deprecated (user: please use the very similar vector.broadcast instead) with the last PR landing about 6 weeks ago. The discourse discussion is at https://discourse.llvm.org/t/rfc-mlir-vector-deprecate-then-remove-vector-splat/87143/1 The last PR was #152230 This PR completely removes vector.splat. In addition to removing vector.splat from VectorOps.td, it - Updates the few remaining places where vector::SplatOp is created (now vector::BroadcastOp is created) - Removes temporary patterns where vector.splat is replaced by vector.broadcast The only place 'vector.splat' appears is now the files https://github.com/llvm/llvm-project/blob/main/mlir/utils/tree-sitter-mlir/test/corpus/op.txt and https://github.com/llvm/llvm-project/blob/main/mlir/utils/tree-sitter-mlir/dialect/vector.js --------- Signed-off-by: James Newling <james.newling@gmail.com> | 9 个月前 | |
[mlir][gpu][vector] Lower Vector dialect to GPU for element-wise ops only (#159091) Current convertVectorToMMAOps starts from vector.contract and finds its dependencies as the targets to convert. In GPU dialect, we have gpu.subgroup_mma_elementwise operation. We should be able to lower element-wise operations to GPU MMA operations without vector.contract. This patch adds this case to the pattern. | 10 个月前 | |
[MLIR][Conversion] Vector to LLVM: Remove unneeded vector shuffle (#162946) if vector.broadcast source is a scalar and target is a single element 1D vector. | 9 个月前 | |
[mlir][vector] Folder: shape_cast(extract) -> extract (#146368) In a later PR more shape_cast ops will appear. Specifically, broadcasts that just prepend ones become shape_cast ops (i.e. volume preserving broadcasts are canonicalized to shape_casts). This PR ensures that broadcast-like shape_cast ops fold at least as well as broadcast ops. This is done by modifying patterns that target broadcast ops, to target 'broadcast-like' ops. No new patterns are added, the patterns that exist are just made to match on shape_casts where appropriate. This PR also includes minor code simplifications: use isBroadcastableTo to simplify ExtractOpFromBroadcast and simplify how broadcast dims are detected in foldExtractFromBroadcast. These are NFC. --------- Co-authored-by: Andrzej Warzyński <andrzej.warzynski@gmail.com> | 1 年前 | |
[mlir][spirv][vector] Use adaptor.getElements() in FromElements lowering. (#156972) Signed-off-by: hanhanW <hanhan0912@gmail.com> | 10 个月前 | |
[mlir][XeGPU][VectorToXeGPU] Use 'xegpu.load' to lower 1D 'vector.transfer_read' for PVC & BMG (#168910) The PR changes the TransferReadLowering to always use xegpu.load (and not xegpu.load_nd) for 1D cases as it has more developed interface (e.g. layouts capabilites). Signed-off-by: dchigarev <dmitry.chigarev@intel.com> | 8 个月前 | |
[MLIR][Conversion] XeGPU to XeVM: Use adaptor for getting base address from memref. (#168610) adaptor already lowers memref to base address. Conversion patterns should use it instead of generating code to get base address from memref. | 8 个月前 | |
[mlir][LLVM] Resync memory effect attribute with LLVM IR (#168568) - Add missing locations, namely 'ErrnoMem', 'TargetMem0', and 'TargetMem1'. | 8 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 8 个月前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 11 个月前 | ||
| 2 年前 | ||
| 10 个月前 | ||
| 3 年前 | ||
| 9 个月前 | ||
| 11 个月前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 9 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 11 个月前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 11 个月前 | ||
| 9 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 2 年前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 9 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 9 个月前 | ||
| 10 个月前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 |