Namespaces
namespace	detail
namespace	impl

Classes
struct	ACCImplicitDataOptions
struct	ACCImplicitRoutineOptions
struct	ACCLoopTilingOptions
struct	ConstructResource
struct	CurrentDeviceIdResource
struct	LegalizeDataValuesInRegionOptions
class	OpenACCSupport
struct	RuntimeCounters

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

Functions
bool	isValidSymbolUse (Operation user, SymbolRefAttr symbol, Operation *definingOpPtr)
remark::detail::InFlightRemark	emitRemark (Operation *op, const Twine &message, llvm::StringRef category)
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 ()
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.
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 llvm::Twine &message, llvm::StringRef category="openacc")
	Emit an OpenACC remark for the given operation with the given message.
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.
std::unique_ptr<::mlir::Pass >	createACCImplicitData ()
std::unique_ptr<::mlir::Pass >	createACCImplicitData (ACCImplicitDataOptions options)
std::unique_ptr<::mlir::Pass >	createACCImplicitDeclare ()
std::unique_ptr<::mlir::Pass >	createACCImplicitRoutine ()
std::unique_ptr<::mlir::Pass >	createACCImplicitRoutine (ACCImplicitRoutineOptions options)
std::unique_ptr<::mlir::Pass >	createACCLegalizeSerial ()
std::unique_ptr<::mlir::Pass >	createACCLoopTiling ()
std::unique_ptr<::mlir::Pass >	createACCLoopTiling (ACCLoopTilingOptions options)
std::unique_ptr<::mlir::Pass >	createLegalizeDataValuesInRegion ()
std::unique_ptr<::mlir::Pass >	createLegalizeDataValuesInRegion (LegalizeDataValuesInRegionOptions options)
void	registerACCImplicitData ()
void	registerACCImplicitDataPass ()
void	registerACCImplicitDeclare ()
void	registerACCImplicitDeclarePass ()
void	registerACCImplicitRoutine ()
void	registerACCImplicitRoutinePass ()
void	registerACCLegalizeSerial ()
void	registerACCLegalizeSerialPass ()
void	registerACCLoopTiling ()
void	registerACCLoopTilingPass ()
void	registerLegalizeDataValuesInRegion ()
void	registerLegalizeDataValuesInRegionPass ()
void	registerOpenACCPasses ()

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 85 of file OpenACC.h.

Function Documentation

◆ createACCImplicitData() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitData ( )

Definition at line 102 of file ACCImplicitData.cpp.

◆ createACCImplicitData() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitData ( ACCImplicitDataOptions options )

Definition at line 106 of file ACCImplicitData.cpp.

◆ createACCImplicitDeclare()

std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitDeclare ( )

We declare an explicit private instantiation because Pass classes should only be visible by the current library.

Definition at line 181 of file ACCImplicitDeclare.cpp.

◆ createACCImplicitRoutine() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitRoutine ( )

Definition at line 278 of file ACCImplicitRoutine.cpp.

◆ createACCImplicitRoutine() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitRoutine ( ACCImplicitRoutineOptions options )

Definition at line 282 of file ACCImplicitRoutine.cpp.

◆ createACCLegalizeSerial()

std::unique_ptr<::mlir::Pass > mlir::acc::createACCLegalizeSerial ( )

We declare an explicit private instantiation because Pass classes should only be visible by the current library.

Definition at line 358 of file ACCLegalizeSerial.cpp.

◆ createACCLoopTiling() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCLoopTiling ( )

Definition at line 450 of file ACCLoopTiling.cpp.

◆ createACCLoopTiling() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCLoopTiling ( ACCLoopTilingOptions options )

Definition at line 454 of file ACCLoopTiling.cpp.

◆ createLegalizeDataValuesInRegion() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createLegalizeDataValuesInRegion ( )

Definition at line 548 of file LegalizeDataValues.cpp.

◆ createLegalizeDataValuesInRegion() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createLegalizeDataValuesInRegion ( LegalizeDataValuesInRegionOptions options )

Definition at line 552 of file LegalizeDataValues.cpp.

◆ emitRemark() [1/2]

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

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 271 of file OpenACCUtils.cpp.

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

◆ emitRemark() [2/2]

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

Referenced by mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::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 4823 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 4835 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 4890 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 4868 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 4882 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 154 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 4853 of file OpenACC.cpp.

◆ getCombinedConstructsAttrName()

constexpr StringLiteral mlir::acc::getCombinedConstructsAttrName ( )

staticconstexpr

Definition at line 207 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 4908 of file OpenACC.cpp.

Referenced by checkDeclareOperands().

◆ getDataOperands()

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

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

Definition at line 4926 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

◆ getDeclareActionAttrName()

constexpr StringLiteral mlir::acc::getDeclareActionAttrName ( )

staticconstexpr

Definition at line 175 of file OpenACC.h.

◆ getDeclareAttrName()

constexpr StringLiteral mlir::acc::getDeclareAttrName ( )

staticconstexpr

Used to obtain the attribute name for declare.

Definition at line 171 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 55 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 212 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 21 of file OpenACCUtils.cpp.

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

◆ getFromDefaultClauseAttrName()

constexpr StringLiteral mlir::acc::getFromDefaultClauseAttrName ( )

staticconstexpr

Definition at line 199 of file OpenACC.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 4918 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 4936 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

◆ getRecipe()

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

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

Definition at line 4945 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 116 of file OpenACCUtils.cpp.

References mlir::Type::print().

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

◆ getRoutineInfoAttrName()

constexpr StringLiteral mlir::acc::getRoutineInfoAttrName ( )

staticconstexpr

Definition at line 179 of file OpenACC.h.

Referenced by isAccRoutine(), and isValidSymbolUse().

