| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[MLIR][Affine] Fix affine loop permute validity check for iter_args loops (#155077) We cannot guarantee the validity of the interchange if the loops have iter_args, since the dependence analysis does not take them into account. Conservatively return false in such cases. Add an option to check permutation validity in test-loop-permutation pass to test this change. Signed-off-by: Prathamesh Tagore <prathamesh+1@polymagelabs.com> | 11 个月前 | |
[mlir][NFC] update mlir/Dialect create APIs (28/n) (#150641) See https://github.com/llvm/llvm-project/pull/147168 for more info. | 1 年前 | |
[MLIR][NFC] Retire let constructor for passes in Conversion directory (part1) (#127403) let constructor is deprecated since the table gen backend emits most of the glue logic to build a pass. This PR retires the td method for most (I need another pass) passes in the Conversion directory. | 1 年前 | |
[mlir][bufferization] Test tensor encoding -> memref layout conversion (#161166) Support custom types (4/N): test that it is possible to customize memref layout specification for custom operations and function boundaries. This is purely a test setup (no API modifications) to ensure users are able to pass information from tensors to memrefs within bufferization process. To achieve this, a test pass is required (since bufferization options have to be set manually). As there is already a --test-one-shot-module-bufferize pass present, it is extended for the purpose. | 9 个月前 | |
Reapply "[mlir] Link libraries that aren't included in libMLIR to libMLIR" (#123910) Use mlir_target_link_libraries() to link dependencies of libraries that are not included in libMLIR, to ensure that they link to the dylib when they are used in Flang. Otherwise, they implicitly pull in all their static dependencies, effectively causing Flang binaries to simultaneously link to the dylib and to static libraries, which is never a good idea. I have only covered the libraries that are used by Flang. If you wish, I can extend this approach to all non-libMLIR libraries in MLIR, making MLIR itself also link to the dylib consistently. [v3 with more -DBUILD_SHARED_LIBS=ON fixes] | 1 年前 | |
[MLIR][LLVM][DLTI] Handle data layout token 'n32:64' (#141299) | 1 年前 | |
[mlir][NFC] update mlir/Dialect create APIs (28/n) (#150641) See https://github.com/llvm/llvm-project/pull/147168 for more info. | 1 年前 | |
[MLIR][GPU][XeVM] Add XeVM target and XeVM dialect integration tests. (#148286) As part of XeVM dialect upsteaming, covers remaining parts required for XeVM dialect integration and testing. It has two high level components - XeVM target and serialization support - XeVM dialect integration tests using level zero runtime Co-Authored-by: Artem Kroviakov <artem.kroviakov@intel.com> | 11 个月前 | |
[mlir][LLVM] ArithToLLVM: Add 1:N support for arith.select lowering (#153944) Add 1:N support for the arith.select lowering. Only cases where the entire true/false value is selected are supported. | 11 个月前 | |
[MLIR] Apply clang-tidy fixes for readability-simplify-boolean-expr in TestLinalgTransforms.cpp (NFC) | 9 个月前 | |
[mlir][math] Add clampf and clean math ExpandOps API (#151153) This patch adds the clampf operation to the math dialect. The semantics op are defined as: clampf(x, min_v, max_v) = max(min(x, min_v), max_v) The reasoning behind adding this operation is that some GPU vendors offer specialized intrinsics for this operation, or subsets of this operation. For example, [__saturatef](https://docs.nvidia.com/cuda/cuda-math-api/cuda_math_api/group__CUDA__MATH__INTRINSIC__SINGLE.html#group__cuda__math__intrinsic__single_1ga2c84f08e0db7117a14509d21c3aec04e) in NVIDIA GPUs, or __builtin_amdgcn_fmed3f in AMD GPUs. This patch also removes test-expand-math in favor of math-expand-ops. Finally, it removes individual expansion population API calls like populateExpandCoshPattern in favor of: C++ void populateExpansionPatterns(RewritePatternSet &patterns, ArrayRef<StringRef> opMnemonics = {}); | 10 个月前 | |
[MLIR] Make 1-D memref flattening a prerequisite for vector narrow type emulation (#157771) Addresses: https://github.com/llvm/llvm-project/issues/115653 We already have utilities to flatten memrefs into 1-D. This change makes memref flattening a prerequisite for vector narrow type emulation, ensuring that emulation patterns only need to handle 1-D scenarios. | 10 个月前 | |
Reapply "[mlir] Link libraries that aren't included in libMLIR to libMLIR" (#123910) Use mlir_target_link_libraries() to link dependencies of libraries that are not included in libMLIR, to ensure that they link to the dylib when they are used in Flang. Otherwise, they implicitly pull in all their static dependencies, effectively causing Flang binaries to simultaneously link to the dylib and to static libraries, which is never a good idea. I have only covered the libraries that are used by Flang. If you wish, I can extend this approach to all non-libMLIR libraries in MLIR, making MLIR itself also link to the dylib consistently. [v3 with more -DBUILD_SHARED_LIBS=ON fixes] | 1 年前 | |
[acc] Expand OpenACCSupport to provide getRecipeName and emitNYI (#165628) Extends OpenACCSupport utilities to include recipe name generation and error reporting for unsupported features, providing foundation for variable privatization handling. Changes: - Add RecipeKind enum (private, firstprivate, reduction) for APIs that request a specific kind of recipe - Add getRecipeName() API to OpenACCSupport and OpenACCUtils that generates recipe names from types (e.g., "privatization_memref_5x10xf32_") - Add emitNYI() API to OpenACCSupport for graceful handling of not-yet-implemented cases - Generalize MemRefPointerLikeModel template to support UnrankedMemRefType - Add unit tests and integration tests for new APIs | 8 个月前 | |
[mlir][scf] Add parallelLoopUnrollByFactors() (#164958) - In the SCF Utils, add the parallelLoopUnrollByFactors() function to unroll scf::ParallelOp loops according to the specified unroll factors - Add a test pass "TestParallelLoopUnrolling" and the related LIT test - Expose mlir::parallelLoopUnrollByFactors(), mlir::generateUnrolledLoop(), and mlir::scf::computeUbMinusLb() functions in the mlir/Dialect/SCF/Utils/Utils.h and /IR/SCF.h headers to make them available to other passes. - In mlir::generateUnrolledLoop(), add an optional IRMapping *clonedToSrcOpsMap argument to map the new cloned operations to their original ones. In the function body, change the default AnnotateFn type to static const to silence potential warnings about dangling references when a function_ref is assigned to a variable with automatic storage. Signed-off-by: Fabrizio Indirli <Fabrizio.Indirli@arm.com> | 8 个月前 | |
[mlir][spirv] Add support for SPV_ARM_graph extension - part 1 (#151934) This is the first patch to add support for the SPV_ARM_graph SPIR-V extension to MLIR’s SPIR-V dialect. The extension introduces a new Graph abstraction for expressing dataflow computations over full resources. The part 1 implementation includes: - A new GraphType, modeled similarly to FunctionType, for typed graph signatures. - New operations in the spirv.arm namespace: - spirv.arm.Graph - spirv.arm.GraphEntryPoint - spirv.arm.GraphConstant - spirv.arm.GraphOutput - Verifier and VCE updates to properly gate usage under SPV_ARM_graph. - Tests covering parsing and verification. Graphs currently support only SPV_ARM_tensors, but are designed to generalize to other resource types, such as images. Spec: KhronosGroup/SPIRV-Registry#346 RFC: https://discourse.llvm.org/t/rfc-add-support-for-spv-arm-graph-extension-in-mlir-spir-v-dialect/86947 --------- Signed-off-by: Davide Grohmann <davide.grohmann@arm.com> | 10 个月前 | |
Reapply "[mlir] Link libraries that aren't included in libMLIR to libMLIR" (#123910) Use mlir_target_link_libraries() to link dependencies of libraries that are not included in libMLIR, to ensure that they link to the dylib when they are used in Flang. Otherwise, they implicitly pull in all their static dependencies, effectively causing Flang binaries to simultaneously link to the dylib and to static libraries, which is never a good idea. I have only covered the libraries that are used by Flang. If you wish, I can extend this approach to all non-libMLIR libraries in MLIR, making MLIR itself also link to the dylib consistently. [v3 with more -DBUILD_SHARED_LIBS=ON fixes] | 1 年前 | |
[mlir][NFC] update mlir create APIs (34/n) (#150660) See https://github.com/llvm/llvm-project/pull/147168 for more info. | 1 年前 | |
[mlir][NFC] update mlir/Dialect create APIs (28/n) (#150641) See https://github.com/llvm/llvm-project/pull/147168 for more info. | 1 年前 | |
[mlir][bufferization] Refine tensor-buffer compatibility checks (#167705) Generally, to_tensor and to_buffer already perform sufficient verification. However, there are some unnecessarily strict constraints: * builtin tensor requires its buffer counterpart to always be memref * to_buffer on ranked tensor requires to always return memref These checks are assertions (i.e. preconditions), however, they actually prevent an apparently useful bufferization where builtin tensors could become custom buffers. Lift these assertions, maintaining the verification procedure unchanged, to allow builtin -> custom bufferizations at operation boundary level. | 8 个月前 | |
Reapply "[mlir] Link libraries that aren't included in libMLIR to libMLIR" (#123910) Use mlir_target_link_libraries() to link dependencies of libraries that are not included in libMLIR, to ensure that they link to the dylib when they are used in Flang. Otherwise, they implicitly pull in all their static dependencies, effectively causing Flang binaries to simultaneously link to the dylib and to static libraries, which is never a good idea. I have only covered the libraries that are used by Flang. If you wish, I can extend this approach to all non-libMLIR libraries in MLIR, making MLIR itself also link to the dylib consistently. [v3 with more -DBUILD_SHARED_LIBS=ON fixes] | 1 年前 | |
[mlir] [irdl] Add support for regions in irdl-to-cpp (#158540) Fixes https://github.com/llvm/llvm-project/issues/158034 For the input mlir irdl.dialect @conditional_dialect { // A conditional operation with regions irdl.operation @conditional { // Create region constraints %r0 = irdl.region // Unconstrained region %r1 = irdl.region() // Region with no entry block arguments %v0 = irdl.any %r2 = irdl.region(%v0) // Region with one i1 entry block argument irdl.regions(cond: %r2, then: %r0, else: %r1) } } This produces the following cpp: https://gist.github.com/j2kun/d2095f108efbd8d403576d5c460e0c00 Summary of changes: - The op class and adaptor get named accessors to the regions Region &get<RegionName>() and getRegions() - The op now gets OpTrait::NRegions<3> and OpInvariants to trigger the region verification - Support for region block argument constraints is added, but not working for all constraints until codegen for irdl.is is added (filed https://github.com/llvm/llvm-project/issues/161018 and left a TODO). - Helper functions for the individual verification steps are added, following mlir-tblgen's format (in the above gist, __mlir_irdl_local_region_constraint_ConditionalOp_cond and similar), and verifyInvariantsImpl that calls them. - Regions are added in the builder ## Questions for the reviewer ### What is the "correct" interface for verification? I used mlir-tblgen on an analogous version of the example ConditionalOp in this PR, and I see an ::mlir::OpTrait::OpInvariants trait as well as cpp ::llvm::LogicalResult ConditionalOp::verifyInvariantsImpl() { { unsigned index = 0; (void)index; for (auto ®ion : ::llvm::MutableArrayRef((*this)->getRegion(0))) if (::mlir::failed(__mlir_ods_local_region_constraint_test1(*this, region, "cond", index++))) return ::mlir::failure(); for (auto ®ion : ::llvm::MutableArrayRef((*this)->getRegion(1))) if (::mlir::failed(__mlir_ods_local_region_constraint_test1(*this, region, "then", index++))) return ::mlir::failure(); for (auto ®ion : ::llvm::MutableArrayRef((*this)->getRegion(2))) if (::mlir::failed(__mlir_ods_local_region_constraint_test1(*this, region, "else", index++))) return ::mlir::failure(); } return ::mlir::success(); } ::llvm::LogicalResult ConditionalOp::verifyInvariants() { if(::mlir::succeeded(verifyInvariantsImpl()) && ::mlir::succeeded(verify())) return ::mlir::success(); return ::mlir::failure(); } However, OpInvariants only seems to need verifyInvariantsImpl, so it's not clear to me what is the purpose of the verifyInvariants function, or, if I leave out verifyInvariants, whether I need to call verify() in my implementation of verifyInvariantsImpl. In this PR, I omitted verifyInvariants and generated verifyInvariantsImpl. ### Is testing sufficient? I am not certain I implemented the builders properly, and it's unclear to me to what extent the existing tests check this (which look like they compile the generated cpp, but don't actually use it). Did I omit some standard function or overload? --------- Co-authored-by: Jeremy Kun <j2kun@users.noreply.github.com> | 9 个月前 | |
[mlir][tosa] Convert TOSA enumerations from StringBasedAttr to Tosa_I32EnumAttr (#152856) Fixes #152129 Use Tosa_I32EnumAttr instead of StringBasedAttr to represent Tosa enumerations. This PR replaces StringBasedAttr with Tosa_I32EnumAttr to represent Tosa enumerations as per the specification. The intent is to make the IR and C++ APIs more type-safe and prevent fragile string comparisons in passes. Enumerations rewritten are: - Tosa_ResizeTypeAttr - Tosa_NanPropagationAttr - Tosa_RoundingTypeAttr **BREAKING CHANGE**: This commit changes attribute assembly and the C++ API surface for the listed attributes. Code that previously used StringAttr for these fields must now be updated to use the new enum representation. In .mlir files, replace string literals with the enum assembly (e.g. mode = #tosa.resize_type<BILINEAR>). In C++, update call sites to either pass the generated enum (e.g. ::mlir::tosa::RoundingType::SINGLE_ROUND) into builder overloads or construct the typed attribute with tosa::RoundingTypeAttr::get(context, /*enum*/) and pass that. | 10 个月前 | |
[MLIR] Apply clang-tidy fixes for bugprone-argument-comment in TestTransformDialectExtension.cpp (NFC) | 8 个月前 | |
[MLIR][Vector] Add unroll pattern for vector.shape_cast (#167738) This PR adds pattern for unrolling shape_cast given a targetShape. This PR is a follow up of #164010 which was very general and was using inserts and extracts on each element (which is also LowerVectorShapeCast.cpp is doing). After doing some more research on use cases, we (me and @Jianhui-Li ) realized that the previous version in #164010 is unnecessarily generic and doesn't fit our performance needs. Our use case requires that targetShape is contiguous in both source and result vector. This pattern only applies when contiguous slices can be extracted from the source vector and inserted into the result vector such that each slice remains in vector form with targetShape (and not decompose to scalars). In these cases, the unrolling proceeds as: vector.extract_strided_slice -> vector.shape_cast (on the slice unrolled) -> vector.insert_strided_slice | 8 个月前 | |
[MLIR][XeGPU] Matrix load/store subgroup distribution (#165008) | 8 个月前 | |
[mlir][acc] Extend PointerLikeType to provide alloc, dealloc, copy (#162328) A variable in an acc data clause operation must have a type that implements either PointerLikeType or a MappableType interface. These interfaces provide the contract that allows acc dialect and its transform passes to interact with a source dialect. One of these requirements is ability to generate code that creates memory for a private copy and ability to initialize that copy from another variable. Thus, update the PointerLikeType API to provide the means to create allocation, deallocation, and copy. This will be used as a way to fill in privatization and firstprivatization recipes. This new API was implemented for memref along with testing to exercise it via the implementation of a testing pass. | 9 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 9 个月前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 9 个月前 |