[all-commits] [llvm/llvm-project] 987980: [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM ...

[Chris via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 9879807393d3f502d3cac468c5f6451db872aa5f
      https://github.com/llvm/llvm-project/commit/9879807393d3f502d3cac468c5f6451db872aa5f
  Author: Christopher Bate <cbate at nvidia.com>
  Date:   2022-05-08 (Sun, 08 May 2022)

  Changed paths:
    M mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
    M mlir/test/Conversion/NVGPUToNVVM/mma-sync-to-nvvm.mlir

  Log Message:
  -----------
  [mlir][NvGpu] Fix nvgpu.mma.sync lowering to NVVM for f32, tf32 types

Adds missing logic in the lowering from NvGPU to NVVM to support fp32
(in an accumulator operand) and tf32 (in multiplicand operand) types.
Fixes logic in one of the helper functions for converting the result
of a mma.sync operation with multiple 8x256bit output tiles, which is
the case for f32 outputs.

Differential Revision: https://reviews.llvm.org/D124533