CreateAsyncGroups.cpp File Reference

#include "mlir/Dialect/NVGPU/Transforms/Transforms.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h"
#include "mlir/Dialect/NVGPU/Transforms/Utils.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinTypes.h"

Go to the source code of this file.

Functions
template<typename OpTy>
static bool	isContiguousXferOp (OpTy op)
	Return "true" if the given vector transfer op is contiguous and suitable for replacement with an async copy.
static bool	isContiguousStore (Operation *write)
	Return "true" if the given op is a contiguous and suitable vector.transfer_write or vector.store op.
static bool	isContiguousRead (Operation *read)
	Return "true" if the given op is a contiguous and suitable vector.transfer_read or vector.load op.
static FailureOr< TransferMask >	getMaskOp (Operation *loadOp)
	If the given vector load op has a mask that is defined by vector.create_mask, return that op.
static Value	buildNumReadElements (OpBuilder &b, Location loc, Operation *readOp)
	Build an SSA value that represents the number of read elements.
static bool	resultsInSupportedAsyncCopy (MemRefType memrefType, VectorType vecType)
	Return "true" if the conversion to async copy is supported by "async copy".

Function Documentation

buildNumReadElements()

Value buildNumReadElements ( OpBuilder & b, Location loc, Operation * readOp )

static

Build an SSA value that represents the number of read elements.

Definition at line 83 of file CreateAsyncGroups.cpp.

References b, mlir::arith::ConstantIndexOp::create(), and getMaskOp().

Referenced by mlir::nvgpu::createAsyncGroups().

getMaskOp()

FailureOr< TransferMask > getMaskOp ( Operation * loadOp )

static

If the given vector load op has a mask that is defined by vector.create_mask, return that op.

Definition at line 58 of file CreateAsyncGroups.cpp.

Referenced by buildNumReadElements(), and mlir::nvgpu::createAsyncGroups().

isContiguousRead()

bool isContiguousRead ( Operation * read )

static

Return "true" if the given op is a contiguous and suitable vector.transfer_read or vector.load op.

Definition at line 41 of file CreateAsyncGroups.cpp.

References isContiguousXferOp().

Referenced by mlir::nvgpu::createAsyncGroups().

isContiguousStore()

bool isContiguousStore ( Operation * write )

static

Return "true" if the given op is a contiguous and suitable vector.transfer_write or vector.store op.

Definition at line 32 of file CreateAsyncGroups.cpp.

References isContiguousXferOp().

Referenced by mlir::nvgpu::createAsyncGroups().

isContiguousXferOp()

template<typename OpTy>

bool isContiguousXferOp ( OpTy op )

static

Return "true" if the given vector transfer op is contiguous and suitable for replacement with an async copy.

Definition at line 23 of file CreateAsyncGroups.cpp.

References mlir::nvgpu::getMemrefOperand(), and mlir::getType().

Referenced by isContiguousRead(), and isContiguousStore().

resultsInSupportedAsyncCopy()

bool resultsInSupportedAsyncCopy ( MemRefType memrefType, VectorType vecType )

static

Return "true" if the conversion to async copy is supported by "async copy".

Definition at line 126 of file CreateAsyncGroups.cpp.

References mlir::Type::getIntOrFloatBitWidth().

Referenced by mlir::nvgpu::createAsyncGroups().