Skip to content

Commit acee27c

Browse files
Regenerate MLIR Bindings (#1825)
Co-authored-by: enzyme-ci-bot[bot] <78882869+enzyme-ci-bot[bot]@users.noreply.github.com>
1 parent 4f318d5 commit acee27c

File tree

1 file changed

+26
-0
lines changed

1 file changed

+26
-0
lines changed

src/mlir/Dialects/MosaicGPU.jl

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -450,6 +450,32 @@ function slice_smem(offset::Value; result_0::IR.Type, location=Location())
450450
)
451451
end
452452

453+
"""
454+
`slice_tmem`
455+
456+
The principal use case for this op is to do a single TMEM allocation and
457+
slice it into multiple smaller TMEM references. `source` is the large TMEM
458+
allocation and `offset` is the number of columns to start slicing from.
459+
"""
460+
function slice_tmem(source::Value; result_0::IR.Type, offset, location=Location())
461+
op_ty_results = IR.Type[result_0,]
462+
operands = Value[source,]
463+
owned_regions = Region[]
464+
successors = Block[]
465+
attributes = NamedAttribute[namedattribute("offset", offset),]
466+
467+
return create_operation(
468+
"mosaic_gpu.slice_tmem",
469+
location;
470+
operands,
471+
owned_regions,
472+
successors,
473+
attributes,
474+
results=op_ty_results,
475+
result_inference=false,
476+
)
477+
end
478+
453479
"""
454480
`tcgen05_mma`
455481

0 commit comments

Comments
 (0)