summaryrefslogtreecommitdiff
path: root/mlir/lib/Target/LLVMIR/Dialect
AgeCommit message (Collapse)Author
2025-11-20[MLIR] Apply clang-tidy fixes for readability-identifier-naming in ↵Mehdi Amini
LLVMToLLVMIRTranslation.cpp (NFC)
2025-11-19[mlir][LLVM] Resync memory effect attribute with LLVM IR (#168568)darkbuck
- Add missing locations, namely 'ErrnoMem', 'TargetMem0', and 'TargetMem1'.
2025-11-07[mlir][NVVM] Add nvvm.membar operation (#166698)Valentin Clement (バレンタイン クレメン)
Add nvvm.membar operation with level as defined in https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-membar This will be used to replace direct intrinsic call in CUDA Fortran for `threadfence()`, `threadfence_block` and `thread fence_system()` currently lowered here: https://github.com/llvm/llvm-project/blob/e700f157026bf8b4d58f936c5db8f152e269d77f/flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp#L1310 The nvvm membar intrsinsic are also used in CUDA C/C++ (https://github.com/llvm/llvm-project/blob/49f55f4991227f3c7a2b8161bbf45c74b7023944/clang/lib/Headers/__clang_cuda_device_functions.h#L528)
2025-11-04[MLIR][NVVM] Update redux.sync op (#166125)Srinivasa Ravi
This change: - Updates the `redux.sync` NVVM Op input and output type constraints. - Adds a verifier for the Op to prevent stack dumps and hitting `llvm_unreachable` in certain invalid usage scenarios. Instead, we gracefully error out with an informative message now.
2025-10-31[MLIR][OpenMP] Fix and simplify bounds offset calculation for 1-D GEP ↵agozillon
offsets (#165486) Currently this is being calculated incorrectly and will result in incorrect index offsets in more complicated array slices. This PR tries to address it by refactoring and changing the calculation to be more correct.
2025-10-22[Flang][mlir] - Translation of delayed privatization for deferred ↵Pranav Bhandarkar
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.
2025-10-21[Flang][OpenMP][Dialect] Swap to using MLIR dialect enum to encode map flags ↵agozillon
(#164043) This PR shifts from using the LLVM OpenMP enumerator bit flags to an OpenMP dialect specific enumerator. This allows us to better represent map types that wouldn't be of interest to the LLVM backend and runtime in the dialect. Primarily things like ref_ptr/ref_ptee/ref_ptr_ptee/atach_none/attach_always/attach_auto which are of interest to the compiler for certrain transformations (primarily in the FIR transformation passes dealing with mapping), but the runtime has no need to know about them. It also means if another OpenMP implementation comes along they won't need to stick to the same bit flag system LLVM chose/do leg work to address it.
2025-10-17[MLIR] Apply clang-tidy fixes for performance-unnecessary-value-param in ↵Mehdi Amini
OpenMPToLLVMIRTranslation.cpp (NFC)
2025-10-06[mlir] Simplify unreachable type switch cases. NFC. (#162032)Jakub Kuderski
Use `DefaultUnreachable` from https://github.com/llvm/llvm-project/pull/161970.
2025-10-02[mlir][omp] Add omp.tile operation (#160292)Michael Kruse
Add the `omp.tile` loop transformations for the OpenMP dialect. Used for lowering a standalone `!$omp tile` in Flang.
2025-09-26[clang] Use the VFS to create the OpenMP region entry ID (#160918)Jan Svoboda
This PR uses the VFS to create the OpenMP target entry instead of going straight to the real file system. This matches the behavior of other input files of the compiler.
2025-09-26[Flang][OpenMP] Enable no-loop kernels (#155818)Dominik Adamski
Enable the generation of no-loop kernels for Fortran OpenMP code. target teams distribute parallel do pragmas can be promoted to no-loop kernels if the user adds the -fopenmp-assume-teams-oversubscription and -fopenmp-assume-threads-oversubscription flags. If the OpenMP kernel contains reduction or num_teams clauses, it is not promoted to no-loop mode. The global OpenMP device RTL oversubscription flags no longer force no-loop code generation for Fortran.
2025-09-24[mlir][ptr] Add `ptr.ptr_diff` operation (#157354)Fabian Mora
Thi patch introduces the `ptr.ptr_diff` operation for computing pointer differences. The semantics of the operation are given by: ``` The `ptr_diff` operation computes the difference between two pointers, returning an integer or index value representing the number of bytes between them. The operation supports both scalar and shaped types with value semantics: - When both operands are scalar: produces a single difference value - When both are shaped: performs element-wise subtraction, shapes must be the same The operation also supports the following flags: - `none`: No flags are set. - `nuw`: No Unsigned Wrap, if the subtraction causes an unsigned overflow, the result is a poison value. - `nsw`: No Signed Wrap, if the subtraction causes a signed overflow, the result is a poison value. NOTE: The pointer difference is calculated using an integer type specified by the data layout. The final result will be sign-extended or truncated to fit the result type as necessary. ``` This patch also adds translation to LLVM IR hooks for the `ptr_diff` op. This translation uses the `ptrtoaddr` builder to compute only index bits difference. Example: ```mlir llvm.func @ptr_diff_vector_i32(%ptrs1: vector<8x!ptr.ptr<#llvm.address_space<0>>>, %ptrs2: vector<8x!ptr.ptr<#llvm.address_space<0>>>) -> vector<8xi32> { %diffs = ptr.ptr_diff %ptrs1, %ptrs2 : vector<8x!ptr.ptr<#llvm.address_space<0>>> -> vector<8xi32> llvm.return %diffs : vector<8xi32> } ``` Translation to LLVM IR: ```llvm define <8 x i32> @ptr_diff_vector_i32(<8 x ptr> %0, <8 x ptr> %1) { %3 = ptrtoint <8 x ptr> %0 to <8 x i64> %4 = ptrtoint <8 x ptr> %1 to <8 x i64> %5 = sub <8 x i64> %3, %4 %6 = trunc <8 x i64> %5 to <8 x i32> ret <8 x i32> %6 } ``` --------- Co-authored-by: Mehdi Amini <joker.eph@gmail.com>
2025-09-19[OpenMP][MLIR] Preserve to/from flags in mapper base entry for mappers (#159799)Akash Banerjee
With declare mapper, the parent base entry was emitted as `TARGET_PARAM` only. The mapper received a map-type without `to/from`, causing components to degrade to `alloc`-only (no copies), breaking allocatable payload mapping. This PR preserves the map-type bits from the parent. This fixes #156466.
2025-09-14[mlir][ptr] Add ConstantOp with NullAttr and AddressAttr support (#157347)Fabian Mora
This patch introduces the `ptr.constant` operation. It also adds the `NullAttr` and `AddressAttr` for representing null pointers, and integer raw addresses. It also implements LLVM IR translation for `ptr.constant` with `#ptr.null` or `#ptr.address` attributes. Finally, it extends `FieldParser` to support APInt parsing. Example: ```mlir llvm.func @constant_address_op() -> !llvm.struct<(!ptr.ptr<#llvm.address_space<0>>, !ptr.ptr<#llvm.address_space<1>>, !ptr.ptr<#llvm.address_space<2>>)> { %0 = ptr.constant #ptr.null : !ptr.ptr<#llvm.address_space<0>> %1 = ptr.constant #ptr.address<0x1000> : !ptr.ptr<#llvm.address_space<1>> %2 = ptr.constant #ptr.address<3735928559> : !ptr.ptr<#llvm.address_space<2>> %3 = llvm.mlir.poison : !llvm.struct<(!ptr.ptr<#llvm.address_space<0>>, !ptr.ptr<#llvm.address_space<1>>, !ptr.ptr<#llvm.address_space<2>>)> %4 = llvm.insertvalue %0, %3[0] : !llvm.struct<(!ptr.ptr<#llvm.address_space<0>>, !ptr.ptr<#llvm.address_space<1>>, !ptr.ptr<#llvm.address_space<2>>)> %5 = llvm.insertvalue %1, %4[1] : !llvm.struct<(!ptr.ptr<#llvm.address_space<0>>, !ptr.ptr<#llvm.address_space<1>>, !ptr.ptr<#llvm.address_space<2>>)> %6 = llvm.insertvalue %2, %5[2] : !llvm.struct<(!ptr.ptr<#llvm.address_space<0>>, !ptr.ptr<#llvm.address_space<1>>, !ptr.ptr<#llvm.address_space<2>>)> llvm.return %6 : !llvm.struct<(!ptr.ptr<#llvm.address_space<0>>, !ptr.ptr<#llvm.address_space<1>>, !ptr.ptr<#llvm.address_space<2>>)> } ``` Result of translation to LLVM IR: ```llvm define { ptr, ptr addrspace(1), ptr addrspace(2) } @constant_address_op() { ret { ptr, ptr addrspace(1), ptr addrspace(2) } { ptr null, ptr addrspace(1) inttoptr (i64 4096 to ptr addrspace(1)), ptr addrspace(2) inttoptr (i64 3735928559 to ptr addrspace(2)) } } ``` This patch also changes all the `convert*` occurrences in function names or comments to `translate` in the PtrToLLVM file. --------- Co-authored-by: Mehdi Amini <joker.eph@gmail.com>
2025-09-14[mlir][LLVM] Add LLVMAddrSpaceAttrInterface and NVVMMemorySpaceAttr (#157339)Fabian Mora
This patch introduces the `LLVMAddrSpaceAttrInterface` for defining compatible LLVM address space attributes To test this interface, this patch also adds: - Adds NVVMMemorySpaceAttr implementing both LLVMAddrSpaceAttrInterface and MemorySpaceAttrInterface - Converts NVVM memory space constants from enum to MLIR enums - Updates all NVVM memory space references to use new attribute system - Adds support for NVVM memory spaces in ptr dialect translation Example: ```mlir llvm.func @nvvm_ptr_address_space( !ptr.ptr<#nvvm.memory_space<global>>, !ptr.ptr<#nvvm.memory_space<shared>>, !ptr.ptr<#nvvm.memory_space<constant>>, !ptr.ptr<#nvvm.memory_space<local>>, !ptr.ptr<#nvvm.memory_space<tensor>>, !ptr.ptr<#nvvm.memory_space<shared_cluster>> ) -> !ptr.ptr<#nvvm.memory_space<generic>> ``` Translating the above code to LLVM produces: ```llvm declare ptr @nvvm_ptr_address_space(ptr addrspace(1), ptr addrspace(3), ptr addrspace(4), ptr addrspace(5), ptr addrspace(6), ptr addrspace(7)) ``` To convert the memory space enum to the new enum class use: ```bash grep -r . -e "NVVMMemorySpace::kGenericMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kGenericMemorySpace/NVVMMemorySpace::Generic/g" grep -r . -e "NVVMMemorySpace::kGlobalMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kGlobalMemorySpace/NVVMMemorySpace::Global/g" grep -r . -e "NVVMMemorySpace::kSharedMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kSharedMemorySpace/NVVMMemorySpace::Shared/g" grep -r . -e "NVVMMemorySpace::kConstantMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kConstantMemorySpace/NVVMMemorySpace::Constant/g" grep -r . -e "NVVMMemorySpace::kLocalMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kLocalMemorySpace/NVVMMemorySpace::Local/g" grep -r . -e "NVVMMemorySpace::kTensorMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kTensorMemorySpace/NVVMMemorySpace::Tensor/g" grep -r . -e "NVVMMemorySpace::kSharedClusterMemorySpace" -l | xargs sed -i -e "s/NVVMMemorySpace::kSharedClusterMemorySpace/NVVMMemorySpace::SharedCluster/g" ``` NOTE: A future patch will add support for ROCDL, it wasn't added here to keep the patch small.
2025-09-11[clang] Use VFS for `-fopenmp-host-ir-file-path` (#156727)Jan Svoboda
This is a follow-up to #150124. This PR makes it so that the `-fopenmp-host-ir-file-path` respects VFS overlays, like any other input file.
2025-09-10[mlir][LLVMIR] Support memory model relaxation annotations (MMRA) (#157770)Krzysztof Drewniak
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>
2025-09-10[flang][OpenMP] Enable tiling (#143715)Jan Leyonberg
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.
2025-09-09[MLIR][XeVM] Add lowering for llvm load store ops with XeVM cache control ↵Sang Ik Lee
(#156768) Add lowering support for LLVM load / store ops with XeVM cache control attributes.
2025-09-07[MLIR] Apply clang-tidy fixes for misc-use-internal-linkage in ↵Mehdi Amini
OpenMPToLLVMIRTranslation.cpp (NFC)
2025-09-04[MLIR] Apply clang-tidy fixes for readability-container-size-empty in ↵Mehdi Amini
OpenMPToLLVMIRTranslation.cpp (NFC)
2025-09-03[mlir][ptr] Add `gather`, `masked_load`, `masked_store`, and `scatter` ops ↵Fabian Mora
(#156368) This patch adds the `gather`, `masked_load`, `masked_store`, and `scatter` operations to the `ptr` dialect. It also implements translation from these operations to LLVM intrinsics: - ptr.gather -> llvm.masked.gather - ptr.masked_load -> llvm.masked.load - ptr.masked_store -> llvm.masked.store - ptr.scatter -> llvm.masked.scatter Example: ```mlir llvm.func @mixed_masked_ops_address_spaces(%ptr: !ptr.ptr<#llvm.address_space<3>>, %ptrs: vector<4x!ptr.ptr<#llvm.address_space<3>>>, %mask: vector<4xi1>, %value: vector<4xf64>, %passthrough: vector<4xf64>) { %0 = ptr.gather %ptrs, %mask, %passthrough alignment = 8 : vector<4x!ptr.ptr<#llvm.address_space<3>>> -> vector<4xf64> ptr.scatter %value, %ptrs, %mask alignment = 8 : vector<4xf64>, vector<4x!ptr.ptr<#llvm.address_space<3>>> %1 = ptr.masked_load %ptr, %mask, %passthrough alignment = 8 : !ptr.ptr<#llvm.address_space<3>> -> vector<4xf64> ptr.masked_store %value, %ptr, %mask alignment = 8 : vector<4xf64>, !ptr.ptr<#llvm.address_space<3>> llvm.return } ``` Translates to: ```llvm define void @mixed_masked_ops_address_spaces(ptr addrspace(3) %0, <4 x ptr addrspace(3)> %1, <4 x i1> %2, <4 x double> %3, <4 x double> %4) { %6 = call <4 x double> @llvm.masked.gather.v4f64.v4p3(<4 x ptr addrspace(3)> %1, i32 8, <4 x i1> %2, <4 x double> %4) call void @llvm.masked.scatter.v4f64.v4p3(<4 x double> %3, <4 x ptr addrspace(3)> %1, i32 8, <4 x i1> %2) %7 = call <4 x double> @llvm.masked.load.v4f64.p3(ptr addrspace(3) %0, i32 8, <4 x i1> %2, <4 x double> %4) call void @llvm.masked.store.v4f64.p3(<4 x double> %3, ptr addrspace(3) %0, i32 8, <4 x i1> %2) ret void } ```
2025-09-03[mlir][ptr] Add translations to LLVMIR for ptr ops. (#156355)Fabian Mora
Implements translation from ptr dialect to LLVM IR for core pointer operations: - `ptr.ptr_add` -> `getelementptr` - `ptr.load` -> `load` with atomic ordering, volatility, and metadata support - `ptr.store` -> `store` with atomic ordering, volatility, and metadata support - `ptr.type_offset` -> GEP-based size computation Example: ```mlir llvm.func @test(%arg0: !ptr.ptr<#llvm.address_space<0>>) { %0 = ptr.type_offset f64 : i32 %1 = ptr.ptr_add inbounds %arg0, %0 : !ptr.ptr<#llvm.address_space<0>>, i32 %2 = ptr.load volatile %1 : !ptr.ptr<#llvm.address_space<0>> -> f64 ptr.store %2, %arg0 : f64, !ptr.ptr<#llvm.address_space<0>> llvm.return } ``` Translates to: ```llvm define void @test(ptr %0) { %2 = getelementptr inbounds i8, ptr %0, i32 8 %3 = load volatile double, ptr %2, align 8 store double %3, ptr %0, align 8 ret void } ```
2025-09-02[mlir][debug] Inherit DISubprogramAttr from DILocalScopeAttr. (#156081)Abid Qadeer
As mentioned in https://github.com/llvm/llvm-project/pull/154926, `DISubprogramAttr` is inherited from `DIScopeAttr` while in llvm, the `DISubprogram` inherits from `DILocalScope`. This change corrects the hierarchy. Also does the same change for `DILexicalBlockAttr` and `DILexicalBlockFileAttr`.
2025-09-02[mlir][LLVM|ptr] Add the `#llvm.address_space` attribute, and allow `ptr` ↵Fabian Mora
translation (#156333) This commit introduces the `#llvm.address_space` attribute. This attribute implements the `ptr::MemorySpaceAttrInterface`, establishing the semantics of the LLVM address space. This allows making `!ptr.ptr` translatable to LLVM IR as long it uses the `#llvm.address_space` attribute. Concretely, `!ptr.ptr<#llvm.address_space<N>>` now translates to `ptr addrspace(N)`. Additionally, this patch makes `PtrLikeTypes` with no metadata, no element type, and with `#llvm.address_space` memory space, compatible with the LLVM dialect. **Infrastructure Updates:** - Refactor `ptr::MemorySpaceAttrInterface` to include DataLayout parameter for better validation - Add new utility functions `LLVM::isLoadableType()` and `LLVM::isTypeCompatibleWithAtomicOp()` - Update type compatibility checks to support ptr-like types with LLVM address spaces - Splice the `MemorySpaceAttrInterface` to its own library, so the LLVMDialect won't depend on the PtrDialect yet **Translation Support:** - New `PtrToLLVMIRTranslation` module for converting ptr dialect to LLVM IR - Type translation support for ptr types with LLVM address spaces - Proper address space preservation during IR lowering Example: ```mlir llvm.func @llvm_ops_with_ptr_values(%arg0: !llvm.ptr) { %1 = llvm.load %arg0 : !llvm.ptr -> !ptr.ptr<#llvm.address_space<1>> llvm.store %1, %arg0 : !ptr.ptr<#llvm.address_space<1>>, !llvm.ptr llvm.return } ``` Translates to: ```llvmir ; ModuleID = 'LLVMDialectModule' source_filename = "LLVMDialectModule" define void @llvm_ops_with_ptr_values(ptr %0) { %2 = load ptr addrspace(1), ptr %0, align 8 store ptr addrspace(1) %2, ptr %0, align 8 ret void } !llvm.module.flags = !{!0} !0 = !{i32 2, !"Debug Info Version", i32 3} ```
2025-08-31[MLIR] Apply clang-tidy fixes for llvm-else-after-return in ↵Mehdi Amini
OpenMPToLLVMIRTranslation.cpp (NFC)
2025-08-29[mlir][debug] Convert DbgIntrOp to DebugRecords directly. (#154926)Abid Qadeer
Currently the `DbgDeclareOP/DbgValueOP/DbgLabelOp` are first converted to llvm debug intrinsics which are later translated to debug records by a call of `convertToNewDbgValues`. This is not only inefficient but also makes the code that works on intermediate IR unnecessarily complicated. The debug intrinsics are also being phased out. This PR converts these Ops directly to debug records. The conversion is relatively simple but there is a bit of code repetition due to how the APIs in the `DIBuilders` are named. There are few `cast<>` which I would like to do without but could not see a good way around them. Any suggestions welcome here. Also noticed that `DISubprogramAttr` is inherited from `DIScopeAttr` while in llvm, the `DISubprogram` inherits from `DILocalScope`. I am going to fix this separately and then we could use `FusedLocWith<LLVM::DILocalScopeAttr>` and cast to `DILocalScope` will be much safer. As the output remains the same, the existing tests cover this change. I also ran the `GDB` tests with flang and there was no regression. --------- Co-authored-by: Tobias Gysi <tobias.gysi@nextsilicon.com>
2025-08-29[Flang][OpenMP] Fix to resolve the crash with SIMD aligned clause. (#150612)Kaviya Rajendiran
**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
2025-08-27[NVPTX] Auto-upgrade nvvm.grid_constant to param attribute (#155489)Alex MacLean
Upgrade the !"grid_constant" !nvvm.annotation to a "nvvm.grid_constant" attribute. This attribute is much simpler for front-ends to apply and faster and simpler to query.
2025-08-22[MLIR][NVVM] Add "blocksareclusters" kernel attribute support (#154519)Rajat Bajpai
This change adds "nvvm.blocksareclusters" kernel attribute support in NVVM Dialect/MLIR.
2025-08-18[OpenMP] Introduce omp.target_allocmem and omp.target_freemem omp dialect ↵Chaitanya
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)
2025-08-12[MLIR][NVVM] Support generating all the ldmatrix intrinsics from NVVM ops ↵Gao Yanfeng
(#148783) Previously, the NVVM dialect's ldmatrix operation could only generate a limited subset of the available NVVM ldmatrix intrinsics. The intrinsics generating new ops introduced in BlackWell are not accessible through the NVVM ops. This commit extends the ldmatrix operation to support all available ldmatrix intrinsics.
2025-08-11[OMPIRBuilder] Avoid invalid debug location. (#151306)Abid Qadeer
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.
2025-08-07 [flang][OMPIRBuilder][MLIR][llvm] Backend support for atomic control ↵Anchu Rajendran S
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
2025-08-04[mlir][OpenMP][NFC] Fix gcc 14 warning (#151941)Tom Eccles
GCC couldn't tell that the enum is checked exhaustively and so was warning about there being no return on this path from the function.
2025-08-01[MLIR][NVVM] Support stmatrix intrinsics (#148377)Gao Yanfeng
Add support for the `@llvm.nvvm.stmatrix` intrinsic series. These correspond to PTX stmatrix operations, as documented in the [PTX ISA reference](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-store-instruction-stmatrix).
2025-07-30Reapply "[mlir][llvm] Add intrinsic arg and result attribute support … ↵Tobias Gysi
(#151324) …(… (#151099) This reverts commit 2780b8f22058b35a8e70045858b87a1966df8df3 and relands b7bfbc0c4c7b20d6623a5b0b4a7fea8ae08a62da. Adds the following fixes compared to the original PR (https://github.com/llvm/llvm-project/pull/150783): - A bazel fix - Use `let methods` instead of `list<InterfaceMethod> methods` The missing forward declaration has been added in meantime: https://github.com/llvm/llvm-project/commit/9164d206b33d61c93f5fc4628797485f96d654ca.
2025-07-30Revert "Reland "[mlir][llvm] Add intrinsic arg and result attribute support ↵Mehdi Amini
(…" (#151316) Reverts llvm/llvm-project#151125 Broke the gcc-7 build: include/mlir/Target/LLVMIR/ModuleTranslation.h:318:34: error: no type named 'CallBase' in namespace 'llvm' llvm::CallBase *call, ~~~~~~^
2025-07-30Reland "[mlir][llvm] Add intrinsic arg and result attribute support (… ↵Tobias Gysi
(#151125) …… (#151099) This reverts commit 2780b8f22058b35a8e70045858b87a1966df8df3 to reland 59013d44058ef423a117f95092150e16e16fdb09. In addition to the original commit this one includes: - This includes a bazel fix - Use `let methods` instead of `list<InterfaceMethod> methods` The original commit message was: This patch extends the LLVM dialect's intrinsic infra to support argument and result attributes. Initial support is added for the memory intrinsics llvm.intr.memcpy, llvm.intr.memmove, and llvm.intr.memset. Additionally, an ArgAndResultAttrsOpInterface is factored out of CallOpInterface and CallableOpInterface, enabling operations to have argument and result attributes without requiring them to be a call or a callable operation.
2025-07-30[MLIR] Avoid #include OMPIRBuilder.h (#151302)Michael Kruse
`#include "llvm/Frontend/OpenMP/OMPIRBuilder.h"` can be replaced with forward-declarations of `OpenMPIRBuilder` and `CanonicalLoopInfo`. This also avoids a dependency to `omp_gen` of the LLVMFrontendOpenMP component which is included indirectly in `OMPIRBuilder.h`. Since its inclusion in #147069, additional indirect dependencies on headers included by `OMPIRBuilder.h` were introduced as well. These are now included directly. Reported-by: fabrizio-indirli See https://github.com/llvm/llvm-project/pull/147069#issuecomment-3114034973
2025-07-29[MLIR][XeVM] Add XeVM to LLVMIR translation. (#150696)Sang Ik Lee
Add XeVM dialect to LLVMIR translation. Currently no ops are translated. Only xevm.DecorationCacheControl are translated to metadata for spirv decoration - !spirv.DecorationCacheControlINTEL. Co-authored-by: Artem Kroviakov artem.kroviakov@intel.com
2025-07-29[mlir][NFC] Use range-based overload of `llvm::sort` (#150934)Longsheng Mou
Replace explicit begin/end iterator pairs with the range-based overload of `llvm::sort`, which simplifies the code and improves readability.
2025-07-29Revert "[mlir][llvm] Add intrinsic arg and result attribute support (… ↵Tobias Gysi
(#151099) …#150783)" This reverts commit 59013d44058ef423a117f95092150e16e16fdb09. The change breaks a flang build bot: https://lab.llvm.org/buildbot/#/builders/207/builds/4441
2025-07-29[mlir][llvm] Add intrinsic arg and result attribute support (#150783)Tobias Gysi
This patch extends the LLVM dialect's intrinsic infra to support argument and result attributes. Initial support is added for the memory intrinsics `llvm.intr.memcpy`, `llvm.intr.memmove`, and `llvm.intr.memset`. Additionally, an ArgAndResultAttrsOpInterface is factored out of CallOpInterface and CallableOpInterface, enabling operations to have argument and result attributes without requiring them to be a call or a callable operation.
2025-07-25[mlir][NFC] Use `getDefiningOp<OpTy>()` instead of ↵Longsheng Mou
`dyn_cast<OpTy>(getDefiningOp())` (#150428) This PR uses `val.getDefiningOp<OpTy>()` to replace `dyn_cast<OpTy>(val.getDefiningOp())` , `dyn_cast_or_null<OpTy>(val.getDefiningOp())` and `dyn_cast_if_present<OpTy>(val.getDefiningOp())`.
2025-07-24[mlir] Remove unused includes (NFC) (#150476)Kazu Hirata
These are identified by misc-include-cleaner. I've filtered out those that break builds. Also, I'm staying away from llvm-config.h, config.h, and Compiler.h, which likely cause platform- or compiler-specific build failures.
2025-07-24[mlir][NFC] Use `hasOneBlock` instead of `llvm::hasSingleElement(region)` ↵Longsheng Mou
(#149809)
2025-07-17[mlir][LLVMIR] Add IFuncOp to LLVM dialect (#147697)Robert Konicar
Add IFunc to LLVM dialect and add support for lifting/exporting LLVMIR IFunc.
2025-07-15[mlir][nvvm] Fix control reaches end of non-void function warning (#148965)Valentin Clement (バレンタイン クレメン)