<feed xmlns='http://www.w3.org/2005/Atom'>
<title>llvm-project.git/mlir/lib/Dialect/LLVMIR, branch main</title>
<subtitle>Unnamed repository; edit this file 'description' to name the repository.
</subtitle>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/'/>
<entry>
<title>[MLIR][LLVM] Extend DIScopeForLLVMFuncOp to handle cross-file operatio… (#167844)</title>
<updated>2025-11-20T11:14:14+00:00</updated>
<author>
<name>Zichen Lu</name>
<email>mikaovo2000@gmail.com</email>
</author>
<published>2025-11-20T11:14:14+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=0a88e9622891aa111d07928d144e042174a9d570'/>
<id>0a88e9622891aa111d07928d144e042174a9d570</id>
<content type='text'>
The current `DIScopeForLLVMFuncOp` pass handles debug information for
inlined code by processing `CallSiteLoc` attributes. However, some
compilation scenarios compose code from multiple source files directly
into a single function without generating `CallSiteLoc`.

**Scenario:**
```python
# a.py
def kernel_a(tensor):
    print("a: {}", tensor)  # a.py:3
    jit_func_b(tensor)           # Calls b.py code

# b.py
def func_b(tensor):
    print("b: {}", tensor)  # b.py:7
```

The scenario executes Python at compile-time and directly inserts
operations from `b.py` into the kernel function, resulting in MLIR like:

```mlir
@kernel_a(...) {
  print("a: {}", %arg0) loc(#loc_a)  // a.py:3
  print("b: {}", %arg0) loc(#loc_b)  // b.py:7 &lt;- FileLineColLoc, not CallSiteLoc
} loc(#loc_kernel)  // a.py:1

#loc1 = loc("a.py":3:.)
#loc2 = loc("b.py":7:.)
#loc_a = loc("print"(#loc1))
#loc_b = loc("print"(#loc2))
```
```llvm
!6 = !DIFile(filename: "a.py", directory: "...")
!9 = distinct !DISubprogram(name: "...", linkageName: "...", scope: !6, file: !6, line: 13, ...)
!10 = !DILocation(line: 7, column: ., scope: !9)  // Points to kernel's DISubprogram, not correct
```</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
The current `DIScopeForLLVMFuncOp` pass handles debug information for
inlined code by processing `CallSiteLoc` attributes. However, some
compilation scenarios compose code from multiple source files directly
into a single function without generating `CallSiteLoc`.

**Scenario:**
```python
# a.py
def kernel_a(tensor):
    print("a: {}", tensor)  # a.py:3
    jit_func_b(tensor)           # Calls b.py code

# b.py
def func_b(tensor):
    print("b: {}", tensor)  # b.py:7
```

The scenario executes Python at compile-time and directly inserts
operations from `b.py` into the kernel function, resulting in MLIR like:

```mlir
@kernel_a(...) {
  print("a: {}", %arg0) loc(#loc_a)  // a.py:3
  print("b: {}", %arg0) loc(#loc_b)  // b.py:7 &lt;- FileLineColLoc, not CallSiteLoc
} loc(#loc_kernel)  // a.py:1

#loc1 = loc("a.py":3:.)
#loc2 = loc("b.py":7:.)
#loc_a = loc("print"(#loc1))
#loc_b = loc("print"(#loc2))
```
```llvm
!6 = !DIFile(filename: "a.py", directory: "...")
!9 = distinct !DISubprogram(name: "...", linkageName: "...", scope: !6, file: !6, line: 13, ...)
!10 = !DILocation(line: 7, column: ., scope: !9)  // Points to kernel's DISubprogram, not correct
```</pre>
</div>
</content>
</entry>
<entry>
<title>[mlir][LLVM] Resync memory effect attribute with LLVM IR (#168568)</title>
<updated>2025-11-19T16:56:04+00:00</updated>
<author>
<name>darkbuck</name>
<email>michael.hliao@gmail.com</email>
</author>
<published>2025-11-19T16:56:04+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=2f6f045ea8d9342a2c57ea93f6343622499dd87a'/>
<id>2f6f045ea8d9342a2c57ea93f6343622499dd87a</id>
<content type='text'>
- Add missing locations, namely 'ErrnoMem', 'TargetMem0', and
'TargetMem1'.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
- Add missing locations, namely 'ErrnoMem', 'TargetMem0', and
'TargetMem1'.</pre>
</div>
</content>
</entry>
<entry>
<title>Reland "[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)" (#168638)</title>
<updated>2025-11-19T09:21:05+00:00</updated>
<author>
<name>Pradeep Kumar</name>
<email>pradeepku@nvidia.com</email>
</author>
<published>2025-11-19T09:21:05+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=4ab1d06406ea425ac40072c3bb3fd96002ba2b0a'/>
<id>4ab1d06406ea425ac40072c3bb3fd96002ba2b0a</id>
<content type='text'>
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");
   ^
```</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
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");
   ^
```</pre>
</div>
</content>
</entry>
<entry>
<title>Revert "[MLIR][NVVM] Add tcgen05.mma MLIR Ops" (#168583)</title>
<updated>2025-11-18T18:15:37+00:00</updated>
<author>
<name>Mehdi Amini</name>
<email>joker.eph@gmail.com</email>
</author>
<published>2025-11-18T18:15:37+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=5407e62611abfbb359f595d89d9f29adf647be02'/>
<id>5407e62611abfbb359f595d89d9f29adf647be02</id>
<content type='text'>
Reverts llvm/llvm-project#164356

The bots are broken.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Reverts llvm/llvm-project#164356

The bots are broken.</pre>
</div>
</content>
</entry>
<entry>
<title>[MLIR][NVVM] Add tcgen05.mma MLIR Ops (#164356)</title>
<updated>2025-11-18T11:10:31+00:00</updated>
<author>
<name>Pradeep Kumar</name>
<email>pradeepku@nvidia.com</email>
</author>
<published>2025-11-18T11:10:31+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=fb829bf11feeb53f815a3abf539e63ec3a23ed3d'/>
<id>fb829bf11feeb53f815a3abf539e63ec3a23ed3d</id>
<content type='text'>
This commit adds support for tgen05.mma family of instructions in the NVVM MLIR dialect and lowers to LLVM Intrinsics. Please refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-instructions) for information</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
This commit adds support for tgen05.mma family of instructions in the NVVM MLIR dialect and lowers to LLVM Intrinsics. Please refer [PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/#tcgen05-mma-instructions) for information</pre>
</div>
</content>
</entry>
<entry>
<title>Fix build breakage (MLIR LLVM dialect requires MLIRInferIntRangeInterface) (#168440)</title>
<updated>2025-11-17T23:30:39+00:00</updated>
<author>
<name>Jeremy Furtek</name>
<email>jfurtek@nvidia.com</email>
</author>
<published>2025-11-17T23:30:39+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=b48f29356641102a52ac8aa05f007bfce719df24'/>
<id>b48f29356641102a52ac8aa05f007bfce719df24</id>
<content type='text'>
This MR fixes a recent build breakage by this MR:
https://github.com/llvm/llvm-project/pull/166648

(Post-merge build error here:
https://lab.llvm.org/buildbot/#/builders/138/builds/21929)

The `MLIRInferIntRangeInterface` library is now a public dependency of
`MLIRLLVMDialect`.</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
This MR fixes a recent build breakage by this MR:
https://github.com/llvm/llvm-project/pull/166648

(Post-merge build error here:
https://lab.llvm.org/buildbot/#/builders/138/builds/21929)

The `MLIRInferIntRangeInterface` library is now a public dependency of
`MLIRLLVMDialect`.</pre>
</div>
</content>
</entry>
<entry>
<title>Fix side effects for LLVM integer operations (udiv, sdiv) incorrectly marked as Pure (#166648)</title>
<updated>2025-11-17T17:46:56+00:00</updated>
<author>
<name>Jeremy Furtek</name>
<email>jfurtek@nvidia.com</email>
</author>
<published>2025-11-17T17:46:56+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=9349a10f93308a196499d2c80a222476c78f1065'/>
<id>9349a10f93308a196499d2c80a222476c78f1065</id>
<content type='text'>
This MR modifies side effect traits of some integer arithmetic
operations in the LLVM dialect.

Prior to this MR, the LLVM dialect `sdiv` and `udiv` operations were
marked as `Pure` through `tblgen` inheritance of the
`LLVM_ArithmeticOpBase` class. The `Pure` trait allowed incorrect
hoisting of `sdiv`/`udiv` operations by the
`loop-independent-code-motion` pass.

This MR modifies the `sdiv` and `udiv` LLVM operations to have traits
and code motion behavior identical to their counterparts in the `arith`
dialect, which were established by the commit/review below.


https://github.com/llvm/llvm-project/commit/ed39825be48805b174d3177f1d8d41ed84784d18
https://reviews.llvm.org/D137814</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
This MR modifies side effect traits of some integer arithmetic
operations in the LLVM dialect.

Prior to this MR, the LLVM dialect `sdiv` and `udiv` operations were
marked as `Pure` through `tblgen` inheritance of the
`LLVM_ArithmeticOpBase` class. The `Pure` trait allowed incorrect
hoisting of `sdiv`/`udiv` operations by the
`loop-independent-code-motion` pass.

This MR modifies the `sdiv` and `udiv` LLVM operations to have traits
and code motion behavior identical to their counterparts in the `arith`
dialect, which were established by the commit/review below.


https://github.com/llvm/llvm-project/commit/ed39825be48805b174d3177f1d8d41ed84784d18
https://reviews.llvm.org/D137814</pre>
</div>
</content>
</entry>
<entry>
<title>[MLIR][NVVM] Add support for shared::cta destination (#168056)</title>
<updated>2025-11-17T08:28:21+00:00</updated>
<author>
<name>Durgadoss R</name>
<email>durgadossr@nvidia.com</email>
</author>
<published>2025-11-17T08:28:21+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=95aa70cf209ab317d908e38630c92bf62d149247'/>
<id>95aa70cf209ab317d908e38630c92bf62d149247</id>
<content type='text'>
This patch adds support for shared::cta as destination space in
the TMA non-tensor copy Op (from global to shared::cta).

* Appropriate verifier checks are added.
* Unit tests are added to verify the lowering.

The related intrinsic changes were merged through PR #167508.

Signed-off-by: Durgadoss R &lt;durgadossr@nvidia.com&gt;</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
This patch adds support for shared::cta as destination space in
the TMA non-tensor copy Op (from global to shared::cta).

* Appropriate verifier checks are added.
* Unit tests are added to verify the lowering.

The related intrinsic changes were merged through PR #167508.

Signed-off-by: Durgadoss R &lt;durgadossr@nvidia.com&gt;</pre>
</div>
</content>
</entry>
<entry>
<title>[mlir] Add FP software implementation lowering pass: `arith-to-apfloat` (#167848)</title>
<updated>2025-11-13T09:35:30+00:00</updated>
<author>
<name>Matthias Springer</name>
<email>me@m-sp.org</email>
</author>
<published>2025-11-13T09:35:30+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=7a53d33e7c29ce3058eb12d2f85ca0a46b0ea363'/>
<id>7a53d33e7c29ce3058eb12d2f85ca0a46b0ea363</id>
<content type='text'>
Reland pass and fix linker errors.

---------

Co-authored-by: Maksim Levental &lt;maksim.levental@gmail.com&gt;</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Reland pass and fix linker errors.

---------

Co-authored-by: Maksim Levental &lt;maksim.levental@gmail.com&gt;</pre>
</div>
</content>
</entry>
<entry>
<title>Revert "Reland yet again: [mlir] Add FP software implementation lowering pass: `arith-to-apfloat`" (#167834)</title>
<updated>2025-11-13T07:02:21+00:00</updated>
<author>
<name>Maksim Levental</name>
<email>maksim.levental@gmail.com</email>
</author>
<published>2025-11-13T07:02:21+00:00</published>
<link rel='alternate' type='text/html' href='https://git.belthelziquor.com/llvm-project.git/commit/?id=140e07c862ccf02c3c7412187ce79139ac45f017'/>
<id>140e07c862ccf02c3c7412187ce79139ac45f017</id>
<content type='text'>
Reverts llvm/llvm-project#167608

Broken builder https://lab.llvm.org/buildbot/#/builders/52/builds/12781</content>
<content type='xhtml'>
<div xmlns='http://www.w3.org/1999/xhtml'>
<pre>
Reverts llvm/llvm-project#167608

Broken builder https://lab.llvm.org/buildbot/#/builders/52/builds/12781</pre>
</div>
</content>
</entry>
</feed>
