MmaSyncBuilder Struct Reference

Helper struct to provide a simple mapping from matmul operations to the corresponding mma.sync operation. More...

Public Types
using	IndexCalculator = std::function< SmallVector< RowColIndexing >(MLIRContext *)>

Public Member Functions
	MmaSyncBuilder (OpBuilder &b, Location loc, OpFoldResult laneId)
FailureOr< Operation * >	buildMmaSync (LinalgOp linalgOp)
	Create the mma.sync operation corresponding to linalgOp along with all the supporting load/store and vector operations. More...

Detailed Description

Helper struct to provide a simple mapping from matmul operations to the corresponding mma.sync operation.

This is constrained to the case where the matmul matches the mma.sync operation 1-1.

Definition at line 440 of file NVGPUTransformOps.cpp.

Member Typedef Documentation

IndexCalculator

using MmaSyncBuilder::IndexCalculator = std::function<SmallVector<RowColIndexing>(MLIRContext *)>

Definition at line 444 of file NVGPUTransformOps.cpp.

Constructor & Destructor Documentation

MmaSyncBuilder()

MmaSyncBuilder::MmaSyncBuilder ( OpBuilder &  b, Location  loc, OpFoldResult  laneId  )

inline

Definition at line 441 of file NVGPUTransformOps.cpp.

Member Function Documentation

buildMmaSync()

FailureOr< Operation * > MmaSyncBuilder::buildMmaSync

(

LinalgOp

linalgOp

)

Create the mma.sync operation corresponding to linalgOp along with all the supporting load/store and vector operations.

Definition at line 772 of file NVGPUTransformOps.cpp.

References mlir::remark::failed(), mlir::Value::getDefiningOp(), mlir::getElementTypeOrSelf(), and mlir::Value::getType().

---

The documentation for this struct was generated from the following file:

• lib/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp