GPUDialect.cpp File Reference

#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/BuiltinAttributes.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/BuiltinTypes.h"
#include "mlir/IR/DialectImplementation.h"
#include "mlir/IR/FunctionImplementation.h"
#include "mlir/IR/Matchers.h"
#include "mlir/IR/OpImplementation.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/IR/TypeUtilities.h"
#include "mlir/Interfaces/SideEffectInterfaces.h"
#include "mlir/Transforms/InliningUtils.h"
#include "llvm/ADT/TypeSwitch.h"
#include "mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc"
#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
#include "mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc"

Go to the source code of this file.

Classes
struct	FoldLaunchArguments
	Simplify the gpu.launch when the range of a thread or block ID is trivially known to be one. More...

Macros
#define	GET_OP_LIST
#define	GET_ATTRDEF_LIST
#define	GET_ATTRDEF_CLASSES
#define	GET_OP_CLASSES

Enumerations
enum	GPUMemorySpace { kGenericMemorySpace = 0 , kGlobalMemorySpace = 1 , kSharedMemorySpace = 3 }
	GPU memory space identifiers. More...

Functions
static std::string	getSparseHandleKeyword (SparseHandleKind kind)
static ParseResult	parseAsyncDependencies (OpAsmParser &parser, Type &asyncTokenType, SmallVectorImpl< OpAsmParser::UnresolvedOperand > &asyncDependencies)
	Parses an optional list of async operands with an optional leading keyword. More...
static void	printAsyncDependencies (OpAsmPrinter &printer, Operation *op, Type asyncTokenType, OperandRange asyncDependencies)
	Prints optional async dependencies with its leading keyword. More...
static ParseResult	parseAttributions (OpAsmParser &parser, StringRef keyword, SmallVectorImpl< OpAsmParser::Argument > &args)
	Parses a GPU function memory attribution. More...
static void	printAttributions (OpAsmPrinter &p, StringRef keyword, ArrayRef< BlockArgument > values)
	Prints a GPU function memory attribution. More...
static LogicalResult	verifyAttributions (Operation *op, ArrayRef< BlockArgument > attributions, gpu::AddressSpace memorySpace)
	Verifies a GPU function memory attribution. More...
static bool	verifyReduceOpAndType (gpu::AllReduceOperation opName, Type resType)
static bool	canMakeGroupOpUniform (Operation *op)
static ParseResult	parseAllReduceOperation (AsmParser &parser, AllReduceOperationAttr &attr)
static void	printAllReduceOperation (AsmPrinter &printer, Operation *op, AllReduceOperationAttr attr)
static void	printSizeAssignment (OpAsmPrinter &p, KernelDim3 size, KernelDim3 operands, KernelDim3 ids)
static ParseResult	parseSizeAssignment (OpAsmParser &parser, MutableArrayRef< OpAsmParser::UnresolvedOperand > sizes, MutableArrayRef< OpAsmParser::UnresolvedOperand > regionSizes, MutableArrayRef< OpAsmParser::UnresolvedOperand > indices)
static ParseResult	parseLaunchFuncOperands (OpAsmParser &parser, SmallVectorImpl< OpAsmParser::UnresolvedOperand > &argNames, SmallVectorImpl< Type > &argTypes)
static void	printLaunchFuncOperands (OpAsmPrinter &printer, Operation *, OperandRange operands, TypeRange types)
static ParseResult	parseAttributions (OpAsmParser &parser, StringRef keyword, SmallVectorImpl< OpAsmParser::Argument > &args, Attribute &attributionAttrs)
	Parses a GPU function memory attribution. More...
static void	printAttributions (OpAsmPrinter &p, StringRef keyword, ArrayRef< BlockArgument > values, ArrayAttr attributes)
static DictionaryAttr	getAttributionAttrs (GPUFuncOp op, unsigned index, StringAttr attrName)
static void	setAttributionAttrs (GPUFuncOp op, unsigned index, DictionaryAttr value, StringAttr attrName)
static Attribute	getAttributionAttr (GPUFuncOp op, unsigned index, StringAttr name, StringAttr attrsName)
static void	setAttributionAttr (GPUFuncOp op, unsigned index, StringAttr name, Attribute value, StringAttr attrsName)
static LogicalResult	verifyKnownLaunchSizeAttr (gpu::GPUFuncOp op, StringRef attrName)
static bool	isLastMemrefDimUnitStride (MemRefType type)
	Return true if the last dimension of the MemRefType has unit stride. More...

Macro Definition Documentation

