mlir::acc Namespace Reference

Namespaces
namespace	AttributeTrait
namespace	detail
namespace	impl

Classes
struct	ACCBindRoutineOptions
struct	ACCCGToGPUOptions
struct	ACCComputeLoweringOptions
struct	ACCDeclareCtorDtorConversionOptions
class	ACCDeclareEnterOpConversion
	Pattern to erase acc.declare_enter and its associated acc.declare_exit. More...
struct	ACCDeclareGPUModuleInsertionOptions
struct	ACCEmitRemarksLoopOptions
struct	ACCImplicitDataOptions
struct	ACCImplicitRoutineOptions
struct	ACCLoopTilingOptions
struct	ACCMaterializeRoutineBindTargetsOptions
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	ACCRoutineLoweringOptions
struct	ACCRoutineToGPUFuncOptions
class	ACCRuntimeCallConfig
	Configuration for OpenACC to LLVM runtime lowering. More...
struct	ACCSpecializeForDeviceOptions
struct	ACCSpecializeForHostOptions
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...
struct	LegalizeDataValuesInRegionOptions
struct	OffloadTargetVerifierOptions
class	OpenACCSupport
struct	RuntimeCounters
class	SharedMemoryBudget
	Tracks aligned byte consumption against a configurable shared memory cap. More...
class	VariableInfoAttr
	Base attribute class for language-specific variable information carried through the OpenACC type interface helpers. More...
struct	VariableNameConfig
	How the name of a variable is to be rendered. More...

Typedefs
using	ACCToGPUMappingPolicy
	Type alias for the GPU-specific mapping policy.
using	TypeSizeAndAlignment = std::pair<llvm::TypeSize, llvm::TypeSize>
using	TypesForDevice = llvm::SmallSetVector<int64_t, 3>
	Holds information for which integers represent a device type in the runtime.

Enumerations
enum	OpenACCExecMapping { NONE = 0 , VECTOR = 1 , WORKER = 2 , GANG = 4 }
	Enumeration used to encode the execution mapping on a loop construct. More...
enum class	RuntimeFunction
	IDs for OpenACC compiler-to-runtime entry points (__tgt_acc_*). More...
enum class	RuntimeResult : int32_t { Success = 0 , Error = 1 }
	Status returned by the runtime entry points that report whether they completed. More...
enum class	DataDescriptor
	IDs for the argument descriptors of the OpenACC data entry points, declared in OpenACCRuntimeDescriptors.def. More...
enum class	KernelArgumentField : int64_t
	Fields of the argument struct passed to __tgt_acc_kernel, in declaration order. More...

Functions
Location	unfuseLoc (Location loc)
	Unfuses fused locations, returning the last sub-location.
std::optional< FileLineColLoc >	getFileLineColLoc (Location loc, bool errorOnInvalidLocation)
	Returns file:line:column location information when available.
StringRef	getParentFunctionName (Operation *op)
	Returns the symbol name of the function op belongs to, or of op itself when it is a function.
StringRef	getParentFunctionName (Value value)
	Returns the enclosing function symbol name for value's defining op.
StringRef	getParentFunctionName (ValueRange values)
	Returns the first non-empty enclosing function name from values.
std::string	getInternalGlobalName (StringRef kind, StringRef detail)
	Returns the name to give a global that the conversion creates to hold detail of kind, such as the name of a variable or a source position.
Value	getOrCreateGlobalString (Location loc, OpBuilder &builder, StringRef name, StringRef value, Region &globalSymbolRegion, SymbolTable *symbolTable=nullptr)
	Creates or reuses a null-terminated string global in globalSymbolRegion.
Value	createIdent (Location loc, StringRef functionName, OpBuilder &builder, Region &globalSymbolRegion, const ACCRuntimeCallConfig &config, SymbolTable *symbolTable=nullptr)
	Returns a pointer to a constant global holding an ident_t for OpenACC runtime calls.
Value	castToI64 (Location loc, Value value, OpBuilder &builder)
	Sign-extends or truncates value to the i64 the runtime entry points take for values like queue numbers.
Value	constantI64 (Location loc, int64_t value, OpBuilder &builder)
	Creates an i64 constant.
Value	remapOrSelf (Value value, ConversionPatternRewriter &rewriter)
	Returns the value value was converted to, or value itself when the conversion left it as it was.
Value	getAsyncQueue (Location loc, Value asyncOperand, bool asyncOnly, OpBuilder &builder, const ACCRuntimeCallConfig &config)
	Returns the queue an async clause selects: the value of the clause when it has one, the queue standing for an async clause without a value when asyncOnly is set, and the synchronous queue when there is no clause at all.
LogicalResult	emitWaitCall (Location loc, ValueRange waitOperands, Value asyncQueue, OpBuilder &builder, Region &globalSymbolRegion, SymbolTable &symbolTable, const ACCRuntimeCallConfig &config, Value deviceNum={})
	Emits the runtime call that waits for waitOperands on asyncQueue, which is what a wait clause or an acc.wait directive asks for.
FailureOr< Value >	emitGetDevicePtrCall (Operation *clauseOp, Value hostPtr, bool ifPresent, OpBuilder &builder, Region &globalSymbolRegion, SymbolTable &symbolTable, const ACCRuntimeCallConfig &config)
	Emits the runtime call that asks for the device address the object of the data clause clauseOp is mapped to, which is what a use_device clause exposes in the body of its construct.
LogicalResult	emitGuardedByIfCond (Location loc, Value ifCond, RewriterBase &rewriter, function_ref< LogicalResult()> emitFn)
	Runs emitFn guarded by a branch on ifCond, or unguarded when there is no condition.
FailureOr< Value >	emitValueSelectedByIfCond (Location loc, Value ifCond, RewriterBase &rewriter, function_ref< FailureOr< Value >()> thenFn, function_ref< FailureOr< Value >()> elseFn)
	Emits thenFn on the path a branch on ifCond takes and elseFn on the other one, and returns the value that reaches the code following the branch.
void	spliceConstructRegion (Operation *op, Region &region, RewriterBase &rewriter)
	Splices region, the body of a structured construct, into the block holding op, so that the construct itself can be erased.
template<typename OpTy>
Value	getAsyncClauseValue (OpTy op, DeviceType deviceType, bool &asyncOnly)
	Returns the value the async clause of op names for deviceType, and sets asyncOnly when the clause names no queue.
template<typename OpTy>
std::optional< DeviceType >	getWaitClauseValues (OpTy op, DeviceType deviceType, SmallVectorImpl< Value > &waitValues)
	Appends to waitValues the queues the wait clause of op names for deviceType, and returns the device type the clause naming them is given for - a clause naming no queue waits for every one of them.
template<typename OpTy>
bool	hasWaitDevnum (OpTy op, DeviceType deviceType)
	Returns whether the wait clause op gives for deviceType carries a devnum modifier, which selects the device the queues belong to.
template<typename OpTy>
Value	getAsyncQueue (OpTy op, DeviceType deviceType, ConversionPatternRewriter &rewriter, const ACCRuntimeCallConfig &config)
	Returns the queue that the runtime calls of op run on, from the async clause it gives for deviceType.
template<typename OpTy>
LogicalResult	emitWaitClause (OpTy op, DeviceType deviceType, Value asyncQueue, ConversionPatternRewriter &rewriter, OpenACCSupport &accSupport, Region &globalSymbolRegion, SymbolTable &symbolTable, const ACCRuntimeCallConfig &config)
	Emits the wait that a wait clause on op asks for before the runtime calls of the construct, waiting on asyncQueue for the queues the clause names, or for every queue when it names none.
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.
SmallVector< DeviceType, 3 >	getDeviceTypesByPrecedence (DeviceType deviceType)
	The clauses of a construct can be given once per device type.
std::optional< llvm::StringRef >	getVarName (mlir::Operation *accOp)
	Used to obtain the name from an acc operation.
bool	isDataEntryOp (mlir::Operation *op)
	Used to find out whether op is a data entry operation or an acc.map_info.
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.
bool	getSyntheticFlag (mlir::Operation *accDataClauseOp)
	Used to find out whether the implementation created the data operation for its own bookkeeping, instead of to represent a data attribute of the program being compiled.
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	getNoDeviceDescriptorAttrName ()
	Name for a unit attribute on a data entry operation stating that nothing in the region reads the descriptor its variable is reached through, so the mapping supplies the bounds but needs no device copy of the descriptor.
static constexpr StringLiteral	getCollapseCountAttrName ()
	Name for an attribute attached to a loop indicating the number of loops collapsed to create that loop.
static constexpr StringLiteral	getUnstructuredIndependentLoopAttrName ()
	Name for a unit attribute attached to the region an unstructured independent loop was converted to, recording that the loop could not be represented as a structured loop and is therefore not parallelized.
static constexpr StringLiteral	getVarNameAttrName ()
static constexpr StringLiteral	getCombinedConstructsAttrName ()
template<typename ComputeOpT>
bool	isGangWorkerVectorAllOne (ComputeOpT op)
ParLevel	getGangParLevel (int64_t gangDimValue)
	Convert a gang dimension value (1, 2, or 3) to the corresponding ParLevel.
StringRef	getRuntimeFunctionName (RuntimeFunction fn)
	Returns the default runtime symbol name for fn.
LLVM::LLVMFunctionType	getRuntimeFunctionType (MLIRContext *ctx, RuntimeFunction fn)
	Builds the LLVM function type for fn in ctx.
template<typename FieldEnum>
int64_t	getDataDescriptorFieldIndex (FieldEnum field)
	The index an insert or extract of field addresses.
int64_t	getKernelArgumentFieldIndex (KernelArgumentField field)
	The index an insert of field addresses.
LLVM::LLVMStructType	getKernelArgumentType (MLIRContext *ctx)
	The LLVM type of the argument struct passed to __tgt_acc_kernel.
StringRef	getDataDescriptorName (DataDescriptor desc)
	Returns the name of the runtime type desc materializes.
DataDescKind	getDataDescriptorKind (DataDescriptor desc)
	Returns the descriptor kind the runtime reads from the version field of desc.
LLVM::LLVMStructType	getDataDescriptorType (MLIRContext *ctx, DataDescriptor desc, Type baseType={})
	Builds the LLVM type of desc in ctx.
void	populateDialectIdentityDeviceTypeMapping (ACCRuntimeCallConfig &config)
	Install a device-type mapping that uses OpenACC dialect enum ordinals as the runtime encoding.
void	populateDialectIdentityMapFlagsMapping (ACCRuntimeCallConfig &config)
	Install a map-flag mapping that uses the OpenACC dialect bit positions as the runtime encoding, with the same caveat as populateDialectIdentityDeviceTypeMapping.
FailureOr< LLVM::CallOp >	createRuntimeCall (Location loc, OpBuilder &builder, Region &globalSymbolRegion, SymbolTable &symbolTable, RuntimeFunction fn, const ACCRuntimeCallConfig &config, ArrayRef< Value > arguments)
	Declares (if needed) and returns a call to the runtime function identified by fn using the name from config.
mlir::Operation *	getEnclosingComputeOp (mlir::Region &region)
	Used to obtain the enclosing compute construct operation that contains the provided region.
mlir::Value	getACCOperandForBlockArg (mlir::Value v)
	If v is not a block argument of an acc.compute_region body, returns nullptr.
mlir::Operation *	getACCDataClauseOpForBlockArg (mlir::Value v)
	If v is not a block argument of an acc.compute_region body, returns nullptr.
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, the name of a data clause operation, or the symbol a global is addressed through is found.
llvm::StringLiteral	getVarNamePlaceholder ()
	Returns a placeholder string for use as an acc.var_name attribute value when the actual variable name is not yet known at the point of IR construction.
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	isDeviceAccessibleValue (mlir::Value val)
	Check if a value represents device data.
bool	isInDeviceMemoryValue (mlir::Value val)
	Check if a value is backed by memory that is residing in the current device, and therefore requires no runtime mapping or attach.
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.
bool	isInOffTargetBranch (mlir::Operation *op, llvm::ArrayRef< int64_t > deviceTypes)
	Returns true if op is in a branch that does not execute on the target.
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.
void	insertParDim (llvm::SmallVector< GPUParallelDimAttr > &parDims, GPUParallelDimAttr parDim)
	Insert parDim into parDims while preserving dimension ordering.
void	removeParDim (llvm::SmallVector< GPUParallelDimAttr > &parDims, GPUParallelDimAttr parDim)
	Remove parDim from parDims if present.
GPUParallelDimsAttr	getParDimsAttr (Operation *op)
	Obtain the parallel dimensions carried by op, if any.
bool	hasParDimsAttr (Operation *op)
	Return whether op carries parallel dimensions.
bool	hasSeqParDims (Operation *op)
	Return whether op carries sequential parallel dimensions.
void	setParDimsAttr (Operation *op, GPUParallelDimsAttr attr)
	Set parallel dimensions on op.
void	updateParDimsAttr (Operation *op, GPUParallelDimsAttr attr)
	Update parallel dimensions on op.
void	copyParDimsAttr (Operation *from, Operation *to)
	Copy parallel dimensions from from to to.
ActiveParDimsAttr	getActiveParDimsAttr (Operation *op)
	Obtain the active parallel dimensions carried by op, if any.
bool	hasActiveParDimsAttr (Operation *op)
	Return whether op carries active parallel dimensions.
void	setActiveParDimsAttr (Operation *op, ActiveParDimsAttr attr)
	Set active parallel dimensions on op.
void	setActiveParDimsAttr (Operation *op, ArrayRef< GPUParallelDimAttr > dims)
	Set active parallel dimensions on op from a dimension list.
