#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/Dialect/Utils/VerificationUtils.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 <optional>
#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"

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.
static void	printAsyncDependencies (OpAsmPrinter &printer, Operation *op, Type asyncTokenType, OperandRange asyncDependencies)
	Prints optional async dependencies with its leading keyword.
static ParseResult	parseAttributions (OpAsmParser &parser, StringRef keyword, SmallVectorImpl< OpAsmParser::Argument > &args)
	Parses a GPU function memory attribution.
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.
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, StringRef keyword)
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 LogicalResult	eraseRedundantGpuBarrierOps (BarrierOp op, PatternRewriter &rewriter)
	Remove gpu.barrier after gpu.barrier, the threads are already synchronized!
static ParseResult	parseAttributions (OpAsmParser &parser, StringRef keyword, SmallVectorImpl< OpAsmParser::Argument > &args, Attribute &attributionAttrs)
	Parses a GPU function memory attribution.
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.

Macro Definition Documentation

◆ GET_ATTRDEF_CLASSES

#define GET_ATTRDEF_CLASSES

Definition at line 2914 of file GPUDialect.cpp.

◆ GET_ATTRDEF_LIST

#define GET_ATTRDEF_LIST

◆ GET_OP_CLASSES

#define GET_OP_CLASSES

Definition at line 2917 of file GPUDialect.cpp.

◆ GET_OP_LIST

#define GET_OP_LIST

Function Documentation

◆ canMakeGroupOpUniform()

bool canMakeGroupOpUniform ( Operation * op )

static

Definition at line 593 of file GPUDialect.cpp.

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

◆ eraseRedundantGpuBarrierOps()

LogicalResult eraseRedundantGpuBarrierOps	(	BarrierOp	op,
		PatternRewriter &	rewriter )

static

Remove gpu.barrier after gpu.barrier, the threads are already synchronized!

Definition at line 1538 of file GPUDialect.cpp.

References mlir::RewriterBase::eraseOp(), mlir::Builder::getArrayAttr(), mlir::RewriterBase::modifyOpInPlace(), and success().

◆ getAttributionAttr()

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

static

Definition at line 1852 of file GPUDialect.cpp.

References getAttributionAttrs().

◆ getAttributionAttrs()

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

static

Definition at line 1809 of file GPUDialect.cpp.

Referenced by getAttributionAttr(), and setAttributionAttr().

◆ getSparseHandleKeyword()

std::string getSparseHandleKeyword ( SparseHandleKind kind )

static

Definition at line 313 of file GPUDialect.cpp.

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

◆ parseAllReduceOperation()

ParseResult parseAllReduceOperation	(	AsmParser &	parser,
		AllReduceOperationAttr &	attr )

static

Definition at line 615 of file GPUDialect.cpp.

References mlir::AsmParser::emitError(), mlir::AsmParser::getContext(), mlir::AsmParser::getCurrentLocation(), mlir::AsmParser::parseOptionalKeyword(), and success().

◆ parseAsyncDependencies()

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 443 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]

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

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

◆ parseAttributions() [2/2]

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

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

◆ parseLaunchDimType()

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

static

Definition at line 1444 of file GPUDialect.cpp.

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

◆ parseLaunchFuncOperands()

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

static

Definition at line 1466 of file GPUDialect.cpp.

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

◆ parseOffloadingHandler()

ParseResult parseOffloadingHandler	(	OpAsmParser &	parser,
		Attribute &	offloadingHandler )

static

Definition at line 2061 of file GPUDialect.cpp.

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

◆ parseSizeAssignment()

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

static

Definition at line 959 of file GPUDialect.cpp.

References mlir::AsmParser::emitError(), mlir::AsmParser::getNameLoc(), indices, 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()

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

static

Definition at line 628 of file GPUDialect.cpp.

◆ printAsyncDependencies()

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

◆ printAttributions()

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

static

Definition at line 490 of file GPUDialect.cpp.

References ArrayAttr().

◆ printLaunchDimType()

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

static

Definition at line 1459 of file GPUDialect.cpp.

References mlir::Type::isIndex().

◆ printLaunchFuncOperands()

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

static

Definition at line 1482 of file GPUDialect.cpp.

◆ printOffloadingHandler()

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

static

Definition at line 2074 of file GPUDialect.cpp.

References mlir::Operation::getContext().

◆ printSizeAssignment()

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

static

Definition at line 890 of file GPUDialect.cpp.

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

◆ setAttributionAttr()

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

static

Definition at line 1876 of file GPUDialect.cpp.

References getAttributionAttrs(), and setAttributionAttrs().

◆ setAttributionAttrs()

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

static

Definition at line 1825 of file GPUDialect.cpp.

References ArrayAttr().

Referenced by setAttributionAttr().

◆ verifyAttributions()

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

static

Verifies a GPU function memory attribution.

Definition at line 513 of file GPUDialect.cpp.

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

◆ verifyDistributedType()

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

References mlir::Operation::emitOpError(), and success().

◆ verifyKnownLaunchSizeAttr()

LogicalResult verifyKnownLaunchSizeAttr	(	Operation *	op,
		NamedAttribute	attr )

static

Definition at line 407 of file GPUDialect.cpp.

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

◆ verifyReduceOpAndType()

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

static

Definition at line 539 of file GPUDialect.cpp.

References success().

Classes

Macros

Functions

Macro Definition Documentation

◆ GET_ATTRDEF_CLASSES

◆ GET_ATTRDEF_LIST

◆ GET_OP_CLASSES

◆ GET_OP_LIST

Function Documentation

◆ canMakeGroupOpUniform()

◆ eraseRedundantGpuBarrierOps()

◆ getAttributionAttr()

◆ getAttributionAttrs()

◆ getSparseHandleKeyword()

◆ parseAllReduceOperation()

◆ parseAsyncDependencies()

◆ parseAttributions() [1/2]

◆ parseAttributions() [2/2]

◆ parseLaunchDimType()

◆ parseLaunchFuncOperands()

◆ parseOffloadingHandler()

◆ parseSizeAssignment()

◆ printAllReduceOperation()

◆ printAsyncDependencies()

◆ printAttributions()

◆ printLaunchDimType()

◆ printLaunchFuncOperands()

◆ printOffloadingHandler()

◆ printSizeAssignment()

◆ setAttributionAttr()

◆ setAttributionAttrs()

◆ verifyAttributions()

◆ verifyDistributedType()

◆ verifyKnownLaunchSizeAttr()

◆ verifyReduceOpAndType()