MLIR 24.0.0git
mlir::acc Namespace Reference

Namespaces

namespace  AttributeTrait
namespace  detail

Classes

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

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

Functions

mlir::Value getVar (mlir::Operation *accDataClauseOp)
 Used to obtain the var from a data clause operation.
mlir::TypedValue< mlir::acc::PointerLikeType > getVarPtr (mlir::Operation *accDataClauseOp)
 Used to obtain the var from a data clause operation if it implements PointerLikeType.
mlir::Type getVarType (mlir::Operation *accDataClauseOp)
 Used to obtains the varType from a data clause operation which records the type of variable.
mlir::Value getAccVar (mlir::Operation *accDataClauseOp)
 Used to obtain the accVar from a data clause operation.
mlir::TypedValue< mlir::acc::PointerLikeType > getAccPtr (mlir::Operation *accDataClauseOp)
 Used to obtain the accVar from a data clause operation if it implements PointerLikeType.
mlir::Value getVarPtrPtr (mlir::Operation *accDataClauseOp)
 Used to obtain the varPtrPtr from a data clause operation.
mlir::SmallVector< mlir::ValuegetBounds (mlir::Operation *accDataClauseOp)
 Used to obtain bounds from an acc data clause operation.
mlir::SmallVector< mlir::ValuegetAsyncOperands (mlir::Operation *accDataClauseOp)
 Used to obtain async operands from an acc data clause operation.
mlir::ArrayAttr getAsyncOperandsDeviceType (mlir::Operation *accDataClauseOp)
 Returns an array of acc:DeviceTypeAttr attributes attached to an acc data clause operation, that correspond to the device types associated with the async clauses with an async-value.
mlir::ArrayAttr getAsyncOnly (mlir::Operation *accDataClauseOp)
 Returns an array of acc:DeviceTypeAttr attributes attached to an acc data clause operation, that correspond to the device types associated with the async clauses without an async-value.
std::optional< llvm::StringRef > getVarName (mlir::Operation *accOp)
 Used to obtain the name from an acc operation.
std::optional< mlir::acc::DataClause > getDataClause (mlir::Operation *accDataEntryOp)
 Used to obtain the dataClause from a data entry operation.
bool getImplicitFlag (mlir::Operation *accDataEntryOp)
 Used to find out whether data operation is implicit.
mlir::ValueRange getDataOperands (mlir::Operation *accOp)
 Used to get an immutable range iterating over the data operands.
mlir::MutableOperandRange getMutableDataOperands (mlir::Operation *accOp)
 Used to get a mutable range iterating over the data operands.
mlir::SymbolRefAttr getRecipe (mlir::Operation *accOp)
 Used to get the recipe attribute from a data clause operation.
bool isPointerLikeType (mlir::Type type)
 Used to check whether the provided type implements the PointerLikeType interface.
bool isMappableType (mlir::Type type)
 Used to check whether the provided type implements the MappableType interface.
static constexpr StringLiteral getDeclareAttrName ()
 Used to obtain the attribute name for declare.
static constexpr StringLiteral getDeclareActionAttrName ()
static constexpr StringLiteral getRoutineInfoAttrName ()
static constexpr StringLiteral getSpecializedRoutineAttrName ()
bool isAccRoutine (mlir::Operation *op)
 Used to check whether the current operation is marked with acc routine.
bool isSpecializedAccRoutine (mlir::Operation *op)
 Used to check whether this is a specialized accelerator version of acc routine function.
static constexpr StringLiteral getFromDefaultClauseAttrName ()
static constexpr StringLiteral getCollapseCountAttrName ()
 Name for an attribute attached to a loop indicating the number of loops collapsed to create that loop.
static constexpr StringLiteral getVarNameAttrName ()
static constexpr StringLiteral getCombinedConstructsAttrName ()
template<typename ComputeOpT>
static bool isGangWorkerVectorAllOne (ComputeOpT op)
ParLevel getGangParLevel (int64_t gangDimValue)
 Convert a gang dimension value (1, 2, or 3) to the corresponding ParLevel.
