| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[mlir][llvm] Handle debug record import edge cases (#168774) This commit enables the direct import of debug records by default and fixes issues with two edge cases: - Detect early on if the address operand is an argument list (calling getAddress() for argument lists asserts) - Use getAddress() to check if the address operand is null, which means the address operand is an empty metadata node, which currently is not supported. - Add support for debug label records. This is a follow-up to: https://github.com/llvm/llvm-project/pull/167812 | 8 个月前 | |
Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638) Reland commit fb829bf11feeb53f815a3abf539e63ec3a23ed3d with additional fixes relating to post-merge CI failure /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp: In function ‘constexpr llvm::nvvm::CTAGroupKind getNVVMCtaGroupKind(mlir::NVVM::CTAGroupKind)’: /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/llvm/include/llvm/Support/ErrorHandling.h:165:36: error: call to non-constexpr function ‘void llvm::llvm_unreachable_internal(const char*, const char*, unsigned int)’ ::llvm::llvm_unreachable_internal(msg, __FILE__, __LINE__) ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~~~~~~~~~~~~~~~~~~~ /vol/worker/mlir-nvidia/mlir-nvidia-gcc7/llvm.src/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp:73:3: note: in expansion of macro ‘llvm_unreachable’ llvm_unreachable("unsupported cta_group value"); ^ | 8 个月前 | |
[MLIR][LLVM] Implement LLVM dialect support for global aliases (#125295) This includes support for module translation, module import and add tests for both. Fix https://github.com/llvm/llvm-project/issues/115390 ClangIR cannot currently lower global aliases to LLVM because of missing support for this. | 1 年前 | |
[mlir][amx] Optional stride for tile load and store (#159569) Adds an optional stride argument to amx.tile_load and amx.tile_store operations. The stride argument aligns ops closer to the hardware intrinsics. However, stride remains optional to preserve current op behavior. Explicit stride allows greater flexibility in terms of the base buffer shapes and allows different read and write memory patterns. When stride is not provided, it is inferred from the buffer shape as before. Operations documentation is expanded to make ops easier to use. | 9 个月前 | |
[mlir][Pass] Include anchor op in -pass-pipeline In D134622 the printed form of a pass manager is changed to include the name of the op that the pass manager is anchored on. This updates the -pass-pipeline argument format to include the anchor op as well, so that the printed form of a pipeline can be directly passed to -pass-pipeline. In most cases this requires updating -pass-pipeline='pipeline' to -pass-pipeline='builtin.module(pipeline)'. This also fixes an outdated assert that prevented running a PassManager anchored on 'any'. Reviewed By: rriddle Differential Revision: https://reviews.llvm.org/D134900 | 3 年前 | |
[MLIR][ArmNeon] Add an ArmNeon operation which maps to bfmmla (#145038) | 1 年前 | |
[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 个月前 | |
[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][ArmSVE] Add an ArmSVE dialect operation mapping to bfmmla (#145064) | 1 年前 | |
[IR] Convert from nocapture to captures(none) (#123181) This PR removes the old nocapture attribute, replacing it with the new captures attribute introduced in #116990. This change is intended to be essentially NFC, replacing existing uses of nocapture with captures(none) without adding any new analysis capabilities. Making use of non-none values is left for a followup. Some notes: * nocapture will be upgraded to captures(none) by the bitcode reader. * nocapture will also be upgraded by the textual IR reader. This is to make it easier to use old IR files and somewhat reduce the test churn in this PR. * Helper APIs like doesNotCapture() will check for captures(none). * MLIR import will convert captures(none) into an llvm.nocapture attribute. The representation in the LLVM IR dialect should be updated separately. | 1 年前 | |
[MLIR][LLVM] Import dereferenceable metadata from LLVM IR (#130974) Add support for importing dereferenceable and dereferenceable_or_null metadata into LLVM dialect. Add a new attribute which models these two metadata nodes and a new OpInterface. | 1 年前 | |
[mlir][llvm] Translate alias scopes lazily Change the LLVM dialect to LLVM IR translation to convert the alias scope attributes lazily to LLVM IR metadata. Previously, the alias scopes have been translated upfront walking the alias scopes of operations that implement the AliasAnalysisOpInterface. As a result, the translation of a module that contains only a noalias scope intrinsic failed, since its alias scope attribute has not been translated due to the intrinsic not implementing AliasAnalysisOpInterface. Reviewed By: zero9178 Differential Revision: https://reviews.llvm.org/D159187 | 2 年前 | |
[MLIR][LLVM] Fix blockaddress mapping to LLVM blocks (#139814) After each function is translated, both value and block maps are erased, which makes the current mapping of blockaddresses to llvm blocks broken - the patching happens only after *all* functions are translated. Simplify the overall mapping, update comments, variable names and fix the bug. --------- Co-authored-by: Christian Ulmann <christianulmann@gmail.com> | 1 年前 | |
[mlir][LLVM] handle argument and result attributes in llvm.call and llvm.invoke (#123177) Update llvm.call/llvm.invoke pretty printer/parser and the llvm ir import/export to deal with the argument and result attributes. This patch is made on top of PR 123176 that modified the CallOpInterface and added the argument and result attributes to llvm.call and llvm.invoke without doing anything with them. RFC: https://discourse.llvm.org/t/mlir-rfc-adding-argument-and-result-attributes-to-llvm-call/84107 | 1 年前 | |
[mlir][llvm] Add comdat attribute to functions This revision adds comdat support to functions. Additionally, it ensures only comdats that have uses are imported/exported and only non-empty global comdat operations are created. Reviewed By: Dinistro Differential Revision: https://reviews.llvm.org/D153739 | 3 年前 | |
| 1 年前 | ||
[MLIR][LLVM][DLTI] Handle data layout token 'n32:64' (#141299) | 1 年前 | |
[mlir][LLVM] Add disjoint flag (#115855) The implementation is mostly based on the one existing for the exact flag. disjoint means that for each bit, that bit is zero in at least one of the inputs. This allows the Or to be treated as an Add since no carry can occur from any bit. If the disjoint keyword is present, the result value of the or is a [poison value](https://llvm.org/docs/LangRef.html#poisonvalues) if both inputs have a one in the same bit position. For vectors, only the element containing the bit is poison. | 1 年前 | |
[mlir][LLVM] Add exact flag (#115327) The implementation is mostly based on the one existing for the nsw and nuw flags. If the exact flag is present, the corresponding operation returns a poison value when the result is not exact. (For a division, if rounding happens; for a right shift, if a non-zero bit is shifted out.) | 1 年前 | |
[NVPTX] Convert scalar function nvvm.annotations to attributes (#125908) Replace some more nvvm.annotations with function attributes, auto-upgrading the annotations as needed. These new attributes will be more idiomatic and compile-time efficient than the annotations. - !"maxclusterrank" / !"cluster_max_blocks" -> "nvvm.maxclusterrank" - !"minctasm" -> "nvvm.minctasm" - !"maxnreg" -> "nvvm.maxnreg" | 1 年前 | |
[mlir][LLVMIR] Remove "unsafe-fp-math" attribute support (#162782) These global flags block furthur improvements for clang, users should always use fast-math flags see also https://discourse.llvm.org/t/rfc-honor-pragmas-with-ffp-contract-fast/80797 Remove them incrementally, this is the mlir part. | 9 个月前 | |
[MLIR] Add support for frame pointers in MLIR (#72145) Add support for frame pointers in MLIR. --------- Co-authored-by: Markus Böck <markus.boeck02@gmail.com> Co-authored-by: Christian Ulmann <christianulmann@gmail.com> | 2 年前 | |
| 10 个月前 | ||
[mlir][gpu] Change GPU modules to globals (#135478) Load/unload GPU modules in global ctors/dtors instead of each time when launching a kernel. Loading GPU modules is a heavy-weight operation and synchronizes the GPU context. Now that the modules are loaded ahead of time, asynchronously launched kernels can run concurrently, see https://discourse.llvm.org/t/how-to-lower-the-combination-of-async-gpu-ops-in-gpu-dialect. The implementations of embedBinary() and launchKernel() use slightly different mechanics at the moment but I prefer to not change the latter more than necessary as part of this PR. I will prepare a follow-up NFC for launchKernel() to align them again. | 1 年前 | |
[flang][mlir] Add llvm.ident metadata when compiling with flang This brings the behavior of flang in line with clang which also adds this metadata unconditionally. Co-authored-by: Tarun Prabhu <tarun.prabhu@gmail.com> | 1 年前 | |
[mlir][LLVMIR] Add IFuncOp to LLVM dialect (#147697) Add IFunc to LLVM dialect and add support for lifting/exporting LLVMIR IFunc. | 1 年前 | |
Reapply [MLIR][LLVM] Support for indirectbr (#136378) Fix msan issue that caused revert in https://github.com/llvm/llvm-project/pull/135695 ### Original message Now that LLVM dialect has blockaddress support, introduce import/translation for indirectbr instruction. | 1 年前 | |
[mlir][llvm] adds an attribute for the module level assembly (#151318) Adds support for the module level assembly in the LLVM IR dialect. --------- Co-authored-by: Tobias Gysi <tobias.gysi@nextsilicon.com> | 11 个月前 | |
[mlir][LLVM] handle argument and result attributes in llvm.call and llvm.invoke (#123177) Update llvm.call/llvm.invoke pretty printer/parser and the llvm ir import/export to deal with the argument and result attributes. This patch is made on top of PR 123176 that modified the CallOpInterface and added the argument and result attributes to llvm.call and llvm.invoke without doing anything with them. RFC: https://discourse.llvm.org/t/mlir-rfc-adding-argument-and-result-attributes-to-llvm-call/84107 | 1 年前 | |
[mlir] Add splitDebugFilename field in DIComplileUnitAttr. (#160704) Mostly mechanical changes to add the missing field. | 10 个月前 | |
[MLIR][LLVMIR] Add {s,u}cmp intrinsics to LLVM dialect (#167870) | 8 个月前 | |
[MLIR] Add target_specific_attrs attribute to mlir.global (#154706) Adds a target_specific_attrs optional array attribute to mlir.global, as well as conversions to and from LLVM attributes on llvm::GlobalVariable objects. This is necessary to preserve unknown attributes on global variables when converting to and from the LLVM Dialect. Previously, any attributes on an llvm::GlobalVariable not explicitly modeled by mlir.global were dropped during conversion. | 10 个月前 | |
[mlir][test] XFAIL little-endian-only tests on SPARC (#103726) 3 MLIR tests FAIL on SPARC, both Solaris/sparcv9 and Linux/sparc64: MLIR :: Conversion/ArithToSPIRV/arith-to-spirv-le-specific.mlir MLIR :: IR/elements-attr-interface.mlir MLIR :: Target/LLVMIR/llvmir-le-specific.mlir The issue is always the same: the tests in question are little-endian-only currently, so this patch XFAILs them on sparc* as is already done for s390x. Tested on sparcv9-sun-solaris2.11, sparc64-unknown-linux-gnu, amd64-pc-solaris2.11, and x86_64-pc-linux-gnu. | 1 年前 | |
[mlir][llvmir][debug] Correctly generate location for phi nodes. (#105534) In [convertBlockImpl](https://github.com/llvm/llvm-project/blob/87eeed1f0ebe57abffde560c25dd9829dc6038f3/mlir/lib/Target/LLVMIR/ModuleTranslation.cpp#L959), the debug location is set on the builder before the op is processed. This results in correct location being given to corresponding llvm instructions. But same is not done when phi nodes are created a few lines above. This result is phi nodes getting whatever the current debug location of the builder is. It can be nothing or in worst case a stale location. Fixed by calling SetCurrentDebugLocation before generating phi nodes. | 1 年前 | |
[mlir][LLVM] Delete LLVMFixedVectorType and LLVMScalableVectorType (#133286) Since #125690, the MLIR vector type supports !llvm.ptr as an element type. The only remaining element type for LLVMFixedVectorType is now LLVMPPCFP128Type. This commit turns LLVMPPCFP128Type into a proper FP type (by implementing FloatTypeInterface), so that the MLIR vector type accepts it as an element type. This makes LLVMFixedVectorType obsolete. LLVMScalableVectorType is also obsolete. This commit deletes LLVMFixedVectorType and LLVMScalableVectorType. Note for LLVM integration: Use VectorType instead of LLVMFixedVectorType and LLVMScalableVectorType. | 1 年前 | |
[mlir][LLVM] Resync memory effect attribute with LLVM IR (#168568) - Add missing locations, namely 'ErrnoMem', 'TargetMem0', and 'TargetMem1'. | 8 个月前 | |
[MLIR][LLVM] Add distinct identifier to DICompileUnit attribute (#77070) This commit adds a distinct attribute parameter to the DICompileUnit to enable the modeling of distinctness. LLVM requires DICompileUnits to be distinct and there are cases where one gets two equivalent compilation units but LLVM still requires differentiates them. We observed such cases for combinations of LTO and inline functions. This patch also changes the DIScopeForLLVMFuncOp pass to a module pass, to ensure that only one distinct DICompileUnit is created, instead of one for each function. | 2 年前 | |
[mlir][LLVMIR] Support memory model relaxation annotations (MMRA) (#157770) This commit adds support for exportind and importing MMRA data in the LLVM dialect. MMRA is a potentilly-discardable piece of metadata that can be placed on any operation that touches memory (fences, loads, stores, atomics, and intrinsics that operate on memory). It includes one (technically zero) ome more prefix:suffix string pairs which indicate ways in which the LLVM memory model can be relaxed for these annotations. At the MLIR level, each tag is represented with a #llvm.mmra_tag<"prefix":"suffex"> attribute, and the MMRA metadata as a whole is represented as a discardable llvm.mmra attribute. (This discardability both allows us to transparently enable MMRA for wrapper dialects like ROCDL and ensures that MLIR passes which don't know about MMRA combining will, conservatively, discard the annotations, per the LLVM spec). --------- Co-authored-by: Tobias Gysi <tobias.gysi@nextsilicon.com> | 10 个月前 | |
[mlir][llvm] adds an attribute for the module level assembly (#151318) Adds support for the module level assembly in the LLVM IR dialect. --------- Co-authored-by: Tobias Gysi <tobias.gysi@nextsilicon.com> | 11 个月前 | |
[flang] Add support for -mrecip[=<list>] (#143418) This patch adds support for the -mrecip command line option. The parsing of this options is equivalent to Clang's and it is implemented by setting the "reciprocal-estimates" function attribute. Also move the ParseMRecip(...) function to CommonArgs, so that Flang is able to make use of it as well. --------- Co-authored-by: Cameron McInally <cmcinally@nvidia.com> | 1 年前 | |
[mlir][LLVM] Add nneg flag (#115498) This implementation is based on the existing one for the exact flag. If the nneg flag is set and the argument is negative, the result is a poison value. | 1 年前 | |
[mlir][LLVM] Add nsw and nuw flags to trunc (#115509) This implementation is based on the one already existing for the binary operations. If the nuw keyword is present, and any of the truncated bits are non-zero, the result is a poison value. If the nsw keyword is present, and any of the truncated bits are not the same as the top bit of the truncation result, the result is a poison value. | 1 年前 | |
[NVVM] Make nanosleep op duration SSA value (#167331) | 8 个月前 | |
[mlir][NVVM] Add support for barrier0-reduction operation (#167036) Add support for nvvm.barrier0.[popc|and|or] operation. It is added as a separate operation since Barrier0Op has no result. https://docs.nvidia.com/cuda/nvvm-ir-spec/#barrier-and-memory-fence This will be used in CUDA Fortran lowering: https://github.com/llvm/llvm-project/blob/49f55f4991227f3c7a2b8161bbf45c74b7023944/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp#L1081 And could be used later in the CUDA C/C++ with CIR https://github.com/llvm/llvm-project/blob/49f55f4991227f3c7a2b8161bbf45c74b7023944/clang/lib/Headers/__clang_cuda_device_functions.h#L524 --------- Co-authored-by: Guray Ozen <guray.ozen@gmail.com> | 8 个月前 | |
[OpenMP][OMPIRBuilder] Collect users of a value before replacing them in target outlined function (#139064) This PR fixes a crash that curently happens given the following input: fortran subroutine caller() real :: x integer :: i !$omp target x = i call callee(x,x) !$omp end target endsubroutine caller subroutine callee(x1,x2) real :: x1, x2 endsubroutine callee The crash happens because the following sequence of events is taken by the OMPIRBuilder: 1. .... 2. An outlined function for the target region is created. At first the outlined function still refers to the SSA values from the original function of the target region. 3. The builder then iterates over the users of SSA values used in the target region to replace them with the corresponding function arguments of outlined function. 4. If the same instruction references the SSA value more than once (say m), all uses of that SSA value are replaced in the instruction. Deleting all m uses of the value. 5. The next m-1 iterations will still iterate over the same instruction dropping the last m-1 actual users of the value. Hence, we collect all users first before modifying them. | 1 年前 | |
[OpenMP] Introduce omp.target_allocmem and omp.target_freemem omp dialect ops. (#145464) This PR introduces two new ops in omp dialect, omp.target_allocmem and omp.target_freemem. omp.target_allocmem: Allocates heap memory on device. Will be lowered to omp_target_alloc call in llvm. omp.target_freemem: Deallocates heap memory on device. Will be lowered to omp+target_free call in llvm. Example: %1 = omp.target_allocmem %device : i32, i64 omp.target_freemem %device, %1 : i32, i64 The work in this PR is C-P/inspired from @ivanradanov commit from coexecute implementation: [Add fir omp target alloc and free ops](https://github.com/ivanradanov/llvm-project/commit/be860ac8baf24b8405e6f396c75d7f0d26375de5) [Lower omp_target_{alloc,free} to llvm](https://github.com/ivanradanov/llvm-project/commit/6e2d584dc93ff99bb89adc28c7afbc2b21c46d39) | 11 个月前 | |
[MLIR][LLVMIR] Always use TargetFolder in IRBuilder (#126929) This is a followup to https://github.com/llvm/llvm-project/pull/126745, generalizing it to always use TargetFolder, including inside function bodies. This avoids generating non-canonical constant expressions that can be folded away. | 1 年前 | |
[flang][OMPIRBuilder][MLIR][llvm] Backend support for atomic control options (#151579) Adding mlir to llvm support for atomic control options. Atomic Control Options are used to specify architectural characteristics to help lowering of atomic operations. The options used are: -f[no-]atomic-remote-memory, -f[no-]atomic-fine-grained-memory, -f[no-]atomic-ignore-denormal-mode. Legacy option -m[no-]unsafe-fp-atomics is aliased to -f[no-]ignore-denormal-mode. More details can be found in https://github.com/llvm/llvm-project/pull/102569. This PR implements the MLIR to LLVM lowering support of atomic control attributes specified with OpenMP atomicUpdateOp. Initial support can be found in PR: https://github.com/llvm/llvm-project/pull/150860 | 11 个月前 | |
[flang][OMPIRBuilder][MLIR][llvm] Backend support for atomic control options (#151579) Adding mlir to llvm support for atomic control options. Atomic Control Options are used to specify architectural characteristics to help lowering of atomic operations. The options used are: -f[no-]atomic-remote-memory, -f[no-]atomic-fine-grained-memory, -f[no-]atomic-ignore-denormal-mode. Legacy option -m[no-]unsafe-fp-atomics is aliased to -f[no-]ignore-denormal-mode. More details can be found in https://github.com/llvm/llvm-project/pull/102569. This PR implements the MLIR to LLVM lowering support of atomic control attributes specified with OpenMP atomicUpdateOp. Initial support can be found in PR: https://github.com/llvm/llvm-project/pull/150860 | 11 个月前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[MLIR][OpenMP] Normalize representation of entry block arg-defining clauses (#109809) This patch updates printing and parsing of operations including clauses that define entry block arguments to the operation's region. This impacts in_reduction, map, private, reduction and task_reduction. The proposed representation to be used by all such clauses is the following: <clause_name>([byref] [@<sym>] %value -> %block_arg [, ...] : <type>[, ...]) { ... } The byref tag is only allowed for reduction-like clauses and the @<sym> is required and only allowed for the private and reduction-like clauses. The map clause does not accept any of these two. This change fixes some currently broken op representations, like omp.teams or omp.sections reduction: omp.teams reduction([byref] @<sym> -> %value : <type>) { ^bb0(%block_arg : <type>): ... } Additionally, it addresses some redundancy in the representation of the previously mentioned cases, as well as e.g. map in omp.target. The problem is that the block argument name after the arrow is not checked in any way, which makes some misleading representations legal: mlir omp.target map_entries(%x -> %arg1, %y -> %arg0, %z -> %doesnt_exist : !llvm.ptr, !llvm.ptr, !llvm.ptr) { ^bb0(%arg0 : !llvm.ptr, %arg1 : !llvm.ptr, %arg2 : !llvm.ptr): ... } In that case, %x maps to %arg0, contrary to what the representation states, and %z maps to %arg2. %doesnt_exist is not resolved, so it would likely cause issues if used anywhere inside of the operation's region. The solution implemented in this patch makes it so that values introduced after the arrow on the representation of these clauses implicitly define the corresponding entry block arguments, removing the potential for these problematic representations. This is what is already implemented for the private and reduction clauses of omp.parallel. There are a couple of consequences of this change: - Entry block argument-defining clauses must come at the end of the operation's representation and in alphabetical order. This is because they are printed/parsed as part of the region and a standardized ordering is needed to reliably match op arguments with their corresponding entry block arguments via the BlockArgOpenMPOpInterface. - We can no longer define per-clause assembly formats to be reused by all operations that take these clauses, since they must be passed to a custom printer including the region and arguments of all other entry block argument-defining clauses. Code duplication and potential for introducing issues is minimized by providing the generic {print,parse}BlockArgRegion helpers and associated structures. MLIR and Flang lowering unit tests are updated due to changes in the order and formatting of impacted operations. | 1 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[MLIR][OpenMP] Introduce overlapped record type map support (#119588) This PR introduces a new additional type of map lowering for record types that Clang currently supports, in which a user can map a top-level record type and then individual members with different mapping, effectively creating a sort of "overlapping" mapping that we attempt to cut around. This is currently most predominantly used in Fortran, when mapping descriptors and there data, we map the descriptor and its data with separate map modifiers and "cut around" the pointer data, so that wedo not overwrite it unless the runtime deems it a neccesary action based on its reference counting mechanism. However, it is a mechanism that will come in handy/trigger when a user explitily maps a record type (derived type or structure) and then explicitly maps a member with a different map type. These additions were predominantly in the OpenMPToLLVMIRTranslation.cpp file and phase, however, one Flang test that checks end-to-end IR compilation (as far as we care for now at least) was altered. 2/3 required PRs to enable declare target to mapping, should look at PR 3/3 to check for full green passes (this one will fail a number due to some dependencies). Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com | 8 个月前 | |
[OMPIRBuilder] Avoid invalid debug location. (#151306) This fixes #147063. I tried to fix this issue in more general way in https://github.com/llvm/llvm-project/pull/147091 but the reviewer suggested to fix the locations which are causing this issue. So this is a more targeted approach. The restoreIP is frequently used in the OMPIRBuilder to change the insert position. This function eventually calls SetInsertPoint(BasicBlock *TheBB, BasicBlock::iterator IP). This function updates the insert point and the debug location. But if the IP is pointing to the end of the TheBB, then the debug location is not updated and we could have a mismatch between insert point and the debug location. The problem can occur in 2 different code patterns. This code below shows the first scenario. 1. auto curPos = builder.saveIP(); 2. builder.restoreIP(/* some new pos */); 3. // generate some code 4. builder.restoreIP(curPos); If curPos points to the end of basic block, we could have a problem. But it is easy one to handle as we have the location before hand and can save the correct debug location before 2 and then restore it after 3. This can be done either manually or using the llvm::InsertPointGuard as shown below. // manual approach auto curPos = builder.saveIP(); llvm::DebugLoc DbgLoc = builder.getCurrentDebugLocation(); builder.restoreIP(/* some new pos */); // generate some code builder.SetCurrentDebugLocation(DbgLoc); builder.restoreIP(curPos); { // using InsertPointGuard llvm::InsertPointGuard IPG(builder); builder.restoreIP(/* some new pos */); // generate some code } This PR fixes one problematic case using the manual approach. For the 2nd scenario, look at the code below. 1. void fn(InsertPointTy allocIP, InsertPointTy codegenIP) { 2. builder.setInsertPoint(allocIP); 3. // generate some alloca 4. builder.setInsertPoint(codegenIP); 5. } The fn can be called from anywhere and we can't assume the debug location of the builder is valid at the start of the function. So if 4 does not update the debug location because the codegenIP points at the end of the block, the rest of the code can end up using the debug location of the allocaIP. Unlike the first case, we don't have a debug location that we can save before hand and restore afterwards. The solution here is to use the location of the last instruction in that block. I have added a wrapper function over restoreIP that could be called for such cases. This PR uses it to fix one problematic case. | 11 个月前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[MLIR][OpenMP] LLVM IR translation of host_eval (#116052) This patch adds support for processing the host_eval clause of omp.target to populate default and runtime kernel launch attributes. Specifically, these related to the num_teams, thread_limit and num_threads clauses attached to operations nested inside of omp.target. As a result, the thread_limit clause of omp.target is also supported. The implementation of initTargetDefaultAttrs() is intended to reflect clang's own processing of multiple constructs and clauses in order to define a default number of teams and threads to be used as kernel attributes and to populate global variables in the target device module. One side effect of this change is that it is no longer possible to translate to LLVM IR target device MLIR modules unless they have a supported target triple. This is because the local getGridValue() function in the OpenMPIRBuilder only works for certain architectures, and it is called whenever the maximum number of threads has not been explicitly defined. This limitation also matches clang. Evaluating the collapsed loop trip count of SPMD and Generic-SPMD kernels remains unsupported. | 1 年前 | |
[MLIR][LLVMIR] Use TargetFolder when creating globals (#126745) The LLVM dialect lowers globals using IRBuilder, relying on it creating constant expressions where possible. As we remove support for more constant expressions (per https://discourse.llvm.org/t/rfc-remove-most-constant-expressions/63179), this can cause issues for cases where the constant expression is no longer supported, and the operation cannot be constant folded without DataLayout being available. In particular, I ran into this issue with flang and the removal of mul constant expressions. Address this by using TargetFolder when creating globals, which will perform DL-aware constant folding. I think it would make sense to also do this in general, but I'm starting with globals where not doing this can result in translation failures. Ideally, globals with these problematic expressions would never be generated in the first place, but there has been little movement on fixing this (https://github.com/llvm/llvm-project/issues/96047). | 1 年前 | |
[Flang][OpenMP][MLIR] Initial declare target to for variables implementation (#119589) While the infrastructure for declare target to/enter and link for variables exists in the MLIR dialect and at the Flang level, the current lowering from MLIR -> LLVM IR isn't in place, it's only in place for variables that have the link clause applied. This PR aims to extend that lowering to an initial implementation that incorporates declare target to as well, which primarily requires changes in the OpenMPToLLVMIRTranslation phase. However, a minor addition to the OpenMP dialect was required to extend the declare target enumerator to include a default None field as well. This also requires a minor change to the Flang lowering's MapInfoFinlization.cpp pass to alter the map type for descriptors to deal with cases where a variable is marked declare to. Currently, when a descriptor variable is mapped declare target to the descriptor component can become attatched, and cannot be updated, this results in issues when an unusual allocation range is specified (effectively an off-by X error). The current solution is to map the descriptor always, as we always require an up-to-date version of this data. However, this also requires an interlinked PR that adds a more intricate type of mapping of structures/record types that clang currently implements, to circumvent the overwriting of the pointer in the descriptor. 3/3 required PRs to enable declare target to mapping, this PR should pass all tests and provide an all green CI. Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com | 8 个月前 | |
[MLIR][OpenMP] Introduce overlapped record type map support (#119588) This PR introduces a new additional type of map lowering for record types that Clang currently supports, in which a user can map a top-level record type and then individual members with different mapping, effectively creating a sort of "overlapping" mapping that we attempt to cut around. This is currently most predominantly used in Fortran, when mapping descriptors and there data, we map the descriptor and its data with separate map modifiers and "cut around" the pointer data, so that wedo not overwrite it unless the runtime deems it a neccesary action based on its reference counting mechanism. However, it is a mechanism that will come in handy/trigger when a user explitily maps a record type (derived type or structure) and then explicitly maps a member with a different map type. These additions were predominantly in the OpenMPToLLVMIRTranslation.cpp file and phase, however, one Flang test that checks end-to-end IR compilation (as far as we care for now at least) was altered. 2/3 required PRs to enable declare target to mapping, should look at PR 3/3 to check for full green passes (this one will fail a number due to some dependencies). Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com | 8 个月前 | |
[OMPIRBuilder] - Fix emitTargetTaskProxyFunc to not generate empty functions (#126958) This is a fix for https://github.com/llvm/llvm-project/issues/126949 There are two issues being fixed here. First, in some cases, OMPIRBuilder generates empty target task proxy functions. This happens when the target kernel doesn't use any stack-allocated data (either no data or only globals). The second problem is encountered when the target task i.e the code that makes the target call spans a single basic block. This usually happens when we do not generate a target or device kernel launch and instead fall back to the host. In such cases, we end up not outlining the target task entirely. This can cause us to call target kernel twice - once via the target task proxy function and a second time via the host fallback This PR fixes both of these problems and updates some tests to catch these problems should this patch fail. | 1 年前 | |
[OMPIRBuilder] - Make offloading input data persist for deferred target tasks (#133499) When we offload to the target, the pointers to data used by the kernel are passed in arrays created by OMPIRBuilder. These arrays of pointers are allocated on the stack on the host. This is fine for the most part because absent the nowait clause, the default behavior is that target tasks are included tasks. That is, the host is blocked until the offloaded target kernel is done. In turn, this means that the host's stack frame is intact and accessing the array of pointers when offloading is safe. However, when nowait is used on the !$ omp target instance, then the target task is a deferred task meaning, the generating task on the host does not have to wait for the target task to finish. In such cases, it is very likely that the stack frame of the function invoking the target call is wound up thereby leading to memory access errors as shown below. `` AMDGPU error: Error in hsa_amd_memory_pool_allocate: HSA_STATUS_ERROR_INVALID_ALLOCATION: The requested allocation is not valid. AMDGPU error: Error in hsa_amd_memory_pool_allocate: HSA_STATUS_ERROR_INVALID_ALLOCATION: The requested allocation is not valid. "PluginInterface" error: Failure to allocate device memory: Failed to allocate from memory manager fort.cod.out: /llvm/llvm-project/offload/plugins-nextgen/common/src/PluginInterface.cpp:1434: Error llvm::omp::target::plugin::PinnedAllocationMapTy::lockMappedHostBuffer(void *, size_t): Assertion HstPtr && "Invalid pointer"' failed. Aborted (core dumped) `` This PR implements support in OMPIRBuilder` to store these arrays of pointers in the task structure that is passed to the target task thereby ensuring it is available to the target task when the target task is eventually scheduled. --------- Co-authored-by: Sergio Afonso <safonsof@amd.com> | 1 年前 | |
[MLIR][LLVMIR] Always use TargetFolder in IRBuilder (#126929) This is a followup to https://github.com/llvm/llvm-project/pull/126745, generalizing it to always use TargetFolder, including inside function bodies. This avoids generating non-canonical constant expressions that can be folded away. | 1 年前 | |
[OpenMP][IRBuilder] Handle target directives with both if & nowait (#125029) This fixes a bug when a target directive has both an if and a nowait clauses. The bug happens because we tried to emitKernelLaunch for else branch of the if clause. | 1 年前 | |
[OMPIRBuilder][MLIR] Add support for target 'if' clause (#122478) This patch implements support for handling the 'if' clause of OpenMP 'target' constructs in the OMPIRBuilder and updates MLIR to LLVM IR translation of the omp.target MLIR operation to make use of this new feature. | 1 年前 | |
[flang][OpenMP] Skip runtime mapping with no offload targets (#145594) When no offload targets are specified flang will avoid offloading for "target" constructs, but not "target data" constructs. This patch makes the behavior consistent across all offload-related operations. While ignoring "target" may produce semantically incorrect code, it may still be a useful debugging tool. -- This reinstates commits 6ba1955 and 349f8d6, reverted due to compilation failures in the gfortran test suite. These build problems were caused by an unrelated issue (https://github.com/llvm/llvm-project/issues/145558) which is now fixed. Ref: https://github.com/llvm/llvm-project/pull/144534 | 1 年前 | |
[AMDGPU] Set AS8 address width to 48 bits Of the 128-bits of buffer descriptor only 48 bits are address bits, so following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54, the logic conclusion is to set the index width to 48 bits instead of the current value of 128. Most of the test changes are mechanical datalayout updates, but there is one actual change: the ptrmask test now uses .i48 instead of .i128 and I had to update SelectionDAGBuilder to correctly extend the mask. Reviewed By: krzysz00 Pull Request: https://github.com/llvm/llvm-project/pull/139419 | 1 年前 | |
[flang][OpenMP] Support multi-block reduction combiner regions on the GPU (#156837) Fixes a bug related to insertion points when inlining multi-block combiner reduction regions. The IP at the end of the inlined region was not used resulting in emitting BBs with multiple terminators. PR stack: - https://github.com/llvm/llvm-project/pull/155754 - https://github.com/llvm/llvm-project/pull/155987 - https://github.com/llvm/llvm-project/pull/155992 - https://github.com/llvm/llvm-project/pull/155993 - https://github.com/llvm/llvm-project/pull/157638 - https://github.com/llvm/llvm-project/pull/156610 - https://github.com/llvm/llvm-project/pull/156837 ◀️ | 10 个月前 | |
[AMDGPU] Set AS8 address width to 48 bits Of the 128-bits of buffer descriptor only 48 bits are address bits, so following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54, the logic conclusion is to set the index width to 48 bits instead of the current value of 128. Most of the test changes are mechanical datalayout updates, but there is one actual change: the ptrmask test now uses .i48 instead of .i128 and I had to update SelectionDAGBuilder to correctly extend the mask. Reviewed By: krzysz00 Pull Request: https://github.com/llvm/llvm-project/pull/139419 | 1 年前 | |
[MLIR][OpenMP] Improve omp.map.info verification (#132066) This patch makes the map_type and map_capture_type arguments of the omp.map.info operation required, which was already an invariant being verified by its users via verifyMapClause(). This makes it clearer, as getters no longer return misleading std::optional values. Checks for the mapper_id argument are moved to a verifier for the operation, rather than being checked by users. Functionally NFC, but not marked as such due to a reordering of arguments in the assembly format of omp.map.info. | 1 年前 | |
[OpenMP][MLIR] Descriptor explicit member map lowering changes (#113556) This is one of 3 PRs in a PR stack that aims to add support for explicit mapping of allocatable members in derived types. The primary changes in this PR are the OpenMPToLLVMIRTranslation.cpp changes, which are small and seek to alter the current member mapping to add an additional map insertion for pointers. Effectively, if the member is a pointer (currently indicated by having a varPtrPtr field) we add an additional map for the pointer and then alter the subsequent mapping of the member (the data) to utilise the member rather than the parents base pointer. This appears to be necessary in certain cases when mapping pointer data within record types to avoid segfaulting on device (due to incorrect data mapping). In general this record type mapping may be simplifiable in the future. There are also additions of tests which should help to showcase the affect of the changes above. | 1 年前 | |
[OMPIRBuilder] - Fix emitTargetTaskProxyFunc to not generate empty functions (#126958) This is a fix for https://github.com/llvm/llvm-project/issues/126949 There are two issues being fixed here. First, in some cases, OMPIRBuilder generates empty target task proxy functions. This happens when the target kernel doesn't use any stack-allocated data (either no data or only globals). The second problem is encountered when the target task i.e the code that makes the target call spans a single basic block. This usually happens when we do not generate a target or device kernel launch and instead fall back to the host. In such cases, we end up not outlining the target task entirely. This can cause us to call target kernel twice - once via the target task proxy function and a second time via the host fallback This PR fixes both of these problems and updates some tests to catch these problems should this patch fail. | 1 年前 | |
[OMPIRBuilder] - Make offloading input data persist for deferred target tasks (#133499) When we offload to the target, the pointers to data used by the kernel are passed in arrays created by OMPIRBuilder. These arrays of pointers are allocated on the stack on the host. This is fine for the most part because absent the nowait clause, the default behavior is that target tasks are included tasks. That is, the host is blocked until the offloaded target kernel is done. In turn, this means that the host's stack frame is intact and accessing the array of pointers when offloading is safe. However, when nowait is used on the !$ omp target instance, then the target task is a deferred task meaning, the generating task on the host does not have to wait for the target task to finish. In such cases, it is very likely that the stack frame of the function invoking the target call is wound up thereby leading to memory access errors as shown below. `` AMDGPU error: Error in hsa_amd_memory_pool_allocate: HSA_STATUS_ERROR_INVALID_ALLOCATION: The requested allocation is not valid. AMDGPU error: Error in hsa_amd_memory_pool_allocate: HSA_STATUS_ERROR_INVALID_ALLOCATION: The requested allocation is not valid. "PluginInterface" error: Failure to allocate device memory: Failed to allocate from memory manager fort.cod.out: /llvm/llvm-project/offload/plugins-nextgen/common/src/PluginInterface.cpp:1434: Error llvm::omp::target::plugin::PinnedAllocationMapTy::lockMappedHostBuffer(void *, size_t): Assertion HstPtr && "Invalid pointer"' failed. Aborted (core dumped) `` This PR implements support in OMPIRBuilder` to store these arrays of pointers in the task structure that is passed to the target task thereby ensuring it is available to the target task when the target task is eventually scheduled. --------- Co-authored-by: Sergio Afonso <safonsof@amd.com> | 1 年前 | |
[MLIR][OpenMP] Introduce overlapped record type map support (#119588) This PR introduces a new additional type of map lowering for record types that Clang currently supports, in which a user can map a top-level record type and then individual members with different mapping, effectively creating a sort of "overlapping" mapping that we attempt to cut around. This is currently most predominantly used in Fortran, when mapping descriptors and there data, we map the descriptor and its data with separate map modifiers and "cut around" the pointer data, so that wedo not overwrite it unless the runtime deems it a neccesary action based on its reference counting mechanism. However, it is a mechanism that will come in handy/trigger when a user explitily maps a record type (derived type or structure) and then explicitly maps a member with a different map type. These additions were predominantly in the OpenMPToLLVMIRTranslation.cpp file and phase, however, one Flang test that checks end-to-end IR compilation (as far as we care for now at least) was altered. 2/3 required PRs to enable declare target to mapping, should look at PR 3/3 to check for full green passes (this one will fail a number due to some dependencies). Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com | 8 个月前 | |
[MLIR][OpenMP] Improve omp.map.info verification (#132066) This patch makes the map_type and map_capture_type arguments of the omp.map.info operation required, which was already an invariant being verified by its users via verifyMapClause(). This makes it clearer, as getters no longer return misleading std::optional values. Checks for the mapper_id argument are moved to a verifier for the operation, rather than being checked by users. Functionally NFC, but not marked as such due to a reordering of arguments in the assembly format of omp.map.info. | 1 年前 | |
[MLIR][OpenMP] Introduce overlapped record type map support (#119588) This PR introduces a new additional type of map lowering for record types that Clang currently supports, in which a user can map a top-level record type and then individual members with different mapping, effectively creating a sort of "overlapping" mapping that we attempt to cut around. This is currently most predominantly used in Fortran, when mapping descriptors and there data, we map the descriptor and its data with separate map modifiers and "cut around" the pointer data, so that wedo not overwrite it unless the runtime deems it a neccesary action based on its reference counting mechanism. However, it is a mechanism that will come in handy/trigger when a user explitily maps a record type (derived type or structure) and then explicitly maps a member with a different map type. These additions were predominantly in the OpenMPToLLVMIRTranslation.cpp file and phase, however, one Flang test that checks end-to-end IR compilation (as far as we care for now at least) was altered. 2/3 required PRs to enable declare target to mapping, should look at PR 3/3 to check for full green passes (this one will fail a number due to some dependencies). Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com | 8 个月前 | |
[flang][debug] Generate DISubprogramAttr for omp::TargetOp. (#146532) This is combination of https://github.com/llvm/llvm-project/pull/138149 and https://github.com/llvm/llvm-project/pull/138039 which were opened separately for ease of reviewing. Only other change is adjustments in 2 tests which have gone in since. There are DeclareOp present for the variables mapped into target region. That allow us to generate debug information for them. But the TargetOp is still part of parent function and those variables get the parent function's DISubprogram as a scope. In OMPIRBuilder, a new function is created for the TargetOp. We also create a new DISubprogram for it. All the variables that were in the target region now have to be updated to have the correct scope. This after the fact updating of debug information becomes very difficult in certain cases. Take the example of variable arrays. The type of those arrays depend on the artificial DILocalVariable(s) which hold the size(s) of the array. This new function will now require that we generate the new variable and and new types. Similar issue exist for character type variables too. To avoid this after the fact updating, this PR generates a DISubprogramAttr for the TargetOp while generating the debug info in flang. Then we don't need to generate a DISubprogram in OMPIRBuilder. This change is made a bit more complicated by the the fact that in new scheme, the debug location already points to the new DISubprogram by the time it reaches convertOmpTarget. But we need some code generation in the parent function so we have to carefully manage the debug locations. This fixes issue #134991. | 1 年前 | |
[AMDGPU] Set AS8 address width to 48 bits Of the 128-bits of buffer descriptor only 48 bits are address bits, so following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54, the logic conclusion is to set the index width to 48 bits instead of the current value of 128. Most of the test changes are mechanical datalayout updates, but there is one actual change: the ptrmask test now uses .i48 instead of .i128 and I had to update SelectionDAGBuilder to correctly extend the mask. Reviewed By: krzysz00 Pull Request: https://github.com/llvm/llvm-project/pull/139419 | 1 年前 | |
[Offload] Add oneInterationPerThread param to loop device RTL (#151959) Currently, Flang can generate no-loop kernels for all OpenMP target kernels in the program if the flags -fopenmp-assume-teams-oversubscription or -fopenmp-assume-threads-oversubscription are set. If we add an additional parameter, we can choose in the future which OpenMP kernels should be generated as no-loop kernels. This PR doesn't modify current behavior of oversubscription flags. RFC for no-loop kernels: https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517 | 11 个月前 | |
[AMDGPU] Set AS8 address width to 48 bits Of the 128-bits of buffer descriptor only 48 bits are address bits, so following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54, the logic conclusion is to set the index width to 48 bits instead of the current value of 128. Most of the test changes are mechanical datalayout updates, but there is one actual change: the ptrmask test now uses .i48 instead of .i128 and I had to update SelectionDAGBuilder to correctly extend the mask. Reviewed By: krzysz00 Pull Request: https://github.com/llvm/llvm-project/pull/139419 | 1 年前 | |
[MLIR][OpenMP] Normalize representation of entry block arg-defining clauses (#109809) This patch updates printing and parsing of operations including clauses that define entry block arguments to the operation's region. This impacts in_reduction, map, private, reduction and task_reduction. The proposed representation to be used by all such clauses is the following: <clause_name>([byref] [@<sym>] %value -> %block_arg [, ...] : <type>[, ...]) { ... } The byref tag is only allowed for reduction-like clauses and the @<sym> is required and only allowed for the private and reduction-like clauses. The map clause does not accept any of these two. This change fixes some currently broken op representations, like omp.teams or omp.sections reduction: omp.teams reduction([byref] @<sym> -> %value : <type>) { ^bb0(%block_arg : <type>): ... } Additionally, it addresses some redundancy in the representation of the previously mentioned cases, as well as e.g. map in omp.target. The problem is that the block argument name after the arrow is not checked in any way, which makes some misleading representations legal: mlir omp.target map_entries(%x -> %arg1, %y -> %arg0, %z -> %doesnt_exist : !llvm.ptr, !llvm.ptr, !llvm.ptr) { ^bb0(%arg0 : !llvm.ptr, %arg1 : !llvm.ptr, %arg2 : !llvm.ptr): ... } In that case, %x maps to %arg0, contrary to what the representation states, and %z maps to %arg2. %doesnt_exist is not resolved, so it would likely cause issues if used anywhere inside of the operation's region. The solution implemented in this patch makes it so that values introduced after the arrow on the representation of these clauses implicitly define the corresponding entry block arguments, removing the potential for these problematic representations. This is what is already implemented for the private and reduction clauses of omp.parallel. There are a couple of consequences of this change: - Entry block argument-defining clauses must come at the end of the operation's representation and in alphabetical order. This is because they are printed/parsed as part of the region and a standardized ordering is needed to reliably match op arguments with their corresponding entry block arguments via the BlockArgOpenMPOpInterface. - We can no longer define per-clause assembly formats to be reused by all operations that take these clauses, since they must be passed to a custom printer including the region and arguments of all other entry block argument-defining clauses. Code duplication and potential for introducing issues is minimized by providing the generic {print,parse}BlockArgRegion helpers and associated structures. MLIR and Flang lowering unit tests are updated due to changes in the order and formatting of impacted operations. | 1 年前 | |
[MLIR][OpenMP] Introduce overlapped record type map support (#119588) This PR introduces a new additional type of map lowering for record types that Clang currently supports, in which a user can map a top-level record type and then individual members with different mapping, effectively creating a sort of "overlapping" mapping that we attempt to cut around. This is currently most predominantly used in Fortran, when mapping descriptors and there data, we map the descriptor and its data with separate map modifiers and "cut around" the pointer data, so that wedo not overwrite it unless the runtime deems it a neccesary action based on its reference counting mechanism. However, it is a mechanism that will come in handy/trigger when a user explitily maps a record type (derived type or structure) and then explicitly maps a member with a different map type. These additions were predominantly in the OpenMPToLLVMIRTranslation.cpp file and phase, however, one Flang test that checks end-to-end IR compilation (as far as we care for now at least) was altered. 2/3 required PRs to enable declare target to mapping, should look at PR 3/3 to check for full green passes (this one will fail a number due to some dependencies). Co-authored-by: Raghu Maddhipatla raghu.maddhipatla@amd.com | 8 个月前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[MLIR][OpenMP] Normalize representation of entry block arg-defining clauses (#109809) This patch updates printing and parsing of operations including clauses that define entry block arguments to the operation's region. This impacts in_reduction, map, private, reduction and task_reduction. The proposed representation to be used by all such clauses is the following: <clause_name>([byref] [@<sym>] %value -> %block_arg [, ...] : <type>[, ...]) { ... } The byref tag is only allowed for reduction-like clauses and the @<sym> is required and only allowed for the private and reduction-like clauses. The map clause does not accept any of these two. This change fixes some currently broken op representations, like omp.teams or omp.sections reduction: omp.teams reduction([byref] @<sym> -> %value : <type>) { ^bb0(%block_arg : <type>): ... } Additionally, it addresses some redundancy in the representation of the previously mentioned cases, as well as e.g. map in omp.target. The problem is that the block argument name after the arrow is not checked in any way, which makes some misleading representations legal: mlir omp.target map_entries(%x -> %arg1, %y -> %arg0, %z -> %doesnt_exist : !llvm.ptr, !llvm.ptr, !llvm.ptr) { ^bb0(%arg0 : !llvm.ptr, %arg1 : !llvm.ptr, %arg2 : !llvm.ptr): ... } In that case, %x maps to %arg0, contrary to what the representation states, and %z maps to %arg2. %doesnt_exist is not resolved, so it would likely cause issues if used anywhere inside of the operation's region. The solution implemented in this patch makes it so that values introduced after the arrow on the representation of these clauses implicitly define the corresponding entry block arguments, removing the potential for these problematic representations. This is what is already implemented for the private and reduction clauses of omp.parallel. There are a couple of consequences of this change: - Entry block argument-defining clauses must come at the end of the operation's representation and in alphabetical order. This is because they are printed/parsed as part of the region and a standardized ordering is needed to reliably match op arguments with their corresponding entry block arguments via the BlockArgOpenMPOpInterface. - We can no longer define per-clause assembly formats to be reused by all operations that take these clauses, since they must be passed to a custom printer including the region and arguments of all other entry block argument-defining clauses. Code duplication and potential for introducing issues is minimized by providing the generic {print,parse}BlockArgRegion helpers and associated structures. MLIR and Flang lowering unit tests are updated due to changes in the order and formatting of impacted operations. | 1 年前 | |
[MLIR][OpenMP] Normalize representation of entry block arg-defining clauses (#109809) This patch updates printing and parsing of operations including clauses that define entry block arguments to the operation's region. This impacts in_reduction, map, private, reduction and task_reduction. The proposed representation to be used by all such clauses is the following: <clause_name>([byref] [@<sym>] %value -> %block_arg [, ...] : <type>[, ...]) { ... } The byref tag is only allowed for reduction-like clauses and the @<sym> is required and only allowed for the private and reduction-like clauses. The map clause does not accept any of these two. This change fixes some currently broken op representations, like omp.teams or omp.sections reduction: omp.teams reduction([byref] @<sym> -> %value : <type>) { ^bb0(%block_arg : <type>): ... } Additionally, it addresses some redundancy in the representation of the previously mentioned cases, as well as e.g. map in omp.target. The problem is that the block argument name after the arrow is not checked in any way, which makes some misleading representations legal: mlir omp.target map_entries(%x -> %arg1, %y -> %arg0, %z -> %doesnt_exist : !llvm.ptr, !llvm.ptr, !llvm.ptr) { ^bb0(%arg0 : !llvm.ptr, %arg1 : !llvm.ptr, %arg2 : !llvm.ptr): ... } In that case, %x maps to %arg0, contrary to what the representation states, and %z maps to %arg2. %doesnt_exist is not resolved, so it would likely cause issues if used anywhere inside of the operation's region. The solution implemented in this patch makes it so that values introduced after the arrow on the representation of these clauses implicitly define the corresponding entry block arguments, removing the potential for these problematic representations. This is what is already implemented for the private and reduction clauses of omp.parallel. There are a couple of consequences of this change: - Entry block argument-defining clauses must come at the end of the operation's representation and in alphabetical order. This is because they are printed/parsed as part of the region and a standardized ordering is needed to reliably match op arguments with their corresponding entry block arguments via the BlockArgOpenMPOpInterface. - We can no longer define per-clause assembly formats to be reused by all operations that take these clauses, since they must be passed to a custom printer including the region and arguments of all other entry block argument-defining clauses. Code duplication and potential for introducing issues is minimized by providing the generic {print,parse}BlockArgRegion helpers and associated structures. MLIR and Flang lowering unit tests are updated due to changes in the order and formatting of impacted operations. | 1 年前 | |
[MLIR][OpenMP] Normalize representation of entry block arg-defining clauses (#109809) This patch updates printing and parsing of operations including clauses that define entry block arguments to the operation's region. This impacts in_reduction, map, private, reduction and task_reduction. The proposed representation to be used by all such clauses is the following: <clause_name>([byref] [@<sym>] %value -> %block_arg [, ...] : <type>[, ...]) { ... } The byref tag is only allowed for reduction-like clauses and the @<sym> is required and only allowed for the private and reduction-like clauses. The map clause does not accept any of these two. This change fixes some currently broken op representations, like omp.teams or omp.sections reduction: omp.teams reduction([byref] @<sym> -> %value : <type>) { ^bb0(%block_arg : <type>): ... } Additionally, it addresses some redundancy in the representation of the previously mentioned cases, as well as e.g. map in omp.target. The problem is that the block argument name after the arrow is not checked in any way, which makes some misleading representations legal: mlir omp.target map_entries(%x -> %arg1, %y -> %arg0, %z -> %doesnt_exist : !llvm.ptr, !llvm.ptr, !llvm.ptr) { ^bb0(%arg0 : !llvm.ptr, %arg1 : !llvm.ptr, %arg2 : !llvm.ptr): ... } In that case, %x maps to %arg0, contrary to what the representation states, and %z maps to %arg2. %doesnt_exist is not resolved, so it would likely cause issues if used anywhere inside of the operation's region. The solution implemented in this patch makes it so that values introduced after the arrow on the representation of these clauses implicitly define the corresponding entry block arguments, removing the potential for these problematic representations. This is what is already implemented for the private and reduction clauses of omp.parallel. There are a couple of consequences of this change: - Entry block argument-defining clauses must come at the end of the operation's representation and in alphabetical order. This is because they are printed/parsed as part of the region and a standardized ordering is needed to reliably match op arguments with their corresponding entry block arguments via the BlockArgOpenMPOpInterface. - We can no longer define per-clause assembly formats to be reused by all operations that take these clauses, since they must be passed to a custom printer including the region and arguments of all other entry block argument-defining clauses. Code duplication and potential for introducing issues is minimized by providing the generic {print,parse}BlockArgRegion helpers and associated structures. MLIR and Flang lowering unit tests are updated due to changes in the order and formatting of impacted operations. | 1 年前 | |
[OpenMP][OMPIRBuilder] Use runtime CC for runtime calls (#168608) Some targets have a specific calling convention that should be used for generated calls to runtime functions. Pass that down and use it. Signed-off-by: Nick Sarnie <nick.sarnie@intel.com> | 8 个月前 | |
[Flang][MLIR][OpenMP] Use function-attached target attributes for OpenMP lowering (#78291) This patch removes the omp.target module attribute, since the information it held on the target CPU and features is available through the fir.target_cpu and fir.target_features module attributes. Target outlining during the MLIR to LLVM IR translation stage is updated, so that these attributes, at that point available as llvm.func attributes, are passed along to the newly created function. | 2 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[AMDGPU] Set AS8 address width to 48 bits Of the 128-bits of buffer descriptor only 48 bits are address bits, so following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54, the logic conclusion is to set the index width to 48 bits instead of the current value of 128. Most of the test changes are mechanical datalayout updates, but there is one actual change: the ptrmask test now uses .i48 instead of .i128 and I had to update SelectionDAGBuilder to correctly extend the mask. Reviewed By: krzysz00 Pull Request: https://github.com/llvm/llvm-project/pull/139419 | 1 年前 | |
[MLIR][OpenMP] Skip host omp ops when compiling for the target device (#85239) This patch separates the lowering dispatch for host and target devices. For the target device, if the current operation is not a top-level operation (e.g. omp.target) or is inside a target device code region it will be ignored, since it belongs to the host code. This is an alternative approach to #84611, the new test in this PR was taken from there. | 2 年前 | |
[AMDGPU] Set AS8 address width to 48 bits Of the 128-bits of buffer descriptor only 48 bits are address bits, so following the discussion on https://discourse.llvm.org/t/clarifiying-the-semantics-of-ptrtoint/83987/54, the logic conclusion is to set the index width to 48 bits instead of the current value of 128. Most of the test changes are mechanical datalayout updates, but there is one actual change: the ptrmask test now uses .i48 instead of .i128 and I had to update SelectionDAGBuilder to correctly extend the mask. Reviewed By: krzysz00 Pull Request: https://github.com/llvm/llvm-project/pull/139419 | 1 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[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 个月前 | |
[Offload] Add oneInterationPerThread param to loop device RTL (#151959) Currently, Flang can generate no-loop kernels for all OpenMP target kernels in the program if the flags -fopenmp-assume-teams-oversubscription or -fopenmp-assume-threads-oversubscription are set. If we add an additional parameter, we can choose in the future which OpenMP kernels should be generated as no-loop kernels. This PR doesn't modify current behavior of oversubscription flags. RFC for no-loop kernels: https://discourse.llvm.org/t/rfc-no-loop-mode-for-openmp-gpu-kernels/87517 | 11 个月前 | |
[flang][OpenMP] Skip runtime mapping with no offload targets (#145594) When no offload targets are specified flang will avoid offloading for "target" constructs, but not "target data" constructs. This patch makes the behavior consistent across all offload-related operations. While ignoring "target" may produce semantically incorrect code, it may still be a useful debugging tool. -- This reinstates commits 6ba1955 and 349f8d6, reverted due to compilation failures in the gfortran test suite. These build problems were caused by an unrelated issue (https://github.com/llvm/llvm-project/issues/145558) which is now fixed. Ref: https://github.com/llvm/llvm-project/pull/144534 | 1 年前 | |
[mlir][OpenMP] - Honor dependencies in code-generation of the if clause in omp.task correctly (#90891) This patch fixes the code generation of the if clause, specifically when the condition evaluates to false and when the task directive has the depend clause on it. When the if clause of a task construct evaluates to false, then the task is an undeferred task. This undeferred task still has to honor dependencies. Previously, the OpenMPIRbuilder didn't honor dependencies. This patch fixes that. Fixes https://github.com/llvm/llvm-project/issues/90869 | 2 年前 | |
[IR] Replace alignment argument with attribute on masked intrinsics (#163802) The masked.load, masked.store, masked.gather and masked.scatter intrinsics currently accept a separate alignment immarg. Replace this with an align attribute on the pointer / vector of pointers argument. This is the standard representation for alignment information on intrinsics, and is already used by all other memory intrinsics. This means the signatures now match llvm.expandload, llvm.vp.load, etc. (Things like llvm.memcpy used to have a separate alignment argument as well, but were already migrated a long time ago.) It's worth noting that the masked.gather and masked.scatter intrinsics previously accepted a zero alignment to indicate the ABI type alignment of the element type. This special case is gone now: If the align attribute is omitted, the implied alignment is 1, as usual. If ABI alignment is desired, it needs to be explicitly emitted (which the IRBuilder API already requires anyway). | 9 个月前 | |
[MLIR][acc] Introduce varType to acc data clause operations (#119007) The acc data clause operations hold an operand named varPtr. This was intended to hold a pointer to a variable - where the element type of that pointer specifies the type of the variable. However, for both memref and llvm dialects, this assumption is not true. This is because memref element type for cases like memref<10xf32> is simply f32 and for LLVM, after opaque pointers, the variable type is no longer recoverable. Thus, introduce varType to ensure that appropriate semantics are kept. Both the parser and printer for this new type attribute allow it to not be specified in cases where a dialect's getElementType() applied to varPtr's type has a recoverable type. And more specifically, for FIR, no changes are needed in the MLIR unit tests. | 1 年前 | |
| 1 年前 | ||
[OMPIRBuilder] Don't outline DISTRIBUTE on CPUs (#158317) We use different OpenMP runtime functions on CPU and target offload. The one used for DISTRIBUTE on target offload needs a function pointer to an offloaded function, but the one on CPU doesn't. This caused unnessecary overhead on CPUs because SHARED or FIRSTPRIVATE memory from the surrounding context has to be packaged into a context structure just for an ordinary function call (which would hopefully eventually get inlined). This also makes the IR harder to read. | 10 个月前 | |
[mlir][OpenMP] cancel(lation point) taskgroup LLVMIR (#137841) A cancel or cancellation point for taskgroup is always nested inside of a task inside of the taskgroup. For the task which is cancelled, it is that task which needs to be cleaned up: not the owning taskgroup. Therefore the cancellation branch handler is done in the conversion of the task not in conversion of taskgroup. I added a firstprivate clause to the test for cancel taskgroup to demonstrate that the block being branched to is the same block where mandatory cleanup code is added. Cancellation point follows exactly the same code path. | 1 年前 | |
[mlir][OpenMP] cancel(lation point) taskgroup LLVMIR (#137841) A cancel or cancellation point for taskgroup is always nested inside of a task inside of the taskgroup. For the task which is cancelled, it is that task which needs to be cleaned up: not the owning taskgroup. Therefore the cancellation branch handler is done in the conversion of the task not in conversion of taskgroup. I added a firstprivate clause to the test for cancel taskgroup to demonstrate that the block being branched to is the same block where mandatory cleanup code is added. Cancellation point follows exactly the same code path. | 1 年前 | |
[MLIR][OpenMP] Add canonical loop LLVM-IR lowering (#147069) Support for translating the operations introduced in #144785 to LLVM-IR. In order to keep the lowering simple, OpenMPIRBuider::unrollLoopHeuristic is applied when encountering the omp.unroll_heuristic op. As a result, the operation that unrolling is applied to (omp.canonical_loop) must have been emitted before even though logically there is no such requirement. Eventually, all transformations on a loop must be applied directly after emitting omp.canonical_loop, i.e. future transformations must be looked-up when encountering omp.canonical_loop itself. This is because many OpenMPIRBuilder methods (e.g. createParallel) expect all the region code to be emitted withing a callback. In the case of createParallel, the region code is getting outlined into a new function. Therefore, making the operation order a formal requirement would not make the implementation any easier. | 1 年前 | |
[Flang] Add standalone tile support (#160298) Add support for the standalone OpenMP tile construct: f90 !$omp tile sizes(...) DO i = 1, 100 ... This is complementary to #143715 which added support for the tile construct as part of another loop-associated construct such as worksharing-loop, distribute, etc. | 9 个月前 | |
[Flang] Add standalone tile support (#160298) Add support for the standalone OpenMP tile construct: f90 !$omp tile sizes(...) DO i = 1, 100 ... This is complementary to #143715 which added support for the tile construct as part of another loop-associated construct such as worksharing-loop, distribute, etc. | 9 个月前 | |
[MLIR][OpenMP] Add canonical loop LLVM-IR lowering (#147069) Support for translating the operations introduced in #144785 to LLVM-IR. In order to keep the lowering simple, OpenMPIRBuider::unrollLoopHeuristic is applied when encountering the omp.unroll_heuristic op. As a result, the operation that unrolling is applied to (omp.canonical_loop) must have been emitted before even though logically there is no such requirement. Eventually, all transformations on a loop must be applied directly after emitting omp.canonical_loop, i.e. future transformations must be looked-up when encountering omp.canonical_loop itself. This is because many OpenMPIRBuilder methods (e.g. createParallel) expect all the region code to be emitted withing a callback. In the case of createParallel, the region code is getting outlined into a new function. Therefore, making the operation order a formal requirement would not make the implementation any easier. | 1 年前 | |
[MLIR][OpenMP] Add canonical loop LLVM-IR lowering (#147069) Support for translating the operations introduced in #144785 to LLVM-IR. In order to keep the lowering simple, OpenMPIRBuider::unrollLoopHeuristic is applied when encountering the omp.unroll_heuristic op. As a result, the operation that unrolling is applied to (omp.canonical_loop) must have been emitted before even though logically there is no such requirement. Eventually, all transformations on a loop must be applied directly after emitting omp.canonical_loop, i.e. future transformations must be looked-up when encountering omp.canonical_loop itself. This is because many OpenMPIRBuilder methods (e.g. createParallel) expect all the region code to be emitted withing a callback. In the case of createParallel, the region code is getting outlined into a new function. Therefore, making the operation order a formal requirement would not make the implementation any easier. | 1 年前 | |
[mlir][OpenMP] Allow composite SIMD REDUCTION and IF (#147568) Reduction support: https://github.com/llvm/llvm-project/pull/146671 If Support is fixed in this PR The problem for the IF clause in composite constructs was that wsloop and simd both operate on the same CanonicalLoopInfo structure: with the SIMD processed first, followed by the wsloop. Previously the IF clause generated code like if (cond) { while (...) { simd_loop_body; } } else { while (...) { nonsimd_loop_body; } } The problem with this is that this invalidates the CanonicalLoopInfo structure to be processed by the wsloop later. To avoid this, in this patch I preserve the original loop, moving the IF clause inside of the loop: while (...) { if (cond) { simd_loop_body; } else { non_simd_loop_body; } } On simple examples I tried LLVM was able to hoist the if condition outside of the loop at -O3. The disadvantage of this is that we cannot add the llvm.loop.vectorize.enable attribute on either the SIMD or non-SIMD loops because they both share a loop back edge. There's no way of solving this without keeping the old design of having two different loops: which cannot be represented using only one CanonicalLoopInfo structure. I don't think the presence or absence of this attribute makes much difference. In my testing it is the llvm.loop.parallel_access metadata which makes the difference to vectorization. LLVM will vectorize if legal whether or not this attribute is there in the TRUE branch. In the FALSE branch this means the loop might be vectorized even when the condition is false: but I think this is still standards compliant: OpenMP 6.0 says that when the if clause is false that should be treated like the SIMDLEN clause is one. The SIMDLEN clause is defined as a "hint". For the same reason, SIMDLEN and SAFELEN clauses are silently ignored when SIMD IF is used. I think it is better to implement SIMD IF and ignore SIMDLEN and SAFELEN and some vectorization encouragement metadata when combined with IF than to ignore IF because IF could have correctness consequences whereas the rest are optimiztion hints. For example, the user might use the IF clause to disable SIMD programatically when it is known not safe to vectorize the loop. In this case it is not at all safe to add the parallel access or SAFELEN metadata. | 1 年前 | |
[flang][OpenMP] Skip runtime mapping with no offload targets (#145594) When no offload targets are specified flang will avoid offloading for "target" constructs, but not "target data" constructs. This patch makes the behavior consistent across all offload-related operations. While ignoring "target" may produce semantically incorrect code, it may still be a useful debugging tool. -- This reinstates commits 6ba1955 and 349f8d6, reverted due to compilation failures in the gfortran test suite. These build problems were caused by an unrelated issue (https://github.com/llvm/llvm-project/issues/145558) which is now fixed. Ref: https://github.com/llvm/llvm-project/pull/144534 | 1 年前 | |
[OMPIRBuilder] Don't outline DISTRIBUTE on CPUs (#158317) We use different OpenMP runtime functions on CPU and target offload. The one used for DISTRIBUTE on target offload needs a function pointer to an offloaded function, but the one on CPU doesn't. This caused unnessecary overhead on CPUs because SHARED or FIRSTPRIVATE memory from the surrounding context has to be packaged into a context structure just for an ordinary function call (which would hopefully eventually get inlined). This also makes the IR harder to read. | 10 个月前 | |
[OpenMP][flang][MLIR] Decouple alloc, init, and copy regions for omp.private|declare_reduction ops (#125699) This PR changes the emitted block structure of alloc, init, and copy regions for omp.private and omp.declare_reduction ops a little bit. In particular, this decouples init and copy regions from the alloca insertion-point. The main motivation is fix "Instruction does not dominate all uses!" errors that happen specially when an init region uses a value from the OpenMP region it is being inlined into. The issue happens because, previous to this PR, we inline the init region right after the latest alloc block (since we used the alloca IP); which in some cases (see exmaple below), is too early and causes the use dominance issue. Example that would break without this PR (when delayed privatization is enabled for omp.wsloops): fortran subroutine test2 (xyz) integer :: i integer :: xyz(:) !$omp target map(from:xyz) !$omp do private(xyz) do i = 1, 10 xyz(i) = i end do !$omp end target end subroutine | 1 年前 | |
[flang][openacc][openmp] Support implicit casting on the atomic interface (#114390) ACCMP atomics do not support type conversion. Specifically, I have encountered semantically incorrect code for atomic reads. Example: program main implicit none real(8) :: n integer :: x x = 1.0 !$acc atomic capture n = x x = n !$acc end atomic end program main We have this error when compiling it with flang-new: error: loc("rep.f90":6:9): expected three operations in atomic.capture region (one terminator, and two atomic ops) Yet, in the following generated FIR code, we observe three issues. 1. fir.convert intrudes into the capture region. 2. An incorrect temporary (%2) is being updated instead of n. 3. If we allow n in place of %2, the operand types of atomic.read do not match. Introducing a !fir.ref<i32> -> !fir.ref<f64> conversion on x is inaccurate because we need to convert the value of x. %2 = "fir.alloca"() <{in_type = i32, operandSegmentSizes = array<i32: 0, 0>}> : () -> !fir.ref<i32> %3 = "fir.alloca"() <{bindc_name = "n", in_type = f64, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFEn"}> : () -> !fir.ref<f64> %4:2 = "hlfir.declare"(%3) <{operandSegmentSizes = array<i32: 1, 0, 0, 0>, uniq_name = "_QFEn"}> : (!fir.ref<f64>) -> (!fir.ref<f64>, !fir.ref<f64>) %5 = "fir.alloca"() <{bindc_name = "x", in_type = i32, operandSegmentSizes = array<i32: 0, 0>, uniq_name = "_QFEx"}> : () -> !fir.ref<i32> %6:2 = "hlfir.declare"(%5) <{operandSegmentSizes = array<i32: 1, 0, 0, 0>, uniq_name = "_QFEx"}> : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) %7 = "arith.constant"() <{value = 1 : i32}> : () -> i32 "hlfir.assign"(%7, %6#0) : (i32, !fir.ref<i32>) -> () %8 = "fir.load"(%4#0) : (!fir.ref<f64>) -> f64 %9 = "fir.convert"(%8) : (f64) -> i32 "fir.store"(%9, %2) : (i32, !fir.ref<i32>) -> () %10 = "fir.load"(%6#0) : (!fir.ref<i32>) -> i32 %11 = "fir.convert"(%10) : (i32) -> f64 "acc.atomic.capture"() ({ "acc.atomic.read"(%2, %6#1) <{element_type = f64}> : (!fir.ref<i32>, !fir.ref<i32>) -> () %12 = "fir.convert"(%11) : (f64) -> i32 "acc.atomic.write"(%2, %12) : (!fir.ref<i32>, i32) -> () "acc.terminator"() : () -> () }) : () -> () This PR updates flang/lib/Lower/DirectivesCommon.h to solve the issues by taking the following approaches (from top to bottom): 1. Move fir.convert for atomic.write out of the capture region. 2. Remove the !fir.ref<i32> -> !fir.ref<f64> conversion found in genOmpAccAtomicRead. 3. Eliminate unnecessary genExprAddr calls on the RHS, which create an invalid temporary for x = 1.0. 4. When generating a capture operation, refer to the original LHS instead of the type-casted RHS. Here, we have to allow for the cases where the operand types of atomic.read differ from one another. Thus, this PR also removes the AllTypesMatch trait from both acc.atomic.read and omp.atomic.read. The example code is converted as follows: %0 = fir.alloca f64 {bindc_name = "n", uniq_name = "_QFEn"} %1:2 = hlfir.declare %0 {uniq_name = "_QFEn"} : (!fir.ref<f64>) -> (!fir.ref<f64>, !fir.ref<f64>) %2 = fir.alloca i32 {bindc_name = "x", uniq_name = "_QFEx"} %3:2 = hlfir.declare %2 {uniq_name = "_QFEx"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) %c1_i32 = arith.constant 1 : i32 hlfir.assign %c1_i32 to %3#0 : i32, !fir.ref<i32> %4 = fir.load %1#0 : !fir.ref<f64> %5 = fir.convert %4 : (f64) -> i32 acc.atomic.capture { acc.atomic.read %1#1 = %3#1 : !fir.ref<f64>, !fir.ref<i32>, i32 acc.atomic.write %3#1 = %5 : !fir.ref<i32>, i32 } Fixes #112911. | 1 年前 | |
[OMPIRBuilder] Don't outline DISTRIBUTE on CPUs (#158317) We use different OpenMP runtime functions on CPU and target offload. The one used for DISTRIBUTE on target offload needs a function pointer to an offloaded function, but the one on CPU doesn't. This caused unnessecary overhead on CPUs because SHARED or FIRSTPRIVATE memory from the surrounding context has to be packaged into a context structure just for an ordinary function call (which would hopefully eventually get inlined). This also makes the IR harder to read. | 10 个月前 | |
[Utils][mlir] Fix interaction between CodeExtractor and OpenMPIRBuilder (#145051) CodeExtractor can currently erroneously insert an alloca into a different function than it inserts its users into, in cases where code is being extracted out of a function that has already been outlined. Add an assertion that the two blocks being inserted into are actually in the same function. Add a check to findAllocaInsertPoint in OpenMP to LLVMIR translation to prevent the aforementioned scenario from happening. OpenMPIRBuilder relies on a callback mechanism to fix-up a module later on during the finaliser step. In some cases this results in the module being invalid prior to the finalise step running. Remove calls to verifyModule wrapped in LLVM_DEBUG from CodeExtractor, as the presence of those results in the compiler crashing with -mllvm -debug due to premature module verification where it would not crash without -debug. Call ompBuilder->finalize() the end of mlir::translateModuleToLLVMIR, in order to make sure the module has actually been finalized prior to trying to verify it. Resolves https://github.com/llvm/llvm-project/issues/138102. --------- Signed-off-by: Kajetan Puchalski <kajetan.puchalski@arm.com> | 1 年前 | |
[MLIR][OpenMP] Remove terminators from loop wrappers (#112229) This patch simplifies the representation of OpenMP loop wrapper operations by introducing the NoTerminator trait and updating accordingly the verifier for the LoopWrapperInterface. Since loop wrappers are already limited to having exactly one region containing exactly one block, and this block can only hold a single omp.loop_nest or loop wrapper and an omp.terminator that does not return any values, it makes sense to simplify the representation of loop wrappers by removing the terminator. There is an extensive list of Lit tests that needed updating to remove the omp.terminators adding some noise to this patch, but actual changes are limited to the definition of the omp.wsloop, omp.simd, omp.distribute and omp.taskloop loop wrapper ops, Flang lowering for those, LoopWrapperInterface::verifyImpl(), SCF to OpenMP conversion and OpenMP dialect documentation. | 1 年前 | |
[MLIR][OpenMP] Lowering nontemporal clause to LLVM IR for SIMD directive (#118751) This patch, - Added a new attribute nontemporal to fir.load and fir.store operation in the FIR dialect. - Added a pass lower-nontemporal which is called before FIRToLLVM conversion pass and adds the nontemporal attribute to loads and stores on the list items specified in the nontemporal clause of the SIMD directive. - Set the UnitAttr:$nontemporal to llvm.load and llvm.store operations during FIR to LLVM dialect conversion, if the corresponding fir.load or fir.store operations have the nontemporal attribute. - Attached the nontemporal metadata to load and store instructions that have the nontemporal attribute, during LLVM dialect to LLVM IR translation. | 1 年前 | |
[mlir][OpenMP][flang] make private variable allocation implicit in omp.private (#124019) The intention of this work is to give MLIR->LLVMIR conversion freedom to control how the private variable is allocated so that it can be allocated on the stack in ordinary cases or as part of a structure used to give closure context for tasks which might outlive the current stack frame. See RFC: https://discourse.llvm.org/t/rfc-openmp-supporting-delayed-task-execution-with-firstprivate-variables/83084 For example, a privatizer for an integer used to look like mlir omp.private {type = private} @x.privatizer : !fir.ref<i32> alloc { ^bb0(%arg0: !fir.ref<i32>): %0 = ... allocate proper memory for the private clone ... omp.yield(%0 : !fir.ref<i32>) } After this change, allocation become implicit in the operation: mlir omp.private {type = private} @x.privatizer : i32 For more complex types that require initialization after allocation, an init region can be used: mlir omp.private {type = private} @x.privatizer : !some.type init { ^bb0(%arg0: !some.pointer<!some.type>, %arg1: !some.pointer<!some.type>): // initialize %arg1, using %arg0 as a mold for allocations omp.yield(%arg1 : !some.pointer<!some.type>) } dealloc { ^bb0(%arg0: !some.pointer<!some.type>): ... deallocate memory allocated by the init region ... omp.yield } This patch lays the groundwork for delayed task execution but is not enough on its own. After this patch all gfortran tests which previously passed still pass. There are the following changes to the Fujitsu test suite: - 0380_0009 and 0435_0009 are fixed - 0688_0041 now fails at runtime. This patch is testing firstprivate variables with tasks. Previously we got lucky with the undefined behavior and won the race. After these changes we no longer get lucky. This patch lays the groundwork for a proper fix for this issue. In flang the lowering re-uses the existing lowering used for reduction init and dealloc regions. In flang, before this patch we hit a TODO with the same wording when generating the copy region for firstprivate polymorphic variables. After this patch the box-like fir.class is passed by reference into the copy region, leading to a different path that didn't hit that old TODO but the generated code still didn't work so I added a new TODO in DataSharingProcessor. | 1 年前 | |
[mlir][OpenMP] fix crash outlining infinite loop (#129872) Previously an extra block was created by splitting the previous exit block. This produced incorrect results when the outlined region statically never terminated because then there wouldn't be a valid exit block for the outlined region, this caused this newly added block to have an incoming edge from outside of the outlining region, which caused outlining to fail. So far as I can tell this extra block no longer serves any purpose. The comment says it is supposed to collate multiple control flow edges into one place, but the code as it is now does not achieve this. In fact, as can be seen from the changes to lit tests, this block was not actually outlined in the end. This is because there are actually two code extractors: one in the callback for creating a parallel op which is used to find what the input/output variables are (which does have this block added to it), and another one which actually does the outlining (which this block was not added to). Tested with the gfortran and fujitsu test suites. Fixes #112884 | 1 年前 | |
[OpenMP] Enable simd in non-reduction composite constructs (#146097) Despite currently being ignored with a warning, simd as a leaf in composite constructs behaves as expected when the construct does not contain a reduction. Enable it for those non-reduction constructs. --------- Signed-off-by: Kajetan Puchalski <kajetan.puchalski@arm.com> | 1 年前 | |
[MLIR][OpenMP] Normalize representation of entry block arg-defining clauses (#109809) This patch updates printing and parsing of operations including clauses that define entry block arguments to the operation's region. This impacts in_reduction, map, private, reduction and task_reduction. The proposed representation to be used by all such clauses is the following: <clause_name>([byref] [@<sym>] %value -> %block_arg [, ...] : <type>[, ...]) { ... } The byref tag is only allowed for reduction-like clauses and the @<sym> is required and only allowed for the private and reduction-like clauses. The map clause does not accept any of these two. This change fixes some currently broken op representations, like omp.teams or omp.sections reduction: omp.teams reduction([byref] @<sym> -> %value : <type>) { ^bb0(%block_arg : <type>): ... } Additionally, it addresses some redundancy in the representation of the previously mentioned cases, as well as e.g. map in omp.target. The problem is that the block argument name after the arrow is not checked in any way, which makes some misleading representations legal: mlir omp.target map_entries(%x -> %arg1, %y -> %arg0, %z -> %doesnt_exist : !llvm.ptr, !llvm.ptr, !llvm.ptr) { ^bb0(%arg0 : !llvm.ptr, %arg1 : !llvm.ptr, %arg2 : !llvm.ptr): ... } In that case, %x maps to %arg0, contrary to what the representation states, and %z maps to %arg2. %doesnt_exist is not resolved, so it would likely cause issues if used anywhere inside of the operation's region. The solution implemented in this patch makes it so that values introduced after the arrow on the representation of these clauses implicitly define the corresponding entry block arguments, removing the potential for these problematic representations. This is what is already implemented for the private and reduction clauses of omp.parallel. There are a couple of consequences of this change: - Entry block argument-defining clauses must come at the end of the operation's representation and in alphabetical order. This is because they are printed/parsed as part of the region and a standardized ordering is needed to reliably match op arguments with their corresponding entry block arguments via the BlockArgOpenMPOpInterface. - We can no longer define per-clause assembly formats to be reused by all operations that take these clauses, since they must be passed to a custom printer including the region and arguments of all other entry block argument-defining clauses. Code duplication and potential for introducing issues is minimized by providing the generic {print,parse}BlockArgRegion helpers and associated structures. MLIR and Flang lowering unit tests are updated due to changes in the order and formatting of impacted operations. | 1 年前 | |
[flang][OpenMP] Fix reduction init region block management (#122079) Replaces https://github.com/llvm/llvm-project/pull/121886 Fixes https://github.com/llvm/llvm-project/issues/120254 (hopefully 🤞) ## Problem Consider the following example: fortran program test real :: x(1) integer :: i !$omp parallel do reduction(+:x) do i = 1,1 x = 1 end do !$omp end parallel do end program The HLFIR+OMP IR for this example looks like this: mlir func.func @_QQmain() { ... omp.parallel { %5 = fir.embox %4#0(%3) : (!fir.ref<!fir.array<1xf32>>, !fir.shape<1>) -> !fir.box<!fir.array<1xf32>> %6 = fir.alloca !fir.box<!fir.array<1xf32>> ... omp.wsloop private(@_QFEi_private_ref_i32 %1#0 -> %arg0 : !fir.ref<i32>) reduction(byref @add_reduction_byref_box_1xf32 %6 -> %arg1 : !fir.ref<!fir.box<!fir.array<1xf32>>>) { omp.loop_nest (%arg2) : i32 = (%c1_i32) to (%c1_i32_0) inclusive step (%c1_i32_1) { ... omp.yield } } omp.terminator } return } The problem addressed by this PR is related to: the alloca in the omp.parallel region + the related reduction clause on the omp.wsloop op. When we try translate the reduction from MLIR to LLVM, we have to choose an alloca insertion point. This happens in convertOmpWsloop where at entry to that function, this is what the LLVM module looks like: llvm define void @_QQmain() { %tid.addr = alloca i32, align 4 ... entry: %omp_global_thread_num = call i32 @__kmpc_global_thread_num(ptr @1) br label %omp.par.entry omp.par.entry: %tid.addr.local = alloca i32, align 4 ... br label %omp.par.region omp.par.region: br label %omp.par.region1 omp.par.region1: ... %5 = alloca { ptr, i64, i32, i8, i8, i8, i8, [1 x [3 x i64]] }, align 8 Now, when we choose an alloca insertion point for the reduction, this is the chosen block omp.par.entry (without the changes in this PR). The problem is that the allocation needed for the reduction needs to reference the %5 SSA value. This results in inserting allocations in omp.par.entry that reference allocations in a later block omp.par.region1 which causes the Instruction does not dominate all uses! error. ## Possible solution - take 2: This PR contains a more localized solution than https://github.com/llvm/llvm-project/pull/121886. It makes sure that on entry to initReductionVars, the IR builder is at a point where we can starting inserting initialization region; to make things cleaner, we still split the builder insertion point to a dedicated omp.reduction.init. This way we avoid splitting after the latest allocation block; which is what causing the issue. | 1 年前 | |
[mlir][OpenMP] fix crash outlining infinite loop (#129872) Previously an extra block was created by splitting the previous exit block. This produced incorrect results when the outlined region statically never terminated because then there wouldn't be a valid exit block for the outlined region, this caused this newly added block to have an incoming edge from outside of the outlining region, which caused outlining to fail. So far as I can tell this extra block no longer serves any purpose. The comment says it is supposed to collate multiple control flow edges into one place, but the code as it is now does not achieve this. In fact, as can be seen from the changes to lit tests, this block was not actually outlined in the end. This is because there are actually two code extractors: one in the callback for creating a parallel op which is used to find what the input/output variables are (which does have this block added to it), and another one which actually does the outlining (which this block was not added to). Tested with the gfortran and fujitsu test suites. Fixes #112884 | 1 年前 | |
| 1 年前 | ||
[mlir][OpenMP] fix crash outlining infinite loop (#129872) Previously an extra block was created by splitting the previous exit block. This produced incorrect results when the outlined region statically never terminated because then there wouldn't be a valid exit block for the outlined region, this caused this newly added block to have an incoming edge from outside of the outlining region, which caused outlining to fail. So far as I can tell this extra block no longer serves any purpose. The comment says it is supposed to collate multiple control flow edges into one place, but the code as it is now does not achieve this. In fact, as can be seen from the changes to lit tests, this block was not actually outlined in the end. This is because there are actually two code extractors: one in the callback for creating a parallel op which is used to find what the input/output variables are (which does have this block added to it), and another one which actually does the outlining (which this block was not added to). Tested with the gfortran and fujitsu test suites. Fixes #112884 | 1 年前 | |
[mlir][OpenMP] Convert reduction alloc region to LLVMIR (#102524) The intention of this change is to ensure that allocas end up in the entry block not spread out amongst complex reduction variable initialization code. The tests we have are quite minimized for readability and maintainability, making the benefits less obvious. The use case for this is when there are multiple reduction variables each will multiple blocks inside of the init region for that reduction. 2/3 Part 1: https://github.com/llvm/llvm-project/pull/102522 Part 3: https://github.com/llvm/llvm-project/pull/102525 | 1 年前 | |
[MLIR][OpenMP] NFC: Uniformize OpenMP ops names (#85393) This patch proposes the renaming of certain OpenMP dialect operations with the goal of improving readability and following a uniform naming convention for MLIR operations and associated classes. In particular, the following operations are renamed: - omp.map_info -> omp.map.info - omp.target_update_data -> omp.target_update - omp.ordered_region -> omp.ordered.region - omp.cancellationpoint -> omp.cancellation_point - omp.bounds -> omp.map.bounds - omp.reduction.declare -> omp.declare_reduction Also, the following MLIR operation classes have been renamed: - omp::TaskLoopOp -> omp::TaskloopOp - omp::TaskGroupOp -> omp::TaskgroupOp - omp::DataBoundsOp -> omp::MapBoundsOp - omp::DataOp -> omp::TargetDataOp - omp::EnterDataOp -> omp::TargetEnterDataOp - omp::ExitDataOp -> omp::TargetExitDataOp - omp::UpdateDataOp -> omp::TargetUpdateOp - omp::ReductionDeclareOp -> omp::DeclareReductionOp - omp::WsLoopOp -> omp::WsloopOp | 2 年前 | |
[mlir][OpenMP] fix crash outlining infinite loop (#129872) Previously an extra block was created by splitting the previous exit block. This produced incorrect results when the outlined region statically never terminated because then there wouldn't be a valid exit block for the outlined region, this caused this newly added block to have an incoming edge from outside of the outlining region, which caused outlining to fail. So far as I can tell this extra block no longer serves any purpose. The comment says it is supposed to collate multiple control flow edges into one place, but the code as it is now does not achieve this. In fact, as can be seen from the changes to lit tests, this block was not actually outlined in the end. This is because there are actually two code extractors: one in the callback for creating a parallel op which is used to find what the input/output variables are (which does have this block added to it), and another one which actually does the outlining (which this block was not added to). Tested with the gfortran and fujitsu test suites. Fixes #112884 | 1 年前 | |
[mlir][OpenMP] fix crash outlining infinite loop (#129872) Previously an extra block was created by splitting the previous exit block. This produced incorrect results when the outlined region statically never terminated because then there wouldn't be a valid exit block for the outlined region, this caused this newly added block to have an incoming edge from outside of the outlining region, which caused outlining to fail. So far as I can tell this extra block no longer serves any purpose. The comment says it is supposed to collate multiple control flow edges into one place, but the code as it is now does not achieve this. In fact, as can be seen from the changes to lit tests, this block was not actually outlined in the end. This is because there are actually two code extractors: one in the callback for creating a parallel op which is used to find what the input/output variables are (which does have this block added to it), and another one which actually does the outlining (which this block was not added to). Tested with the gfortran and fujitsu test suites. Fixes #112884 | 1 年前 | |
[mlir][OpenMP] Allow composite SIMD REDUCTION and IF (#147568) Reduction support: https://github.com/llvm/llvm-project/pull/146671 If Support is fixed in this PR The problem for the IF clause in composite constructs was that wsloop and simd both operate on the same CanonicalLoopInfo structure: with the SIMD processed first, followed by the wsloop. Previously the IF clause generated code like if (cond) { while (...) { simd_loop_body; } } else { while (...) { nonsimd_loop_body; } } The problem with this is that this invalidates the CanonicalLoopInfo structure to be processed by the wsloop later. To avoid this, in this patch I preserve the original loop, moving the IF clause inside of the loop: while (...) { if (cond) { simd_loop_body; } else { non_simd_loop_body; } } On simple examples I tried LLVM was able to hoist the if condition outside of the loop at -O3. The disadvantage of this is that we cannot add the llvm.loop.vectorize.enable attribute on either the SIMD or non-SIMD loops because they both share a loop back edge. There's no way of solving this without keeping the old design of having two different loops: which cannot be represented using only one CanonicalLoopInfo structure. I don't think the presence or absence of this attribute makes much difference. In my testing it is the llvm.loop.parallel_access metadata which makes the difference to vectorization. LLVM will vectorize if legal whether or not this attribute is there in the TRUE branch. In the FALSE branch this means the loop might be vectorized even when the condition is false: but I think this is still standards compliant: OpenMP 6.0 says that when the if clause is false that should be treated like the SIMDLEN clause is one. The SIMDLEN clause is defined as a "hint". For the same reason, SIMDLEN and SAFELEN clauses are silently ignored when SIMD IF is used. I think it is better to implement SIMD IF and ignore SIMDLEN and SAFELEN and some vectorization encouragement metadata when combined with IF than to ignore IF because IF could have correctness consequences whereas the rest are optimiztion hints. For example, the user might use the IF clause to disable SIMD programatically when it is known not safe to vectorize the loop. In this case it is not at all safe to add the parallel access or SAFELEN metadata. | 1 年前 | |
[Flang][OpenMP] Fix to resolve the crash with SIMD aligned clause. (#150612) **Issue:** When SIMD aligned clause has a alignment value which is not a power of 2, compiler crashes with error Assertion (alignment & (alignment - 1)) == 0 && "alignment is not power of 2" **Fix:** According to LLVM Language Reference manual [[link]](https://llvm.org/docs/LangRef.html#assume-opbundles), the alignment value may be non-power-of-two. In that case, the pointer value must be a null pointer otherwise the behavior is undefined. So instead of emitting llvm.assume intrinsic function with a null pointer having the specified alignment, modified the implementation which ignores the aligned clause which has an alignment value which is not a power of 2. This patch also emits a warning indicating that the aligned clause is ignored if the alignment value is not a power of two. It fixes the issue https://github.com/llvm/llvm-project/issues/149458 | 10 个月前 | |
[MLIR][OpenMP] Normalize lowering of omp.loop_nest (#127217) This patch refactors the translation of omp.loop_nest operations into LLVM IR so that it is handled similarly to other operations. Before this change, the responsibility of translating the loop nest fell into each loop wrapper, causing code duplication. This patch centralizes that handling of the loop. One consequence of this was fixing an issue lowering non-inclusive omp.simd loops. As a result, it is now expected that the handling of composite constructs is performed collaboratively among translating functions for each operation involved. At the moment, only do/for simd is supported by ignoring SIMD information, and this behavior is preserved. The translation of loop wrapper operations needs access to the llvm::CanonicalLoopInfo loop information structure in order to apply transformations to it. This is now created in the nested call to convertOmpLoopNest, so it needs to be passed up to all associated loop wrapper translation functions. This is done via the creation of an OpenMPLoopInfoStackFrame within convertHostOrTargetOperation, associated to the outermost loop wrapper. This structure is updated by convertOmpLoopNest, making the result available to all loop wrappers after their body has been translated. | 1 年前 | |
[mlir][OpenMP] implement SIMD reduction (#146671) This replicates clang's implementation. Basically: - A private copy of the reduction variable is created, initialized to the reduction neutral value (using regions from the reduction declaration op). - The body of the loop is lowered as usual, with accesses to the reduction variable mapped to the private copy. - After the loop, we inline the reduction region from the declaration op to combine the privatized variable into the original variable. - As usual with the SIMD construct, attributes are added to encourage vectorization of the loop and to assert that memory accesses in the loop don't alias across iterations. I have verified that simple scalar examples do vectorize at -O3 and the tests I could find in the Fujitsu test suite produce correct results. I tested on top of #146097 and this seemed to work for composite constructs as well. Fixes #144290 | 1 年前 | |
[mlir][OpenMP] implement SIMD reduction (#146671) This replicates clang's implementation. Basically: - A private copy of the reduction variable is created, initialized to the reduction neutral value (using regions from the reduction declaration op). - The body of the loop is lowered as usual, with accesses to the reduction variable mapped to the private copy. - After the loop, we inline the reduction region from the declaration op to combine the privatized variable into the original variable. - As usual with the SIMD construct, attributes are added to encourage vectorization of the loop and to assert that memory accesses in the loop don't alias across iterations. I have verified that simple scalar examples do vectorize at -O3 and the tests I could find in the Fujitsu test suite produce correct results. I tested on top of #146097 and this seemed to work for composite constructs as well. Fixes #144290 | 1 年前 | |
[MLIR][OpenMP] Set default address space for OpenMPIRBuilder (#158689) Extension of https://github.com/llvm/llvm-project/pull/158152 for MLIR. --------- Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com> | 10 个月前 | |
[OMPIRBuilder] Don't outline DISTRIBUTE on CPUs (#158317) We use different OpenMP runtime functions on CPU and target offload. The one used for DISTRIBUTE on target offload needs a function pointer to an offloaded function, but the one on CPU doesn't. This caused unnessecary overhead on CPUs because SHARED or FIRSTPRIVATE memory from the surrounding context has to be packaged into a context structure just for an ordinary function call (which would hopefully eventually get inlined). This also makes the IR harder to read. | 10 个月前 | |
[flang][OpenMP] Implement HAS_DEVICE_ADDR clause (#128568) The HAS_DEVICE_ADDR indicates that the object(s) listed exists at an address that is a valid device address. Specifically, has_device_addr(x) means that (in C/C++ terms) &x is a device address. When entering a target region, x does not need to be allocated on the device, or have its contents copied over (in the absence of additional mapping clauses). Passing its address verbatim to the region for use is sufficient, and is the intended goal of the clause. Some Fortran objects use descriptors in their in-memory representation. If x had a descriptor, both the descriptor and the contents of x would be located in the device memory. However, the descriptors are managed by the compiler, and can be regenerated at various points as needed. The address of the effective descriptor may change, hence it's not safe to pass the address of the descriptor to the target region. Instead, the descriptor itself is always copied, but for objects like x, no further mapping takes place (as this keeps the storage pointer in the descriptor unchanged). --------- Co-authored-by: Sergio Afonso <safonsof@amd.com> | 1 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[MLIR][OpenMP] LLVM IR translation of host_eval (#116052) This patch adds support for processing the host_eval clause of omp.target to populate default and runtime kernel launch attributes. Specifically, these related to the num_teams, thread_limit and num_threads clauses attached to operations nested inside of omp.target. As a result, the thread_limit clause of omp.target is also supported. The implementation of initTargetDefaultAttrs() is intended to reflect clang's own processing of multiple constructs and clauses in order to define a default number of teams and threads to be used as kernel attributes and to populate global variables in the target device module. One side effect of this change is that it is no longer possible to translate to LLVM IR target device MLIR modules unless they have a supported target triple. This is because the local getGridValue() function in the OpenMPIRBuilder only works for certain architectures, and it is called whenever the maximum number of threads has not been explicitly defined. This limitation also matches clang. Evaluating the collapsed loop trip count of SPMD and Generic-SPMD kernels remains unsupported. | 1 年前 | |
[mlir][OpenMP][flang] make private variable allocation implicit in omp.private (#124019) The intention of this work is to give MLIR->LLVMIR conversion freedom to control how the private variable is allocated so that it can be allocated on the stack in ordinary cases or as part of a structure used to give closure context for tasks which might outlive the current stack frame. See RFC: https://discourse.llvm.org/t/rfc-openmp-supporting-delayed-task-execution-with-firstprivate-variables/83084 For example, a privatizer for an integer used to look like mlir omp.private {type = private} @x.privatizer : !fir.ref<i32> alloc { ^bb0(%arg0: !fir.ref<i32>): %0 = ... allocate proper memory for the private clone ... omp.yield(%0 : !fir.ref<i32>) } After this change, allocation become implicit in the operation: mlir omp.private {type = private} @x.privatizer : i32 For more complex types that require initialization after allocation, an init region can be used: mlir omp.private {type = private} @x.privatizer : !some.type init { ^bb0(%arg0: !some.pointer<!some.type>, %arg1: !some.pointer<!some.type>): // initialize %arg1, using %arg0 as a mold for allocations omp.yield(%arg1 : !some.pointer<!some.type>) } dealloc { ^bb0(%arg0: !some.pointer<!some.type>): ... deallocate memory allocated by the init region ... omp.yield } This patch lays the groundwork for delayed task execution but is not enough on its own. After this patch all gfortran tests which previously passed still pass. There are the following changes to the Fujitsu test suite: - 0380_0009 and 0435_0009 are fixed - 0688_0041 now fails at runtime. This patch is testing firstprivate variables with tasks. Previously we got lucky with the undefined behavior and won the race. After these changes we no longer get lucky. This patch lays the groundwork for a proper fix for this issue. In flang the lowering re-uses the existing lowering used for reduction init and dealloc regions. In flang, before this patch we hit a TODO with the same wording when generating the copy region for firstprivate polymorphic variables. After this patch the box-like fir.class is passed by reference into the copy region, leading to a different path that didn't hit that old TODO but the generated code still didn't work so I added a new TODO in DataSharingProcessor. | 1 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[MLIR][OpenMP] - Fix translation of omp.target when private variables need cleaning up (#129205) This is a simple fix that ensures that the InsertPoint is properly fixed up after we have translated the dealloc region of all privatized variables during translation of omp.target from MLIR to LLVMIR. Fix for https://github.com/llvm/llvm-project/issues/129202 | 1 年前 | |
[Flang][MLIR][OpenMP] - Add support for firstprivate when translating omp.target ops from MLIR to LLVMIR (#131213) This patch adds support to translate firstprivate clauses on omp.target ops when translating from MLIR to LLVMIR. Presently, this PR is restricted to supporting only included tasks, i.e #omp target nowait firstprivate(some_variable) will likely not work correctly even if it produces object code. | 1 年前 | |
[mlir][OpenMP][flang] make private variable allocation implicit in omp.private (#124019) The intention of this work is to give MLIR->LLVMIR conversion freedom to control how the private variable is allocated so that it can be allocated on the stack in ordinary cases or as part of a structure used to give closure context for tasks which might outlive the current stack frame. See RFC: https://discourse.llvm.org/t/rfc-openmp-supporting-delayed-task-execution-with-firstprivate-variables/83084 For example, a privatizer for an integer used to look like mlir omp.private {type = private} @x.privatizer : !fir.ref<i32> alloc { ^bb0(%arg0: !fir.ref<i32>): %0 = ... allocate proper memory for the private clone ... omp.yield(%0 : !fir.ref<i32>) } After this change, allocation become implicit in the operation: mlir omp.private {type = private} @x.privatizer : i32 For more complex types that require initialization after allocation, an init region can be used: mlir omp.private {type = private} @x.privatizer : !some.type init { ^bb0(%arg0: !some.pointer<!some.type>, %arg1: !some.pointer<!some.type>): // initialize %arg1, using %arg0 as a mold for allocations omp.yield(%arg1 : !some.pointer<!some.type>) } dealloc { ^bb0(%arg0: !some.pointer<!some.type>): ... deallocate memory allocated by the init region ... omp.yield } This patch lays the groundwork for delayed task execution but is not enough on its own. After this patch all gfortran tests which previously passed still pass. There are the following changes to the Fujitsu test suite: - 0380_0009 and 0435_0009 are fixed - 0688_0041 now fails at runtime. This patch is testing firstprivate variables with tasks. Previously we got lucky with the undefined behavior and won the race. After these changes we no longer get lucky. This patch lays the groundwork for a proper fix for this issue. In flang the lowering re-uses the existing lowering used for reduction init and dealloc regions. In flang, before this patch we hit a TODO with the same wording when generating the copy region for firstprivate polymorphic variables. After this patch the box-like fir.class is passed by reference into the copy region, leading to a different path that didn't hit that old TODO but the generated code still didn't work so I added a new TODO in DataSharingProcessor. | 1 年前 | |
[OMPIRBuilder] Don't outline DISTRIBUTE on CPUs (#158317) We use different OpenMP runtime functions on CPU and target offload. The one used for DISTRIBUTE on target offload needs a function pointer to an offloaded function, but the one on CPU doesn't. This caused unnessecary overhead on CPUs because SHARED or FIRSTPRIVATE memory from the surrounding context has to be packaged into a context structure just for an ordinary function call (which would hopefully eventually get inlined). This also makes the IR harder to read. | 10 个月前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[flang][OpenMP] Enable delayed privatization by default omp.wsloop (#125732) Reapplies #122471 This is based on https://github.com/llvm/llvm-project/pull/125699, only the latest commit is relevant. With changes in this PR and the parent one, the previously reported failures in the Fujitsu(*) test suite should hopefully be resolved (I verified all the 14 reported failures and they pass now). (*) https://linaro.atlassian.net/browse/LLVM-1521 | 1 年前 | |
| 1 年前 | ||
[mlir][OpenMP] Pack task private variables into a heap-allocated context struct (#125307) See RFC: https://discourse.llvm.org/t/rfc-openmp-supporting-delayed-task-execution-with-firstprivate-variables/83084 The aim here is to ensure that tasks which are not executed for a while after they are created do not try to reference any data which are now out of scope. This is done by packing the data referred to by the task into a heap allocated structure (freed at the end of the task). I decided to create the task context structure in OpenMPToLLVMIRTranslation instead of adapting how it is done CodeExtractor (via OpenMPIRBuilder] because CodeExtractor is (at least in theory) generic code which could have other unrelated uses. | 1 年前 | |
[NFC][MLIR][OpenMP] Modify tests to have allocas in the correct address space for AMDGPU | 1 年前 | |
[OMPIRBuilder] Don't outline DISTRIBUTE on CPUs (#158317) We use different OpenMP runtime functions on CPU and target offload. The one used for DISTRIBUTE on target offload needs a function pointer to an offloaded function, but the one on CPU doesn't. This caused unnessecary overhead on CPUs because SHARED or FIRSTPRIVATE memory from the surrounding context has to be packaged into a context structure just for an ordinary function call (which would hopefully eventually get inlined). This also makes the IR harder to read. | 10 个月前 | |
[MLIR][OpenMP] Add codegen for teams reductions (#133310) This patch adds the lowering of teams reductions from the omp dialect to LLVM-IR. Some minor cleanup was done in clang to remove an unused parameter. | 1 年前 | |
[MLIR][OpenMP] Add codegen for teams reductions (#133310) This patch adds the lowering of teams reductions from the omp dialect to LLVM-IR. Some minor cleanup was done in clang to remove an unused parameter. | 1 年前 | |
[OpenMPIRBuilder][MLIR] Pass target-cpu and target-features to outlined functions (#80283) This patch adds support for forwarding the target-cpu and target-features attributes to functions outlined in the OpenMPIRBuilder. This, in turn, results in the addition of these attributes for functions created during the translation of the omp.parallel, omp.task and omp.teams operations, and for the omp.wsloop operation when doing codegen for an OpenMP target device. | 2 年前 | |
[Flang][mlir] - Translation of delayed privatization for deferred target-tasks (#155348) This PR adds support for translation of the private clause on deferred target tasks - that is omp.target operations with the nowait clause. An offloading call for a deferred target-task is not blocking - the offloading (target-generating) host task continues its execution after issuing the offloading call. Therefore, the key problem we need to solve is to ensure that the data needed for private variables to be initialized in the target task persists even after the host task has completed. We do this in a new pass called PrepareForOMPOffloadPrivatizationPass. For a privatized variable that needs its host counterpart for initialization (such as the shape of the data from the descriptor when an allocatable is privatized or the value of the data when an allocatable is firstprivatized), - the pass allocates memory on the heap. - it then initializes this memory by using the init and copy (for firstprivate) regions of the corresponding omp::PrivateClauseOp. - Finally the memory allocated on the heap is freed using the dealloc region of the same omp::PrivateClauseOp instance. This step is not straightforward though, because we cannot simply free the memory that's going to be used by another thread without any synchronization. So, for deallocation, we create a omp.task after the omp.target and synchronize the two with a dummy dependency (using the depend clause). In this newly created omp.task we do the deallocation. | 9 个月前 | |
[OpenMP][flang][MLIR] Decouple alloc, init, and copy regions for omp.private|declare_reduction ops (#125699) This PR changes the emitted block structure of alloc, init, and copy regions for omp.private and omp.declare_reduction ops a little bit. In particular, this decouples init and copy regions from the alloca insertion-point. The main motivation is fix "Instruction does not dominate all uses!" errors that happen specially when an init region uses a value from the OpenMP region it is being inlined into. The issue happens because, previous to this PR, we inline the init region right after the latest alloc block (since we used the alloca IP); which in some cases (see exmaple below), is too early and causes the use dominance issue. Example that would break without this PR (when delayed privatization is enabled for omp.wsloops): fortran subroutine test2 (xyz) integer :: i integer :: xyz(:) !$omp target map(from:xyz) !$omp do private(xyz) do i = 1, 10 xyz(i) = i end do !$omp end target end subroutine | 1 年前 | |
[mlir][OpenMP] Add translation of private_barrier attr to LLVMIR (#140090) Part of a series to fix https://github.com/llvm/llvm-project/issues/136357 | 1 年前 | |
[OpenMP][OMPIRBuilder] Refactor reduction initialization logic into one util (#118447) This refactors the logic needed to emit init logic for reductions by moving some duplicated code into a shared util. The logic for doing is quite involved and is needed for any construct that has reductions. Moreover, when a construct has both private and reduction clauses, both sets of clauses need to cooperate with each other when emitting the logic needed for allocation and initialization. Therefore, this PR clearly sets the boundaries for the logic needed to initialize reductions. | 1 年前 | |
[flang][OpenMP] Unconditionally create after_alloca block in allocatePrivateVars (#123168) While https://github.com/llvm/llvm-project/pull/122866 fixed some issues, it introduced a regression in worksharing loops. The new bug comes from the fact that we now conditionally created the after_alloca block based on the number of sucessors of the alloca insertion point. This is unneccessary, we can just alway create the block. If we do this, we respect the post condtions expected after calling allocatePrivateVars (i.e. that the afterAlloca block has a single predecessor. | 1 年前 | |
[flang] Add support for -mprefer-vector-width=<value> (#142073) This patch adds support for the -mprefer-vector-width= command line option. The parsing of this options is equivalent to Clang's and it is implemented by setting the "prefer-vector-width" function attribute. Co-authored-by: Cameron McInally <cmcinally@nvidia.com> | 1 年前 | |
[MLIR] Fix test after ptrtoaddr change b6bbc4b1940006884c49bad7c93b2a949928fe4c fixed IRBuilder::CreatePtrToAddr to produce the correct instruction. Update the test for ptr_diff lowering accordingly. | 9 个月前 | |
[mlir][ROCDL] Adds wmma scaled intrinsics for gfx1250 (#165915) Signed-off-by: Muzammiluddin Syed <muzasyed@amd.com> | 8 个月前 | |
[MLIR][LLVM] Add explicit target_cpu attribute to llvm.func (#78287) This patch adds the target_cpu attribute to llvm.func MLIR operations and updates the translation to/from LLVM IR to match "target-cpu" function attributes. | 2 年前 | |
[MLIR][LLVM] Support named barrier as a global variable type in llvm dialect (#169194) Enables amdgcn.named.barrier target extension type as a global variable type in MLIR. | 8 个月前 | |
[mlir][llvm] Add llvm.target_features features attribute (#71510) This patch adds a target_features (TargetFeaturesAttr) to the LLVM dialect to allow setting and querying the features in use on a function. The motivation for this comes from the Arm SME dialect where we would like a convenient way to check what variants of an operation are available based on the CPU features. Intended usage: The target_features attribute is populated manually or by a pass: mlir func.func @example() attributes { target_features = #llvm.target_features<["+sme", "+sve", "+sme-f64f64"]> } { // ... } Then within a later rewrite the attribute can be checked, and used to make lowering decisions. c++ // Finds the "target_features" attribute on the parent // FunctionOpInterface. auto targetFeatures = LLVM::TargetFeaturesAttr::featuresAt(op); // Check a feature. // Returns false if targetFeatures is null or the feature is not in // the list. if (!targetFeatures.contains("+sme-f64f64")) return failure(); For now, this is rather simple just checks if the exact feature is in the list, though it could be possible to extend with implied features using information from LLVM. | 2 年前 | |
[MLIR][LLVMIR][DLTI] Pass to update #llvm.target's features per relevant backend (#154938) Modifies #llvm.target<..., features = $FEATURES> so that $FEATURES is now an #llvm.target_features<[...]> attribute (rather than a StringAttr). This enables the attribute to respond to DLTI queries for the different target features. The pass updates the $FEATURES attribute of the target attr at name llvm.target in accordance with the (Sub)Target's features that the relevant LLVM backend knows about. --- DEMO: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake"> } { } by way of -llvm-target-to-target-features turns into: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake", features = <["+64bit", "+64bit-mode", "+adx", "+aes", "+allow-light-256-bit", "+avx", "+avx2", "+bmi", "+bmi2", "+clflushopt", "+cmov", "+crc32", "+cx16", "+cx8", "+ermsb", "+f16c", "+false-deps-popcnt", "+fast-15bytenop", "+fast-gather", "+fast-scalar-fsqrt", "+fast-shld-rotate", "+fast-variable-crosslane-shuffle", "+fast-variable-perlane-shuffle", "+fast-vector-fsqrt", "+fma", "+fsgsbase", "+fxsr", "+idivq-to-divl", "+invpcid", "+lzcnt", "+macrofusion", "+mmx", "+movbe", "+no-bypass-delay-blend", "+no-bypass-delay-mov", "+no-bypass-delay-shuffle", "+nopl", "+pclmul", "+popcnt", "+prfchw", "+rdrnd", "+rdseed", "+sahf", "+slow-3ops-lea", "+sse", "+sse2", "+sse3", "+sse4.1", "+sse4.2", "+ssse3", "+vzeroupper", "+x87", "+xsave", "+xsavec", "+xsaveopt", "+xsaves"]>>} { } | 11 个月前 | |
[MLIR][LLVMIR][DLTI] Pass to update #llvm.target's features per relevant backend (#154938) Modifies #llvm.target<..., features = $FEATURES> so that $FEATURES is now an #llvm.target_features<[...]> attribute (rather than a StringAttr). This enables the attribute to respond to DLTI queries for the different target features. The pass updates the $FEATURES attribute of the target attr at name llvm.target in accordance with the (Sub)Target's features that the relevant LLVM backend knows about. --- DEMO: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake"> } { } by way of -llvm-target-to-target-features turns into: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake", features = <["+64bit", "+64bit-mode", "+adx", "+aes", "+allow-light-256-bit", "+avx", "+avx2", "+bmi", "+bmi2", "+clflushopt", "+cmov", "+crc32", "+cx16", "+cx8", "+ermsb", "+f16c", "+false-deps-popcnt", "+fast-15bytenop", "+fast-gather", "+fast-scalar-fsqrt", "+fast-shld-rotate", "+fast-variable-crosslane-shuffle", "+fast-variable-perlane-shuffle", "+fast-vector-fsqrt", "+fma", "+fsgsbase", "+fxsr", "+idivq-to-divl", "+invpcid", "+lzcnt", "+macrofusion", "+mmx", "+movbe", "+no-bypass-delay-blend", "+no-bypass-delay-mov", "+no-bypass-delay-shuffle", "+nopl", "+pclmul", "+popcnt", "+prfchw", "+rdrnd", "+rdseed", "+sahf", "+slow-3ops-lea", "+sse", "+sse2", "+sse3", "+sse4.1", "+sse4.2", "+ssse3", "+vzeroupper", "+x87", "+xsave", "+xsavec", "+xsaveopt", "+xsaves"]>>} { } | 11 个月前 | |
[MLIR][LLVMIR][DLTI] Pass to update #llvm.target's features per relevant backend (#154938) Modifies #llvm.target<..., features = $FEATURES> so that $FEATURES is now an #llvm.target_features<[...]> attribute (rather than a StringAttr). This enables the attribute to respond to DLTI queries for the different target features. The pass updates the $FEATURES attribute of the target attr at name llvm.target in accordance with the (Sub)Target's features that the relevant LLVM backend knows about. --- DEMO: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake"> } { } by way of -llvm-target-to-target-features turns into: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake", features = <["+64bit", "+64bit-mode", "+adx", "+aes", "+allow-light-256-bit", "+avx", "+avx2", "+bmi", "+bmi2", "+clflushopt", "+cmov", "+crc32", "+cx16", "+cx8", "+ermsb", "+f16c", "+false-deps-popcnt", "+fast-15bytenop", "+fast-gather", "+fast-scalar-fsqrt", "+fast-shld-rotate", "+fast-variable-crosslane-shuffle", "+fast-variable-perlane-shuffle", "+fast-vector-fsqrt", "+fma", "+fsgsbase", "+fxsr", "+idivq-to-divl", "+invpcid", "+lzcnt", "+macrofusion", "+mmx", "+movbe", "+no-bypass-delay-blend", "+no-bypass-delay-mov", "+no-bypass-delay-shuffle", "+nopl", "+pclmul", "+popcnt", "+prfchw", "+rdrnd", "+rdseed", "+sahf", "+slow-3ops-lea", "+sse", "+sse2", "+sse3", "+sse4.1", "+sse4.2", "+ssse3", "+vzeroupper", "+x87", "+xsave", "+xsavec", "+xsaveopt", "+xsaves"]>>} { } | 11 个月前 | |
[MLIR][LLVMIR][DLTI] Pass to update #llvm.target's features per relevant backend (#154938) Modifies #llvm.target<..., features = $FEATURES> so that $FEATURES is now an #llvm.target_features<[...]> attribute (rather than a StringAttr). This enables the attribute to respond to DLTI queries for the different target features. The pass updates the $FEATURES attribute of the target attr at name llvm.target in accordance with the (Sub)Target's features that the relevant LLVM backend knows about. --- DEMO: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake"> } { } by way of -llvm-target-to-target-features turns into: mlir module attributes {llvm.target = #llvm.target<triple = "x86_64-unknown-linux", chip = "skylake", features = <["+64bit", "+64bit-mode", "+adx", "+aes", "+allow-light-256-bit", "+avx", "+avx2", "+bmi", "+bmi2", "+clflushopt", "+cmov", "+crc32", "+cx16", "+cx8", "+ermsb", "+f16c", "+false-deps-popcnt", "+fast-15bytenop", "+fast-gather", "+fast-scalar-fsqrt", "+fast-shld-rotate", "+fast-variable-crosslane-shuffle", "+fast-variable-perlane-shuffle", "+fast-vector-fsqrt", "+fma", "+fsgsbase", "+fxsr", "+idivq-to-divl", "+invpcid", "+lzcnt", "+macrofusion", "+mmx", "+movbe", "+no-bypass-delay-blend", "+no-bypass-delay-mov", "+no-bypass-delay-shuffle", "+nopl", "+pclmul", "+popcnt", "+prfchw", "+rdrnd", "+rdseed", "+sahf", "+slow-3ops-lea", "+sse", "+sse2", "+sse3", "+sse4.1", "+sse4.2", "+ssse3", "+vzeroupper", "+x87", "+xsave", "+xsavec", "+xsaveopt", "+xsaves"]>>} { } | 11 个月前 | |
[mlir] support dialect attribute translation to LLVM IR (#75309) Extend the amendOperation mechanism for translating dialect attributes attached to operations from another dialect when translating MLIR to LLVM IR. Previously, this mechanism would have no knowledge of the LLVM IR instructions created for the given operation, making it impossible for it to perform local modifications such as attaching operation-level metadata. Collect instructions inserted by the LLVM IR builder and pass them to amendOperation. | 2 年前 | |
| 2 年前 | ||
[MLIR:LLVM] Add UWTableKind attribute (#135811) Add UWTableKind enum and corresponding attribute to llvm.func including translation to llvm::Function attribute. | 1 年前 | |
[MLIR][VCIX] Support VCIX intrinsics in LLVMIR dialect (#75875) The changeset extends LLVMIR intrinsics with VCIX intrinsics. The VCIX intrinsics allow MLIR users to interact with RISC-V co-processors that are compatible with XSfvcp extension Source: https://www.sifive.com/document-file/sifive-vector-coprocessor-interface-vcix-software | 2 年前 | |
[MLIR][VCIX] Support VCIX intrinsics in LLVMIR dialect (#75875) The changeset extends LLVMIR intrinsics with VCIX intrinsics. The VCIX intrinsics allow MLIR users to interact with RISC-V co-processors that are compatible with XSfvcp extension Source: https://www.sifive.com/document-file/sifive-vector-coprocessor-interface-vcix-software | 2 年前 | |
[mlir][Pass] Include anchor op in -pass-pipeline In D134622 the printed form of a pass manager is changed to include the name of the op that the pass manager is anchored on. This updates the -pass-pipeline argument format to include the anchor op as well, so that the printed form of a pipeline can be directly passed to -pass-pipeline. In most cases this requires updating -pass-pipeline='pipeline' to -pass-pipeline='builtin.module(pipeline)'. This also fixes an outdated assert that prevented running a PassManager anchored on 'any'. Reviewed By: rriddle Differential Revision: https://reviews.llvm.org/D134900 | 3 年前 | |
Add support for MLIR to llvm vscale attribute (#67012) The vscale_range is used for scalabale vector functionality in Arm Scalable Vector Extension to select the size of vector operation (and I thnk RISCV has something similar). This patch adds the base support for the vscale_range attribute to the LLVM::FuncOp, and the marshalling for translation to LLVM-IR and import from LLVM-IR to LLVM dialect. This attribute is intended to be used at higher level MLIR, specified either by command-line options to the compiler or using compiler directives (e.g. pragmas or function attributes in the source code) to indicate the desired range. | 2 年前 | |
[mlir][x86vector] AVX2 I8 Dot Op (#147908) Adds AVX2 i8 dot-product operation and defines lowering to LLVM intrinsics. Target assembly instruction: vpdpbssd.128/256 | 1 年前 | |
[MLIR][XeVM] Add lowering for llvm load store ops with XeVM cache control (#156768) Add lowering support for LLVM load / store ops with XeVM cache control attributes. | 10 个月前 |
| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
| 8 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 3 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 3 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 2 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 8 个月前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 2 年前 | ||
| 10 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 8 个月前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 11 个月前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 10 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 9 个月前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 1 年前 | ||
| 9 个月前 | ||
| 8 个月前 | ||
| 2 年前 | ||
| 8 个月前 | ||
| 2 年前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 11 个月前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 2 年前 | ||
| 2 年前 | ||
| 3 年前 | ||
| 2 年前 | ||
| 1 年前 | ||
| 10 个月前 |