GPUDialect.cpp File Reference

#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/Arith/IR/Arith.h"
#include "mlir/Dialect/Bufferization/IR/BufferDeallocationOpInterface.h"
#include "mlir/Dialect/Math/IR/Math.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/Diagnostics.h"
#include "mlir/IR/DialectImplementation.h"
#include "mlir/IR/Matchers.h"
#include "mlir/IR/OpImplementation.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/IR/SymbolTable.h"
#include "mlir/IR/TypeUtilities.h"
#include "mlir/Interfaces/FunctionImplementation.h"
#include "mlir/Interfaces/SideEffectInterfaces.h"
#include "mlir/Interfaces/ValueBoundsOpInterface.h"
#include "mlir/Transforms/InliningUtils.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/TypeSwitch.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/FormatVariadic.h"
#include "llvm/Support/InterleavedRange.h"
#include "llvm/Support/StringSaver.h"
#include <cassert>
#include <numeric>
#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"
#include "mlir/Dialect/GPU/IR/CompilationAttrInterfaces.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

Functions
static std::string	getSparseHandleKeyword (SparseHandleKind kind)
static LogicalResult	verifyKnownLaunchSizeAttr (Operation *op, NamedAttribute attr)
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, ArrayAttr attributes={})
static LogicalResult	verifyAttributions (Operation *op, ArrayRef< BlockArgument > attributions, gpu::AddressSpace memorySpace)
	Verifies a GPU function memory attribution. More...
static LogicalResult	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	parseLaunchDimType (OpAsmParser &parser, Type &dimTy, std::optional< OpAsmParser::UnresolvedOperand > clusterValue, Type &clusterXTy, Type &clusterYTy, Type &clusterZTy)
static void	printLaunchDimType (OpAsmPrinter &printer, Operation *op, Type dimTy, Value clusterValue, Type clusterXTy, Type clusterYTy, Type clusterZTy)
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 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 ParseResult	parseOffloadingHandler (OpAsmParser &parser, Attribute &offloadingHandler)
static void	printOffloadingHandler (OpAsmPrinter &printer, Operation *op, Attribute offloadingHandler)
static LogicalResult	verifyDistributedType (Type expanded, Type distributed, int64_t warpSize, Operation *op)
	Helper check if the distributed vector type is consistent with the expanded type and distributed size. More...

Macro Definition Documentation

GET_ATTRDEF_CLASSES

#define GET_ATTRDEF_CLASSES

Definition at line 2771 of file GPUDialect.cpp.

GET_ATTRDEF_LIST

#define GET_ATTRDEF_LIST

GET_OP_CLASSES

#define GET_OP_CLASSES

Definition at line 2774 of file GPUDialect.cpp.

GET_OP_LIST

#define GET_OP_LIST

Function Documentation

canMakeGroupOpUniform()

static bool canMakeGroupOpUniform ( Operation *  op )

static

Definition at line 641 of file GPUDialect.cpp.

References mlir::Region::empty(), mlir::Region::front(), mlir::Operation::getBlock(), and mlir::Operation::getParentOp().

getAttributionAttr()

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

static

Definition at line 1730 of file GPUDialect.cpp.

References getAttributionAttrs().

getAttributionAttrs()

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

static

Definition at line 1687 of file GPUDialect.cpp.

Referenced by getAttributionAttr(), and setAttributionAttr().

getSparseHandleKeyword()

static std::string getSparseHandleKeyword ( SparseHandleKind  kind )

static

Definition at line 291 of file GPUDialect.cpp.

References mlir::gpu::DnTensor, kind, mlir::gpu::SpGEMMOp, and mlir::gpu::SpMat.

parseAllReduceOperation()

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

static

Definition at line 663 of file GPUDialect.cpp.

References mlir::AsmParser::emitError(), mlir::get(), mlir::AsmParser::getContext(), mlir::AsmParser::getCurrentLocation(), and 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 491 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(), and mlir::AsmParser::parseOptionalKeyword().

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 528 of file GPUDialect.cpp.

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

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 1545 of file GPUDialect.cpp.

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

parseLaunchDimType()

static ParseResult parseLaunchDimType ( OpAsmParser &  parser, Type &  dimTy, std::optional< OpAsmParser::UnresolvedOperand >  clusterValue, Type &  clusterXTy, Type &  clusterYTy, Type &  clusterZTy  )

static

Definition at line 1385 of file GPUDialect.cpp.

References mlir::get(), mlir::AsmParser::getContext(), mlir::AsmParser::parseOptionalColon(), and mlir::AsmParser::parseType().

parseLaunchFuncOperands()

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

static

Definition at line 1407 of file GPUDialect.cpp.

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

parseOffloadingHandler()

static ParseResult parseOffloadingHandler ( OpAsmParser &  parser, Attribute &  offloadingHandler  )

static

Definition at line 1939 of file GPUDialect.cpp.

References mlir::Builder::getAttr(), mlir::AsmParser::getBuilder(), mlir::AsmParser::parseAttribute(), mlir::AsmParser::parseGreater(), and mlir::AsmParser::parseOptionalLess().

parseSizeAssignment()

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

static

Definition at line 997 of file GPUDialect.cpp.

References 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 676 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 507 of file GPUDialect.cpp.

printAttributions()

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

static

Definition at line 538 of file GPUDialect.cpp.

printLaunchDimType()

static void printLaunchDimType ( OpAsmPrinter &  printer, Operation *  op, Type  dimTy, Value  clusterValue, Type  clusterXTy, Type  clusterYTy, Type  clusterZTy  )

static

Definition at line 1400 of file GPUDialect.cpp.

References mlir::Type::isIndex().

printLaunchFuncOperands()

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

static

Definition at line 1423 of file GPUDialect.cpp.

printOffloadingHandler()

static void printOffloadingHandler ( OpAsmPrinter &  printer, Operation *  op, Attribute  offloadingHandler  )

static

Definition at line 1952 of file GPUDialect.cpp.

References mlir::get(), and mlir::Operation::getContext().

printSizeAssignment()

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

static

Definition at line 932 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 1754 of file GPUDialect.cpp.

References getAttributionAttrs(), and setAttributionAttrs().

setAttributionAttrs()

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

static

Definition at line 1703 of file GPUDialect.cpp.

References mlir::get().

Referenced by setAttributionAttr().

verifyAttributions()

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

static

Verifies a GPU function memory attribution.

Definition at line 561 of file GPUDialect.cpp.

References mlir::Operation::emitOpError(), and mlir::Value::getType().

verifyDistributedType()

static LogicalResult verifyDistributedType ( Type  expanded, Type  distributed, int64_t  warpSize, Operation *  op  )

static

Helper check if the distributed vector type is consistent with the expanded type and distributed size.

Definition at line 2436 of file GPUDialect.cpp.

verifyKnownLaunchSizeAttr()

static LogicalResult verifyKnownLaunchSizeAttr ( Operation *  op, NamedAttribute  attr  )

static

Definition at line 381 of file GPUDialect.cpp.

References mlir::Operation::emitOpError(), mlir::NamedAttribute::getName(), and mlir::NamedAttribute::getValue().

verifyReduceOpAndType()

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

static

Definition at line 587 of file GPUDialect.cpp.

References MINUI.