mlir::OperationgetEnclosingComputeOp (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::OperationgetACCDataClauseOpForBlockArg (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 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 isDeviceValue (mlir::Value val)
 Check if a value represents device data.
bool isValidValueUse (mlir::Value val, mlir::Region &region)
 Check if a value use is valid in an OpenACC region.
llvm::SmallVector< mlir::ValuegetDominatingDataClauses (mlir::Operation *computeConstructOp, mlir::DominanceInfo &domInfo, mlir::PostDominanceInfo &postDomInfo)
 Collects all data clauses that dominate the compute construct.
remark::detail::InFlightRemark emitRemark (mlir::Operation *op, const std::function< std::string()> &messageFn, llvm::StringRef category="openacc")
 Emit an OpenACC remark with lazy message generation.
remark::detail::InFlightRemark emitRemark (mlir::Operation *op, const llvm::Twine &message, llvm::StringRef category="openacc")
 Emit an OpenACC remark for the given operation with the given message.
std::optional< DataLayoutgetDataLayout (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.
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.
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< boolisPrivateLocalSharedMemoryCandidate (PrivateLocalOp privateLocal, ComputeRegionOp computeRegion, ModuleOp module, const ACCToGPUMappingPolicy &policy, OpenACCSupport *support=nullptr)
 True when privateLocal may be placed in shared memory.
std::optional< int64_tgetPrivateLocalSharedMemoryUpperBoundBytes (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.
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::iteratorcloneACCRegionInto (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.
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.
llvm::SmallVector< mlir::acc::LoopOp > uncollapseLoops (mlir::acc::LoopOp origLoop, unsigned tileCount, unsigned collapseCount, mlir::RewriterBase &rewriter)
 Uncollapse tile loops with multiple IVs and collapseCount < tileCount.
mlir::acc::LoopOp tileACCLoops (llvm::SmallVector< mlir::acc::LoopOp > &tileLoops, const llvm::SmallVector< mlir::Value > &tileSizes, int32_t defaultTileSize, mlir::RewriterBase &rewriter)
 Tile ACC loops according to the given tile sizes.
std::optional< TypeSizeAndAlignmentgetTypeSizeAndAlignment (Type ty, ModuleOp module, const DataLayout &dl, OpenACCSupport *support=nullptr)
 Returns the size and ABI alignment in bytes.
std::optional< TypeSizeAndAlignmentgetTypeSizeAndAlignment (Type ty, ModuleOp module, OpenACCSupport *support=nullptr)
 Same as above, obtaining dl from module via getDataLayout.
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).
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 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< TypeSizeAndAlignmentgetTypeSizeAndAlignmentHelper (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

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

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 29 of file OpenACCUtilsType.h.

Enumeration Type Documentation

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

Function Documentation

◆ buildComputeRegion()

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

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

Creates a new acc.compute_region with the given launch arguments and origin string, then clones the operations from regionToClone into its body. 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().

◆ cloneACCRegionInto()

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

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

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

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

◆ collectPrivateLocalParDims()

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

Collect parallel dimensions that govern privatization of privateLocal.

Definition at line 277 of file OpenACCUtilsCG.cpp.

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

Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().

◆ 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
loopOpThe acc.loop operation to convert (must not be unstructured)
rewriterRewriterBase for creating operations
enableCollapseWhether 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 183 of file OpenACCUtilsLoop.cpp.

References mlir::coalesceLoops(), mlir::arith::ConstantIndexOp::create(), mlir::OpBuilder::getInsertionPoint(), mlir::IRMapping::map(), 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
loopOpThe acc.loop operation to convert
rewriterRewriterBase 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 263 of file OpenACCUtilsLoop.cpp.

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

◆ convertUnstructuredACCLoopToSCFExecuteRegion()

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

Convert an unstructured acc.loop to scf.execute_region.

Parameters
loopOpThe acc.loop operation to convert (must be unstructured)
rewriterRewriterBase 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 327 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 221 of file OpenACCUtilsCG.cpp.

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

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

◆ 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
opThe operation to emit the remark for.
messageThe remark message.
categoryOptional 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 150 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
opThe operation to emit the remark for.
messageFnA callable that returns the remark message.
categoryOptional 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 371 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().

◆ findParDim()

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

Definition at line 143 of file OpenACCUtilsCG.cpp.

References lhs, and rhs.

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, mlir::arith::getReductionOp(), lhs, and rhs.

◆ 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 43 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 28 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 5249 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 5261 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

◆ getAsyncOnly()

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

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

Definition at line 5316 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 5294 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 5308 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.

◆ getBaseEntity()

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

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

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

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

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

Referenced by getCollapseCount(), and setCollapseCountAttr().

◆ getCombinedConstructsAttrName()

constexpr StringLiteral mlir::acc::getCombinedConstructsAttrName ( )
staticconstexpr

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

Referenced by checkDeclareOperands().

◆ getDataLayout()

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

Get the data layout for an operation.

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

Parameters
opThe operation to get the data layout for.
allowDefaultIf 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 getTypeSizeAndAlignment().

◆ getDataOperands()

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

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

Definition at line 5352 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

◆ getDeclareActionAttrName()

constexpr StringLiteral mlir::acc::getDeclareActionAttrName ( )
staticconstexpr

Definition at line 181 of file OpenACC.h.

◆ getDeclareAttrName()

constexpr StringLiteral mlir::acc::getDeclareAttrName ( )
staticconstexpr

Used to obtain the attribute name for declare.

Definition at line 177 of file OpenACC.h.

Referenced by isDeviceValue(), 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 75 of file OpenACCUtils.cpp.

References mlir::Operation::getParentOp().

◆ getDominatingDataClauses()

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

Collects all data clauses that dominate the compute construct.

This includes data clauses from:

  • The compute construct itself
  • Enclosing data constructs
  • Applicable declare directives (those that dominate and post-dominate) This is used to determine if a variable is already covered by an existing data clause.
    Parameters
    computeConstructOpThe compute construct operation
    domInfoDominance information
    postDomInfoPost-dominance information
    Returns
    Vector of data clause values that dominate the compute construct

Definition at line 312 of file OpenACCUtils.cpp.

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

◆ getEnclosingComputeOp()

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

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

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

Definition at line 23 of file OpenACCUtils.cpp.

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

◆ getFromDefaultClauseAttrName()

constexpr StringLiteral mlir::acc::getFromDefaultClauseAttrName ( )
staticconstexpr

Definition at line 205 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 110 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

◆ 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

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

References ACC_DATA_ENTRY_OPS.

◆ getMutableDataOperands()

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

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

Definition at line 5362 of file OpenACC.cpp.

References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.

◆ getOrCreateGPUModule()

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

Get or create a GPU module in the given module.

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

Parameters
modThe module to search or create the GPU module in.
createIf true (default), create the GPU module if it doesn't exist.
nameThe 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().

◆ getParDimsAttr()

GPUParallelDimsAttr mlir::acc::getParDimsAttr ( Operation * op)

◆ 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 270 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 411 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 251 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 5371 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 144 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

◆ getRoutineInfoAttrName()

constexpr StringLiteral mlir::acc::getRoutineInfoAttrName ( )
staticconstexpr

Definition at line 185 of file OpenACC.h.

Referenced by isAccRoutine(), and isValidSymbolUse().

◆ getSeqParDimsAttr()

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

Create a sequential GPUParallelDimsAttr based on the mapping policy.

Definition at line 116 of file OpenACCUtilsCG.h.

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

◆ getSpecializedRoutineAttrName()

constexpr StringLiteral mlir::acc::getSpecializedRoutineAttrName ( )
staticconstexpr

Definition at line 189 of file OpenACC.h.

Referenced by isSpecializedAccRoutine().

◆ getTypeCategory()

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

Get the type category of an OpenACC variable.

Definition at line 96 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 )

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.

When support is provided, aggregate element types are sized by recursing through support so dialect-specific implementations can handle nested types.

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

Definition at line 30 of file OpenACCUtilsType.cpp.

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

Referenced by 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 )

Same as above, obtaining dl from module via getDataLayout.

Definition at line 84 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

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

References ACC_DATA_ENTRY_OPS.

Referenced by checkDeclareOperands().

◆ getVariableName()

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

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

Returns empty string if no name is found.

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

Referenced by getVariableName().

◆ getVarNameAttrName()

constexpr StringLiteral mlir::acc::getVarNameAttrName ( )
staticconstexpr

Definition at line 215 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 109 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 5216 of file OpenACC.cpp.

◆ getVarPtrPtr()

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

Used to obtain the varPtrPtr from a data clause operation.

Returns empty value if not a data clause operation.

Definition at line 5269 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

◆ getVarType()

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

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

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

Definition at line 5238 of file OpenACC.cpp.

References ACC_DATA_ENTRY_OPS.

◆ getWorkerPrivateSharedMemoryNumCopies()

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

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

References mlir::Operation::hasAttrOfType().

◆ hasParDimsAttr()

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

Return whether op carries parallel dimensions.

Definition at line 183 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 185 of file OpenACCUtilsCG.cpp.

References getParDimsAttr().

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

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

◆ isDeviceValue()

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

Check if a value represents device data.

This checks if the value represents device data via the MappableType, PointerLikeType, and GlobalVariableOpInterface interfaces, and whether the defining operation carries acc.declare with the deviceptr clause.

Parameters
valThe value to check
Returns
true if the value is device data, false otherwise

Definition at line 246 of file OpenACCUtils.cpp.

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

Referenced by isDeviceValue(), 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)
static

Definition at line 242 of file OpenACC.h.

References mlir::isConstantIntValue().

◆ isInsideACCSpecializedRoutine()

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

◆ isMappableType()

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

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

Definition at line 172 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 64 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 69 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 166 of file OpenACC.h.

◆ isPrivateLocalSharedMemoryCandidate()

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

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

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

Referenced by isInsideACCSpecializedRoutine().

◆ isThreadXPrivatize()

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

Definition at line 263 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
userThe operation using the symbol
symbolThe symbol reference being used
definingOpPtrOptional output parameter to receive the defining op
Returns
true if the symbol use is valid, false otherwise

Definition at line 192 of file OpenACCUtils.cpp.

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

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

◆ isValidValueUse()

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

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

This is true if:

  • The value is produced by an ACC data entry operation
  • The value is device data
  • The value is only used by private clauses in the region
    Parameters
    valThe value to check
    regionThe OpenACC region
    Returns
    true if the value use is valid, false otherwise

Definition at line 289 of file OpenACCUtils.cpp.

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

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

◆ populateACCHostFallbackPatterns()

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

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

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

Parameters
enableLoopConversionWhether to convert orphan acc.loop operations.

Definition at line 404 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
enableLoopConversionWhether to convert orphan acc.loop operations.

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

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

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

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

Referenced by convertACCLoopToSCFFor().

◆ setGPUBlockRedundantAttr()

void mlir::acc::setGPUBlockRedundantAttr ( Operation * op)

Mark op with the acc.gpu_block_redundant attribute.

Definition at line 216 of file OpenACCUtilsCG.cpp.

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

◆ setParDimsAttr()

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

Set parallel dimensions on op.

Definition at line 191 of file OpenACCUtilsCG.cpp.

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

Referenced by copyParDimsAttr().

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

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

◆ tileACCLoops()

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

Tile ACC loops according to the given tile sizes.

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

Before Tiling:

#pragma acc loop tile(tile_size1, tile_size2)
for (i = lb1; i < ub1; i += step1) { // original loop
for (j = lb2; j < ub2; j += step2) {
a[i,j] = i + j;
}
}
Eliminates variable at the specified position using Fourier-Motzkin variable elimination.

After Tiling:

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

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

Parameters
tileLoopsThe loops to tile (outermost first).
tileSizesThe tile sizes for each dimension. Values of -1 are treated as unknown and resolved to defaultTileSize.
defaultTileSizeThe default tile size to use for unknown (*) tiles.
rewriterThe rewriter to use for modifications.
Returns
The outermost loop after tiling.

Definition at line 163 of file OpenACCUtilsTiling.cpp.

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

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

◆ uncollapseLoops()

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

Uncollapse tile loops with multiple IVs and collapseCount < tileCount.

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

Parameters
origLoopThe original loop operation to uncollapse.
tileCountThe number of tile dimensions.
collapseCountThe collapse count from the original loop.
rewriterThe rewriter to use for modifications.
Returns
A vector of uncollapsed loop operations.

Definition at line 280 of file OpenACCUtilsTiling.cpp.

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

◆ updateParDimsAttr()

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

Update parallel dimensions on op.

Definition at line 200 of file OpenACCUtilsCG.cpp.

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

◆ 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
regionThe region to wrap (cloned into the execute_region; not modified).
mappingIR mapping for the clone; updated with block and value mappings.
locLocation for the created execute_region op.
rewriterRewriterBase for creating and erasing operations.
Returns
The created scf.execute_region operation.

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