This is an archive of the discontinued LLVM Phabricator instance.

[mlir][NVVM] Add support for nvvm mma.sync ops
ClosedPublic

Authored by christopherbate on Mar 24 2022, 9:34 AM.

Details

Summary

This patch adds MLIR NVVM support for the various NVPTX mma.sync
operations. There are a number of possible data type, shape,
and other attribute combinations supported by the operation, so a
custom assebmly format is added and attributes are inferred where
possible.

Diff Detail

Event Timeline

Herald added a project: Restricted Project. · View Herald Transcript

Fix MmaOp::parse rebase issues

christopherbate published this revision for review.Mar 24 2022, 10:49 AM

Looks good to me overall. Added some minor comments.

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
927

Is there any advantage having this optional? Always having it looks more consistent to me.

mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
73

I would remove the static and add a some basic doc explaining what the function does,

115

nit: missing space.

181

I'd remove that, I don't see it anywhere else in mlir

373

remove?

christopherbate marked 3 inline comments as done.

Address comments and add missing f64 mma cases.

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
927

I could make it non-optional, but it would increase the verbosity in the case where the two attributes can be easily inferred (for all types above s8/u8).

mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
181

The problem is that if you try to build this op using the default builders, it's a pain because of all the details. I'm using it in lowering pathways I've implemented already, but I didn't see an easy place to insert a unit test for builders like this

373

Yes and it looks like I'm missing f64 tests/support below for m=8, I'll fix that

ThomasRaoux accepted this revision.Mar 24 2022, 9:25 PM
ThomasRaoux added inline comments.
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
927

your call, I feel like consistency is more convenient but I can see both side.

mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
181

I just meant the comment :)

This revision is now accepted and ready to land.Mar 24 2022, 9:25 PM
This revision was automatically updated to reflect the committed changes.