Namespaces
namespace	detail

Classes
class	ACCDeclareEnterOpConversion
	Pattern to erase acc.declare_enter and its associated acc.declare_exit. More...
class	ACCOpEraseConversion
	Pattern to simply erase an ACC op (for ops with no results). More...
class	ACCOpReplaceWithVarConversion
	Pattern to replace an ACC op with its var operand. More...
class	ACCParMappingPolicy
	Policy class that defines how OpenACC parallelism levels map to target-specific parallel dimension attributes. More...
class	ACCRegionUnwrapConversion
	Pattern to unwrap a region from an ACC op and erase the wrapper. More...
struct	ConstructResource
struct	CurrentDeviceIdResource
class	DefaultACCToGPUMappingPolicy
	Default policy that provides the standard GPU mapping: gang(dim:1) -> BlockX (gridDim.x / blockIdx.x) gang(dim:2) -> BlockY (gridDim.y / blockIdx.y) gang(dim:3) -> BlockZ (gridDim.z / blockIdx.z) worker -> ThreadY (blockDim.y / threadIdx.y) vector -> ThreadX (blockDim.x / threadIdx.x) seq -> Sequential. More...
class	OpenACCSupport
struct	RuntimeCounters

Typedefs
using	ACCToGPUMappingPolicy
	Type alias for the GPU-specific mapping policy.

Enumerations
enum	OpenACCExecMapping { NONE = 0 , VECTOR = 1 , WORKER = 2 , GANG = 4 }
	Enumeration used to encode the execution mapping on a loop construct. More...

Functions
mlir::Value	getVar (mlir::Operation *accDataClauseOp)
	Used to obtain the var from a data clause operation.
mlir::TypedValue< mlir::acc::PointerLikeType >	getVarPtr (mlir::Operation *accDataClauseOp)
	Used to obtain the var from a data clause operation if it implements PointerLikeType.
mlir::Type	getVarType (mlir::Operation *accDataClauseOp)
	Used to obtains the varType from a data clause operation which records the type of variable.
mlir::Value	getAccVar (mlir::Operation *accDataClauseOp)
	Used to obtain the accVar from a data clause operation.
mlir::TypedValue< mlir::acc::PointerLikeType >	getAccPtr (mlir::Operation *accDataClauseOp)
	Used to obtain the accVar from a data clause operation if it implements PointerLikeType.
mlir::Value	getVarPtrPtr (mlir::Operation *accDataClauseOp)
	Used to obtain the varPtrPtr from a data clause operation.
mlir::SmallVector< mlir::Value >	getBounds (mlir::Operation *accDataClauseOp)
	Used to obtain bounds from an acc data clause operation.
mlir::SmallVector< mlir::Value >	getAsyncOperands (mlir::Operation *accDataClauseOp)
	Used to obtain async operands from an acc data clause operation.
mlir::ArrayAttr	getAsyncOperandsDeviceType (mlir::Operation *accDataClauseOp)
	Returns an array of acc:DeviceTypeAttr attributes attached to an acc data clause operation, that correspond to the device types associated with the async clauses with an async-value.
mlir::ArrayAttr	getAsyncOnly (mlir::Operation *accDataClauseOp)
	Returns an array of acc:DeviceTypeAttr attributes attached to an acc data clause operation, that correspond to the device types associated with the async clauses without an async-value.
std::optional< llvm::StringRef >	getVarName (mlir::Operation *accOp)
	Used to obtain the name from an acc operation.
std::optional< mlir::acc::DataClause >	getDataClause (mlir::Operation *accDataEntryOp)
	Used to obtain the dataClause from a data entry operation.
bool	getImplicitFlag (mlir::Operation *accDataEntryOp)
	Used to find out whether data operation is implicit.
mlir::ValueRange	getDataOperands (mlir::Operation *accOp)
	Used to get an immutable range iterating over the data operands.
mlir::MutableOperandRange	getMutableDataOperands (mlir::Operation *accOp)
	Used to get a mutable range iterating over the data operands.
mlir::SymbolRefAttr	getRecipe (mlir::Operation *accOp)
	Used to get the recipe attribute from a data clause operation.
bool	isPointerLikeType (mlir::Type type)
	Used to check whether the provided type implements the PointerLikeType interface.
bool	isMappableType (mlir::Type type)
	Used to check whether the provided type implements the MappableType interface.
static constexpr StringLiteral	getDeclareAttrName ()
	Used to obtain the attribute name for declare.
static constexpr StringLiteral	getDeclareActionAttrName ()
static constexpr StringLiteral	getRoutineInfoAttrName ()
static constexpr StringLiteral	getSpecializedRoutineAttrName ()
bool	isAccRoutine (mlir::Operation *op)
	Used to check whether the current operation is marked with acc routine.
bool	isSpecializedAccRoutine (mlir::Operation *op)
	Used to check whether this is a specialized accelerator version of acc routine function.
static constexpr StringLiteral	getFromDefaultClauseAttrName ()
static constexpr StringLiteral	getVarNameAttrName ()
static constexpr StringLiteral	getCombinedConstructsAttrName ()
ParLevel	getGangParLevel (int64_t gangDimValue)
	Convert a gang dimension value (1, 2, or 3) to the corresponding ParLevel.
mlir::Operation *	getEnclosingComputeOp (mlir::Region &region)
	Used to obtain the enclosing compute construct operation that contains the provided region.