GET_ATTRDEF_CLASSES

#define GET_ATTRDEF_CLASSES

Definition at line 1784 of file GPUDialect.cpp.

GET_ATTRDEF_LIST

#define GET_ATTRDEF_LIST

GET_OP_CLASSES

#define GET_OP_CLASSES

Definition at line 1787 of file GPUDialect.cpp.

GET_OP_LIST

#define GET_OP_LIST

Enumeration Type Documentation

GPUMemorySpace

enum GPUMemorySpace

GPU memory space identifiers.

Enumerator
kGenericMemorySpace	Generic memory space identifier.
kGlobalMemorySpace	Global memory space identifier.
kSharedMemorySpace	Shared memory space identifier.

Definition at line 117 of file GPUDialect.cpp.

Function Documentation

canMakeGroupOpUniform()

static bool canMakeGroupOpUniform ( Operation *  op )

static

Definition at line 471 of file GPUDialect.cpp.

getAttributionAttr()

static Attribute getAttributionAttr ( GPUFuncOp  op, unsigned  index, StringAttr  name, StringAttr  attrsName  )

static

Definition at line 1302 of file GPUDialect.cpp.

getAttributionAttrs()

static DictionaryAttr getAttributionAttrs ( GPUFuncOp  op, unsigned  index, StringAttr  attrName  )

static

Definition at line 1259 of file GPUDialect.cpp.

getSparseHandleKeyword()

static std::string getSparseHandleKeyword ( SparseHandleKind  kind )

static

Definition at line 164 of file GPUDialect.cpp.

isLastMemrefDimUnitStride()

static bool isLastMemrefDimUnitStride ( MemRefType  type )

static

Return true if the last dimension of the MemRefType has unit stride.

Also return true for memrefs with no strides.

Definition at line 1564 of file GPUDialect.cpp.

References mlir::failed(), and mlir::getStridesAndOffset().

parseAllReduceOperation()

static ParseResult parseAllReduceOperation ( AsmParser &  parser, AllReduceOperationAttr &  attr  )

static

Definition at line 493 of file GPUDialect.cpp.

References mlir::AsmParser::parseOptionalKeyword().

parseAsyncDependencies()

static ParseResult parseAsyncDependencies ( OpAsmParser &  parser, Type &  asyncTokenType, SmallVectorImpl< OpAsmParser::UnresolvedOperand > &  asyncDependencies  )

static

Parses an optional list of async operands with an optional leading keyword.

(async)? ([ ssa-id-list ])?

This method is used by the tablegen assembly format for async ops as well.

Definition at line 342 of file GPUDialect.cpp.

References mlir::AsmParser::emitError(), mlir::AsmParser::getBuilder(), mlir::AsmParser::getCurrentLocation(), mlir::OpAsmParser::getNumResults(), mlir::Builder::getType(), mlir::AsmParser::OptionalSquare, mlir::OpAsmParser::parseOperandList(), mlir::AsmParser::parseOptionalKeyword(), and mlir::succeeded().

parseAttributions() [1/2]

static ParseResult parseAttributions ( OpAsmParser &  parser, StringRef  keyword, SmallVectorImpl< OpAsmParser::Argument > &  args  )

static

Parses a GPU function memory attribution.

memory-attribution ::= (workgroup ( ssa-id-and-type-list ))? (private ( ssa-id-and-type-list ))?

Note that this function parses only one of the two similar parts, with the keyword provided as argument.

Definition at line 381 of file GPUDialect.cpp.

References mlir::failed(), mlir::AsmParser::Paren, mlir::OpAsmParser::parseArgumentList(), mlir::AsmParser::parseOptionalKeyword(), and mlir::success().

parseAttributions() [2/2]

static ParseResult parseAttributions ( OpAsmParser &  parser, StringRef  keyword, SmallVectorImpl< OpAsmParser::Argument > &  args, Attribute &  attributionAttrs  )

static

Parses a GPU function memory attribution.

memory-attribution ::= (workgroup ( ssa-id-and-type-list ))? (private ( ssa-id-and-type-list ))?

Note that this function parses only one of the two similar parts, with the keyword provided as argument.

Definition at line 1095 of file GPUDialect.cpp.

References mlir::OpAsmParser::Argument::attrs, mlir::failed(), mlir::Builder::getArrayAttr(), mlir::AsmParser::getBuilder(), mlir::Builder::getDictionaryAttr(), mlir::AsmParser::Paren, mlir::OpAsmParser::parseArgumentList(), mlir::AsmParser::parseOptionalKeyword(), and mlir::success().