bool	hasGPUBlockRedundantAttr (Operation *op)
	Return whether op is marked with the acc.gpu_block_redundant attribute, i.e.
void	setGPUBlockRedundantAttr (Operation *op)
	Mark op with the acc.gpu_block_redundant attribute.
ChunkSizeAttr	getChunkSizeAttr (Operation *op)
	Obtain the acc.chunk_size attribute carried by op, if any.
bool	hasChunkSizeAttr (Operation *op)
	Return whether op carries an acc.chunk_size attribute.
void	setChunkSizeAttr (Operation *op, ChunkSizeAttr attr)
	Set the acc.chunk_size attribute on op.
void	setChunkSizeAttr (Operation *op, int64_t chunkSize)
	Set the acc.chunk_size attribute on op from a constant size.
std::optional< int64_t >	getChunkSize (Operation *op)
	Return the chunk size carried by op, if any.
GPUParallelDimsAttr	getGangDim1ParDimsAttr (MLIRContext *ctx, ACCToGPUMappingPolicy &policy)
	Create a gang dim 1 GPUParallelDimsAttr based on the mapping policy.
GPUParallelDimsAttr	getSeqParDimsAttr (MLIRContext *ctx, ACCToGPUMappingPolicy &policy)
	Create a sequential GPUParallelDimsAttr based on the mapping policy.
int64_t	sumExistingSharedMemoryBytes (Region &region)
	Sum aligned static_upper_bound_bytes for all acc.gpu_shared_memory in region.
PrivatizeOp	getPrivatizeOp (PrivateLocalOp privateLocal, ComputeRegionOp computeRegion)
	Resolve the acc.privatize operation associated with a private local.
MemRefType	getPrivateBaseMemRefType (Type baseTy, ModuleOp module)
	Returns the ranked MemRef type used to allocate privatized storage.
SmallVector< GPUParallelDimAttr >	collectPrivateLocalParDims (PrivateLocalOp privateLocal, ComputeRegionOp computeRegion)
	Collect parallel dimensions that govern privatization of privateLocal.
FailureOr< bool >	isPrivateLocalSharedMemoryCandidate (PrivateLocalOp privateLocal, ComputeRegionOp computeRegion, ModuleOp module, const ACCToGPUMappingPolicy &policy, OpenACCSupport *support=nullptr)
	True when privateLocal may be placed in shared memory.
std::optional< int64_t >	getPrivateLocalSharedMemoryUpperBoundBytes (PrivateLocalOp privateLocal, ComputeRegionOp computeRegion, ModuleOp module, const ACCToGPUMappingPolicy &policy, OpenACCSupport *support=nullptr)
	Upper-bound byte size for a shared-memory private_local candidate, or std::nullopt when not eligible or not statically computable.
bool	hasAttachPoint (Operation *mapEntryOp)
	Returns true when mapEntryOp carries an attach point (varPtrPtr).
DataDescKind	getDataDescKind (Operation *mapEntryOp)
	Returns descriptor kind from acc.map_info, or none for other ops.
Value	getDesc (Operation *mapEntryOp)
	Returns descriptor value from acc.map_info.
std::optional< int64_t >	getMapElementSize (Operation *mapEntryOp)
	Returns element size in bytes from acc.map_info, if present.
Value	getMapSize (Operation *mapEntryOp)
	Returns the optional size operand from acc.map_info, or null.
std::optional< MapFlags >	getMapFlags (Operation *mapEntryOp)
	Returns offload map-type flags from acc.map_info, if present.
MapFlags	computePrivatizeMapFlags (PrivatizeOp privatizeOp, const ACCToGPUMappingPolicy &policy)
	Compute the private and parallel-level map flags for privatized storage.
MapFlags	computeDataClauseMapFlags (Operation *entryOp, bool ptrAndObj)
	Fold enter (+ paired exit) data-clause semantics into offload map flags.
bool	hasCopyOutSibling (Operation *entryOp)
	True when another data clause of the same construct maps the same variable with a copy-back and no copy-in.
SmallVector< Operation * >	getPairedDataExitOps (Value entryResult)
	Returns the data exit operations paired with the data entry result entryResult, which take it as their accVar.
std::optional< Location >	getMappingExitLoc (ValueRange dataClauseOperands)
	Returns where the mappings of dataClauseOperands end, taken from the first of them that says.
int64_t	computeMapInfoSizeBytes (Value var, Type varType, DataDescKind descKind, ValueRange bounds, const DataLayout &dataLayout, OpenACCSupport *support=nullptr)
	Compute total mapped byte size for acc.map_info.
int64_t	computeMapInfoSizeBytes (Value var, Type varType, DataDescKind descKind, ValueRange bounds, OpenACCSupport *support=nullptr)
	Same as above, obtaining dataLayout from the module var lives in.
void	populateSourceExtents (ValueRange bounds, ArrayRef< int64_t > shape, OpBuilder &builder)
	Record known extents of the source array on bounds that may describe a section.
std::optional< gpu::GPUModuleOp >	getOrCreateGPUModule (ModuleOp mod, bool create=true, llvm::StringRef name=kDefaultGPUModuleName)
	Get or create a GPU module in the given module.
Value	getGPUSize (gpu::Processor processor, gpu::LaunchOp launch, const llvm::DenseMap< gpu::Processor, Value > &dimensionOps)
	Return the launch dimension for processor from launch, or from dimensionOps when launch is null.
Value	getGPUThreadId (gpu::Processor processor, gpu::LaunchOp launch, const llvm::DenseMap< gpu::Processor, Value > &indexOps)
	Return the thread/block index for processor from launch, or from indexOps when launch is null.
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)
	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.
Value	calculateTripCount (OpBuilder &b, Location loc, Value lb, Value ub, Value step, bool inclusiveUpperbound)
	Calculate trip count for a loop: (ub - lb + step) / step.
void	normalizeIVUses (OpBuilder &b, Location loc, Value iv, Value origLB, Value origStep)
	Normalize IV uses after converting to normalized loop form (lb=0, step=1).
void	setCollapseCountAttr (Operation *op, uint64_t count)
	Record on a collapsed loop how many original loops were folded into it.
uint64_t	getCollapseCount (Operation *op)
	Number of original loops collapsed into op, or 1 when op carries no collapse_count attribute.
SmallVector< GPUParallelDimAttr >	getReductionCombineParDims (ReductionCombineOp op)
	Returns the parallel dimensions that participate in op's combine step.
SmallVector< GPUParallelDimAttr >	getReductionCombineParDims (ReductionCombineRegionOp op)
	Returns the parallel dimensions that participate in op's combine step.
ReductionOperator	translateAtomicRMWKind (arith::AtomicRMWKind kind)
	Maps an arith atomic RMW kind to the corresponding acc reduction operator.
std::optional< arith::AtomicRMWKind >	translateACCReductionOperator (ReductionOperator redOp, Type type)
	Maps an acc reduction operator to the arith atomic RMW kind for type.
Value	createIdentityValue (OpBuilder &b, Location loc, Type type, arith::AtomicRMWKind kind, bool useOnlyFiniteValue=true)
	Creates the identity (neutral) value for a reduction of type and kind.
Value	generateReductionOp (OpBuilder &b, Location loc, Value lhs, Value rhs, arith::AtomicRMWKind kind)
	Combines two reduction partial values using the operator for kind.
mlir::acc::LoopOp	tileACCLoops (mlir::acc::LoopOp tileLoop, const llvm::SmallVector< mlir::Value > &tileSizes, int32_t defaultTileSize, mlir::RewriterBase &rewriter)
	Tile a single fused acc.loop that carries all associated induction variables (one IV per tile dimension).
std::optional< TypeSizeAndAlignment >	getTypeSizeAndAlignment (Type ty, ModuleOp module, const DataLayout &dl, OpenACCSupport *support=nullptr, Value var={})
	Returns the size and ABI alignment in bytes.
std::optional< TypeSizeAndAlignment >	getTypeSizeAndAlignment (Type ty, ModuleOp module, OpenACCSupport *support=nullptr, Value var={})
	Same as above, obtaining dl from module via getDataLayout.
Value	castPointerLikeTypeIfNeeded (OpBuilder &builder, Location loc, Value value, Type resultType)
	Cast value to resultType via PointerLikeType::genCast when needed.
std::unique_ptr<::mlir::Pass >	createACCBindRoutine ()
std::unique_ptr<::mlir::Pass >	createACCBindRoutine (ACCBindRoutineOptions options)
std::unique_ptr<::mlir::Pass >	createACCCGToGPU ()
std::unique_ptr<::mlir::Pass >	createACCCGToGPU (ACCCGToGPUOptions options)
std::unique_ptr<::mlir::Pass >	createACCComputeLowering ()
std::unique_ptr<::mlir::Pass >	createACCComputeLowering (ACCComputeLoweringOptions options)
std::unique_ptr<::mlir::Pass >	createACCDeclareCtorDtorConversion ()
std::unique_ptr<::mlir::Pass >	createACCDeclareCtorDtorConversion (ACCDeclareCtorDtorConversionOptions options)
std::unique_ptr<::mlir::Pass >	createACCDeclareGPUModuleInsertion ()
std::unique_ptr<::mlir::Pass >	createACCDeclareGPUModuleInsertion (ACCDeclareGPUModuleInsertionOptions options)
std::unique_ptr<::mlir::Pass >	createACCEmitRemarksData ()
std::unique_ptr<::mlir::Pass >	createACCEmitRemarksLoop ()
std::unique_ptr<::mlir::Pass >	createACCEmitRemarksLoop (ACCEmitRemarksLoopOptions options)
std::unique_ptr<::mlir::Pass >	createACCEmitRemarksPrivate ()
std::unique_ptr<::mlir::Pass >	createACCIfClauseLowering ()
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 >	createACCMaterializeRoutineBindTargets ()
std::unique_ptr<::mlir::Pass >	createACCMaterializeRoutineBindTargets (ACCMaterializeRoutineBindTargetsOptions options)
std::unique_ptr<::mlir::Pass >	createACCRecipeMaterialization ()
std::unique_ptr<::mlir::Pass >	createACCRoutineLowering ()
std::unique_ptr<::mlir::Pass >	createACCRoutineLowering (ACCRoutineLoweringOptions options)
std::unique_ptr<::mlir::Pass >	createACCRoutineToGPUFunc ()
std::unique_ptr<::mlir::Pass >	createACCRoutineToGPUFunc (ACCRoutineToGPUFuncOptions options)
std::unique_ptr<::mlir::Pass >	createACCSpecializeForDevice ()
std::unique_ptr<::mlir::Pass >	createACCSpecializeForDevice (ACCSpecializeForDeviceOptions options)
std::unique_ptr<::mlir::Pass >	createACCSpecializeForHost ()
std::unique_ptr<::mlir::Pass >	createACCSpecializeForHost (ACCSpecializeForHostOptions options)
std::unique_ptr<::mlir::Pass >	createLegalizeDataValuesInRegion ()
std::unique_ptr<::mlir::Pass >	createLegalizeDataValuesInRegion (LegalizeDataValuesInRegionOptions options)
std::unique_ptr<::mlir::Pass >	createOffloadLiveInValueCanonicalization ()
std::unique_ptr<::mlir::Pass >	createOffloadTargetVerifier ()
std::unique_ptr<::mlir::Pass >	createOffloadTargetVerifier (OffloadTargetVerifierOptions options)
void	populateACCSpecializeForDevicePatterns (RewritePatternSet &patterns, const TypesForDevice &theDeviceTypes)
	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).
void	registerACCBindRoutine ()
void	registerACCBindRoutinePass ()
void	registerACCCGToGPU ()
void	registerACCCGToGPUPass ()
void	registerACCComputeLowering ()
void	registerACCComputeLoweringPass ()
void	registerACCDeclareCtorDtorConversion ()
void	registerACCDeclareCtorDtorConversionPass ()
void	registerACCDeclareGPUModuleInsertion ()
void	registerACCDeclareGPUModuleInsertionPass ()
void	registerACCEmitRemarksData ()
void	registerACCEmitRemarksDataPass ()
void	registerACCEmitRemarksLoop ()
void	registerACCEmitRemarksLoopPass ()
void	registerACCEmitRemarksPrivate ()
void	registerACCEmitRemarksPrivatePass ()
void	registerACCIfClauseLowering ()
void	registerACCIfClauseLoweringPass ()
void	registerACCImplicitData ()
void	registerACCImplicitDataPass ()
void	registerACCImplicitDeclare ()
void	registerACCImplicitDeclarePass ()
void	registerACCImplicitRoutine ()
void	registerACCImplicitRoutinePass ()
void	registerACCLegalizeSerial ()
void	registerACCLegalizeSerialPass ()
void	registerACCLoopTiling ()
void	registerACCLoopTilingPass ()
void	registerACCMaterializeRoutineBindTargets ()
void	registerACCMaterializeRoutineBindTargetsPass ()
void	registerACCRecipeMaterialization ()
void	registerACCRecipeMaterializationPass ()
void	registerACCRoutineLowering ()
void	registerACCRoutineLoweringPass ()
void	registerACCRoutineToGPUFunc ()
void	registerACCRoutineToGPUFuncPass ()
void	registerACCSpecializeForDevice ()
void	registerACCSpecializeForDevicePass ()
void	registerACCSpecializeForHost ()
void	registerACCSpecializeForHostPass ()
void	registerLegalizeDataValuesInRegion ()
void	registerLegalizeDataValuesInRegionPass ()
void	registerOffloadLiveInValueCanonicalization ()
void	registerOffloadLiveInValueCanonicalizationPass ()
void	registerOffloadTargetVerifier ()
void	registerOffloadTargetVerifierPass ()
void	registerOpenACCPasses ()
static SmallVector< GPUParallelDimAttr >::iterator	findParDim (SmallVector< GPUParallelDimAttr > &parDims, GPUParallelDimAttr parDim)
static bool	isThreadXPrivatize (PrivatizeOp privatize)
static FailureOr< std::optional< int64_t > >	getWorkerPrivateSharedMemoryNumCopies (PrivateLocalOp privateLocal, ComputeRegionOp computeRegion, bool isWorkerPrivate, OpenACCSupport *support)
static bool	isInsideACCSpecializedRoutine (Operation *op)
static Operation *	findCorrespondingDataExit (Value entryResult)
static std::optional< Location >	getMappingExitLoc (Value entryResult)
static std::optional< DataClause >	getExitDataClause (Operation *exitOp)
static DataClauseModifier	getEntryModifiers (Operation *entryOp)
static ModuleOp	getEnclosingModule (Value var)
	Returns the module var lives in.
