KKazu Hirata[mlir] Use value_or (NFC)
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[mlir][AMDGPU] Add lds_barrier op The lds_barrier op allows workgroups to wait at a barrier for operations to/from their local data store (LDS) to complete without incurring the performance penalties of a full memory fence. Reviewed By: nirvedhmeshram Differential Revision: https://reviews.llvm.org/D129522 | 3 年前 | |
[mlir] Flip more uses to prefixed accessor form (NFC). Try to keep the final flip small. Need to flip MemRef as there are many templated cases with it and Tensor. | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir][spirv] Rename spv.GLSL ops to spv.GL. NFC. This is to improve consistency within the SPIR-V dialect and make these ops a bit shorter. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D130280 | 3 年前 | |
[mlir] Update accessors to prefixed form (NFC) Follow up from flipping dialects to both, flip accessor used to prefixed variant ahead to flipping from _Both to _Prefixed. This just flips to the accessors introduced in the preceding change which are just prefixed forms of the existing accessor changed from. Mechanical change using helper script https://github.com/jpienaar/llvm-project/blob/main/clang-tools-extra/clang-tidy/misc/AddGetterCheck.cpp and clang-format. | 3 年前 | |
Revert "[MLIR] Generic 'malloc', 'aligned_alloc' and 'free' functions" This reverts commit 3e21fb616d9a1b29bf9d1a1ba484add633d6d5b3. A lot of integration tests are failing on the bot. | 3 年前 | |
[mlir] (NFC) run clang-format on all files | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir] Use has_value instead of hasValue (NFC) | 3 年前 | |
[mlir] (NFC) run clang-format on all files | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir][LLVMIR] Add support for va_start/copy/end intrinsics This patch adds three new LLVM intrinsic operations: llvm.intr.vastart/copy/end. And its translation from LLVM IR. This effectively removes a restriction, imposed by 0126dcf1f0a1, where non-external functions in LLVM dialect cannot be variadic. At that time it was not clear how LLVM intrinsics are going to be modeled, which indirectly affects va_start/copy/end, the core intrinsics used in variadic functions. But since we have LLVM intrinsics as normal MLIR operations, it's not a problem anymore. Differential Revision: https://reviews.llvm.org/D127540 | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir] (NFC) run clang-format on all files | 3 年前 | |
[mlir] replace 'emit_c_wrappers' func->llvm conversion option with a pass The 'emit_c_wrappers' option in the FuncToLLVM conversion requests C interface wrappers to be emitted for every builtin function in the module. While this has been useful to bootstrap the interface, it is problematic in the longer term as it may unintentionally affect the functions that should retain their existing interface, e.g., libm functions obtained by lowering math operations (see D126964 for an example). Since D77314, we have a finer-grain control over interface generation via an attribute that avoids the problem entirely. Remove the 'emit_c_wrappers' option. Introduce the '-llvm-request-c-wrappers' pass that can be run in any pipeline that needs blanket emission of functions to annotate all builtin functions with the attribute before performing the usual lowering that accounts for the attribute. Reviewed By: chelini Differential Revision: https://reviews.llvm.org/D127952 | 3 年前 | |
[mlir] Remove VectorToROCDL Between issues such as https://github.com/llvm/llvm-project/issues/56323, the fact that this lowering (unlike the code in amdgpu-to-rocdl) does not correctly set up bounds checks (and thus will cause page faults on reads that might need to be padded instead), and that fixing these problems would, essentially, involve replicating amdgpu-to-rocdl, remove --vector-to-rocdl for being broken. In addition, the lowering does not support many aspects of transfer_{read,write}, like supervectors, and may not work correctly in their presence. We (the MLIR-based convolution generator at AMD) do not use this conversion pass, nor are we aware of any other clients. Migration strategies: - Use VectorToLLVM - If buffer ops are particularly needed in your application, use amdgpu.raw_buffer_{load,store} A VectorToAMDGPU pass may be introduced in the future. Reviewed By: ThomasRaoux Differential Revision: https://reviews.llvm.org/D129308 | 3 年前 | |
[mlir] Don't use Optional::getValue (NFC) | 3 年前 | |
[mlir] Don't use Optional::getValue (NFC) | 3 年前 | |
Use callables directly in any_of, count_if, etc (NFC) | 3 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[mlir] Flip LinAlg dialect to _Both This one required more changes than ideal due to overlapping generated name with different return types. Changed getIndexingMaps to getIndexingMapsArray to move it out of the way/highlight that it returns (more expensively) a SmallVector and uses the prefixed name for the Attribute. Differential Revision: https://reviews.llvm.org/D129919 | 3 年前 | |
[mlir] (NFC) run clang-format on all files | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir][math] Added math::tan operation. Differential Revision: https://reviews.llvm.org/D129539 | 3 年前 | |
[mlir][spirv] Rename spv.GLSL ops to spv.GL. NFC. This is to improve consistency within the SPIR-V dialect and make these ops a bit shorter. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D130280 | 3 年前 | |
[mlir] Use value_or (NFC) | 3 年前 | |
[mlir] Flip accessors to prefixed form (NFC) Another mechanical sweep to keep diff small for flip to _Prefixed. | 3 年前 | |
[mlir] Use value_or (NFC) | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[Flang][OpenMP] Initial support for integer reduction in worksharing-loop Lower the Flang parse-tree containing OpenMP reductions to the OpenMP dialect. The OpenMP dialect models reductions with, 1) A reduction declaration operation that specifies how to initialize, combine, and atomically combine private reduction variables. 2) The OpenMP operation (like wsloop) that supports reductions has an array of reduction accumulator variables (operands) and an array attribute of the same size that points to the reduction declaration to be used for the reduction accumulation. 3) The OpenMP reduction operation that takes a value and an accumulator. This operation replaces the original reduction operation in the source. (1) is implemented by the createReductionDecl in OpenMP.cpp, (2) is implemented while creating the OpenMP operation, (3) is implemented by the genOpenMPReduction function in OpenMP.cpp, and called from Bridge.cpp. The implementation of (3) is not very robust. NOTE 1: The patch currently supports only reductions for integer type addition. NOTE 2: Only supports reduction in the worksharing loop. NOTE 3: Does not generate atomic combination region. NOTE 4: Other options for creating the reduction operation include a) having the reduction operation as a construct containing an assignment and then handling it appropriately in the Bridge. b) we can modify genAssignment or genFIR(AssignmentStmt) in the Bridge to handle OpenMP reduction but so far we have tried not to mix OpenMP and non-OpenMP code and this will break that. I will try (b) in a separate patch. NOTE 5: OpenMP dialect gained support for reduction with the patches: D105358, D107343. See https://discourse.llvm.org/t/rfc-openmp-reduction-support/3367 for more details. Reviewed By: awarzynski Differential Revision: https://reviews.llvm.org/D130077 Co-authored-by: Peixin-Qiao <qiaopeixin@huawei.com> | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir][NFC] Update the Builtin dialect to use "Both" accessors Differential Revision: https://reviews.llvm.org/D121189 | 4 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[mlir][spirv] Allow converting from index type in unsigned ops index type is converted to i32 in SPIR-V. This is fine to support for all signed/unsigned ops. Reviewed By: hanchung Differential Revision: https://reviews.llvm.org/D124451 | 4 年前 | |
[mlir][spirv] Rename spv.GLSL ops to spv.GL. NFC. This is to improve consistency within the SPIR-V dialect and make these ops a bit shorter. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D130280 | 3 年前 | |
[mlir] move SCF headers to SCF/{IR,Transforms} respectively This aligns the SCF dialect file layout with the majority of the dialects. Reviewed By: jpienaar Differential Revision: https://reviews.llvm.org/D128049 | 3 年前 | |
[mlir] (NFC) Clean up bazel and CMake target names All dialect targets in bazel have been named *Dialect and all dialect targets in CMake have been named MLIR*Dialect. | 3 年前 | |
[mlir] Update accessors to prefixed form (NFC) Follow up from flipping dialects to both, flip accessor used to prefixed variant ahead to flipping from _Both to _Prefixed. This just flips to the accessors introduced in the preceding change which are just prefixed forms of the existing accessor changed from. Mechanical change using helper script https://github.com/jpienaar/llvm-project/blob/main/clang-tools-extra/clang-tidy/misc/AddGetterCheck.cpp and clang-format. | 3 年前 | |
[mlir][tosa] Flip accessors used to prefixed form (NFC) Follow up from dialect flip, just flipping accessors. Both forms still generated. | 3 年前 | |
[mlir][tosa] Flip accessors used to prefixed form (NFC) Follow up from dialect flip, just flipping accessors. Both forms still generated. | 3 年前 | |
[mlir][tosa] Flip accessors used to prefixed form (NFC) Follow up from dialect flip, just flipping accessors. Both forms still generated. | 3 年前 | |
[mlir][tosa] Flip accessors used to prefixed form (NFC) Follow up from dialect flip, just flipping accessors. Both forms still generated. | 3 年前 | |
[mlir] Flip LinAlg dialect to _Both This one required more changes than ideal due to overlapping generated name with different return types. Changed getIndexingMaps to getIndexingMapsArray to move it out of the way/highlight that it returns (more expensively) a SmallVector and uses the prefixed name for the Attribute. Differential Revision: https://reviews.llvm.org/D129919 | 3 年前 | |
Use callables directly in any_of, count_if, etc (NFC) | 3 年前 | |
Convert for_each to range-based for loops (NFC) | 3 年前 | |
[mlir][spirv] Rename spv.GLSL ops to spv.GL. NFC. This is to improve consistency within the SPIR-V dialect and make these ops a bit shorter. Reviewed By: antiagainst Differential Revision: https://reviews.llvm.org/D130280 | 3 年前 | |
[mlir] Remove VectorToROCDL Between issues such as https://github.com/llvm/llvm-project/issues/56323, the fact that this lowering (unlike the code in amdgpu-to-rocdl) does not correctly set up bounds checks (and thus will cause page faults on reads that might need to be padded instead), and that fixing these problems would, essentially, involve replicating amdgpu-to-rocdl, remove --vector-to-rocdl for being broken. In addition, the lowering does not support many aspects of transfer_{read,write}, like supervectors, and may not work correctly in their presence. We (the MLIR-based convolution generator at AMD) do not use this conversion pass, nor are we aware of any other clients. Migration strategies: - Use VectorToLLVM - If buffer ops are particularly needed in your application, use amdgpu.raw_buffer_{load,store} A VectorToAMDGPU pass may be introduced in the future. Reviewed By: ThomasRaoux Differential Revision: https://reviews.llvm.org/D129308 | 3 年前 | |
[MLIR][GPU] Add NvGpu mma.sync path to the VectorToGPU pass This changes adds the option to lower to NvGpu dialect ops during the VectorToGPU convsersion pass. Because this transformation reuses existing VectorToGPU logic, a seperate VectorToNvGpu conversion pass is not created. The option use-nvgpu is added to the VectorToGPU pass. When this is true, the pass will attempt to convert slices rooted at vector.contract operations into nvgpu.mma.sync ops, and vector.transfer_read ops are converted to either nvgpu.ldmatrix or one or more vector.load operations. The specific data loaded will depend on the thread id within a subgroup (warp). These index calculations depend on data type and shape of the MMA op according to the downstream PTX specification. The code for supporting these details is separated into NvGpuSupport.cpp|h. Differential Revision: https://reviews.llvm.org/D122940 | 4 年前 |