bool	isOnlyUsedByPrivateClauses (mlir::Value val, mlir::Region &region)
	Returns true if this value is only used by acc.private operations in the region.
bool	isOnlyUsedByReductionClauses (mlir::Value val, mlir::Region &region)
	Returns true if this value is only used by acc.reduction operations in the region.
std::optional< ClauseDefaultValue >	getDefaultAttr (mlir::Operation *op)
	Looks for an OpenACC default attribute on the current operation op or in a parent operation which encloses op.
mlir::acc::VariableTypeCategory	getTypeCategory (mlir::Value var)
	Get the type category of an OpenACC variable.
std::string	getVariableName (mlir::Value v)
	Attempts to extract the variable name from a value by walking through view-like operations until an acc.var_name attribute is found.
std::string	getRecipeName (mlir::acc::RecipeKind kind, mlir::Type type)
	Get the recipe name for a given recipe kind and type.
mlir::Value	getBaseEntity (mlir::Value val)
bool	isValidSymbolUse (mlir::Operation user, mlir::SymbolRefAttr symbol, mlir::Operation *definingOpPtr=nullptr)
	Check if a symbol use is valid for use in an OpenACC region.
bool	isDeviceValue (mlir::Value val)
	Check if a value represents device data.
bool	isValidValueUse (mlir::Value val, mlir::Region &region)
	Check if a value use is valid in an OpenACC region.
llvm::SmallVector< mlir::Value >	getDominatingDataClauses (mlir::Operation *computeConstructOp, mlir::DominanceInfo &domInfo, mlir::PostDominanceInfo &postDomInfo)
	Collects all data clauses that dominate the compute construct.
remark::detail::InFlightRemark	emitRemark (mlir::Operation *op, const std::function< std::string()> &messageFn, llvm::StringRef category="openacc")
	Emit an OpenACC remark with lazy message generation.
remark::detail::InFlightRemark	emitRemark (mlir::Operation *op, const llvm::Twine &message, llvm::StringRef category="openacc")
	Emit an OpenACC remark for the given operation with the given message.
std::optional< DataLayout >	getDataLayout (Operation *op, bool allowDefault=true)
	Get the data layout for an operation.
ComputeRegionOp	buildComputeRegion (Location loc, ValueRange launchArgs, ValueRange inputArgs, llvm::StringRef origin, Region &regionToClone, RewriterBase &rewriter, IRMapping &mapping, ValueRange output={}, FlatSymbolRefAttr kernelFuncName={}, FlatSymbolRefAttr kernelModuleName={}, Value stream={}, ValueRange inputArgsToMap={})
	Build an acc.compute_region operation by cloning a source region.
std::optional< gpu::GPUModuleOp >	getOrCreateGPUModule (ModuleOp mod, bool create=true, llvm::StringRef name=kDefaultGPUModuleName)
	Get or create a GPU module in the given module.
std::pair< llvm::SmallVector< Value >, Block::iterator >	cloneACCRegionInto (Region src, Block dest, Block::iterator inlinePoint, IRMapping &mapping, ValueRange resultsToReplace)
	Clone an ACC region into a destination block at the given insertion point.
scf::ExecuteRegionOp	wrapMultiBlockRegionWithSCFExecuteRegion (Region &region, IRMapping &mapping, Location loc, RewriterBase &rewriter, bool convertFuncReturn=false)
	Wrap a multi-block region in an scf.execute_region.
scf::ForOp	convertACCLoopToSCFFor (LoopOp loopOp, RewriterBase &rewriter, bool enableCollapse)
	Convert a structured acc.loop to scf.for.
scf::ParallelOp	convertACCLoopToSCFParallel (LoopOp loopOp, RewriterBase &rewriter)
	Convert acc.loop to scf.parallel.
scf::ExecuteRegionOp	convertUnstructuredACCLoopToSCFExecuteRegion (LoopOp loopOp, RewriterBase &rewriter)
	Convert an unstructured acc.loop to scf.execute_region.
llvm::SmallVector< mlir::acc::LoopOp >	uncollapseLoops (mlir::acc::LoopOp origLoop, unsigned tileCount, unsigned collapseCount, mlir::RewriterBase &rewriter)
	Uncollapse tile loops with multiple IVs and collapseCount < tileCount.
mlir::acc::LoopOp	tileACCLoops (llvm::SmallVector< mlir::acc::LoopOp > &tileLoops, const llvm::SmallVector< mlir::Value > &tileSizes, int32_t defaultTileSize, mlir::RewriterBase &rewriter)
	Tile ACC loops according to the given tile sizes.
void	populateACCSpecializeForDevicePatterns (RewritePatternSet &patterns)
	Populates all patterns for device specialization.
void	populateACCOrphanToHostPatterns (RewritePatternSet &patterns, OpenACCSupport &accSupport, bool enableLoopConversion=true)
	Populates patterns for converting orphan ACC operations to host.
void	populateACCHostFallbackPatterns (RewritePatternSet &patterns, OpenACCSupport &accSupport, bool enableLoopConversion=true)
	Populates all patterns for host fallback path (when if clause evaluates to false).

Variables
constexpr llvm::StringLiteral	kDefaultGPUModuleName = "acc_gpu_module"
	Default GPU module name used by OpenACC.

Typedef Documentation

◆ ACCToGPUMappingPolicy

using mlir::acc::ACCToGPUMappingPolicy

Initial value:

ACCParMappingPolicy<mlir::acc::GPUParallelDimAttr>