static Value	getGPUSizeFromLaunch (gpu::LaunchOp launch, gpu::Processor processor)
static Value	getGPUThreadIdFromLaunch (gpu::LaunchOp launch, gpu::Processor processor)
static bool	isFloatOrComplexType (Type ty)
static TypedAttr	getReductionIdentityValueAttr (arith::AtomicRMWKind kind, Type type, OpBuilder &builder, Location loc, bool useOnlyFiniteValue)
static std::optional< TypeSizeAndAlignment >	getTypeSizeAndAlignmentHelper (Type ty, ModuleOp module, const DataLayout &dl, OpenACCSupport *support)

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>

Type alias for the GPU-specific mapping policy.

Definition at line 158 of file OpenACCParMapping.h.

TypesForDevice

using mlir::acc::TypesForDevice = llvm::SmallSetVector<int64_t, 3>

Holds information for which integers represent a device type in the runtime.

Definition at line 33 of file Passes.h.

TypeSizeAndAlignment

using mlir::acc::TypeSizeAndAlignment = std::pair<llvm::TypeSize, llvm::TypeSize>

Definition at line 33 of file OpenACCUtilsType.h.

Enumeration Type Documentation

DataDescriptor

enum class mlir::acc::DataDescriptor

strong

IDs for the argument descriptors of the OpenACC data entry points, declared in OpenACCRuntimeDescriptors.def.

Definition at line 56 of file OpenACCRuntimeUtils.h.

KernelArgumentField

enum class mlir::acc::KernelArgumentField : int64_t

strong

Fields of the argument struct passed to __tgt_acc_kernel, in declaration order.

Declared in OpenACCRuntimeKernelArgs.def.

Definition at line 77 of file OpenACCRuntimeUtils.h.

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

RuntimeFunction

enum class mlir::acc::RuntimeFunction

strong

IDs for OpenACC compiler-to-runtime entry points (__tgt_acc_*).

Definition at line 35 of file OpenACCRuntimeUtils.h.

RuntimeResult

enum class mlir::acc::RuntimeResult : int32_t

strong

Status returned by the runtime entry points that report whether they completed.

The values are part of the runtime ABI.

Enumerator
Success
Error

Definition at line 49 of file OpenACCRuntimeUtils.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. Launch operands should be acc.par_width results (index); the region entry block gets matching index block arguments first, then arguments for each ins operand. 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 69 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::getIndexType(), mlir::Block::getOperations(), mlir::IRMapping::lookup(), mlir::IRMapping::map(), mlir::OpBuilder::setInsertionPointToEnd(), mlir::OpBuilder::setInsertionPointToStart(), and wrapMultiBlockRegionWithSCFExecuteRegion().

calculateTripCount()

Value mlir::acc::calculateTripCount	(	OpBuilder &	b,
		Location	loc,
		Value	lb,
		Value	ub,
		Value	step,
		bool	inclusiveUpperbound )

Calculate trip count for a loop: (ub - lb + step) / step.

Calculate trip count for a loop: (ub - lb + step) / step If inclusiveUpperbound is false, subtracts 1 from ub first.

If inclusiveUpperbound is false, subtracts 1 from ub first. Operands are cast to index type.

Definition at line 30 of file OpenACCUtilsLoop.cpp.

References add, b, mlir::arith::ConstantIndexOp::create(), and mlir::getValueOrCreateCastToIndexLike().

Referenced by convertACCLoopToSCFFor(), and convertACCLoopToSCFParallel().

castPointerLikeTypeIfNeeded()

Value mlir::acc::castPointerLikeTypeIfNeeded	(	OpBuilder &	builder,
		Location	loc,
		Value	value,
		Type	resultType )

Cast value to resultType via PointerLikeType::genCast when needed.

Returns value unchanged if types already match. Emits an error and returns value if no cast can be generated.

Definition at line 108 of file OpenACCUtilsType.cpp.

References mlir::emitError(), and mlir::Value::getType().

Referenced by getPointer().

castToI64()

Value mlir::acc::castToI64	(	Location	loc,
		Value	value,
		OpBuilder &	builder )

Sign-extends or truncates value to the i64 the runtime entry points take for values like queue numbers.

Definition at line 215 of file ACCToLLVMUtils.cpp.

References mlir::Builder::getI64Type(), mlir::Type::getIntOrFloatBitWidth(), and mlir::Value::getType().

Referenced by createKernelArgumentStruct(), emitWaitCall(), getAsyncQueue(), and getBoundValue().

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 leading operands of the cloned region's acc.yield (1:1). Erases acc.yield/terminator and merges blocks. Returns all yielded values, including any values not used as replacements, and the insertion point after the clone.

Definition at line 113 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().

Referenced by cloneACCRegionIntoForLoop().

collectPrivateLocalParDims()

SmallVector< GPUParallelDimAttr > mlir::acc::collectPrivateLocalParDims	(	PrivateLocalOp	privateLocal,
		ComputeRegionOp	computeRegion )

Collect parallel dimensions that govern privatization of privateLocal.

Definition at line 319 of file OpenACCUtilsCG.cpp.

References getParDimsAttr(), getReductionCombineParDims(), mlir::Operation::getResult(), mlir::Value::getUsers(), and insertParDim().

Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().

computeDataClauseMapFlags()

MapFlags mlir::acc::computeDataClauseMapFlags	(	Operation *	entryOp,
		bool	ptrAndObj )

Fold enter (+ paired exit) data-clause semantics into offload map flags.

ptrAndObj is supplied by the caller from type-specific attach discovery.

Definition at line 645 of file OpenACCUtilsCG.cpp.

References findCorrespondingDataExit(), getDataClause(), getEntryModifiers(), getExitDataClause(), getImplicitFlag(), getNoDeviceDescriptorAttrName(), mlir::Operation::getResult(), mlir::Value::getUses(), hasCopyOutSibling(), and mlir::Operation::hasDiscardableAttr().

Referenced by computePackedMapFlags(), and getKernelDataClauseMapFlags().

computeMapInfoSizeBytes() [1/2]

int64_t mlir::acc::computeMapInfoSizeBytes	(	Value	var,
		Type	varType,
		DataDescKind	descKind,
		ValueRange	bounds,
		const DataLayout &	dataLayout,
		OpenACCSupport *	support = nullptr )

Compute total mapped byte size for acc.map_info.

Returns 0 when bounds or a non-none descriptor kind carry size, the mappable size when statically known, and -1 when the size cannot be determined statically. support sizes the members of aggregate types that belong to a dialect, such as a tuple holding dialect-specific references.

Definition at line 773 of file OpenACCUtilsCG.cpp.

References getEnclosingModule(), mlir::Value::getType(), and getTypeSizeAndAlignment().

Referenced by computeMapInfoSizeBytes().

computeMapInfoSizeBytes() [2/2]

int64_t mlir::acc::computeMapInfoSizeBytes	(	Value	var,
		Type	varType,
		DataDescKind	descKind,
		ValueRange	bounds,
		OpenACCSupport *	support = nullptr )

Same as above, obtaining dataLayout from the module var lives in.

Definition at line 800 of file OpenACCUtilsCG.cpp.

References computeMapInfoSizeBytes(), getDataLayout(), and getEnclosingModule().

computePrivatizeMapFlags()

MapFlags mlir::acc::computePrivatizeMapFlags	(	PrivatizeOp	privatizeOp,
		const ACCToGPUMappingPolicy &	policy )

Compute the private and parallel-level map flags for privatized storage.

Storage that names no parallel dimension is private without being replicated per level, so only the private flag is set.

Definition at line 624 of file OpenACCUtilsCG.cpp.

References mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isGang(), mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isVector(), and mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isWorker().

constantI64()

Value mlir::acc::constantI64	(	Location	loc,
		int64_t	value,
		OpBuilder &	builder )

Creates an i64 constant.

Definition at line 227 of file ACCToLLVMUtils.cpp.

References mlir::Builder::getI64Type().

Referenced by createKernelArgumentStruct(), mlir::emitACCDataRuntimeArgs(), emitKernelCall(), emitKernelRuntimeArgs(), getBoundValue(), getElementSize(), and getKernelArgDescriptor().

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 191 of file OpenACCUtilsLoop.cpp.

References calculateTripCount(), cloneACCRegionIntoForLoop(), mlir::coalesceLoops(), copyLoopAnnotationAttr(), mlir::arith::ConstantIndexOp::create(), mlir::OpBuilder::getInsertionPoint(), mlir::IRMapping::map(), mapACCLoopIVsToSCFIVs(), normalizeIVUses(), setCollapseCountAttr(), 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 272 of file OpenACCUtilsLoop.cpp.

References calculateTripCount(), cloneACCRegionIntoForLoop(), copyLoopAnnotationAttr(), mlir::arith::ConstantIndexOp::create(), mlir::RewriterBase::eraseOp(), mlir::OpBuilder::getInsertionBlock(), mlir::OpBuilder::getInsertionPoint(), mlir::Block::getParentOp(), mapACCLoopIVsToSCFIVs(), normalizeIVUses(), setCollapseCountAttr(), mlir::OpBuilder::setInsertionPointToStart(), 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 338 of file OpenACCUtilsLoop.cpp.

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

copyParDimsAttr()

void mlir::acc::copyParDimsAttr	(	Operation *	from,
		Operation *	to )

Copy parallel dimensions from from to to.

Definition at line 246 of file OpenACCUtilsCG.cpp.

References getParDimsAttr(), hasParDimsAttr(), and setParDimsAttr().

createACCBindRoutine() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCBindRoutine

(

)

Definition at line 123 of file ACCBindRoutine.cpp.

createACCBindRoutine() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCBindRoutine

(

ACCBindRoutineOptions

options

)

Definition at line 127 of file ACCBindRoutine.cpp.

References getRoutineInfoAttrName().

createACCCGToGPU() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCCGToGPU

(

)

Definition at line 238 of file ACCCGToGPU.cpp.

createACCCGToGPU() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCCGToGPU

(

ACCCGToGPUOptions

options

)

Definition at line 242 of file ACCCGToGPU.cpp.

createACCComputeLowering() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCComputeLowering

(

)

Definition at line 340 of file ACCComputeLowering.cpp.

createACCComputeLowering() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCComputeLowering

(

ACCComputeLoweringOptions

options

)

Definition at line 344 of file ACCComputeLowering.cpp.

createACCDeclareCtorDtorConversion() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareCtorDtorConversion

(

)

Definition at line 450 of file ACCDeclareCtorDtorConversion.cpp.

createACCDeclareCtorDtorConversion() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareCtorDtorConversion

(

ACCDeclareCtorDtorConversionOptions

options

)

Definition at line 454 of file ACCDeclareCtorDtorConversion.cpp.

createACCDeclareGPUModuleInsertion() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareGPUModuleInsertion

(

)

Definition at line 545 of file ACCDeclareGPUModuleInsertion.cpp.

createACCDeclareGPUModuleInsertion() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareGPUModuleInsertion

(

ACCDeclareGPUModuleInsertionOptions

options

)

Definition at line 549 of file ACCDeclareGPUModuleInsertion.cpp.

createACCEmitRemarksData()

std::unique_ptr<::mlir::Pass > mlir::acc::createACCEmitRemarksData

(

)

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

Definition at line 624 of file ACCEmitRemarksData.cpp.

createACCEmitRemarksLoop() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCEmitRemarksLoop

(

)

Definition at line 715 of file ACCEmitRemarksLoop.cpp.

createACCEmitRemarksLoop() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCEmitRemarksLoop

(

ACCEmitRemarksLoopOptions

options

)

Definition at line 719 of file ACCEmitRemarksLoop.cpp.

createACCEmitRemarksPrivate()

std::unique_ptr<::mlir::Pass > mlir::acc::createACCEmitRemarksPrivate

(

)

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

Definition at line 794 of file ACCEmitRemarksPrivate.cpp.

createACCIfClauseLowering()

std::unique_ptr<::mlir::Pass > mlir::acc::createACCIfClauseLowering

(

)

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

Definition at line 869 of file ACCIfClauseLowering.cpp.

createACCImplicitData() [1/2]

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

(

)

Definition at line 964 of file ACCImplicitData.cpp.

createACCImplicitData() [2/2]

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

(

ACCImplicitDataOptions

options

)

Definition at line 968 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 1043 of file ACCImplicitDeclare.cpp.

createACCImplicitRoutine() [1/2]

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

(

)

Definition at line 1143 of file ACCImplicitRoutine.cpp.

createACCImplicitRoutine() [2/2]

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