parseLaunchFuncOperands()

static ParseResult parseLaunchFuncOperands ( OpAsmParser &  parser, SmallVectorImpl< OpAsmParser::UnresolvedOperand > &  argNames, SmallVectorImpl< Type > &  argTypes  )

static

Definition at line 994 of file GPUDialect.cpp.

References mlir::failure(), mlir::AsmParser::Paren, mlir::AsmParser::parseColonType(), mlir::AsmParser::parseCommaSeparatedList(), mlir::OpAsmParser::parseOperand(), mlir::AsmParser::parseOptionalKeyword(), and mlir::success().

parseSizeAssignment()

static ParseResult parseSizeAssignment ( OpAsmParser &  parser, MutableArrayRef< OpAsmParser::UnresolvedOperand >  sizes, MutableArrayRef< OpAsmParser::UnresolvedOperand >  regionSizes, MutableArrayRef< OpAsmParser::UnresolvedOperand >  indices  )

static

Definition at line 725 of file GPUDialect.cpp.

References mlir::failure(), mlir::AsmParser::Paren, mlir::AsmParser::parseComma(), mlir::AsmParser::parseEqual(), mlir::AsmParser::parseKeyword(), mlir::AsmParser::parseLParen(), mlir::OpAsmParser::parseOperand(), mlir::OpAsmParser::parseOperandList(), and mlir::AsmParser::parseRParen().

printAllReduceOperation()

static void printAllReduceOperation ( AsmPrinter &  printer, Operation *  op, AllReduceOperationAttr  attr  )

static

Definition at line 506 of file GPUDialect.cpp.

printAsyncDependencies()

static void printAsyncDependencies ( OpAsmPrinter &  printer, Operation *  op, Type  asyncTokenType, OperandRange  asyncDependencies  )

static

Prints optional async dependencies with its leading keyword.

(async)? ([ ssa-id-list ])?

Definition at line 358 of file GPUDialect.cpp.

printAttributions() [1/2]

static void printAttributions ( OpAsmPrinter &  p, StringRef  keyword, ArrayRef< BlockArgument >  values  )

static

Prints a GPU function memory attribution.

Definition at line 392 of file GPUDialect.cpp.

References mlir::Value::getType().

printAttributions() [2/2]

static void printAttributions ( OpAsmPrinter &  p, StringRef  keyword, ArrayRef< BlockArgument >  values, ArrayAttr  attributes  )

static

Definition at line 1211 of file GPUDialect.cpp.

References mlir::detail::enumerate(), mlir::Value::getType(), and mlir::OpAsmPrinter::printOptionalAttrDict().

printLaunchFuncOperands()

static void printLaunchFuncOperands ( OpAsmPrinter &  printer, Operation *  , OperandRange  operands, TypeRange  types  )

static

Definition at line 1010 of file GPUDialect.cpp.

References mlir::OpAsmPrinter::printOperand(), and mlir::AsmPrinter::printType().

printSizeAssignment()

static void printSizeAssignment ( OpAsmPrinter &  p, KernelDim3  size, KernelDim3  operands, KernelDim3  ids  )

static

Definition at line 682 of file GPUDialect.cpp.

References mlir::gpu::KernelDim3::x, mlir::gpu::KernelDim3::y, and mlir::gpu::KernelDim3::z.

setAttributionAttr()

static void setAttributionAttr ( GPUFuncOp  op, unsigned  index, StringAttr  name, Attribute  value, StringAttr  attrsName  )

static

Definition at line 1326 of file GPUDialect.cpp.

setAttributionAttrs()

static void setAttributionAttrs ( GPUFuncOp  op, unsigned  index, DictionaryAttr  value, StringAttr  attrName  )

static

Definition at line 1275 of file GPUDialect.cpp.

verifyAttributions()

static LogicalResult verifyAttributions ( Operation *  op, ArrayRef< BlockArgument >  attributions, gpu::AddressSpace  memorySpace  )

static

Verifies a GPU function memory attribution.

Definition at line 404 of file GPUDialect.cpp.

verifyKnownLaunchSizeAttr()

static LogicalResult verifyKnownLaunchSizeAttr ( gpu::GPUFuncOp  op, StringRef  attrName  )

static

Definition at line 1414 of file GPUDialect.cpp.

verifyReduceOpAndType()

static bool verifyReduceOpAndType ( gpu::AllReduceOperation  opName, Type  resType  )

static

Definition at line 430 of file GPUDialect.cpp.