mlir::acc::ACCParMappingPolicy

Policy class that defines how OpenACC parallelism levels map to target-specific parallel dimension at...

Definition OpenACCParMapping.h:64

Type alias for the GPU-specific mapping policy.

Definition at line 158 of file OpenACCParMapping.h.

Enumeration Type Documentation

◆ OpenACCExecMapping

enum mlir::acc::OpenACCExecMapping

Enumeration used to encode the execution mapping on a loop construct.

They refer directly to the OpenACC 3.3 standard: 2.9.2. gang 2.9.3. worker 2.9.4. vector

Value can be combined bitwise to reflect the mapping applied to the construct. e.g. acc.loop gang vector, the gang and vector could be combined and the final mapping value would be 5 (4 | 1).

Enumerator
NONE
VECTOR
WORKER
GANG

Definition at line 87 of file OpenACC.h.

Function Documentation

◆ buildComputeRegion()

ComputeRegionOp mlir::acc::buildComputeRegion	(	Location	loc,
		ValueRange	launchArgs,
		ValueRange	inputArgs,
		llvm::StringRef	origin,
		Region &	regionToClone,
		RewriterBase &	rewriter,
		IRMapping &	mapping,
		ValueRange	output = {},
		FlatSymbolRefAttr	kernelFuncName = {},
		FlatSymbolRefAttr	kernelModuleName = {},
		Value	stream = {},
		ValueRange	inputArgsToMap = {} )

Build an acc.compute_region operation by cloning a source region.

Creates a new acc.compute_region with the given launch arguments and origin string, then clones the operations from regionToClone into its body. Multi-block regions are wrapped with scf.execute_region.

The mapping is used and updated during cloning, allowing callers to track value correspondences. Optional output, kernelFuncName, kernelModuleName, and stream arguments are forwarded to the op.