(

ACCImplicitRoutineOptions

options

)

Definition at line 1147 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 1222 of file ACCLegalizeSerial.cpp.

createACCLoopTiling() [1/2]

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

(

)

Definition at line 1313 of file ACCLoopTiling.cpp.

createACCLoopTiling() [2/2]

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

(

ACCLoopTilingOptions

options

)

Definition at line 1317 of file ACCLoopTiling.cpp.

createACCMaterializeRoutineBindTargets() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCMaterializeRoutineBindTargets

(

)

Definition at line 1414 of file ACCBindRoutine.cpp.

createACCMaterializeRoutineBindTargets() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCMaterializeRoutineBindTargets

(

ACCMaterializeRoutineBindTargetsOptions

options

)

Definition at line 1418 of file ACCBindRoutine.cpp.

createACCRecipeMaterialization()

std::unique_ptr<::mlir::Pass > mlir::acc::createACCRecipeMaterialization

(

)

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

Definition at line 1493 of file ACCRecipeMaterialization.cpp.

createACCRoutineLowering() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineLowering

(

)

Definition at line 1591 of file ACCRoutineLowering.cpp.

createACCRoutineLowering() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineLowering

(

ACCRoutineLoweringOptions

options

)

Definition at line 1595 of file ACCRoutineLowering.cpp.

createACCRoutineToGPUFunc() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineToGPUFunc

(

)

Definition at line 1692 of file ACCRoutineToGPUFunc.cpp.

createACCRoutineToGPUFunc() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineToGPUFunc

(

ACCRoutineToGPUFuncOptions

options

)

Definition at line 1696 of file ACCRoutineToGPUFunc.cpp.

createACCSpecializeForDevice() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForDevice

(

)

Definition at line 1787 of file ACCSpecializeForDevice.cpp.

createACCSpecializeForDevice() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForDevice

(

ACCSpecializeForDeviceOptions

options

)

Definition at line 1791 of file ACCSpecializeForDevice.cpp.

createACCSpecializeForHost() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForHost

(

)

Definition at line 1882 of file ACCSpecializeForHost.cpp.

createACCSpecializeForHost() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForHost

(

ACCSpecializeForHostOptions

options

)

Definition at line 1886 of file ACCSpecializeForHost.cpp.

createIdent()

Value mlir::acc::createIdent	(	Location	loc,
		StringRef	functionName,
		OpBuilder &	builder,
		Region &	globalSymbolRegion,
		const ACCRuntimeCallConfig &	config,
		SymbolTable *	symbolTable = nullptr )

Returns a pointer to a constant global holding an ident_t for OpenACC runtime calls.

globalSymbolRegion and symbolTable are as in getOrCreateGlobalString; the ident and the source string it points to are named after the position they describe, so leaving out the table creates a set of them per call.

Definition at line 145 of file ACCToLLVMUtils.cpp.

References mlir::OpBuilder::createBlock(), mlir::Builder::getContext(), getFileLineColLoc(), mlir::acc::ACCRuntimeCallConfig::getFunctionDisplayName(), mlir::Builder::getI32Type(), mlir::Builder::getI64IntegerAttr(), mlir::Builder::getI64Type(), getInternalGlobalName(), getOrCreateGlobalStringOp(), mlir::SymbolTable::insert(), mlir::SymbolTable::lookup(), mlir::OpBuilder::setInsertionPointAfter(), and mlir::OpBuilder::setInsertionPointToStart().

Referenced by mlir::emitACCDataRuntimeArgs(), emitGetDevicePtrCall(), emitKernelCall(), and emitWaitCall().

createIdentityValue()

Value mlir::acc::createIdentityValue	(	OpBuilder &	b,
		Location	loc,
		Type	type,
		arith::AtomicRMWKind	kind,
		bool	useOnlyFiniteValue = true )

Creates the identity (neutral) value for a reduction of type and kind.

When useOnlyFiniteValue is true, floating-point identities avoid non-finite sentinel values where applicable.

Definition at line 186 of file OpenACCUtilsReduction.cpp.

References b, and getReductionIdentityValueAttr().

createLegalizeDataValuesInRegion() [1/2]

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

(

)

Definition at line 1980 of file LegalizeDataValues.cpp.

createLegalizeDataValuesInRegion() [2/2]

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

(

LegalizeDataValuesInRegionOptions

options

)

Definition at line 1984 of file LegalizeDataValues.cpp.

createOffloadLiveInValueCanonicalization()

std::unique_ptr<::mlir::Pass > mlir::acc::createOffloadLiveInValueCanonicalization

(

)

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

Definition at line 2059 of file OffloadLiveInValueCanonicalization.cpp.

