CreateAsyncGroups.cpp File Reference

#include "mlir/Dialect/NVGPU/Transforms/Transforms.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.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. More...
static bool	isContiguousStore (Operation *write)
	Return "true" if the given op is a contiguous and suitable vector.transfer_write or vector.store op. More...
static bool	isContiguousRead (Operation *read)
	Return "true" if the given op is a contiguous and suitable vector.transfer_read or vector.load op. More...
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. More...
static Value	buildNumReadElements (OpBuilder &b, Location loc, Operation *readOp)
	Build an SSA value that represents the number of read elements. More...
static bool	resultsInSupportedAsyncCopy (MemRefType memrefType, VectorType vecType)
	Return "true" if the conversion to async copy is supported by "async copy". More...

Function Documentation

buildNumReadElements()

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

static

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

Definition at line 84 of file CreateAsyncGroups.cpp.

References mlir::OpBuilder::create(), and getMaskOp().

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

getMaskOp()

static 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 59 of file CreateAsyncGroups.cpp.

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

isContiguousRead()

static 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 42 of file CreateAsyncGroups.cpp.

References isContiguousXferOp().

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

isContiguousStore()

static 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 33 of file CreateAsyncGroups.cpp.

References isContiguousXferOp().

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

isContiguousXferOp()

template<typename OpTy >

static 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 24 of file CreateAsyncGroups.cpp.

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

Referenced by isContiguousRead(), and isContiguousStore().

resultsInSupportedAsyncCopy()

static bool resultsInSupportedAsyncCopy ( MemRefType  memrefType, VectorType  vecType  )

static

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

Definition at line 127 of file CreateAsyncGroups.cpp.

References mlir::Type::getIntOrFloatBitWidth().

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