This is an archive of the discontinued LLVM Phabricator instance.

Revert "Revert "[mlir][Transform] Add support for mma.sync m16n8k16 f16 rewrite." and "[mlir][Transform] Introduce nvgpu transform extensions""
ClosedPublic

Authored by nicolasvasilache on Jun 27 2023, 2:02 AM.

Diff Detail

Event Timeline

Herald added a project: Restricted Project. · View Herald Transcript
nicolasvasilache requested review of this revision.Jun 27 2023, 2:02 AM
Herald added a project: Restricted Project. · View Herald Transcript

Help figuring out the proper lit incantations to filter these tests out on non-sm80 would be most welcome.

I could not find what's the error, but setting ptx version might solve.

mlir/test/Integration/GPU/CUDA/TensorCore/transform-mma-sync-matmul-f16-f16-accum.mlir
15 ↗(On Diff #534874)

Have you tried setting the PTX version? gpu-to-cubin sets StringRef properties="+ptx60" by default. This version does not have mma.sync instructions.
You can try this ...gpu-to-cubin{chip=sm_80 properties=+ptx76}.

Add properties=+ptx76 as suggested

Harbormaster says green but I have no idea whether these are the buildbots that @mehdi_amini mentioned.
Could you please share a link that I should remember to look at when doing this kind of work?

nicolasvasilache edited the summary of this revision. (Show Details)

Update tests behind an sm80 flag.

This revision was not accepted when it landed; it landed in state Needs Review.Jun 27 2023, 11:50 PM
This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.

link for monitoring post-submit NVIDIA buildbots: https://lab.llvm.org/buildbot/#/builders/61