| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
Switch member calls to isa/dyn_cast/cast/... to free function calls. (#89356) This change cleans up call sites. Next step is to mark the member functions deprecated. See https://mlir.llvm.org/deprecation and https://discourse.llvm.org/t/preferred-casting-style-going-forward. | 2 年前 | |
[MLIR][AMX] Remove llvm.ptr<i8> bitcasts from legalize for LLVM pass This commit removes the no longer required llvm.ptr<i8> bitcasts from AMX's legalize-for-llvm-export pass. 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][Affine] NFC. Expose affine loop tiling validity utility (#99459) Move the utility to check for the validity of tiling affine loop nests to affine loop utils and expose for users outside the loop tiling pass or downstream users. | 2 年前 | |
[MLIR] Add f8E4M3 IEEE 754 type (#97118) This PR adds f8E4M3 type to mlir. f8E4M3 type follows IEEE 754 convention c f8E4M3 (IEEE 754) - Exponent bias: 7 - Maximum stored exponent value: 14 (binary 1110) - Maximum unbiased exponent value: 14 - 7 = 7 - Minimum stored exponent value: 1 (binary 0001) - Minimum unbiased exponent value: 1 − 7 = −6 - Precision specifies the total number of bits used for the significand (mantisa), including implicit leading integer bit = 3 + 1 = 4 - Follows IEEE 754 conventions for representation of special values - Has Positive and Negative zero - Has Positive and Negative infinity - Has NaNs Additional details: - Max exp (unbiased): 7 - Min exp (unbiased): -6 - Infinities (+/-): S.1111.000 - Zeros (+/-): S.0000.000 - NaNs: S.1111.{001, 010, 011, 100, 101, 110, 111} - Max normal number: S.1110.111 = +/-2^(7) x (1 + 0.875) = +/-240 - Min normal number: S.0001.000 = +/-2^(-6) - Max subnormal number: S.0000.111 = +/-2^(-6) x 0.875 = +/-2^(-9) x 7 - Min subnormal number: S.0000.001 = +/-2^(-6) x 0.125 = +/-2^(-9) Related PRs: - [PR-97179](https://github.com/llvm/llvm-project/pull/97179) [APFloat] Add support for f8E4M3 IEEE 754 type | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
| 2 年前 | ||
[mlir][ArmSVE] Add arm_sve.psel operation (#95764) This adds a new operation for the SME/SVE2.1 psel instruction. This allows selecting a predicate based on a bit within another predicate, essentially allowing for 2-D predication. Informally, the semantics are: mlir %pd = arm_sve.psel %p1, %p2[%index] : vector<[4]xi1>, vector<[8]xi1> => if p2[index % num_elements(p2)] == 1: pd = p1 : type(p1) else: pd = all-false : type(p1) | 2 年前 | |
[mlir] Use OpBuilder::createBlock in op builders and patterns (#82770) When creating a new block in (conversion) rewrite patterns, OpBuilder::createBlock must be used. Otherwise, no notifyBlockInserted notification is sent to the listener. Note: The dialect conversion relies on listener notifications to keep track of IR modifications. Creating blocks without the builder API can lead to memory leaks during rollback. | 2 年前 | |
[mlir] Fix block merging (#97697) With this PR I am trying to address: https://github.com/llvm/llvm-project/issues/63230. What changed: - While merging identical blocks, don't add a block argument if it is "identical" to another block argument. I.e., if the two block arguments refer to the same Value. The operations operands in the block will point to the argument we already inserted. This needs to happen to all the arguments we pass to the different successors of the parent block - After merged the blocks, get rid of "unnecessary" arguments. I.e., if all the predecessors pass the same block argument, there is no need to pass it as an argument. - This last simplification clashed with BufferDeallocationSimplification. The reason, I think, is that the two simplifications are clashing. I.e., BufferDeallocationSimplification contains an analysis based on the block structure. If we simplify the block structure (by merging and/or dropping block arguments) the analysis is invalid . The solution I found is to do a more prudent simplification when running that pass. **Note**: this a rework of #96871 . I ran all the integration tests (-DMLIR_INCLUDE_INTEGRATION_TESTS=ON) and they passed. | 2 年前 | |
Switch member calls to isa/dyn_cast/cast/... to free function calls. (#89356) This change cleans up call sites. Next step is to mark the member functions deprecated. See https://mlir.llvm.org/deprecation and https://discourse.llvm.org/t/preferred-casting-style-going-forward. | 2 年前 | |
Bufferization with ControlFlow Asserts (#95868) Fixed incorrect bufferization interaction with cf.assert - reordered bufferization condition checking - fixed hasNeitherAllocateNorFreeSideEffect checking bug - implemented memory interface for cf.assert --------- Co-authored-by: McCowan Zhang <mccowan.z@ssi.samsung.com> | 2 年前 | |
DLTI: Simplifying getDevicePropertyValue API by returning Attribute type value (#96706) **Rationale** - With the current flexibility of supporting any type of value, we will need to offer type-specific APIs to fetch a value (e.g., getDevicePropertyValueAsInt for integer type, getDevicePropertyValueAsFloat for float type, etc.) A single type of value will eliminate this need. - Current flexibility can also lead to typing errors when a user fetches the value of a property using an API that is not consistent with the type of the value. **What is the change** For following system description, module attributes { dlti.target_system_spec = #dlti.target_system_spec< "CPU": #dlti.target_device_spec< #dlti.dl_entry<"max_vector_op_width", 64.0 : f32>>, "GPU": #dlti.target_device_spec< #dlti.dl_entry<"max_vector_op_width", 128 : ui32>> >} {} a user no longer needs to use getDevicePropertyValueAsInt for retrieving GPU's max_vector_op_width and getDevicePropertyValueAsFloat for retrieving CPU's max_vector_op_width. Instead it can be done with a uniform API of getDevicePropertyValue. | 2 年前 | |
[mlir][EmitC] Add member access ops (#98460) This adds an emitc.member and emitc.member_of_ptr operation for the corresponding member access operators. Furthermore, emitc.assign is adjusted to be used with the member access operators. | 2 年前 | |
[mlir][side effect] refactor(*): Include more precise side effects (#94213) This patch adds more precise side effects to the current ops with memory effects, allowing us to determine which OpOperand/OpResult/BlockArgument the operation reads or writes, rather than just recording the reading and writing of values. This allows for convenient use of precise side effects to achieve analysis and optimization. Related discussions: https://discourse.llvm.org/t/rfc-add-operandindex-to-sideeffect-instance/79243 | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
[mlir][intrange] Use nsw,nuw flags in inference (#92642) This patch includes the "no signed wrap" and "no unsigned wrap" flags, which can be used to annotate some Ops in the arith dialect and also in LLVMIR, in the integer range inference. The general approach is to use saturating arithmetic operations to infer bounds which are assumed to not wrap and use overflowing arithmetic operations in the normal case. If overflow is detected in the normal case, special handling makes sure that we don't underestimate the result range. | 2 年前 | |
[MLIR][LLVM] Always print variadic callee type (#99293) This commit updates the LLVM dialect CallOp and InvokeOp to always print the variadic callee type (previously callee type) if present. An additional verifier checks that only variadic calls have a non-null variadic callee type, and the builders are adapted accordingly to set the variadic callee type for variadic calls only. Finally, the CallOp and InvokeOp verifiers are strengthened to check that the variadic callee type matches the call argument and result types. The motivation of this change is that CallOp and InvokeOp don't have hidden state that is not pretty printed, but used during the export to LLVM IR. Previously, it could happen that a call looked correct in MLIR, but the return type changed after exporting to LLVM IR (since it has been taken from the hidden callee type attribute). After landing this change, this is not possible anymore since the variadic callee type is always printed if present. | 2 年前 | |
[mlir] [linalg] Add pattern to swap transpose with broadcast (#97063) Add a pattern that implement: transpose(broadcast(input)) -> broadcast(transpose(input)) | 2 年前 | |
[mlir] Make the ml_program dialect allow all of its operations to be inlined. (#85479) | 2 年前 | |
[mlir] Reland "Initial patch to add an MPI dialect" (#81975) This patch introduces the new MPI dialect into MLIR. The Message Passing Interface (MPI) is a widely-used standard for distributed programs to exchange data. This PR goes together with a talk later at today's LLVM Dev Meeting. This is just a first, small patch to get going and add the necessary base files, so that we can add more operations in further patches. Here's the documentation as generated by ninja mlir-doc: # 'mpi' Dialect This dialect models the Message Passing Interface (MPI), version 4.0. It is meant to serve as an interfacing dialect that is targeted by higher-level dialects. The MPI dialect itself can be lowered to multiple MPI implementations and hide differences in ABI. The dialect models the functions of the MPI specification as close to 1:1 as possible while preserving SSA value semantics where it makes sense, and uses memref types instead of bare pointers. This dialect is under active development, and while stability is an eventual goal, it is not guaranteed at this juncture. Given the early state, it is recommended to inquire further prior to using this dialect. For an in-depth documentation of the MPI library interface, please refer to official documentation such as the [OpenMPI online documentation](https://www.open-mpi.org/doc/current/). [TOC] ## Operation definition ### mpi.comm_rank (mpi::CommRankOp) _Get the current rank, equivalent to MPI_Comm_rank(MPI_COMM_WORLD, &rank)_ Syntax: `` operation ::= mpi.comm_rank attr-dict : type(results) ` Communicators other than MPI_COMM_WORLD are not supported for now. This operation can optionally return an !mpi.retval value that can be used to check for errors. #### Results: | Result | Description | | :----: | ----------- | | retval | MPI function call return value | rank | 32-bit signless integer ### mpi.error_class (mpi::ErrorClassOp) _Get the error class from an error code, equivalent to the MPI_Error_class function_ Syntax: ` operation ::= mpi.error_class $val attr-dict : type($val) ` MPI_Error_class maps return values from MPI calls to a set of well-known MPI error classes. #### Operands: | Operand | Description | | :-----: | ----------- | | val | MPI function call return value #### Results: | Result | Description | | :----: | ----------- | | errclass | MPI function call return value ### mpi.finalize (mpi::FinalizeOp) _Finalize the MPI library, equivalent to MPI_Finalize()_ Syntax: ` operation ::= mpi.finalize attr-dict (: type($retval)^)? ` This function cleans up the MPI state. Afterwards, no MPI methods may be invoked (excpet for MPI_Get_version, MPI_Initialized, and MPI_Finalized). Notably, MPI_Init cannot be called again in the same program. This operation can optionally return an !mpi.retval value that can be used to check for errors. #### Results: | Result | Description | | :----: | ----------- | | retval | MPI function call return value ### mpi.init (mpi::InitOp) _Initialize the MPI library, equivalent to MPI_Init(NULL, NULL)_ Syntax: ` operation ::= mpi.init attr-dict (: type($retval)^)? ` This operation must preceed most MPI calls (except for very few exceptions, please consult with the MPI specification on these). Passing &argc, &argv is not supported currently. This operation can optionally return an !mpi.retval value that can be used to check for errors. #### Results: | Result | Description | | :----: | ----------- | | retval | MPI function call return value ### mpi.recv (mpi::RecvOp) _Equivalent to MPI_Recv(ptr, size, dtype, dest, tag, MPI_COMM_WORLD, MPI_STATUS_IGNORE)_ Syntax: ` operation ::= mpi.recv ( $ref , $tag , $rank ) attr-dict : type($ref) , type($tag) , type($rank)(-> type($retval)^)? ` MPI_Recv performs a blocking receive of size elements of type dtype from rank dest. The tag value and communicator enables the library to determine the matching of multiple sends and receives between the same ranks. Communicators other than MPI_COMM_WORLD are not supprted for now. The MPI_Status is set to MPI_STATUS_IGNORE, as the status object is not yet ported to MLIR. This operation can optionally return an !mpi.retval value that can be used to check for errors. #### Operands: | Operand | Description | | :-----: | ----------- | | ref | memref of any type values | tag | 32-bit signless integer | rank | 32-bit signless integer #### Results: | Result | Description | | :----: | ----------- | | retval | MPI function call return value ### mpi.retval_check (mpi::RetvalCheckOp) _Check an MPI return value against an error class_ Syntax: ` operation ::= mpi.retval_check $val = $errclass attr-dict : type($res) ` This operation compares MPI status codes to known error class constants such as MPI_SUCCESS, or MPI_ERR_COMM. #### Attributes: <table> <tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr> <tr><td><code>errclass</code></td><td>::mlir::mpi::MPI_ErrorClassEnumAttr</td><td><details><summary>MPI error class name</summary>{{% markdown %}}Enum cases: * MPI_SUCCESS (MPI_SUCCESS) * MPI_ERR_ACCESS (MPI_ERR_ACCESS) * MPI_ERR_AMODE (MPI_ERR_AMODE) * MPI_ERR_ARG (MPI_ERR_ARG) * MPI_ERR_ASSERT (MPI_ERR_ASSERT) * MPI_ERR_BAD_FILE (MPI_ERR_BAD_FILE) * MPI_ERR_BASE (MPI_ERR_BASE) * MPI_ERR_BUFFER (MPI_ERR_BUFFER) * MPI_ERR_COMM (MPI_ERR_COMM) * MPI_ERR_CONVERSION (MPI_ERR_CONVERSION) * MPI_ERR_COUNT (MPI_ERR_COUNT) * MPI_ERR_DIMS (MPI_ERR_DIMS) * MPI_ERR_DISP (MPI_ERR_DISP) * MPI_ERR_DUP_DATAREP (MPI_ERR_DUP_DATAREP) * MPI_ERR_ERRHANDLER (MPI_ERR_ERRHANDLER) * MPI_ERR_FILE (MPI_ERR_FILE) * MPI_ERR_FILE_EXISTS (MPI_ERR_FILE_EXISTS) * MPI_ERR_FILE_IN_USE (MPI_ERR_FILE_IN_USE) * MPI_ERR_GROUP (MPI_ERR_GROUP) * MPI_ERR_INFO (MPI_ERR_INFO) * MPI_ERR_INFO_KEY (MPI_ERR_INFO_KEY) * MPI_ERR_INFO_NOKEY (MPI_ERR_INFO_NOKEY) * MPI_ERR_INFO_VALUE (MPI_ERR_INFO_VALUE) * MPI_ERR_IN_STATUS (MPI_ERR_IN_STATUS) * MPI_ERR_INTERN (MPI_ERR_INTERN) * MPI_ERR_IO (MPI_ERR_IO) * MPI_ERR_KEYVAL (MPI_ERR_KEYVAL) * MPI_ERR_LOCKTYPE (MPI_ERR_LOCKTYPE) * MPI_ERR_NAME (MPI_ERR_NAME) * MPI_ERR_NO_MEM (MPI_ERR_NO_MEM) * MPI_ERR_NO_SPACE (MPI_ERR_NO_SPACE) * MPI_ERR_NO_SUCH_FILE (MPI_ERR_NO_SUCH_FILE) * MPI_ERR_NOT_SAME (MPI_ERR_NOT_SAME) * MPI_ERR_OP (MPI_ERR_OP) * MPI_ERR_OTHER (MPI_ERR_OTHER) * MPI_ERR_PENDING (MPI_ERR_PENDING) * MPI_ERR_PORT (MPI_ERR_PORT) * MPI_ERR_PROC_ABORTED (MPI_ERR_PROC_ABORTED) * MPI_ERR_QUOTA (MPI_ERR_QUOTA) * MPI_ERR_RANK (MPI_ERR_RANK) * MPI_ERR_READ_ONLY (MPI_ERR_READ_ONLY) * MPI_ERR_REQUEST (MPI_ERR_REQUEST) * MPI_ERR_RMA_ATTACH (MPI_ERR_RMA_ATTACH) * MPI_ERR_RMA_CONFLICT (MPI_ERR_RMA_CONFLICT) * MPI_ERR_RMA_FLAVOR (MPI_ERR_RMA_FLAVOR) * MPI_ERR_RMA_RANGE (MPI_ERR_RMA_RANGE) * MPI_ERR_RMA_SHARED (MPI_ERR_RMA_SHARED) * MPI_ERR_RMA_SYNC (MPI_ERR_RMA_SYNC) * MPI_ERR_ROOT (MPI_ERR_ROOT) * MPI_ERR_SERVICE (MPI_ERR_SERVICE) * MPI_ERR_SESSION (MPI_ERR_SESSION) * MPI_ERR_SIZE (MPI_ERR_SIZE) * MPI_ERR_SPAWN (MPI_ERR_SPAWN) * MPI_ERR_TAG (MPI_ERR_TAG) * MPI_ERR_TOPOLOGY (MPI_ERR_TOPOLOGY) * MPI_ERR_TRUNCATE (MPI_ERR_TRUNCATE) * MPI_ERR_TYPE (MPI_ERR_TYPE) * MPI_ERR_UNKNOWN (MPI_ERR_UNKNOWN) * MPI_ERR_UNSUPPORTED_DATAREP (MPI_ERR_UNSUPPORTED_DATAREP) * MPI_ERR_UNSUPPORTED_OPERATION (MPI_ERR_UNSUPPORTED_OPERATION) * MPI_ERR_VALUE_TOO_LARGE (MPI_ERR_VALUE_TOO_LARGE) * MPI_ERR_WIN (MPI_ERR_WIN) * MPI_ERR_LASTCODE (MPI_ERR_LASTCODE){{% /markdown %}}</details></td></tr> </table> #### Operands: | Operand | Description | | :-----: | ----------- | | val | MPI function call return value #### Results: | Result | Description | | :----: | ----------- | | res | 1-bit signless integer ### mpi.send (mpi::SendOp) _Equivalent to MPI_Send(ptr, size, dtype, dest, tag, MPI_COMM_WORLD)_ Syntax: ` operation ::= mpi.send ( $ref , $tag , $rank ) attr-dict : type($ref) , type($tag) , type($rank)(-> type($retval)^)? ` MPI_Send performs a blocking send of size elements of type dtype to rank dest. The tag value and communicator enables the library to determine the matching of multiple sends and receives between the same ranks. Communicators other than MPI_COMM_WORLD are not supprted for now. This operation can optionally return an !mpi.retval value that can be used to check for errors. #### Operands: | Operand | Description | | :-----: | ----------- | | ref | memref of any type values | tag | 32-bit signless integer | rank | 32-bit signless integer #### Results: | Result | Description | | :----: | ----------- | | retval | MPI function call return value ## Attribute definition ### MPI_ErrorClassEnumAttr MPI error class name Syntax: #mpi.errclass< ::mlir::mpi::MPI_ErrorClassEnum # value > Enum cases: * MPI_SUCCESS (MPI_SUCCESS) * MPI_ERR_ACCESS (MPI_ERR_ACCESS) * MPI_ERR_AMODE (MPI_ERR_AMODE) * ... *all other MPI error codes* #### Parameters: | Parameter | C++ type | Description | | :-------: | :-------: | ----------- | | value | ::mlir::mpi::MPI_ErrorClassEnum | an enum of type MPI_ErrorClassEnum | ## Type definition ### RetvalType MPI function call return value Syntax: !mpi.retval This type represents a return value from an MPI function vall. This value can be MPI_SUCCESS, MPI_ERR_IN_STATUS, or any error code. This return value can be compared agains the known MPI error classes represented by #mpi.errclass using the mpi.retval_check` operation. | 2 年前 | |
[MLIR][Arith][Resubmit] add fastMathAttr on arith::extf and arith::truncf (#95346) Add an fastMathAttr on arith::extf and arith::truncf. If these two ops are inserted by some promotion passes (like legalize-to-f32 / emulate-unsupported-floats), they will be labeled as FastMathFlags::contract, denoting that they can be then eliminated by canonicalizer. The elimination can help improve performance, while may introduce some numerical differences. | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
[mlir][openacc] Added custom builder for acc::ParallelOp (#98191) This change adds a custom builder for acc::ParallelOp. This enables users to only specify the operands they would need for their acc::ParallelOp while building it. They can specify nothing to create an empty acc.parallel, or all of the 11 operands listed [here](https://mlir.llvm.org/docs/Dialects/OpenACCDialect/#operands-27), or anywhere in between following the specified order in this custom builder. Unspecified operands are left empty. Additionally, users can later set the optional attributes such as numGangsDeviceType using the available attribute setters for acc::ParallelOp. | 2 年前 | |
[openacc][openmp] Add dialect representation for acc atomic operations (#65493) The OpenACC standard specifies an atomic construct in section 2.12 (of 3.3 spec), used to ensure that a specific location is accessed or updated atomically. Four different clauses are allowed: read, write, update, or capture. If no clause appears, it is as if update is used. The OpenMP specification defines the same clauses for omp atomic. The types of expression and the clauses in the OpenACC spec match the OpenMP spec exactly. The main difference is that the OpenMP specification is a superset - it includes clauses for hint and memory order. It also allows conditional expression statements. But otherwise, the expression definition matches. Thus, for OpenACC, we refactor and reuse the OpenMP implementation as follows: * The atomic operations are duplicated in OpenACC dialect. This is preferable so that each language's semantics are precisely represented even if specs have divergence. * However, since semantics overlap, a common interface between the atomic operations is being added. The semantics for the interfaces are not generic enough to be used outside of OpenACC and OpenMP, and thus new folders were added to hold common pieces of the two dialects. * The atomic interfaces define common accessors (such as getting x or v) which match the OpenMP and OpenACC specs. It also adds common verifiers intended to be called by each dialect's operation verifier. * The OpenMP write operation was updated to use x and expr to be consistent with its other operations (that use naming based on spec). The frontend lowering necessary to generate the dialect can also be reused. This will be done in a follow up change. | 2 年前 | |
| 2 年前 | ||
Reapply "[mlir][PDL] Add support for native constraints with results (#82760)" with a small stack-use-after-scope fix in getConstraintPredicates() This reverts commit c80e6edba4a9593f0587e27fa0ac825ebe174afd. | 2 年前 | |
[mlir] Move FunctionInterfaces to Interfaces directory and inherit from CallableOpInterface Functions are always callable operations and thus every operation implementing the FunctionOpInterface also implements the CallableOpInterface. The only exception was the FuncOp in the toy example. To make implementation of the FunctionOpInterface easier, this commit lets FunctionOpInterface inherit from CallableOpInterface and merges some of their methods. More precisely, the CallableOpInterface has methods to get the argument and result attributes and a method to get the result types of the callable region. These methods are always implemented the same way as their analogues in FunctionOpInterface and thus this commit moves all the argument and result attribute handling methods to the callable interface as well as the methods to get the argument and result types. The FuntionOpInterface then does not have to declare them as well, but just inherits them from the CallableOpInterface. Adding the inheritance relation also required to move the FunctionOpInterface from the IR directory to the Interfaces directory since IR should not depend on Interfaces. Reviewed By: jpienaar, springerm Differential Revision: https://reviews.llvm.org/D157988 | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
[mlir][Ptr] Init the Ptr dialect with the !ptr.ptr type. (#86860) This patch initializes the ptr dialect directories and base files, adding the !ptr.ptr type and the #ptr.spec<...> data layout spec attribute. The !ptr.ptr type is an opaque pointer type optionally parameterized by a memory space. This type typically represents a handle to an object in memory or target-dependent values like nullptr. The implementation of the DataLayoutTypeInterface interface for !ptr.ptr was adapted from !llvm.ptr's implementation. This implementation uses the #ptr.spec<...> attribute for defining the data layout specification. See [[RFC] ptr dialect & modularizing ptr ops in the LLVM dialect](https://discourse.llvm.org/t/rfc-ptr-dialect-modularizing-ptr-ops-in-the-llvm-dialect/75142) for rationale and roadmap. | 2 年前 | |
mlir/LogicalResult: move into llvm (#97309) This patch is part of a project to move the Presburger library into LLVM. | 2 年前 | |
[MLIR][SCF] fix scf.index_switch fold convergence (#98535) (#98680) If the scf.index_switch op has no result, the current fold logic results in an infinite loop (see #98535). The is because fold mechanism does not support *erasing* zero-result ops. This PR moves the fold logic to a canonicalizer and fix the issue. | 2 年前 | |
[mlir][spirv] Restructure code in SPIRVConversion.cpp. NFC. (#99393) | 2 年前 | |
[mlir] New canonicalization patterns for shape.shape_of and tensor.reshape (#98531) This PR includes 3 new canonicalization patterns: - Operation shape.shape_of: shape of reshape // Before func.func @f(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>) -> tensor<?xindex> { %reshape = tensor.reshape %arg0(%arg1) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> %0 = shape.shape_of %reshape : tensor<*xf32> -> tensor<?xindex> return %0 : tensor<?xindex> } // After func.func @f(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>) -> tensor<?xindex> { return %arg1 : tensor<?xindex> } - Operation tensor.reshape: reshape of reshape // Before func.func @fold_tensor_reshape(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>, %arg2: tensor<?xindex>) -> tensor<*xf32> { %0 = tensor.reshape %arg0(%arg1) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> %1 = tensor.reshape %0(%arg2) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> return %1 : tensor<*xf32> } // After func.func @fold_tensor_reshape(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>, %arg2: tensor<?xindex>) -> tensor<*xf32> { %reshape = tensor.reshape %arg0(%arg2) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> return %reshape : tensor<*xf32> } - Operation tensor.reshape: reshape 1D to 1D // Before func.func @fold_reshape_1d(%input: tensor<?xf32>, %shape: tensor<1xindex>) -> tensor<?xf32> { %0 = tensor.reshape %input(%shape) : (tensor<?xf32>, tensor<1xindex>) -> tensor<?xf32> return %0 : tensor<?xf32> } // After func.func @fold_reshape_1d(%arg0: tensor<?xf32>, %arg1: tensor<1xindex>) -> tensor<?xf32> { return %arg0 : tensor<?xf32> } These three canonicalization patterns cooperate to simplify the IR structure emerging from the lowering of certain element-wise ops with unranked tensor inputs. See file unranked-tensor-lowering.mlir in the proposed change list for a detailed example and description. For context, this PR is meant to enable code optimizations for the code generated while lowering ops quant.qcast and quant.dcast with unranked tensors, as proposed in https://discourse.llvm.org/t/rfc-improvements-in-the-quant-dialect/79942 (implementation currently in progress). | 2 年前 | |
[mlir][sparse] Use vector.step for index vector generation (#97692) | 2 年前 | |
[mlir] New canonicalization patterns for shape.shape_of and tensor.reshape (#98531) This PR includes 3 new canonicalization patterns: - Operation shape.shape_of: shape of reshape // Before func.func @f(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>) -> tensor<?xindex> { %reshape = tensor.reshape %arg0(%arg1) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> %0 = shape.shape_of %reshape : tensor<*xf32> -> tensor<?xindex> return %0 : tensor<?xindex> } // After func.func @f(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>) -> tensor<?xindex> { return %arg1 : tensor<?xindex> } - Operation tensor.reshape: reshape of reshape // Before func.func @fold_tensor_reshape(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>, %arg2: tensor<?xindex>) -> tensor<*xf32> { %0 = tensor.reshape %arg0(%arg1) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> %1 = tensor.reshape %0(%arg2) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> return %1 : tensor<*xf32> } // After func.func @fold_tensor_reshape(%arg0: tensor<*xf32>, %arg1: tensor<?xindex>, %arg2: tensor<?xindex>) -> tensor<*xf32> { %reshape = tensor.reshape %arg0(%arg2) : (tensor<*xf32>, tensor<?xindex>) -> tensor<*xf32> return %reshape : tensor<*xf32> } - Operation tensor.reshape: reshape 1D to 1D // Before func.func @fold_reshape_1d(%input: tensor<?xf32>, %shape: tensor<1xindex>) -> tensor<?xf32> { %0 = tensor.reshape %input(%shape) : (tensor<?xf32>, tensor<1xindex>) -> tensor<?xf32> return %0 : tensor<?xf32> } // After func.func @fold_reshape_1d(%arg0: tensor<?xf32>, %arg1: tensor<1xindex>) -> tensor<?xf32> { return %arg0 : tensor<?xf32> } These three canonicalization patterns cooperate to simplify the IR structure emerging from the lowering of certain element-wise ops with unranked tensor inputs. See file unranked-tensor-lowering.mlir in the proposed change list for a detailed example and description. For context, this PR is meant to enable code optimizations for the code generated while lowering ops quant.qcast and quant.dcast with unranked tensors, as proposed in https://discourse.llvm.org/t/rfc-improvements-in-the-quant-dialect/79942 (implementation currently in progress). | 2 年前 | |
[mlir][tosa] Use roundeven in TOSA cast splat constant op folding (#99484) The behavior of TOSA Cast operation for floating-point to integers is to round to the nearest even. This commit aligns the behavior of folding a TOSA Cast of a float splat to int, so it also uses roundeven. | 2 年前 | |
[mlir] Makes zip_shortest an optional keyword in transform.foreach (#98492) This PR addresses a [comment] made by @ftynse about the syntax for ForeachOp. The syntax was modified by @muneebkhan85 in #82792, where the attribute dictionary was moved to the middle. This patch moves it back to its original place at the end. And introduces an optional keyword for zip_shortest. [comment]: https://github.com/llvm/llvm-project/pull/82792#pullrequestreview-2132814144 | 2 年前 | |
[NFC][mlir] Reorder declarePromisedInterface() operands (#86628) Reorder the template operands of declarePromisedInterface() to match declarePromisedInterfaces(). | 2 年前 | |
[mlir] [linalg] Add pattern to swap transpose with broadcast (#97063) Add a pattern that implement: transpose(broadcast(input)) -> broadcast(transpose(input)) | 2 年前 | |
[MLIR][Vector] Generalize DropUnitDimFromElementwiseOps to non leading / trailing dimensions. (#98455) Generalizes DropUnitDimFromElementwiseOps to support inner unit dimensions. This change stems from improving lowering of contractionOps for Arm SME. Where we end up with inner unit dimensions on MulOp, BroadcastOp and TransposeOp, preventing the generation of outerproducts. discussed [here](https://discourse.llvm.org/t/on-improving-arm-sme-lowering-resilience-in-mlir/78543/17?u=nujaa). Fix after : https://github.com/llvm/llvm-project/pull/97652 showed an unhandled edge case when all dimensions are one. The generated target VectorType would be vector<f32> which is apparently not supported by the mulf. In case all dimensions are dropped, the target vectorType is vector<1xf32> --------- Co-authored-by: Benjamin Maxwell <macdue@dueutil.tech> | 2 年前 | |
[mlir][Conversion] Store const type converter in ConversionPattern ConversionPatterns do not (and should not) modify the type converter that they are using. * Make ConversionPattern::typeConverter const. * Make member functions of the LLVMTypeConverter const. * Conversion patterns take a const type converter. * Various helper functions (that are called from patterns) now also take a const type converter. Differential Revision: https://reviews.llvm.org/D157601 | 2 年前 | |
[mlir][xegpu] Patch dynamic descriptor creation (#93580) fixes the bug in XeGPU's CreateNdDescOp tensor creation with dynamic offset and strides. | 2 年前 | |
[mlir][Ptr] Init the Ptr dialect with the !ptr.ptr type. (#86860) This patch initializes the ptr dialect directories and base files, adding the !ptr.ptr type and the #ptr.spec<...> data layout spec attribute. The !ptr.ptr type is an opaque pointer type optionally parameterized by a memory space. This type typically represents a handle to an object in memory or target-dependent values like nullptr. The implementation of the DataLayoutTypeInterface interface for !ptr.ptr was adapted from !llvm.ptr's implementation. This implementation uses the #ptr.spec<...> attribute for defining the data layout specification. See [[RFC] ptr dialect & modularizing ptr ops in the LLVM dialect](https://discourse.llvm.org/t/rfc-ptr-dialect-modularizing-ptr-ops-in-the-llvm-dialect/75142) for rationale and roadmap. | 2 年前 | |
[mlir][NFC] Simplify type checks with isa predicates (#87183) For more context on isa predicates, see: https://github.com/llvm/llvm-project/pull/83753. | 2 年前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 2 年前 |