◆ getSpecializedRoutineAttrName()

constexpr StringLiteral mlir::acc::getSpecializedRoutineAttrName ( )

staticconstexpr

Definition at line 183 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 76 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 4804 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 89 of file OpenACCUtils.cpp.

References 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 4897 of file OpenACC.cpp.

Referenced by getVariableName().

◆ getVarNameAttrName()

constexpr StringLiteral mlir::acc::getVarNameAttrName ( )

staticconstexpr

Definition at line 203 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 4790 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 4843 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 4812 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 189 of file OpenACC.h.

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

◆ isMappableType()

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

inline

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

Definition at line 166 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 44 of file OpenACCUtils.cpp.

References isOnlyUsedByOpClauses().

◆ 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 49 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 160 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 195 of file OpenACC.h.

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

◆ isValidSymbolUse() [1/2]

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 164 of file OpenACCUtils.cpp.

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

◆ isValidSymbolUse() [2/2]

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

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

◆ registerACCImplicitData()

void mlir::acc::registerACCImplicitData ( )

inline

Definition at line 572 of file Passes.h.

◆ registerACCImplicitDataPass()

void mlir::acc::registerACCImplicitDataPass ( )

inline

Definition at line 579 of file Passes.h.

◆ registerACCImplicitDeclare()

void mlir::acc::registerACCImplicitDeclare ( )

inline

Definition at line 593 of file Passes.h.

◆ registerACCImplicitDeclarePass()

void mlir::acc::registerACCImplicitDeclarePass ( )

inline

Definition at line 600 of file Passes.h.

◆ registerACCImplicitRoutine()

void mlir::acc::registerACCImplicitRoutine ( )

inline

Definition at line 614 of file Passes.h.

◆ registerACCImplicitRoutinePass()

void mlir::acc::registerACCImplicitRoutinePass ( )

inline

Definition at line 621 of file Passes.h.

◆ registerACCLegalizeSerial()

void mlir::acc::registerACCLegalizeSerial ( )

inline

Definition at line 635 of file Passes.h.

◆ registerACCLegalizeSerialPass()

void mlir::acc::registerACCLegalizeSerialPass ( )

inline

Definition at line 642 of file Passes.h.

◆ registerACCLoopTiling()

void mlir::acc::registerACCLoopTiling ( )

inline

Definition at line 656 of file Passes.h.

◆ registerACCLoopTilingPass()

void mlir::acc::registerACCLoopTilingPass ( )

inline

Definition at line 663 of file Passes.h.

◆ registerLegalizeDataValuesInRegion()

void mlir::acc::registerLegalizeDataValuesInRegion ( )

inline

Definition at line 677 of file Passes.h.

◆ registerLegalizeDataValuesInRegionPass()

void mlir::acc::registerLegalizeDataValuesInRegionPass ( )

inline

Definition at line 684 of file Passes.h.

◆ registerOpenACCPasses()

void mlir::acc::registerOpenACCPasses ( )

inline

Definition at line 698 of file Passes.h.

Referenced by mlir::registerAllPasses().

◆ 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().

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

Namespaces

Classes

Enumerations

Functions

Enumeration Type Documentation

◆ OpenACCExecMapping

Function Documentation

◆ createACCImplicitData() [1/2]

◆ createACCImplicitData() [2/2]

◆ createACCImplicitDeclare()

◆ createACCImplicitRoutine() [1/2]

◆ createACCImplicitRoutine() [2/2]

◆ createACCLegalizeSerial()

◆ createACCLoopTiling() [1/2]

◆ createACCLoopTiling() [2/2]

◆ createLegalizeDataValuesInRegion() [1/2]

◆ createLegalizeDataValuesInRegion() [2/2]

◆ emitRemark() [1/2]

◆ emitRemark() [2/2]

◆ getAccPtr()

◆ getAccVar()

◆ getAsyncOnly()

◆ getAsyncOperands()

◆ getAsyncOperandsDeviceType()

◆ getBaseEntity()

◆ getBounds()

◆ getCombinedConstructsAttrName()

◆ getDataClause()

◆ getDataOperands()

◆ getDeclareActionAttrName()

◆ getDeclareAttrName()

◆ getDefaultAttr()

◆ getDominatingDataClauses()

◆ getEnclosingComputeOp()

◆ getFromDefaultClauseAttrName()

◆ getImplicitFlag()

◆ getMutableDataOperands()

◆ getRecipe()

◆ getRecipeName()

◆ getRoutineInfoAttrName()

◆ getSpecializedRoutineAttrName()

◆ getTypeCategory()

◆ getVar()

◆ getVariableName()

◆ getVarName()

◆ getVarNameAttrName()

◆ getVarPtr()

◆ getVarPtrPtr()

◆ getVarType()

◆ isAccRoutine()

◆ isMappableType()

◆ isOnlyUsedByPrivateClauses()

◆ isOnlyUsedByReductionClauses()

◆ isPointerLikeType()

◆ isSpecializedAccRoutine()

◆ isValidSymbolUse() [1/2]

◆ isValidSymbolUse() [2/2]

◆ registerACCImplicitData()

◆ registerACCImplicitDataPass()

◆ registerACCImplicitDeclare()

◆ registerACCImplicitDeclarePass()

◆ registerACCImplicitRoutine()

◆ registerACCImplicitRoutinePass()

◆ registerACCLegalizeSerial()

◆ registerACCLegalizeSerialPass()

◆ registerACCLoopTiling()

◆ registerACCLoopTilingPass()

◆ registerLegalizeDataValuesInRegion()

◆ registerLegalizeDataValuesInRegionPass()

◆ registerOpenACCPasses()

◆ tileACCLoops()

◆ uncollapseLoops()