When inputArgsToMap is non-empty, it is used as the key set for the clone mapping (instead of inputArgs). Use this when cloning a region that references one set of values (e.g. the source function's args) while the op's operands are another set (e.g. the current block's args). inputArgsToMap must have the same size as inputArgs when provided.

Definition at line 58 of file OpenACCUtilsCG.cpp.

References mlir::Block::addArgument(), mlir::OpBuilder::clone(), mlir::OpBuilder::createBlock(), mlir::RewriterBase::eraseOp(), mlir::Region::front(), mlir::Block::getArgument(), mlir::Region::getBlocks(), mlir::Builder::getContext(), mlir::Block::getOperations(), mlir::IRMapping::lookup(), mlir::IRMapping::map(), mlir::OpBuilder::setInsertionPointToEnd(), mlir::OpBuilder::setInsertionPointToStart(), and wrapMultiBlockRegionWithSCFExecuteRegion().

◆ cloneACCRegionInto()

std::pair< SmallVector< Value >, Block::iterator > mlir::acc::cloneACCRegionInto	(	Region *	src,
		Block *	dest,
		Block::iterator	inlinePoint,
		IRMapping &	mapping,
		ValueRange	resultsToReplace )

Clone an ACC region into a destination block at the given insertion point.

Requires a single-block source region. Maps block arguments and optional result replacement: values in resultsToReplace are replaced with the operands of the cloned region's acc.yield (1:1). Erases acc.yield/terminator and merges blocks. Returns (replacement values, insertion point after clone).

Definition at line 107 of file OpenACCUtilsLoop.cpp.

References mlir::Region::cloneInto(), mlir::Block::end(), mlir::Block::erase(), mlir::Block::getOperations(), mlir::Block::getParent(), mlir::Block::getTerminator(), mlir::Region::hasOneBlock(), mlir::replaceAllUsesInRegionWith(), replacement(), and mlir::Block::splitBlock().

◆ convertACCLoopToSCFFor()

scf::ForOp mlir::acc::convertACCLoopToSCFFor	(	LoopOp	loopOp,
		RewriterBase &	rewriter,
		bool	enableCollapse )

Convert a structured acc.loop to scf.for.

The loop arguments are converted to index type. If enableCollapse is true, nested loops are collapsed into a single loop.

Parameters

loopOp	The acc.loop operation to convert (must not be unstructured)
rewriter	RewriterBase for creating operations
enableCollapse	Whether to collapse nested loops into one

Returns: The created scf.for operation or nullptr on creation error. An InFlightDiagnostic is emitted on creation error.

Definition at line 186 of file OpenACCUtilsLoop.cpp.

References mlir::coalesceLoops(), mlir::arith::ConstantIndexOp::create(), mlir::OpBuilder::getInsertionPoint(), mlir::IRMapping::map(), mlir::OpBuilder::setInsertionPoint(), mlir::OpBuilder::setInsertionPointToStart(), and mlir::Value::use_empty().

◆ convertACCLoopToSCFParallel()

scf::ParallelOp mlir::acc::convertACCLoopToSCFParallel	(	LoopOp	loopOp,
		RewriterBase &	rewriter )

Convert acc.loop to scf.parallel.

The loop induction variables are converted to index types.

Parameters

loopOp	The acc.loop operation to convert
rewriter	RewriterBase for creating and erasing operations

Returns: The created scf.parallel operation or nullptr on creation error. An InFlightDiagnostic is emitted on creation error.

Definition at line 262 of file OpenACCUtilsLoop.cpp.

References mlir::arith::ConstantIndexOp::create(), mlir::RewriterBase::eraseOp(), mlir::OpBuilder::getInsertionBlock(), mlir::OpBuilder::getInsertionPoint(), mlir::Block::getParentOp(), mlir::OpBuilder::setInsertionPointToStart(), mlir::Value::use_empty(), and wrapMultiBlockRegionWithSCFExecuteRegion().

◆ convertUnstructuredACCLoopToSCFExecuteRegion()

scf::ExecuteRegionOp mlir::acc::convertUnstructuredACCLoopToSCFExecuteRegion	(	LoopOp	loopOp,
		RewriterBase &	rewriter )

Convert an unstructured acc.loop to scf.execute_region.

Parameters

loopOp	The acc.loop operation to convert (must be unstructured)
rewriter	RewriterBase for creating and erasing operations

Returns: The created scf.execute_region operation or nullptr on creation error. An InFlightDiagnostic is emitted on creation error.

Definition at line 326 of file OpenACCUtilsLoop.cpp.

References mlir::OpBuilder::getInsertionBlock(), mlir::Block::getParentOp(), and wrapMultiBlockRegionWithSCFExecuteRegion().

◆ emitRemark() [1/2]

remark::detail::InFlightRemark mlir::acc::emitRemark	(	mlir::Operation *	op,
		const llvm::Twine &	message,
		llvm::StringRef	category = "openacc" )

inline

Emit an OpenACC remark for the given operation with the given message.

Parameters

op	The operation to emit the remark for.
message	The remark message.
category	Optional category for the remark. Defaults to "openacc".

Returns: An in-flight remark object that can be used to append additional information to the remark.

Definition at line 129 of file OpenACCUtils.h.

References emitRemark().

◆ emitRemark() [2/2]

mlir::remark::detail::InFlightRemark mlir::acc::emitRemark	(	mlir::Operation *	op,
		const std::function< std::string()> &	messageFn,
		llvm::StringRef	category = "openacc" )

Emit an OpenACC remark with lazy message generation.

The messageFn is only invoked if remarks are enabled, allowing callers to avoid constructing expensive messages when remarks are disabled.

Parameters

op	The operation to emit the remark for.
messageFn	A callable that returns the remark message.
category	Optional category for the remark. Defaults to "openacc".

Returns: An in-flight remark object that can be used to append additional information to the remark.

Definition at line 334 of file OpenACCUtils.cpp.

References mlir::Attribute::getContext(), mlir::Operation::getLoc(), mlir::Operation::getParentOfType(), mlir::MLIRContext::getRemarkEngine(), and mlir::remark::RemarkOpts::name().

Referenced by mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::emitRemark(), emitRemark(), and mlir::acc::OpenACCSupport::emitRemark().

◆ getAccPtr()

mlir::TypedValue< mlir::acc::PointerLikeType > mlir::acc::getAccPtr ( mlir::Operation * accDataClauseOp )

Used to obtain the accVar from a data clause operation if it implements PointerLikeType.

Definition at line 5069 of file OpenACC.cpp.

◆ getAccVar()

mlir::Value mlir::acc::getAccVar ( mlir::Operation * accDataClauseOp )

Used to obtain the accVar from a data clause operation.

When a data entry operation, it obtains its result accVar value. If a data exit operation, it obtains its operand accVar value. Returns empty value if not a data clause operation.

Definition at line 5081 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

◆ getAsyncOnly()

mlir::ArrayAttr mlir::acc::getAsyncOnly ( mlir::Operation * accDataClauseOp )

Returns an array of acc:DeviceTypeAttr attributes attached to an acc data clause operation, that correspond to the device types associated with the async clauses without an async-value.

Definition at line 5136 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

◆ getAsyncOperands()

mlir::SmallVector< mlir::Value > mlir::acc::getAsyncOperands ( mlir::Operation * accDataClauseOp )

Used to obtain async operands from an acc data clause operation.

Returns an empty vector if there are no such operands.

Definition at line 5114 of file OpenACC.cpp.

◆ getAsyncOperandsDeviceType()

mlir::ArrayAttr mlir::acc::getAsyncOperandsDeviceType ( mlir::Operation * accDataClauseOp )

Returns an array of acc:DeviceTypeAttr attributes attached to an acc data clause operation, that correspond to the device types associated with the async clauses with an async-value.

Definition at line 5128 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

◆ getBaseEntity()

mlir::Value mlir::acc::getBaseEntity ( mlir::Value val )

Definition at line 155 of file OpenACCUtils.cpp.

References mlir::Value::getDefiningOp().

◆ getBounds()

mlir::SmallVector< mlir::Value > mlir::acc::getBounds ( mlir::Operation * accDataClauseOp )

Used to obtain bounds from an acc data clause operation.

Returns an empty vector if there are no bounds.

Definition at line 5099 of file OpenACC.cpp.

◆ getCombinedConstructsAttrName()

constexpr StringLiteral mlir::acc::getCombinedConstructsAttrName ( )

staticconstexpr

Definition at line 209 of file OpenACC.h.

◆ getDataClause()

std::optional< mlir::acc::DataClause > mlir::acc::getDataClause ( mlir::Operation * accDataEntryOp )

Used to obtain the dataClause from a data entry operation.

Returns empty optional if not a data entry operation.

Definition at line 5154 of file OpenACC.cpp.

Referenced by checkDeclareOperands().

◆ getDataLayout()

std::optional< DataLayout > mlir::acc::getDataLayout	(	Operation *	op,
		bool	allowDefault = true )

Get the data layout for an operation.

Attempts to get the data layout from the operation or its parent module. If allowDefault is true (default), a default data layout may be constructed when no explicit data layout spec is found.

Parameters

op	The operation to get the data layout for.
allowDefault	If true, allow returning a default data layout.

Returns: The data layout if available, std::nullopt otherwise.

Definition at line 23 of file OpenACCUtilsCG.cpp.

References mlir::Operation::getParentOfType(), and mlir::Operation::getParentOp().

◆ getDataOperands()

mlir::ValueRange mlir::acc::getDataOperands ( mlir::Operation * accOp )

Used to get an immutable range iterating over the data operands.

Definition at line 5172 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

◆ getDeclareActionAttrName()

constexpr StringLiteral mlir::acc::getDeclareActionAttrName ( )

staticconstexpr

Definition at line 177 of file OpenACC.h.

◆ getDeclareAttrName()

constexpr StringLiteral mlir::acc::getDeclareAttrName ( )

staticconstexpr

Used to obtain the attribute name for declare.

Definition at line 173 of file OpenACC.h.

Referenced by isValidSymbolUse().

◆ getDefaultAttr()

std::optional< mlir::acc::ClauseDefaultValue > mlir::acc::getDefaultAttr ( mlir::Operation * op )

Looks for an OpenACC default attribute on the current operation op or in a parent operation which encloses op.

This is useful because OpenACC specification notes that a visible default clause is the nearest default clause appearing on the compute construct or a lexically containing data construct.

Definition at line 52 of file OpenACCUtils.cpp.

References mlir::Operation::getParentOp().

◆ getDominatingDataClauses()

llvm::SmallVector< mlir::Value > mlir::acc::getDominatingDataClauses	(	mlir::Operation *	computeConstructOp,
		mlir::DominanceInfo &	domInfo,
		mlir::PostDominanceInfo &	postDomInfo )

Collects all data clauses that dominate the compute construct.

This includes data clauses from:

The compute construct itself
Enclosing data constructs
Applicable declare directives (those that dominate and post-dominate) This is used to determine if a variable is already covered by an existing data clause.
Parameters

computeConstructOp The compute construct operation

domInfo Dominance information

postDomInfo Post-dominance information

Returns
Vector of data clause values that dominate the compute construct

Definition at line 275 of file OpenACCUtils.cpp.

References mlir::DominanceInfo::dominates(), mlir::Operation::getParentOfType(), and mlir::Operation::getParentOp().

◆ getEnclosingComputeOp()

mlir::Operation * mlir::acc::getEnclosingComputeOp ( mlir::Region & region )

Used to obtain the enclosing compute construct operation that contains the provided region.

Returns nullptr if no compute construct operation is found. The returned operation is one of types defined by ACC_COMPUTE_CONSTRUCT_OPS.

Definition at line 23 of file OpenACCUtils.cpp.

References ACC_COMPUTE_CONSTRUCT_OPS, and mlir::Region::getParentOfType().

◆ getFromDefaultClauseAttrName()

constexpr StringLiteral mlir::acc::getFromDefaultClauseAttrName ( )

staticconstexpr

Definition at line 201 of file OpenACC.h.

◆ getGangParLevel()

ParLevel mlir::acc::getGangParLevel ( int64_t gangDimValue )

inline

Convert a gang dimension value (1, 2, or 3) to the corresponding ParLevel.

Asserts if the value is not a valid gang dimension.

Definition at line 33 of file OpenACCParMapping.h.

◆ getImplicitFlag()

bool mlir::acc::getImplicitFlag ( mlir::Operation * accDataEntryOp )

Used to find out whether data operation is implicit.

Returns false if not a data operation or if it is a data operation without implicit flag.

Definition at line 5164 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

◆ getMutableDataOperands()

mlir::MutableOperandRange mlir::acc::getMutableDataOperands ( mlir::Operation * accOp )

Used to get a mutable range iterating over the data operands.

Definition at line 5182 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

◆ getOrCreateGPUModule()

std::optional< gpu::GPUModuleOp > mlir::acc::getOrCreateGPUModule	(	ModuleOp	mod,
		bool	create = true,
		llvm::StringRef	name = kDefaultGPUModuleName )

Get or create a GPU module in the given module.

If a GPU module with the specified name already exists, it is returned. If create is true and no GPU module exists, one is created. If create is false and no GPU module exists, std::nullopt is returned.

Parameters

mod	The module to search or create the GPU module in.
create	If true (default), create the GPU module if it doesn't exist.
name	The name for the GPU module. If empty, uses kDefaultGPUModuleName.

Returns: The GPU module if found or created, std::nullopt otherwise.

Definition at line 20 of file OpenACCUtilsGPU.cpp.

References mlir::SymbolTable::insert(), kDefaultGPUModuleName, and mlir::SymbolTable::lookup().

Referenced by mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::getOrCreateGPUModule(), and mlir::acc::OpenACCSupport::getOrCreateGPUModule().

◆ getRecipe()

mlir::SymbolRefAttr mlir::acc::getRecipe ( mlir::Operation * accOp )

Used to get the recipe attribute from a data clause operation.

Definition at line 5191 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

◆ getRecipeName()

std::string mlir::acc::getRecipeName	(	mlir::acc::RecipeKind	kind,
		mlir::Type	type )

Get the recipe name for a given recipe kind and type.

Returns an empty string if not possible to generate a recipe name.

Definition at line 117 of file OpenACCUtils.cpp.

References mlir::Type::print().

Referenced by mlir::acc::OpenACCSupport::getRecipeName().

◆ getRoutineInfoAttrName()

constexpr StringLiteral mlir::acc::getRoutineInfoAttrName ( )

staticconstexpr

Definition at line 181 of file OpenACC.h.

Referenced by isAccRoutine(), and isValidSymbolUse().

◆ getSpecializedRoutineAttrName()

constexpr StringLiteral mlir::acc::getSpecializedRoutineAttrName ( )

staticconstexpr

Definition at line 185 of file OpenACC.h.

Referenced by isSpecializedAccRoutine().

◆ getTypeCategory()

mlir::acc::VariableTypeCategory mlir::acc::getTypeCategory ( mlir::Value var )

Get the type category of an OpenACC variable.

Definition at line 73 of file OpenACCUtils.cpp.

References mlir::Value::getType().

◆ getVar()

mlir::Value mlir::acc::getVar ( mlir::Operation * accDataClauseOp )

Used to obtain the var from a data clause operation.

Returns empty value if not a data clause operation or is a data exit operation with no var.

Definition at line 5050 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

Referenced by checkDeclareOperands().

◆ getVariableName()

std::string mlir::acc::getVariableName ( mlir::Value v )

Attempts to extract the variable name from a value by walking through view-like operations until an acc.var_name attribute is found.

Returns empty string if no name is found.

Definition at line 86 of file OpenACCUtils.cpp.

References getConstantIntValue(), mlir::Value::getDefiningOp(), getVarName(), and getVarNameAttrName().

Referenced by mlir::acc::OpenACCSupport::getVariableName().

◆ getVarName()

std::optional< llvm::StringRef > mlir::acc::getVarName ( mlir::Operation * accOp )

Used to obtain the name from an acc operation.

Definition at line 5143 of file OpenACC.cpp.

Referenced by getVariableName().

◆ getVarNameAttrName()

constexpr StringLiteral mlir::acc::getVarNameAttrName ( )

staticconstexpr

Definition at line 205 of file OpenACC.h.

Referenced by getVariableName().

◆ getVarPtr()

mlir::TypedValue< mlir::acc::PointerLikeType > mlir::acc::getVarPtr ( mlir::Operation * accDataClauseOp )

Used to obtain the var from a data clause operation if it implements PointerLikeType.

Definition at line 5036 of file OpenACC.cpp.

◆ getVarPtrPtr()

mlir::Value mlir::acc::getVarPtrPtr ( mlir::Operation * accDataClauseOp )

Used to obtain the varPtrPtr from a data clause operation.

Returns empty value if not a data clause operation.

Definition at line 5089 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

◆ getVarType()

mlir::Type mlir::acc::getVarType ( mlir::Operation * accDataClauseOp )

Used to obtains the varType from a data clause operation which records the type of variable.

When var is PointerLikeType, this returns the type of the pointer target.

Definition at line 5058 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

◆ isAccRoutine()

bool mlir::acc::isAccRoutine ( mlir::Operation * op )

inline

Used to check whether the current operation is marked with acc routine.

The operation passed in should be a function.

Definition at line 191 of file OpenACC.h.

References getRoutineInfoAttrName(), and mlir::Operation::hasAttr().

◆ isDeviceValue()

bool mlir::acc::isDeviceValue ( mlir::Value val )

Check if a value represents device data.

This checks if the value represents device data via the MappableType, PointerLikeType, and GlobalVariableOpInterface interfaces.

Parameters

val	The value to check

Returns: true if the value is device data, false otherwise

Definition at line 219 of file OpenACCUtils.cpp.

References mlir::Value::getDefiningOp(), mlir::Value::getType(), isDeviceValue(), and mlir::SymbolTable::lookupNearestSymbolFrom().

Referenced by isDeviceValue(), and isValidValueUse().

◆ isMappableType()

bool mlir::acc::isMappableType ( mlir::Type type )

inline

Used to check whether the provided type implements the MappableType interface.

Definition at line 168 of file OpenACC.h.

Referenced by checkRecipe().

◆ isOnlyUsedByPrivateClauses()

bool mlir::acc::isOnlyUsedByPrivateClauses	(	mlir::Value	val,
		mlir::Region &	region )

Returns true if this value is only used by acc.private operations in the region.

Definition at line 41 of file OpenACCUtils.cpp.

References isOnlyUsedByOpClauses().

Referenced by isValidValueUse().

◆ isOnlyUsedByReductionClauses()

bool mlir::acc::isOnlyUsedByReductionClauses	(	mlir::Value	val,
		mlir::Region &	region )

Returns true if this value is only used by acc.reduction operations in the region.

Definition at line 46 of file OpenACCUtils.cpp.

References isOnlyUsedByOpClauses().

◆ isPointerLikeType()

bool mlir::acc::isPointerLikeType ( mlir::Type type )

inline

Used to check whether the provided type implements the PointerLikeType interface.

Definition at line 162 of file OpenACC.h.

◆ isSpecializedAccRoutine()

bool mlir::acc::isSpecializedAccRoutine ( mlir::Operation * op )

inline

Used to check whether this is a specialized accelerator version of acc routine function.

Definition at line 197 of file OpenACC.h.

References getSpecializedRoutineAttrName(), and mlir::Operation::hasAttr().

◆ isValidSymbolUse()

bool mlir::acc::isValidSymbolUse	(	mlir::Operation *	user,
		mlir::SymbolRefAttr	symbol,
		mlir::Operation **	definingOpPtr = nullptr )

Check if a symbol use is valid for use in an OpenACC region.

This includes looking for various attributes such as acc.routine_info and acc.declare attributes.

Parameters

user	The operation using the symbol
symbol	The symbol reference being used
definingOpPtr	Optional output parameter to receive the defining op

Returns: true if the symbol use is valid, false otherwise

Definition at line 165 of file OpenACCUtils.cpp.

References getDeclareAttrName(), getRoutineInfoAttrName(), mlir::Operation::hasAttr(), mlir::SymbolTable::lookupNearestSymbolFrom(), and mlir::SymbolTable::Private.

Referenced by mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::isValidSymbolUse(), and mlir::acc::OpenACCSupport::isValidSymbolUse().

◆ isValidValueUse()

bool mlir::acc::isValidValueUse	(	mlir::Value	val,
		mlir::Region &	region )

Check if a value use is valid in an OpenACC region.

This is true if:

The value is produced by an ACC data entry operation
The value is device data
The value is only used by private clauses in the region
Parameters

val The value to check

region The OpenACC region

Returns
true if the value use is valid, false otherwise

Definition at line 252 of file OpenACCUtils.cpp.

References mlir::Value::getDefiningOp(), mlir::Value::getType(), isDeviceValue(), mlir::Type::isIntOrIndexOrFloat(), and isOnlyUsedByPrivateClauses().

Referenced by mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::isValidValueUse(), and mlir::acc::OpenACCSupport::isValidValueUse().

◆ populateACCHostFallbackPatterns()

void mlir::acc::populateACCHostFallbackPatterns	(	RewritePatternSet &	patterns,
		OpenACCSupport &	accSupport,
		bool	enableLoopConversion = true )

Populates all patterns for host fallback path (when if clause evaluates to false).

In this mode, ALL ACC operations should be converted or removed.

Parameters

enableLoopConversion Whether to convert orphan acc.loop operations.

Definition at line 397 of file ACCSpecializeForHost.cpp.

References mlir::RewritePatternSet::getContext(), and mlir::RewritePatternSet::insert().

◆ populateACCOrphanToHostPatterns()

void mlir::acc::populateACCOrphanToHostPatterns	(	RewritePatternSet &	patterns,
		OpenACCSupport &	accSupport,
		bool	enableLoopConversion = true )

Populates patterns for converting orphan ACC operations to host.

All patterns check that the operation is NOT inside or associated with a compute region before converting.

Parameters

enableLoopConversion Whether to convert orphan acc.loop operations.

Definition at line 360 of file ACCSpecializeForHost.cpp.

References mlir::RewritePatternSet::getContext(), and mlir::RewritePatternSet::insert().

◆ populateACCSpecializeForDevicePatterns()

void mlir::acc::populateACCSpecializeForDevicePatterns ( RewritePatternSet & patterns )

Populates all patterns for device specialization.

In specialized device code (such as specialized acc routine), many ACC operations do not make sense because they are host-side constructs. This function adds patterns to remove or transform them.

Definition at line 124 of file ACCSpecializeForDevice.cpp.

References mlir::RewritePatternSet::getContext(), and mlir::RewritePatternSet::insert().

◆ tileACCLoops()

mlir::acc::LoopOp mlir::acc::tileACCLoops	(	llvm::SmallVector< mlir::acc::LoopOp > &	tileLoops,
		const llvm::SmallVector< mlir::Value > &	tileSizes,
		int32_t	defaultTileSize,
		mlir::RewriterBase &	rewriter )

Tile ACC loops according to the given tile sizes.

Tiling a 2-level nested loop will create two 'tile' loops containing two 'element' loops. The transformation looks like:

Before Tiling:

#pragma acc loop tile(tile_size1, tile_size2)
 for (i = lb1; i < ub1; i += step1) { // original loop
   for (j = lb2; j < ub2; j += step2) {
     a[i,j] = i + j;
   }
 }

After Tiling:

for (i = lb1; i < ub1; i += (step1 * tile_size1)) { // tile loop 1
  for (j = lb2; j < ub2; j += (step2 * tile_size2)) { // tile loop 2
    for (ii = i; ii < min(ub1, (step1 * tile_size1) + i); ii += step1) {
    // element loop 1
      for (jj = j; jj < min(ub2, (step2 * tile_size2) + j); jj += step2)
      { // element loop 2
        a[ii,jj] = i + j;
      }
    }
  }
}

Unknown tile sizes (represented as -1 in acc dialect for tile(*)) are resolved to the provided default tile size.

Parameters

tileLoops	The loops to tile (outermost first).
tileSizes	The tile sizes for each dimension. Values of -1 are treated as unknown and resolved to defaultTileSize.
defaultTileSize	The default tile size to use for unknown (*) tiles.
rewriter	The rewriter to use for modifications.

Returns: The outermost loop after tiling.

Definition at line 136 of file OpenACCUtilsTiling.cpp.

References createInnerLoop(), mlir::RewriterBase::finalizeOpModification(), mlir::Builder::getDenseBoolArrayAttr(), mlir::Builder::getIntegerAttr(), mlir::getType(), moveOpsAndReplaceIVs(), removeWorkerVectorFromLoop(), resolveAndCastTileSize(), mlir::OpBuilder::setInsertionPoint(), and mlir::RewriterBase::startOpModification().

◆ uncollapseLoops()

llvm::SmallVector< mlir::acc::LoopOp > mlir::acc::uncollapseLoops	(	mlir::acc::LoopOp	origLoop,
		unsigned	tileCount,
		unsigned	collapseCount,
		mlir::RewriterBase &	rewriter )

Uncollapse tile loops with multiple IVs and collapseCount < tileCount.

This is used to prepare loops for tiling when the collapse count is less than the tile count.

Parameters

origLoop	The original loop operation to uncollapse.
tileCount	The number of tile dimensions.
collapseCount	The collapse count from the original loop.
rewriter	The rewriter to use for modifications.

Returns: A vector of uncollapsed loop operations.

Definition at line 253 of file OpenACCUtilsTiling.cpp.

References createACCLoopFromOriginal(), mlir::OpBuilder::createBlock(), createInnerLoop(), mlir::Builder::getDenseBoolArrayAttr(), moveOpsAndReplaceIVs(), mlir::OpBuilder::setInsertionPoint(), and mlir::OpBuilder::setInsertionPointToEnd().

◆ wrapMultiBlockRegionWithSCFExecuteRegion()

scf::ExecuteRegionOp mlir::acc::wrapMultiBlockRegionWithSCFExecuteRegion	(	Region &	region,
		IRMapping &	mapping,
		Location	loc,
		RewriterBase &	rewriter,
		bool	convertFuncReturn = false )

Wrap a multi-block region in an scf.execute_region.

Wrap a multi-block region with scf.execute_region.

Clones the given region into a new scf.execute_region. Replaces acc.yield with scf.yield; when convertFuncReturn is true, also replaces func.return with scf.yield. Use this to convert unstructured control flow (e.g. multiple blocks with branches) into a single SCF region.

Parameters

region	The region to wrap (cloned into the execute_region; not modified).
mapping	IR mapping for the clone; updated with block and value mappings.
loc	Location for the created execute_region op.
rewriter	RewriterBase for creating and erasing operations.
convertFuncReturn	When true, replace func.return with scf.yield in addition to acc.yield. Default is false.

Returns: The created scf.execute_region operation, or nullptr if any replaced terminator has operands (results not yet supported).

Definition at line 151 of file OpenACCUtilsLoop.cpp.

References mlir::OpBuilder::cloneRegionBefore(), mlir::Region::end(), mlir::RewriterBase::eraseOp(), mlir::Region::getBlocks(), mlir::Operation::getLoc(), mlir::Operation::getOperands(), mlir::IRMapping::lookup(), mlir::OpBuilder::setInsertionPoint(), and TypeRange.

Referenced by buildComputeRegion(), convertACCLoopToSCFParallel(), and convertUnstructuredACCLoopToSCFExecuteRegion().

Variable Documentation

◆ kDefaultGPUModuleName

llvm::StringLiteral mlir::acc::kDefaultGPUModuleName = "acc_gpu_module"

constexpr

Default GPU module name used by OpenACC.

Definition at line 25 of file OpenACCUtilsGPU.h.

Referenced by getOrCreateGPUModule().

computeConstructOp	The compute construct operation
domInfo	Dominance information
postDomInfo	Post-dominance information

Namespaces

Classes

Typedefs

Enumerations

Functions

Variables

Typedef Documentation

◆ ACCToGPUMappingPolicy

Enumeration Type Documentation

◆ OpenACCExecMapping

Function Documentation

◆ buildComputeRegion()

◆ cloneACCRegionInto()

◆ convertACCLoopToSCFFor()

◆ convertACCLoopToSCFParallel()

◆ convertUnstructuredACCLoopToSCFExecuteRegion()

◆ emitRemark() [1/2]

◆ emitRemark() [2/2]

◆ getAccPtr()

◆ getAccVar()

◆ getAsyncOnly()

◆ getAsyncOperands()

◆ getAsyncOperandsDeviceType()

◆ getBaseEntity()

◆ getBounds()

◆ getCombinedConstructsAttrName()

◆ getDataClause()

◆ getDataLayout()

◆ getDataOperands()

◆ getDeclareActionAttrName()

◆ getDeclareAttrName()

◆ getDefaultAttr()

◆ getDominatingDataClauses()

◆ getEnclosingComputeOp()

◆ getFromDefaultClauseAttrName()

◆ getGangParLevel()

◆ getImplicitFlag()

◆ getMutableDataOperands()

◆ getOrCreateGPUModule()

◆ getRecipe()

◆ getRecipeName()

◆ getRoutineInfoAttrName()

◆ getSpecializedRoutineAttrName()

◆ getTypeCategory()

◆ getVar()

◆ getVariableName()

◆ getVarName()

◆ getVarNameAttrName()

◆ getVarPtr()

◆ getVarPtrPtr()

◆ getVarType()

◆ isAccRoutine()

◆ isDeviceValue()

◆ isMappableType()

◆ isOnlyUsedByPrivateClauses()

◆ isOnlyUsedByReductionClauses()

◆ isPointerLikeType()

◆ isSpecializedAccRoutine()

◆ isValidSymbolUse()

◆ isValidValueUse()

◆ populateACCHostFallbackPatterns()

◆ populateACCOrphanToHostPatterns()

◆ populateACCSpecializeForDevicePatterns()

◆ tileACCLoops()

◆ uncollapseLoops()

◆ wrapMultiBlockRegionWithSCFExecuteRegion()

Variable Documentation

◆ kDefaultGPUModuleName