createOffloadTargetVerifier() [1/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createOffloadTargetVerifier

(

)

Definition at line 2159 of file OffloadTargetVerifier.cpp.

createOffloadTargetVerifier() [2/2]

std::unique_ptr<::mlir::Pass > mlir::acc::createOffloadTargetVerifier

(

OffloadTargetVerifierOptions

options

)

Definition at line 2163 of file OffloadTargetVerifier.cpp.

createRuntimeCall()

FailureOr< LLVM::CallOp > mlir::acc::createRuntimeCall	(	Location	loc,
		OpBuilder &	builder,
		Region &	globalSymbolRegion,
		SymbolTable &	symbolTable,
		RuntimeFunction	fn,
		const ACCRuntimeCallConfig &	config,
		ArrayRef< Value >	arguments )

Declares (if needed) and returns a call to the runtime function identified by fn using the name from config.

Fails and emits a diagnostic if the symbol is already declared with a signature the runtime cannot be called through. The declaration is created in globalSymbolRegion and registered in symbolTable.

Definition at line 243 of file OpenACCRuntimeUtils.cpp.

References mlir::OpBuilder::atBlockEnd(), mlir::emitError(), mlir::Region::front(), mlir::Builder::getContext(), mlir::acc::ACCRuntimeCallConfig::getName(), getRuntimeFunctionType(), mlir::SymbolTable::insert(), and mlir::SymbolTable::lookup().

Referenced by mlir::emitACCDataRuntimeCall(), emitGetDevicePtrCall(), emitKernelCall(), and emitWaitCall().

emitGetDevicePtrCall()

FailureOr< Value > mlir::acc::emitGetDevicePtrCall	(	Operation *	clauseOp,
		Value	hostPtr,
		bool	ifPresent,
		OpBuilder &	builder,
		Region &	globalSymbolRegion,
		SymbolTable &	symbolTable,
		const ACCRuntimeCallConfig &	config )

Emits the runtime call that asks for the device address the object of the data clause clauseOp is mapped to, which is what a use_device clause exposes in the body of its construct.

hostPtr is the address of that object, already converted to the LLVM dialect. With ifPresent, an object that is not mapped keeps its host address instead of being reported.

Bounds on the clause say which part of the object it names, which the address asked about does not have to state: the result stands for the object, so whatever reads it addresses the part it wants as it would on the host.

Definition at line 301 of file ACCToLLVMUtils.cpp.

References createIdent(), createRuntimeCall(), mlir::Builder::getI64Type(), mlir::Operation::getLoc(), mlir::acc::ACCRuntimeCallConfig::getMapFlagsRuntimeValue(), and getParentFunctionName().

emitGuardedByIfCond()

LogicalResult mlir::acc::emitGuardedByIfCond	(	Location	loc,
		Value	ifCond,
		RewriterBase &	rewriter,
		function_ref< LogicalResult()>	emitFn )

Runs emitFn guarded by a branch on ifCond, or unguarded when there is no condition.

Leaves the insertion point after the guarded code, so that a caller can keep emitting into the same block either way.

Definition at line 325 of file ACCToLLVMUtils.cpp.

References mlir::OpBuilder::createBlock(), mlir::OpBuilder::getInsertionBlock(), mlir::OpBuilder::getInsertionPoint(), mlir::Block::getParent(), result, mlir::OpBuilder::setInsertionPointToEnd(), mlir::OpBuilder::setInsertionPointToStart(), and mlir::RewriterBase::splitBlock().

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 173 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 509 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().

emitValueSelectedByIfCond()

FailureOr< Value > mlir::acc::emitValueSelectedByIfCond	(	Location	loc,
		Value	ifCond,
		RewriterBase &	rewriter,
		function_ref< FailureOr< Value >()>	thenFn,
		function_ref< FailureOr< Value >()>	elseFn )

Emits thenFn on the path a branch on ifCond takes and elseFn on the other one, and returns the value that reaches the code following the branch.

Both have to produce a value, and of the same type. With no condition only thenFn is emitted and its value returned. As with emitGuardedByIfCond, the insertion point is left after the branch.

Definition at line 350 of file ACCToLLVMUtils.cpp.

References mlir::Block::addArgument(), mlir::OpBuilder::createBlock(), mlir::OpBuilder::getInsertionBlock(), mlir::OpBuilder::getInsertionPoint(), mlir::Block::getParent(), mlir::OpBuilder::setInsertionPointToEnd(), mlir::OpBuilder::setInsertionPointToStart(), and mlir::RewriterBase::splitBlock().

emitWaitCall()

LogicalResult mlir::acc::emitWaitCall	(	Location	loc,
		ValueRange	waitOperands,
		Value	asyncQueue,
		OpBuilder &	builder,
		Region &	globalSymbolRegion,
		SymbolTable &	symbolTable,
		const ACCRuntimeCallConfig &	config,
		Value	deviceNum = {} )

Emits the runtime call that waits for waitOperands on asyncQueue, which is what a wait clause or an acc.wait directive asks for.

An empty waitOperands waits for every queue, as a wait clause without values does. The values must already be converted to the LLVM dialect. If provided, deviceNum must be an i32 value; otherwise the device number defaults to zero.

Definition at line 252 of file ACCToLLVMUtils.cpp.

References castToI64(), createIdent(), createRuntimeCall(), mlir::Builder::getContext(), mlir::acc::ACCRuntimeCallConfig::getDeviceTypeRuntimeValue(), mlir::Builder::getI32Type(), mlir::Builder::getI64Type(), mlir::OpBuilder::getInsertionBlock(), and getParentFunctionName().

Referenced by emitWaitClause().

emitWaitClause()

template<typename OpTy>

LogicalResult mlir::acc::emitWaitClause	(	OpTy	op,
		DeviceType	deviceType,
		Value	asyncQueue,
		ConversionPatternRewriter &	rewriter,
		OpenACCSupport &	accSupport,
		Region &	globalSymbolRegion,
		SymbolTable &	symbolTable,
		const ACCRuntimeCallConfig &	config )

Emits the wait that a wait clause on op asks for before the runtime calls of the construct, waiting on asyncQueue for the queues the clause names, or for every queue when it names none.

Nothing is emitted when there is no such clause for deviceType. A devnum modifier, which selects the device the queues belong to, is reported through accSupport as not yet implemented.

Definition at line 263 of file ACCToLLVMUtils.h.

References mlir::acc::OpenACCSupport::emitNYI(), emitWaitCall(), getWaitClauseValues(), hasWaitDevnum(), and success().

findCorrespondingDataExit()

Operation * mlir::acc::findCorrespondingDataExit ( Value entryResult )

static

Definition at line 555 of file OpenACCUtilsCG.cpp.

References getPairedDataExitOps(), and nullptr.

Referenced by computeDataClauseMapFlags(), and getMappingExitLoc().

findParDim()

SmallVector< GPUParallelDimAttr >::iterator mlir::acc::findParDim ( SmallVector< GPUParallelDimAttr > & parDims, GPUParallelDimAttr parDim )

static

Definition at line 143 of file OpenACCUtilsCG.cpp.

Referenced by insertParDim(), and removeParDim().

generateReductionOp()

Value mlir::acc::generateReductionOp	(	OpBuilder &	b,
		Location	loc,
		Value	lhs,
		Value	rhs,
		arith::AtomicRMWKind	kind )

Combines two reduction partial values using the operator for kind.

Definition at line 203 of file OpenACCUtilsReduction.cpp.

References b, and mlir::arith::getReductionOp().

getACCDataClauseOpForBlockArg()

mlir::Operation * mlir::acc::getACCDataClauseOpForBlockArg

(

mlir::Value

v

)

If v is not a block argument of an acc.compute_region body, returns nullptr.

Otherwise maps the block argument to its operand and returns the defining operation if it is one of ACC_DATA_ENTRY_OPS.

Definition at line 47 of file OpenACCUtils.cpp.

References getACCOperandForBlockArg(), and mlir::Value::getDefiningOp().

getACCOperandForBlockArg()

mlir::Value mlir::acc::getACCOperandForBlockArg

(

mlir::Value

v

)

If v is not a block argument of an acc.compute_region body, returns nullptr.

Otherwise maps the block argument to its operand and returns it.

Definition at line 32 of file OpenACCUtils.cpp.

References mlir::Block::getParentOp().

Referenced by getACCDataClauseOpForBlockArg().

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 5422 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 5435 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

getActiveParDimsAttr()

ActiveParDimsAttr mlir::acc::getActiveParDimsAttr

(

Operation *

op

)

Obtain the active parallel dimensions carried by op, if any.

Definition at line 252 of file OpenACCUtilsCG.cpp.

References mlir::Operation::getDiscardableAttrOfType().

Referenced by hasActiveParDimsAttr().

getAsyncClauseValue()

template<typename OpTy>

Value mlir::acc::getAsyncClauseValue	(	OpTy	op,
		DeviceType	deviceType,
		bool &	asyncOnly )

Returns the value the async clause of op names for deviceType, and sets asyncOnly when the clause names no queue.

The value is the one the operation holds, so a caller in a conversion has to remap it.

Definition at line 168 of file ACCToLLVMUtils.h.

References getDeviceTypesByPrecedence().

Referenced by getAsyncQueue().

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 5493 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 5471 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 5485 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

getAsyncQueue() [1/2]

Value mlir::acc::getAsyncQueue	(	Location	loc,
		Value	asyncOperand,
		bool	asyncOnly,
		OpBuilder &	builder,
		const ACCRuntimeCallConfig &	config )

Returns the queue an async clause selects: the value of the clause when it has one, the queue standing for an async clause without a value when asyncOnly is set, and the synchronous queue when there is no clause at all.

asyncOperand must already be converted to the LLVM dialect.

Definition at line 239 of file ACCToLLVMUtils.cpp.

References castToI64(), mlir::acc::ACCRuntimeCallConfig::getAsyncNoValueRuntimeValue(), mlir::acc::ACCRuntimeCallConfig::getAsyncSyncRuntimeValue(), and mlir::Builder::getI64Type().

Referenced by getAsyncQueue().

getAsyncQueue() [2/2]

template<typename OpTy>

Value mlir::acc::getAsyncQueue	(	OpTy	op,
		DeviceType	deviceType,
		ConversionPatternRewriter &	rewriter,
		const ACCRuntimeCallConfig &	config )

Returns the queue that the runtime calls of op run on, from the async clause it gives for deviceType.

Definition at line 245 of file ACCToLLVMUtils.h.

References getAsyncClauseValue(), and getAsyncQueue().

getBaseEntity()

mlir::Value mlir::acc::getBaseEntity

(

mlir::Value

val

)

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

Referenced by mlir::createACCArgumentDescriptor().

getChunkSize()

std::optional< int64_t > mlir::acc::getChunkSize

(

Operation *

op

)

Return the chunk size carried by op, if any.

Definition at line 240 of file OpenACCUtilsCG.cpp.

References getChunkSizeAttr().

getChunkSizeAttr()

ChunkSizeAttr mlir::acc::getChunkSizeAttr

(

Operation *

op

)

Obtain the acc.chunk_size attribute carried by op, if any.

Definition at line 226 of file OpenACCUtilsCG.cpp.

References mlir::Operation::getDiscardableAttrOfType().

Referenced by getChunkSize(), and hasChunkSizeAttr().

getCollapseCount()

uint64_t mlir::acc::getCollapseCount

(

Operation *

op

)

Number of original loops collapsed into op, or 1 when op carries no collapse_count attribute.

Definition at line 358 of file OpenACCUtilsLoop.cpp.

References getCollapseCountAttrName(), and mlir::Operation::getDiscardableAttrOfType().

getCollapseCountAttrName()

constexpr StringLiteral mlir::acc::getCollapseCountAttrName ( )

staticconstexpr

Name for an attribute attached to a loop indicating the number of loops collapsed to create that loop.

Definition at line 238 of file OpenACC.h.

Referenced by getCollapseCount(), and setCollapseCountAttr().

getCombinedConstructsAttrName()

constexpr StringLiteral mlir::acc::getCombinedConstructsAttrName ( )

staticconstexpr

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

Referenced by checkDeclareOperands(), computeDataClauseMapFlags(), getKernelDataClauseMapFlags(), hasAttachPoint(), hasCopyOutSibling(), and isDevicePtrMapping().

getDataDescKind()

DataDescKind mlir::acc::getDataDescKind

(

Operation *

mapEntryOp

)

Returns descriptor kind from acc.map_info, or none for other ops.

Definition at line 501 of file OpenACCUtilsCG.cpp.

Referenced by mlir::createACCArgumentDescriptor().

getDataDescriptorFieldIndex()

template<typename FieldEnum>

int64_t mlir::acc::getDataDescriptorFieldIndex

(

FieldEnum

field

)

The index an insert or extract of field addresses.

Definition at line 71 of file OpenACCRuntimeUtils.h.

Referenced by mlir::createACCArgumentDescriptor(), mlir::createACCDataDescriptor(), and mlir::createACCMemRefDescriptorWrapperArg().

getDataDescriptorKind()

DataDescKind mlir::acc::getDataDescriptorKind

(

DataDescriptor

desc

)

Returns the descriptor kind the runtime reads from the version field of desc.

An overlay contributes its kind to the descriptor it nests.

Definition at line 62 of file OpenACCRuntimeUtils.cpp.

Referenced by mlir::createACCArgumentDescriptor(), mlir::createACCDataDescriptor(), and mlir::createACCMemRefDescriptorWrapperArg().

getDataDescriptorName()

StringRef mlir::acc::getDataDescriptorName

(

DataDescriptor

desc

)

Returns the name of the runtime type desc materializes.

Definition at line 52 of file OpenACCRuntimeUtils.cpp.

getDataDescriptorType()

LLVM::LLVMStructType mlir::acc::getDataDescriptorType	(	MLIRContext *	ctx,
		DataDescriptor	desc,
		Type	baseType = {} )

Builds the LLVM type of desc in ctx.

A descriptor that nests another one takes it as baseType, which is ignored otherwise.

Definition at line 72 of file OpenACCRuntimeUtils.cpp.

Referenced by mlir::createACCArgumentDescriptor(), mlir::createACCDataDescriptor(), and mlir::createACCMemRefDescriptorWrapperArg().

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 34 of file OpenACCUtilsCG.cpp.

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

Referenced by computeMapInfoSizeBytes(), and getTypeSizeAndAlignment().

getDataOperands()

mlir::ValueRange mlir::acc::getDataOperands

(

mlir::Operation *

accOp

)

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

Definition at line 5553 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

getDeclareActionAttrName()

constexpr StringLiteral mlir::acc::getDeclareActionAttrName ( )

staticconstexpr

Definition at line 200 of file OpenACC.h.

Referenced by getEnclosingDeclareAction().

getDeclareAttrName()

constexpr StringLiteral mlir::acc::getDeclareAttrName ( )

staticconstexpr

Used to obtain the attribute name for declare.

Definition at line 196 of file OpenACC.h.

Referenced by isDeviceAccessibleValue(), isInDeviceMemoryValue(), and 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 79 of file OpenACCUtils.cpp.

References mlir::Operation::getParentOp().

getDesc()

Value mlir::acc::getDesc

(

Operation *

mapEntryOp

)

Returns descriptor value from acc.map_info.

When desc is omitted and descKind is not none, returns var (the mapped object is the descriptor). Returns null for other ops.

Definition at line 507 of file OpenACCUtilsCG.cpp.

Referenced by mlir::createACCArgumentDescriptor().

getDeviceTypesByPrecedence()

SmallVector< DeviceType, 3 > mlir::acc::getDeviceTypesByPrecedence

(

DeviceType

deviceType

)

The clauses of a construct can be given once per device type.

Of the values that reach a given device type, the ones naming it are the most specific, then the ones naming every device type, then the ones given before any device_type clause.

Definition at line 5501 of file OpenACC.cpp.

Referenced by getAsyncClauseValue(), and getWaitClauseValues().

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

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

getEnclosingModule()

ModuleOp mlir::acc::getEnclosingModule ( Value var )

static

Returns the module var lives in.

Definition at line 764 of file OpenACCUtilsCG.cpp.

References mlir::Value::getDefiningOp(), mlir::Operation::getParentOp(), and mlir::Value::getParentRegion().

Referenced by computeMapInfoSizeBytes(), and computeMapInfoSizeBytes().

getEntryModifiers()

DataClauseModifier mlir::acc::getEntryModifiers ( Operation * entryOp )

static

Definition at line 617 of file OpenACCUtilsCG.cpp.

References ACC_DATA_ENTRY_OPS.

Referenced by computeDataClauseMapFlags().

getExitDataClause()

std::optional< DataClause > mlir::acc::getExitDataClause ( Operation * exitOp )

static

Definition at line 576 of file OpenACCUtilsCG.cpp.

Referenced by computeDataClauseMapFlags().

getFileLineColLoc()

std::optional< FileLineColLoc > mlir::acc::getFileLineColLoc	(	Location	loc,
		bool	errorOnInvalidLocation )

Returns file:line:column location information when available.

Definition at line 32 of file ACCToLLVMUtils.cpp.

References getFileLineColLoc(), and unfuseLoc().

Referenced by createIdent(), and getFileLineColLoc().

getFromDefaultClauseAttrName()

constexpr StringLiteral mlir::acc::getFromDefaultClauseAttrName ( )

staticconstexpr

Definition at line 225 of file OpenACC.h.

getGangDim1ParDimsAttr()

GPUParallelDimsAttr mlir::acc::getGangDim1ParDimsAttr ( MLIRContext * ctx, ACCToGPUMappingPolicy & policy )

inline

Create a gang dim 1 GPUParallelDimsAttr based on the mapping policy.

Definition at line 138 of file OpenACCUtilsCG.h.

References mlir::acc::ACCParMappingPolicy< ParDimAttrT >::gangDim().

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.

getGPUSize()

Value mlir::acc::getGPUSize	(	gpu::Processor	processor,
		gpu::LaunchOp	launch,
		const llvm::DenseMap< gpu::Processor, Value > &	dimensionOps )

Return the launch dimension for processor from launch, or from dimensionOps when launch is null.

Definition at line 68 of file OpenACCUtilsGPU.cpp.

References getGPUSizeFromLaunch().

getGPUSizeFromLaunch()

Value mlir::acc::getGPUSizeFromLaunch ( gpu::LaunchOp launch, gpu::Processor processor )

static

Definition at line 46 of file OpenACCUtilsGPU.cpp.

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

Referenced by getGPUSize().

getGPUThreadId()

Value mlir::acc::getGPUThreadId	(	gpu::Processor	processor,
		gpu::LaunchOp	launch,
		const llvm::DenseMap< gpu::Processor, Value > &	indexOps )

Return the thread/block index for processor from launch, or from indexOps when launch is null.

Definition at line 98 of file OpenACCUtilsGPU.cpp.

References getGPUThreadIdFromLaunch().

getGPUThreadIdFromLaunch()

Value mlir::acc::getGPUThreadIdFromLaunch ( gpu::LaunchOp launch, gpu::Processor processor )

static

Definition at line 76 of file OpenACCUtilsGPU.cpp.

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

Referenced by getGPUThreadId().

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

References ACC_DATA_ENTRY_OPS.

Referenced by computeDataClauseMapFlags(), and getKernelArgMapFlags().

getInternalGlobalName()

std::string mlir::acc::getInternalGlobalName	(	StringRef	kind,
		StringRef	detail )

Returns the name to give a global that the conversion creates to hold detail of kind, such as the name of a variable or a source position.

The dots make a name no identifier of the program can carry, so these globals are reachable by name without walking the symbols of the module. detail keeps letters, digits, _, $ and . and folds every other character to an underscore. Distinct details can therefore collide on the same name; getOrCreateGlobalString is what then keeps the globals apart.

Definition at line 75 of file ACCToLLVMUtils.cpp.

Referenced by mlir::createACCMappedObjectName(), createIdent(), emitHostFallbackNotImplemented(), and emitKernelCall().

getKernelArgumentFieldIndex()

int64_t mlir::acc::getKernelArgumentFieldIndex ( KernelArgumentField field )

inline

The index an insert of field addresses.

Definition at line 83 of file OpenACCRuntimeUtils.h.

Referenced by createKernelArgumentStruct().

getKernelArgumentType()

LLVM::LLVMStructType mlir::acc::getKernelArgumentType

(

MLIRContext *

ctx

)

The LLVM type of the argument struct passed to __tgt_acc_kernel.

Definition at line 96 of file OpenACCRuntimeUtils.cpp.

Referenced by createKernelArgumentStruct().

getMapElementSize()

std::optional< int64_t > mlir::acc::getMapElementSize

(

Operation *

mapEntryOp

)

Returns element size in bytes from acc.map_info, if present.

Definition at line 520 of file OpenACCUtilsCG.cpp.

Referenced by getElementSize().

getMapFlags()

std::optional< MapFlags > mlir::acc::getMapFlags

(

Operation *

mapEntryOp

)

Returns offload map-type flags from acc.map_info, if present.

Definition at line 533 of file OpenACCUtilsCG.cpp.

Referenced by hasCopyOutSibling(), and isDevicePtrMapping().

getMappingExitLoc() [1/2]

std::optional< Location > mlir::acc::getMappingExitLoc ( Value entryResult )

static

Definition at line 560 of file OpenACCUtilsCG.cpp.

References findCorrespondingDataExit(), and mlir::Value::getDefiningOp().

getMappingExitLoc() [2/2]

std::optional< Location > mlir::acc::getMappingExitLoc

(

ValueRange

dataClauseOperands

)

Returns where the mappings of dataClauseOperands end, taken from the first of them that says.

This is where a structured construct tears its mappings down, which is past the end of its region and therefore not a position the construct itself carries. A mapping that is never closed has none.

Definition at line 569 of file OpenACCUtilsCG.cpp.

References getMappingExitLoc().

Referenced by getMappingExitLoc().

getMapSize()

Value mlir::acc::getMapSize

(

Operation *

mapEntryOp

)

Returns the optional size operand from acc.map_info, or null.

Definition at line 527 of file OpenACCUtilsCG.cpp.

getMutableDataOperands()

mlir::MutableOperandRange mlir::acc::getMutableDataOperands

(

mlir::Operation *

accOp

)

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

Definition at line 5564 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

getNoDeviceDescriptorAttrName()

constexpr StringLiteral mlir::acc::getNoDeviceDescriptorAttrName ( )

staticconstexpr

Name for a unit attribute on a data entry operation stating that nothing in the region reads the descriptor its variable is reached through, so the mapping supplies the bounds but needs no device copy of the descriptor.

Definition at line 232 of file OpenACC.h.

Referenced by computeDataClauseMapFlags().

getOrCreateGlobalString()

Value mlir::acc::getOrCreateGlobalString	(	Location	loc,
		OpBuilder &	builder,
		StringRef	name,
		StringRef	value,
		Region &	globalSymbolRegion,
		SymbolTable *	symbolTable = nullptr )

Creates or reuses a null-terminated string global in globalSymbolRegion.

With symbolTable, a global already going by name is reused when it holds value, and a name that two different values arrive under gets a suffix to tell the globals apart. Without it the global is created under name as given, which only a caller whose names are unique by construction can ask for.

Definition at line 129 of file ACCToLLVMUtils.cpp.

References mlir::Builder::getContext(), mlir::Builder::getI64IntegerAttr(), mlir::Builder::getI64Type(), and getOrCreateGlobalStringOp().

Referenced by mlir::createACCMappedObjectName(), emitHostFallbackNotImplemented(), and emitKernelCall().

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

getPairedDataExitOps()

SmallVector< Operation * > mlir::acc::getPairedDataExitOps

(

Value

entryResult

)

Returns the data exit operations paired with the data entry result entryResult, which take it as their accVar.

Definition at line 539 of file OpenACCUtilsCG.cpp.

References ACC_DATA_EXIT_OPS, and mlir::Value::getUses().

Referenced by findCorrespondingDataExit().

getParDimsAttr()

GPUParallelDimsAttr mlir::acc::getParDimsAttr

(

Operation *

op

)

Obtain the parallel dimensions carried by op, if any.

Definition at line 170 of file OpenACCUtilsCG.cpp.

References ACC_OP_WITH_PAR_DIMS_LIST, and mlir::Operation::getDiscardableAttr().

Referenced by collectPrivateLocalParDims(), copyParDimsAttr(), getReductionCombineParDims(), getReductionCombineParDims(), hasParDimsAttr(), and hasSeqParDims().

getParentFunctionName() [1/3]

StringRef mlir::acc::getParentFunctionName

(

Operation *

op

)

Returns the symbol name of the function op belongs to, or of op itself when it is a function.

Definition at line 52 of file ACCToLLVMUtils.cpp.

References mlir::Operation::getParentOfType().

Referenced by mlir::emitACCDataRuntimeArgs(), emitGetDevicePtrCall(), emitKernelCall(), emitWaitCall(), getParentFunctionName(), and getParentFunctionName().

getParentFunctionName() [2/3]

StringRef mlir::acc::getParentFunctionName

(

Value

value

)

Returns the enclosing function symbol name for value's defining op.

Definition at line 61 of file ACCToLLVMUtils.cpp.

References mlir::Value::getDefiningOp(), and getParentFunctionName().

getParentFunctionName() [3/3]

StringRef mlir::acc::getParentFunctionName

(

ValueRange

values

)

Returns the first non-empty enclosing function name from values.

Definition at line 67 of file ACCToLLVMUtils.cpp.

References getParentFunctionName().

getPrivateBaseMemRefType()

MemRefType mlir::acc::getPrivateBaseMemRefType	(	Type	baseTy,
		ModuleOp	module )

Returns the ranked MemRef type used to allocate privatized storage.

baseTy is the baseTy parameter of acc.private_type (the privatized variable's type).

Definition at line 312 of file OpenACCUtilsCG.cpp.

Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().

getPrivateLocalSharedMemoryUpperBoundBytes()

std::optional< int64_t > mlir::acc::getPrivateLocalSharedMemoryUpperBoundBytes	(	PrivateLocalOp	privateLocal,
		ComputeRegionOp	computeRegion,
		ModuleOp	module,
		const ACCToGPUMappingPolicy &	policy,
		OpenACCSupport *	support = nullptr )

Upper-bound byte size for a shared-memory private_local candidate, or std::nullopt when not eligible or not statically computable.

Definition at line 453 of file OpenACCUtilsCG.cpp.

References collectPrivateLocalParDims(), getPrivateBaseMemRefType(), getTypeSizeAndAlignment(), getWorkerPrivateSharedMemoryNumCopies(), isPrivateLocalSharedMemoryCandidate(), and mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isWorker().

getPrivatizeOp()

PrivatizeOp mlir::acc::getPrivatizeOp	(	PrivateLocalOp	privateLocal,
		ComputeRegionOp	computeRegion )

Resolve the acc.privatize operation associated with a private local.

Definition at line 293 of file OpenACCUtilsCG.cpp.

References mlir::Value::getDefiningOp().

Referenced by isPrivateLocalSharedMemoryCandidate().

getRecipe()

mlir::SymbolRefAttr mlir::acc::getRecipe

(

mlir::Operation *

accOp

)

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

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

References mlir::Type::print().

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

getReductionCombineParDims() [1/2]

SmallVector< GPUParallelDimAttr > mlir::acc::getReductionCombineParDims

(

ReductionCombineOp

op

)

Returns the parallel dimensions that participate in op's combine step.

Used when lowering reductions to determine which GPU parallelism levels must be synchronized before combining partial results.

Definition at line 25 of file OpenACCUtilsReduction.cpp.

References getParDimsAttr().

Referenced by collectPrivateLocalParDims().

getReductionCombineParDims() [2/2]

SmallVector< GPUParallelDimAttr > mlir::acc::getReductionCombineParDims

(

ReductionCombineRegionOp

op

)

Returns the parallel dimensions that participate in op's combine step.

Prefers dimensions from an acc.reduction_accumulate user of the source variable; otherwise falls back to op's acc.par_dims attribute.

Definition at line 33 of file OpenACCUtilsReduction.cpp.

References getParDimsAttr(), and mlir::Operation::getUsers().

getReductionIdentityValueAttr()

TypedAttr mlir::acc::getReductionIdentityValueAttr ( arith::AtomicRMWKind kind, Type type, OpBuilder & builder, Location loc, bool useOnlyFiniteValue )

static

Definition at line 141 of file OpenACCUtilsReduction.cpp.

References mlir::emitError(), mlir::arith::getIdentityValueAttr(), and mlir::Type::isIntOrIndexOrFloat().

Referenced by createIdentityValue().

getRoutineInfoAttrName()

constexpr StringLiteral mlir::acc::getRoutineInfoAttrName ( )

staticconstexpr

Definition at line 204 of file OpenACC.h.

Referenced by createACCBindRoutine(), mlir::acc::impl::ACCCGToGPUBase< DerivedT >::getDependentDialects(), isAccRoutine(), and isValidSymbolUse().

getRuntimeFunctionName()

StringRef mlir::acc::getRuntimeFunctionName

(

RuntimeFunction

fn

)

Returns the default runtime symbol name for fn.

Definition at line 25 of file OpenACCRuntimeUtils.cpp.

Referenced by mlir::acc::ACCRuntimeCallConfig::getName().

getRuntimeFunctionType()

LLVM::LLVMFunctionType mlir::acc::getRuntimeFunctionType	(	MLIRContext *	ctx,
		RuntimeFunction	fn )

Builds the LLVM function type for fn in ctx.

Definition at line 35 of file OpenACCRuntimeUtils.cpp.

Referenced by createRuntimeCall().

getSeqParDimsAttr()

GPUParallelDimsAttr mlir::acc::getSeqParDimsAttr ( MLIRContext * ctx, ACCToGPUMappingPolicy & policy )

inline

Create a sequential GPUParallelDimsAttr based on the mapping policy.

Definition at line 144 of file OpenACCUtilsCG.h.

References mlir::acc::ACCParMappingPolicy< ParDimAttrT >::seqDim().

getSpecializedRoutineAttrName()

constexpr StringLiteral mlir::acc::getSpecializedRoutineAttrName ( )

staticconstexpr

Definition at line 208 of file OpenACC.h.

Referenced by mlir::acc::impl::ACCCGToGPUBase< DerivedT >::getDependentDialects(), isSpecializedAccRoutine(), and isValidSymbolUse().

getSyntheticFlag()

bool mlir::acc::getSyntheticFlag

(

mlir::Operation *

accDataClauseOp

)

Used to find out whether the implementation created the data operation for its own bookkeeping, instead of to represent a data attribute of the program being compiled.

Such an operation has no counterpart in that program, so it must not be described to the user. Returns false if not a data clause operation or if it is a data clause operation without the synthetic flag.

Definition at line 5546 of file OpenACC.cpp.

References ACC_DATA_CLAUSE_OPS.

getTypeCategory()

mlir::acc::VariableTypeCategory mlir::acc::getTypeCategory

(

mlir::Value

var

)

Get the type category of an OpenACC variable.

Definition at line 100 of file OpenACCUtils.cpp.

References mlir::Value::getType().

getTypeSizeAndAlignment() [1/2]

std::optional< TypeSizeAndAlignment > mlir::acc::getTypeSizeAndAlignment	(	Type	ty,
		ModuleOp	module,
		const DataLayout &	dl,
		OpenACCSupport *	support = nullptr,
		Value	var = {} )

Returns the size and ABI alignment in bytes.

For aggregate structures and arrays, padding between members or elements is not taken into account. The result is a close estimate suitable for early OpenACC layout decisions, but not a complete ABI guarantee. For final size computations, use LLVM materialized types.

ty itself is sized dialect-agnostically; when support is provided it sizes aggregate element types, so that nested dialect types are handled. Callers that hold an OpenACCSupport should therefore ask it directly - OpenACCSupport::getTypeSizeAndAlignment covers dialect types and falls back to this utility - and call this utility directly only for a type this utility is expected to know.

When var is provided, MappableType sizes the mapped object rather than the type's storage alone.

Returns std::nullopt when the size is not statically computable or the type is not supported.

Definition at line 32 of file OpenACCUtilsType.cpp.

References mlir::Type::getContext(), mlir::DataLayout::getTypeABIAlignment(), mlir::DataLayout::getTypeSize(), getTypeSizeAndAlignmentHelper(), and mlir::Type::isIntOrIndexOrFloat().

Referenced by computeMapInfoSizeBytes(), getElementSize(), getKernelArgDescriptor(), getPrivateLocalSharedMemoryUpperBoundBytes(), mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::getTypeSizeAndAlignment(), getTypeSizeAndAlignment(), mlir::acc::OpenACCSupport::getTypeSizeAndAlignment(), and getTypeSizeAndAlignmentHelper().

getTypeSizeAndAlignment() [2/2]

std::optional< TypeSizeAndAlignment > mlir::acc::getTypeSizeAndAlignment	(	Type	ty,
		ModuleOp	module,
		OpenACCSupport *	support = nullptr,
		Value	var = {} )

Same as above, obtaining dl from module via getDataLayout.

Definition at line 100 of file OpenACCUtilsType.cpp.

References getDataLayout(), and getTypeSizeAndAlignment().

getTypeSizeAndAlignmentHelper()

std::optional< TypeSizeAndAlignment > mlir::acc::getTypeSizeAndAlignmentHelper ( Type ty, ModuleOp module, const DataLayout & dl, OpenACCSupport * support )

static

Definition at line 24 of file OpenACCUtilsType.cpp.

References getTypeSizeAndAlignment(), and mlir::acc::OpenACCSupport::getTypeSizeAndAlignment().

Referenced by getTypeSizeAndAlignment().

getUnstructuredIndependentLoopAttrName()

constexpr StringLiteral mlir::acc::getUnstructuredIndependentLoopAttrName ( )

staticconstexpr

Name for a unit attribute attached to the region an unstructured independent loop was converted to, recording that the loop could not be represented as a structured loop and is therefore not parallelized.

Definition at line 245 of file OpenACC.h.

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

References ACC_DATA_ENTRY_OPS.

Referenced by checkDeclareOperands(), mlir::createACCArgumentDescriptor(), mlir::emitACCDataRuntimeArgs(), and hasCopyOutSibling().

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, the name of a data clause operation, or the symbol a global is addressed through is found.

Returns empty string if no name is found.

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

Referenced by getVariableName().

getVarNameAttrName()

constexpr StringLiteral mlir::acc::getVarNameAttrName ( )

staticconstexpr

Definition at line 249 of file OpenACC.h.

Referenced by getVariableName().

getVarNamePlaceholder()

llvm::StringLiteral mlir::acc::getVarNamePlaceholder

(

)

Returns a placeholder string for use as an acc.var_name attribute value when the actual variable name is not yet known at the point of IR construction.

The placeholder is meant to be replaced with the real name at a later lowering stage. For example, recipe init regions may attach this to ops at recipe-generation time, and ACCRecipeMaterialization will subsequently replace the placeholder with the actual variable name on all marked ops after inlining the recipe into the compute construct.

Definition at line 113 of file OpenACCUtils.cpp.

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

Referenced by hasCopyOutSibling().

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

References ACC_DATA_ENTRY_OPS.

Referenced by hasAttachPoint().

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

References ACC_DATA_ENTRY_OPS.

Referenced by getElementSize().

getWaitClauseValues()

template<typename OpTy>

std::optional< DeviceType > mlir::acc::getWaitClauseValues	(	OpTy	op,
		DeviceType	deviceType,
		SmallVectorImpl< Value > &	waitValues )

Appends to waitValues the queues the wait clause of op names for deviceType, and returns the device type the clause naming them is given for - a clause naming no queue waits for every one of them.

Returns std::nullopt when op gives no such clause for deviceType. The values are the ones the operation holds, so a caller in a conversion has to remap them.

Definition at line 199 of file ACCToLLVMUtils.h.

References getDeviceTypesByPrecedence().

Referenced by emitWaitClause().

getWorkerPrivateSharedMemoryNumCopies()

FailureOr< std::optional< int64_t > > mlir::acc::getWorkerPrivateSharedMemoryNumCopies ( PrivateLocalOp privateLocal, ComputeRegionOp computeRegion, bool isWorkerPrivate, OpenACCSupport * support )

static

Definition at line 363 of file OpenACCUtilsCG.cpp.

References mlir::ValueBoundsConstraintSet::computeConstantBound(), mlir::acc::OpenACCSupport::emitNYI(), and mlir::presburger::EQ.

Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().

hasActiveParDimsAttr()

bool mlir::acc::hasActiveParDimsAttr

(

Operation *

op

)

Return whether op carries active parallel dimensions.

Definition at line 257 of file OpenACCUtilsCG.cpp.

References getActiveParDimsAttr().

hasAttachPoint()

bool mlir::acc::hasAttachPoint

(

Operation *

mapEntryOp

)

Returns true when mapEntryOp carries an attach point (varPtrPtr).

Definition at line 487 of file OpenACCUtilsCG.cpp.

References getDataClause(), and getVarPtrPtr().

Referenced by computePackedMapFlags(), and getKernelDataClauseMapFlags().

hasChunkSizeAttr()

bool mlir::acc::hasChunkSizeAttr

(

Operation *

op

)

Return whether op carries an acc.chunk_size attribute.

Definition at line 230 of file OpenACCUtilsCG.cpp.

References getChunkSizeAttr().

hasCopyOutSibling()

bool mlir::acc::hasCopyOutSibling

(

Operation *

entryOp

)

True when another data clause of the same construct maps the same variable with a copy-back and no copy-in.

One entry operation is emitted per clause, so a construct such as create(a) copyout(a) or copyin(a) copyout(a) maps a twice. The runtime keeps one mapping per variable and acts on it once, which makes the copy-back depend on whichever entry it processes: the entry that has no copy-back of its own must therefore carry from as well.

Definition at line 582 of file OpenACCUtilsCG.cpp.

References getDataClause(), getMapFlags(), mlir::Operation::getResult(), mlir::Value::getUsers(), getVar(), and getVarPtr().

Referenced by computeDataClauseMapFlags().

hasGPUBlockRedundantAttr()

bool mlir::acc::hasGPUBlockRedundantAttr

(

Operation *

op

)

Return whether op is marked with the acc.gpu_block_redundant attribute, i.e.

it executes redundantly across all thread blocks. Such an op must not be assigned block/grid-level work-sharing; only thread-level parallelism may apply, and its enclosing block dimensions are treated as active (not predicated).

Definition at line 216 of file OpenACCUtilsCG.cpp.

References mlir::Operation::hasDiscardableAttrOfType().

hasParDimsAttr()

bool mlir::acc::hasParDimsAttr

(

Operation *

op

)

Return whether op carries parallel dimensions.

Definition at line 185 of file OpenACCUtilsCG.cpp.

References getParDimsAttr().

Referenced by copyParDimsAttr(), setParDimsAttr(), and updateParDimsAttr().

hasSeqParDims()

bool mlir::acc::hasSeqParDims

(

Operation *

op

)

Return whether op carries sequential parallel dimensions.

Definition at line 187 of file OpenACCUtilsCG.cpp.

References getParDimsAttr().

hasWaitDevnum()

template<typename OpTy>

bool mlir::acc::hasWaitDevnum	(	OpTy	op,
		DeviceType	deviceType )

Returns whether the wait clause op gives for deviceType carries a devnum modifier, which selects the device the queues belong to.

Only the clause given for that device type is asked, so a caller passes the device type getWaitClauseValues took the queues from.

Definition at line 233 of file ACCToLLVMUtils.h.

Referenced by emitWaitClause(), getWaitDevnumValue(), and getWaitValuesWithoutDevnum().

insertParDim()

void mlir::acc::insertParDim	(	llvm::SmallVector< GPUParallelDimAttr > &	parDims,
		GPUParallelDimAttr	parDim )

Insert parDim into parDims while preserving dimension ordering.

If the dimension is already present, this is a no-op.

Definition at line 152 of file OpenACCUtilsCG.cpp.

References findParDim().

Referenced by collectPrivateLocalParDims().

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

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

Referenced by isValidValueUse().

isDataEntryOp()

bool mlir::acc::isDataEntryOp

(

mlir::Operation *

op

)

Used to find out whether op is a data entry operation or an acc.map_info.

op may be null.

Definition at line 5532 of file OpenACC.cpp.

Referenced by mlir::emitACCDataRuntimeArgs(), emitKernelRuntimeArgs(), getKernelArgDescriptor(), getKernelArgMapFlags(), and isDataEntryInsideEnvironment().

isDeviceAccessibleValue()

bool mlir::acc::isDeviceAccessibleValue

(

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, and whether the defining operation carries acc.declare with the deviceptr clause.

Parameters

val	The value to check

Returns

true if the value is device data, false otherwise

Definition at line 280 of file OpenACCUtils.cpp.

References getDeclareAttrName(), mlir::Value::getDefiningOp(), mlir::Operation::getDiscardableAttrOfType(), mlir::Value::getType(), isDeviceAccessibleValue(), and mlir::SymbolTable::lookupNearestSymbolFrom().

Referenced by isDeviceAccessibleValue(), and isValidValueUse().

isFloatOrComplexType()

bool mlir::acc::isFloatOrComplexType ( Type ty )

static

Definition at line 20 of file OpenACCUtilsReduction.cpp.

Referenced by translateACCReductionOperator().

isGangWorkerVectorAllOne()

template<typename ComputeOpT>

bool mlir::acc::isGangWorkerVectorAllOne

(

ComputeOpT

op

)

Definition at line 276 of file OpenACC.h.

References mlir::isConstantIntValue().

isInDeviceMemoryValue()

bool mlir::acc::isInDeviceMemoryValue

(

mlir::Value

val

)

Check if a value is backed by memory that is residing in the current device, and therefore requires no runtime mapping or attach.

This is stricter than isDeviceAccessibleValue: isDeviceAccessibleValue answers device accessibility (whether the current device can reach the storage, regardless of where it physically resides), whereas this answers device residence (whether the storage physically lives in device memory). Storage that is accessible but physically shared with the host may migrate on demand, so it is not in device memory: it must still be mapped so the runtime can attach rather than be treated as already resident.

Parameters

val	The value to check

Returns

true if the value is in device memory, false otherwise

Definition at line 324 of file OpenACCUtils.cpp.

References getDeclareAttrName(), mlir::Value::getDefiningOp(), mlir::Operation::getDiscardableAttrOfType(), mlir::Value::getType(), isInDeviceMemoryValue(), and mlir::SymbolTable::lookupNearestSymbolFrom().

Referenced by isInDeviceMemoryValue().

isInOffTargetBranch()

bool mlir::acc::isInOffTargetBranch	(	mlir::Operation *	op,
		llvm::ArrayRef< int64_t >	deviceTypes )

Returns true if op is in a branch that does not execute on the target.

deviceTypes lists the acc_device_t values for which acc.on_device is true. That branch is the then of a constant acc.on_device outside the list, or the else of one in the list. A non-constant operand is ignored. An empty list returns false.

Definition at line 472 of file OpenACCUtils.cpp.

References ACC_COMPUTE_CONSTRUCT_OPS, mlir::getConstantIntValue(), mlir::Value::getDefiningOp(), getIfCondition(), mlir::Operation::getParentOp(), and mlir::Operation::getParentRegion().

isInsideACCSpecializedRoutine()

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

static

Definition at line 394 of file OpenACCUtilsCG.cpp.

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

Referenced by isPrivateLocalSharedMemoryCandidate().

isMappableType()

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

inline

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

Definition at line 191 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 68 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 73 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 185 of file OpenACC.h.

isPrivateLocalSharedMemoryCandidate()

FailureOr< bool > mlir::acc::isPrivateLocalSharedMemoryCandidate	(	PrivateLocalOp	privateLocal,
		ComputeRegionOp	computeRegion,
		ModuleOp	module,
		const ACCToGPUMappingPolicy &	policy,
		OpenACCSupport *	support = nullptr )

True when privateLocal may be placed in shared memory.

Definition at line 399 of file OpenACCUtilsCG.cpp.

References collectPrivateLocalParDims(), getPrivateBaseMemRefType(), getPrivatizeOp(), getWorkerPrivateSharedMemoryNumCopies(), mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isGang(), isInsideACCSpecializedRoutine(), isThreadXPrivatize(), mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isVector(), and mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isWorker().

Referenced by getPrivateLocalSharedMemoryUpperBoundBytes().

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

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

Referenced by mlir::acc::impl::ACCCGToGPUBase< DerivedT >::getDependentDialects(), isInsideACCSpecializedRoutine(), and isValidValueUse().

isThreadXPrivatize()

bool mlir::acc::isThreadXPrivatize ( PrivatizeOp privatize )

static

Definition at line 305 of file OpenACCUtilsCG.cpp.

Referenced by isPrivateLocalSharedMemoryCandidate().

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

References getDeclareAttrName(), getRoutineInfoAttrName(), getSpecializedRoutineAttrName(), mlir::Operation::hasDiscardableAttr(), mlir::SymbolTable::lookupNearestSymbolFrom(), lookupSymbolInGPUModules(), 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 368 of file OpenACCUtils.cpp.

References mlir::Value::getDefiningOp(), mlir::Region::getParentOp(), mlir::Value::getType(), isAccRoutine(), isDeviceAccessibleValue(), mlir::Type::isIntOrIndexOrFloat(), isOnlyUsedByPrivateClauses(), and isSpecializedAccRoutine().

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

normalizeIVUses()

void mlir::acc::normalizeIVUses	(	OpBuilder &	b,
		Location	loc,
		Value	iv,
		Value	origLB,
		Value	origStep )

Normalize IV uses after converting to normalized loop form (lb=0, step=1).

Normalize IV uses after converting to normalized loop form.

Replaces uses of iv with iv * origStep + origLB.

For normalized loops (lb=0, step=1), we need to denormalize the IV: original_iv = new_iv * orig_step + orig_lb

Definition at line 68 of file OpenACCUtilsLoop.cpp.

References b, mlir::Value::getDefiningOp(), mlir::getValueOrCreateCastToIndexLike(), and mlir::Value::replaceAllUsesExcept().

Referenced by convertACCLoopToSCFFor(), and convertACCLoopToSCFParallel().

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 408 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 371 of file ACCSpecializeForHost.cpp.

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

populateACCSpecializeForDevicePatterns()

void mlir::acc::populateACCSpecializeForDevicePatterns	(	RewritePatternSet &	patterns,
		const TypesForDevice &	theDeviceTypes )

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. theDeviceTypes hold the device types for which acc_on_device should return true.

Definition at line 161 of file ACCSpecializeForDevice.cpp.

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

populateDialectIdentityDeviceTypeMapping()

void mlir::acc::populateDialectIdentityDeviceTypeMapping

(

ACCRuntimeCallConfig &

config

)

Install a device-type mapping that uses OpenACC dialect enum ordinals as the runtime encoding.

This is only correct when the target runtime happens to use the same numbering; runtimes with a different ABI must install their own mapping via setDeviceTypeRuntimeValue.

Definition at line 227 of file OpenACCRuntimeUtils.cpp.

References mlir::acc::ACCRuntimeCallConfig::setDeviceTypeRuntimeValue().

Referenced by mlir::acc::ACCRuntimeCallConfig::ACCRuntimeCallConfig().

populateDialectIdentityMapFlagsMapping()

void mlir::acc::populateDialectIdentityMapFlagsMapping

(

ACCRuntimeCallConfig &

config

)

Install a map-flag mapping that uses the OpenACC dialect bit positions as the runtime encoding, with the same caveat as populateDialectIdentityDeviceTypeMapping.

Definition at line 234 of file OpenACCRuntimeUtils.cpp.

References mlir::acc::ACCRuntimeCallConfig::setMapFlagRuntimeValue().

Referenced by mlir::acc::ACCRuntimeCallConfig::ACCRuntimeCallConfig().

populateSourceExtents()

void mlir::acc::populateSourceExtents	(	ValueRange	bounds,
		ArrayRef< int64_t >	shape,
		OpBuilder &	builder )

Record known extents of the source array on bounds that may describe a section.

Dynamic / unknown extents are left unset.

Definition at line 812 of file OpenACCUtilsCG.cpp.

References mlir::arith::ConstantIndexOp::create(), and mlir::OpBuilder::setInsertionPoint().

registerACCBindRoutine()

void mlir::acc::registerACCBindRoutine ( )

inline

Definition at line 2200 of file Passes.h.

registerACCBindRoutinePass()

void mlir::acc::registerACCBindRoutinePass ( )

inline

Definition at line 2207 of file Passes.h.

registerACCCGToGPU()

void mlir::acc::registerACCCGToGPU ( )

inline

Definition at line 2221 of file Passes.h.

registerACCCGToGPUPass()

void mlir::acc::registerACCCGToGPUPass ( )

inline

Definition at line 2228 of file Passes.h.

registerACCComputeLowering()

void mlir::acc::registerACCComputeLowering ( )

inline

Definition at line 2242 of file Passes.h.

registerACCComputeLoweringPass()

void mlir::acc::registerACCComputeLoweringPass ( )

inline

Definition at line 2249 of file Passes.h.

registerACCDeclareCtorDtorConversion()

void mlir::acc::registerACCDeclareCtorDtorConversion ( )

inline

Definition at line 2263 of file Passes.h.

registerACCDeclareCtorDtorConversionPass()

void mlir::acc::registerACCDeclareCtorDtorConversionPass ( )

inline

Definition at line 2270 of file Passes.h.

registerACCDeclareGPUModuleInsertion()

void mlir::acc::registerACCDeclareGPUModuleInsertion ( )

inline

Definition at line 2284 of file Passes.h.

registerACCDeclareGPUModuleInsertionPass()

void mlir::acc::registerACCDeclareGPUModuleInsertionPass ( )

inline

Definition at line 2291 of file Passes.h.

registerACCEmitRemarksData()

void mlir::acc::registerACCEmitRemarksData ( )

inline

Definition at line 2305 of file Passes.h.

registerACCEmitRemarksDataPass()

void mlir::acc::registerACCEmitRemarksDataPass ( )

inline

Definition at line 2312 of file Passes.h.

registerACCEmitRemarksLoop()

void mlir::acc::registerACCEmitRemarksLoop ( )

inline

Definition at line 2326 of file Passes.h.

registerACCEmitRemarksLoopPass()

void mlir::acc::registerACCEmitRemarksLoopPass ( )

inline

Definition at line 2333 of file Passes.h.

registerACCEmitRemarksPrivate()

void mlir::acc::registerACCEmitRemarksPrivate ( )

inline

Definition at line 2347 of file Passes.h.

registerACCEmitRemarksPrivatePass()

void mlir::acc::registerACCEmitRemarksPrivatePass ( )

inline

Definition at line 2354 of file Passes.h.

registerACCIfClauseLowering()

void mlir::acc::registerACCIfClauseLowering ( )

inline

Definition at line 2368 of file Passes.h.

registerACCIfClauseLoweringPass()

void mlir::acc::registerACCIfClauseLoweringPass ( )

inline

Definition at line 2375 of file Passes.h.

registerACCImplicitData()

void mlir::acc::registerACCImplicitData ( )

inline

Definition at line 2389 of file Passes.h.

registerACCImplicitDataPass()

void mlir::acc::registerACCImplicitDataPass ( )

inline

Definition at line 2396 of file Passes.h.

registerACCImplicitDeclare()

void mlir::acc::registerACCImplicitDeclare ( )

inline

Definition at line 2410 of file Passes.h.

registerACCImplicitDeclarePass()

void mlir::acc::registerACCImplicitDeclarePass ( )

inline

Definition at line 2417 of file Passes.h.

registerACCImplicitRoutine()

void mlir::acc::registerACCImplicitRoutine ( )

inline

Definition at line 2431 of file Passes.h.

registerACCImplicitRoutinePass()

void mlir::acc::registerACCImplicitRoutinePass ( )

inline

Definition at line 2438 of file Passes.h.

registerACCLegalizeSerial()

void mlir::acc::registerACCLegalizeSerial ( )

inline

Definition at line 2452 of file Passes.h.

registerACCLegalizeSerialPass()

void mlir::acc::registerACCLegalizeSerialPass ( )

inline

Definition at line 2459 of file Passes.h.

registerACCLoopTiling()

void mlir::acc::registerACCLoopTiling ( )

inline

Definition at line 2473 of file Passes.h.

registerACCLoopTilingPass()

void mlir::acc::registerACCLoopTilingPass ( )

inline

Definition at line 2480 of file Passes.h.

registerACCMaterializeRoutineBindTargets()

void mlir::acc::registerACCMaterializeRoutineBindTargets ( )

inline

Definition at line 2494 of file Passes.h.

registerACCMaterializeRoutineBindTargetsPass()

void mlir::acc::registerACCMaterializeRoutineBindTargetsPass ( )

inline

Definition at line 2501 of file Passes.h.

registerACCRecipeMaterialization()

void mlir::acc::registerACCRecipeMaterialization ( )

inline

Definition at line 2515 of file Passes.h.

registerACCRecipeMaterializationPass()

void mlir::acc::registerACCRecipeMaterializationPass ( )

inline

Definition at line 2522 of file Passes.h.

registerACCRoutineLowering()

void mlir::acc::registerACCRoutineLowering ( )

inline

Definition at line 2536 of file Passes.h.

registerACCRoutineLoweringPass()

void mlir::acc::registerACCRoutineLoweringPass ( )

inline

Definition at line 2543 of file Passes.h.

registerACCRoutineToGPUFunc()

void mlir::acc::registerACCRoutineToGPUFunc ( )

inline

Definition at line 2557 of file Passes.h.

registerACCRoutineToGPUFuncPass()

void mlir::acc::registerACCRoutineToGPUFuncPass ( )

inline

Definition at line 2564 of file Passes.h.

registerACCSpecializeForDevice()

void mlir::acc::registerACCSpecializeForDevice ( )

inline

Definition at line 2578 of file Passes.h.

registerACCSpecializeForDevicePass()

void mlir::acc::registerACCSpecializeForDevicePass ( )

inline

Definition at line 2585 of file Passes.h.

registerACCSpecializeForHost()

void mlir::acc::registerACCSpecializeForHost ( )

inline

Definition at line 2599 of file Passes.h.

registerACCSpecializeForHostPass()

void mlir::acc::registerACCSpecializeForHostPass ( )

inline

Definition at line 2606 of file Passes.h.

registerLegalizeDataValuesInRegion()

void mlir::acc::registerLegalizeDataValuesInRegion ( )

inline

Definition at line 2620 of file Passes.h.

registerLegalizeDataValuesInRegionPass()

void mlir::acc::registerLegalizeDataValuesInRegionPass ( )

inline

Definition at line 2627 of file Passes.h.

registerOffloadLiveInValueCanonicalization()

void mlir::acc::registerOffloadLiveInValueCanonicalization ( )

inline

Definition at line 2641 of file Passes.h.

registerOffloadLiveInValueCanonicalizationPass()

void mlir::acc::registerOffloadLiveInValueCanonicalizationPass ( )

inline

Definition at line 2648 of file Passes.h.

registerOffloadTargetVerifier()

void mlir::acc::registerOffloadTargetVerifier ( )

inline

Definition at line 2662 of file Passes.h.

registerOffloadTargetVerifierPass()

void mlir::acc::registerOffloadTargetVerifierPass ( )

inline

Definition at line 2669 of file Passes.h.

registerOpenACCPasses()

void mlir::acc::registerOpenACCPasses ( )

inline

Definition at line 2683 of file Passes.h.

Referenced by mlir::registerAllPasses().

remapOrSelf()

Value mlir::acc::remapOrSelf	(	Value	value,
		ConversionPatternRewriter &	rewriter )

Returns the value value was converted to, or value itself when the conversion left it as it was.

Returns a null value for a null value.

Definition at line 231 of file ACCToLLVMUtils.cpp.

Referenced by mlir::createACCArgumentDescriptor(), and getBoundValue().

removeParDim()

void mlir::acc::removeParDim	(	llvm::SmallVector< GPUParallelDimAttr > &	parDims,
		GPUParallelDimAttr	parDim )

Remove parDim from parDims if present.

Definition at line 159 of file OpenACCUtilsCG.cpp.

References findParDim().

setActiveParDimsAttr() [1/2]

void mlir::acc::setActiveParDimsAttr	(	Operation *	op,
		ActiveParDimsAttr	attr )

Set active parallel dimensions on op.

Definition at line 261 of file OpenACCUtilsCG.cpp.

References mlir::Operation::setDiscardableAttr().

Referenced by setActiveParDimsAttr().

setActiveParDimsAttr() [2/2]

void mlir::acc::setActiveParDimsAttr	(	Operation *	op,
		ArrayRef< GPUParallelDimAttr >	dims )

Set active parallel dimensions on op from a dimension list.

Definition at line 265 of file OpenACCUtilsCG.cpp.

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

setChunkSizeAttr() [1/2]

void mlir::acc::setChunkSizeAttr	(	Operation *	op,
		ChunkSizeAttr	attr )

Set the acc.chunk_size attribute on op.

Definition at line 232 of file OpenACCUtilsCG.cpp.

References mlir::Operation::setDiscardableAttr().

Referenced by setChunkSizeAttr().

setChunkSizeAttr() [2/2]

void mlir::acc::setChunkSizeAttr	(	Operation *	op,
		int64_t	chunkSize )

Set the acc.chunk_size attribute on op from a constant size.

Definition at line 236 of file OpenACCUtilsCG.cpp.

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

setCollapseCountAttr()

void mlir::acc::setCollapseCountAttr	(	Operation *	op,
		uint64_t	count )

Record on a collapsed loop how many original loops were folded into it.

Definition at line 352 of file OpenACCUtilsLoop.cpp.

References getCollapseCountAttrName(), mlir::Operation::getContext(), and mlir::Operation::setDiscardableAttr().

Referenced by convertACCLoopToSCFFor(), and convertACCLoopToSCFParallel().

setGPUBlockRedundantAttr()

void mlir::acc::setGPUBlockRedundantAttr

(

Operation *

op

)

Mark op with the acc.gpu_block_redundant attribute.

Definition at line 221 of file OpenACCUtilsCG.cpp.

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

setParDimsAttr()

void mlir::acc::setParDimsAttr	(	Operation *	op,
		GPUParallelDimsAttr	attr )

Set parallel dimensions on op.

Definition at line 193 of file OpenACCUtilsCG.cpp.

References ACC_OP_WITH_PAR_DIMS_LIST, hasParDimsAttr(), and mlir::Operation::setDiscardableAttr().

Referenced by copyParDimsAttr().

spliceConstructRegion()

void mlir::acc::spliceConstructRegion	(	Operation *	op,
		Region &	region,
		RewriterBase &	rewriter )

Splices region, the body of a structured construct, into the block holding op, so that the construct itself can be erased.

The code that follows op is branched to where the region ends, whether it ends in an acc.terminator or, as a region holding no terminator does, at the end of its blocks.

Definition at line 391 of file ACCToLLVMUtils.cpp.

References mlir::RewriterBase::eraseOp(), mlir::Operation::getBlock(), mlir::Region::getBlocks(), mlir::Operation::getLoc(), mlir::RewriterBase::inlineRegionBefore(), mlir::OpBuilder::setInsertionPoint(), mlir::OpBuilder::setInsertionPointToEnd(), and mlir::RewriterBase::splitBlock().

sumExistingSharedMemoryBytes()

int64_t mlir::acc::sumExistingSharedMemoryBytes

(

Region &

region

)

Sum aligned static_upper_bound_bytes for all acc.gpu_shared_memory in region.

Definition at line 284 of file OpenACCUtilsCG.cpp.

References mlir::acc::SharedMemoryBudget::alignOffset(), and mlir::Region::walk().

tileACCLoops()

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

Tile a single fused acc.loop that carries all associated induction variables (one IV per tile dimension).

This produces exactly two multi-IV loops, each carrying all of the tiled induction variables:

• a "tile group" loop that steps over tiles: each step is the original step multiplied by the tile size. It keeps the original loop's gang attribute.
• an "element group" loop that walks the iterations inside one tile: it keeps the original step, its lower bound is the current tile's starting index, and its upper bound is clamped to min(original upper bound, tile start + tile extent). It keeps the original loop's vector (or worker) attribute.

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 (each group is one multi-IV loop over all tiled IVs):

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

The element bounds above assume ascending steps. For a descending dimension, the tile's far edge lies below the original bound, so the clamp uses max instead of min.

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

Parameters

tileLoop	The fused loop to tile.
tileSizes	The tile sizes for each tiled 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 tile group loop that is modified in place.

Definition at line 155 of file OpenACCUtilsTiling.cpp.

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

translateACCReductionOperator()

std::optional< arith::AtomicRMWKind > mlir::acc::translateACCReductionOperator	(	ReductionOperator	redOp,
		Type	type )

Maps an acc reduction operator to the arith atomic RMW kind for type.

Returns std::nullopt when redOp is not supported for type.

Definition at line 75 of file OpenACCUtilsReduction.cpp.

References mlir::Type::isFloat(), isFloatOrComplexType(), mlir::Type::isInteger(), and mlir::Type::isUnsignedInteger().

translateAtomicRMWKind()

ReductionOperator mlir::acc::translateAtomicRMWKind

(

arith::AtomicRMWKind

kind

)

Maps an arith atomic RMW kind to the corresponding acc reduction operator.

Definition at line 44 of file OpenACCUtilsReduction.cpp.

unfuseLoc()

Location mlir::acc::unfuseLoc

(

Location

loc

)

Unfuses fused locations, returning the last sub-location.

Definition at line 25 of file ACCToLLVMUtils.cpp.

Referenced by getFileLineColLoc().

updateParDimsAttr()

void mlir::acc::updateParDimsAttr	(	Operation *	op,
		GPUParallelDimsAttr	attr )

Update parallel dimensions on op.

Definition at line 203 of file OpenACCUtilsCG.cpp.

References ACC_OP_WITH_PAR_DIMS_LIST, hasParDimsAttr(), and mlir::Operation::setDiscardableAttr().

wrapMultiBlockRegionWithSCFExecuteRegion()

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

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. Terminators with no successors (i.e., region exit points) are replaced 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.

Returns

The created scf.execute_region operation.

Definition at line 158 of file OpenACCUtilsLoop.cpp.

References mlir::OpBuilder::cloneRegionBefore(), mlir::Region::end(), mlir::RewriterBase::eraseOp(), mlir::Region::getBlocks(), mlir::Operation::getLoc(), mlir::Operation::getNumSuccessors(), 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 27 of file OpenACCUtilsGPU.h.

Referenced by getOrCreateGPUModule().