From ef403b6951852b8dba8523ccfb25ef74565d3d64 Mon Sep 17 00:00:00 2001
From: "enzyme-ci-bot[bot]"
 <78882869+enzyme-ci-bot[bot]@users.noreply.github.com>
Date: Wed, 5 Nov 2025 00:40:36 +0000
Subject: [PATCH] Regenerate MLIR Bindings

---
 src/mlir/Dialects/MosaicGPU.jl | 26 ++++++++++++++++++++++++++
 1 file changed, 26 insertions(+)

diff --git a/src/mlir/Dialects/MosaicGPU.jl b/src/mlir/Dialects/MosaicGPU.jl
index 4ae46ff27a..d8b466bce5 100644
--- a/src/mlir/Dialects/MosaicGPU.jl
+++ b/src/mlir/Dialects/MosaicGPU.jl
@@ -450,6 +450,32 @@ function slice_smem(offset::Value; result_0::IR.Type, location=Location())
     )
 end
 
+"""
+`slice_tmem`
+
+The principal use case for this op is to do a single TMEM allocation and
+slice it into multiple smaller TMEM references. `source` is the large TMEM
+allocation and `offset` is the number of columns to start slicing from.
+"""
+function slice_tmem(source::Value; result_0::IR.Type, offset, location=Location())
+    op_ty_results = IR.Type[result_0,]
+    operands = Value[source,]
+    owned_regions = Region[]
+    successors = Block[]
+    attributes = NamedAttribute[namedattribute("offset", offset),]
+
+    return create_operation(
+        "mosaic_gpu.slice_tmem",
+        location;
+        operands,
+        owned_regions,
+        successors,
+        attributes,
+        results=op_ty_results,
+        result_inference=false,
+    )
+end
+
 """
 `tcgen05_mma`