|
MLIR 24.0.0git
|
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 |
| 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 | DataDescriptor |
| IDs for the argument descriptors of the OpenACC data entry points, declared in OpenACCRuntimeDescriptors.def. 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 | 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) |
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 ®ion, RewriterBase &rewriter) |
Splices region, the body of a structured construct, into the block holding op, so that the construct itself can be erased. | |
| SmallVector< DeviceType, 3 > | getDeviceTypesByPrecedence (DeviceType deviceType) |
| The clauses of a construct can be given once per device type. | |
| 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. | |
| 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> | |
| 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. | |
| 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 ®ion) |
| 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 ®ion) |
| Returns true if this value is only used by acc.private operations in the region. | |
| bool | isOnlyUsedByReductionClauses (mlir::Value val, mlir::Region ®ion) |
| 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 | isDeviceValue (mlir::Value val) |
| Check if a value represents device data. | |
| bool | isValidValueUse (mlir::Value val, mlir::Region ®ion) |
| Check if a value use is valid in an OpenACC region. | |
| llvm::SmallVector< mlir::Value > | getDominatingDataClauses (mlir::Operation *computeConstructOp, mlir::DominanceInfo &domInfo, mlir::PostDominanceInfo &postDomInfo) |
| Collects all data clauses that dominate the compute construct. | |
| remark::detail::InFlightRemark | emitRemark (mlir::Operation *op, const std::function< std::string()> &messageFn, llvm::StringRef category="openacc") |
| Emit an OpenACC remark with lazy message generation. | |
| remark::detail::InFlightRemark | emitRemark (mlir::Operation *op, const llvm::Twine &message, llvm::StringRef category="openacc") |
| Emit an OpenACC remark for the given operation with the given message. | |
| std::optional< DataLayout > | getDataLayout (Operation *op, bool allowDefault=true) |
| Get the data layout for an operation. | |
| ComputeRegionOp | buildComputeRegion (Location loc, ValueRange launchArgs, ValueRange inputArgs, llvm::StringRef origin, Region ®ionToClone, 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. | |
| 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 ®ion) |
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 ®ion, 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 > | 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 | 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. | |
Type alias for the GPU-specific mapping policy.
Definition at line 158 of file OpenACCParMapping.h.
| using mlir::acc::TypesForDevice = llvm::SmallSetVector<int64_t, 3> |
| using mlir::acc::TypeSizeAndAlignment = std::pair<llvm::TypeSize, llvm::TypeSize> |
Definition at line 33 of file OpenACCUtilsType.h.
|
strong |
IDs for the argument descriptors of the OpenACC data entry points, declared in OpenACCRuntimeDescriptors.def.
Definition at line 48 of file OpenACCRuntimeUtils.h.
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 | |
|
strong |
IDs for OpenACC compiler-to-runtime entry points (__tgt_acc_*).
Definition at line 34 of file OpenACCRuntimeUtils.h.
| 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().
| 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().
| 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().
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 emitWaitCall(), getAsyncQueue(), getBoundValue(), and getMapSize().
| 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().
| SmallVector< GPUParallelDimAttr > mlir::acc::collectPrivateLocalParDims | ( | PrivateLocalOp | privateLocal, |
| ComputeRegionOp | computeRegion ) |
Collect parallel dimensions that govern privatization of privateLocal.
Definition at line 299 of file OpenACCUtilsCG.cpp.
References getParDimsAttr(), getReductionCombineParDims(), mlir::Operation::getResult(), mlir::Value::getUsers(), and insertParDim().
Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().
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 625 of file OpenACCUtilsCG.cpp.
References findCorrespondingDataExit(), getDataClause(), getEntryModifiers(), getExitDataClause(), getImplicitFlag(), mlir::Operation::getResult(), mlir::Value::getUses(), and hasCopyOutSibling().
Referenced by computePackedMapFlags().
| 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 751 of file OpenACCUtilsCG.cpp.
References getEnclosingModule(), mlir::Value::getType(), and getTypeSizeAndAlignment().
Referenced by computeMapInfoSizeBytes(), and getMapSize().
| 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 778 of file OpenACCUtilsCG.cpp.
References computeMapInfoSizeBytes(), getDataLayout(), and getEnclosingModule().
| 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 604 of file OpenACCUtilsCG.cpp.
References mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isGang(), mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isVector(), and mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isWorker().
| 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.
| loopOp | The acc.loop operation to convert (must not be unstructured) |
| rewriter | RewriterBase for creating operations |
| enableCollapse | Whether to collapse nested loops into one |
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().
Referenced by mlir::acc::impl::ACCComputeLoweringBase< DerivedT >::createACCComputeLowering.
| scf::ParallelOp mlir::acc::convertACCLoopToSCFParallel | ( | LoopOp | loopOp, |
| RewriterBase & | rewriter ) |
Convert acc.loop to scf.parallel.
The loop induction variables are converted to index types.
| loopOp | The acc.loop operation to convert |
| rewriter | RewriterBase for creating and erasing operations |
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().
| scf::ExecuteRegionOp mlir::acc::convertUnstructuredACCLoopToSCFExecuteRegion | ( | LoopOp | loopOp, |
| RewriterBase & | rewriter ) |
Convert an unstructured acc.loop to scf.execute_region.
| loopOp | The acc.loop operation to convert (must be unstructured) |
| rewriter | RewriterBase for creating and erasing operations |
Definition at line 338 of file OpenACCUtilsLoop.cpp.
References mlir::OpBuilder::getInsertionBlock(), mlir::Block::getParentOp(), and wrapMultiBlockRegionWithSCFExecuteRegion().
Copy parallel dimensions from from to to.
Definition at line 226 of file OpenACCUtilsCG.cpp.
References getParDimsAttr(), hasParDimsAttr(), and setParDimsAttr().
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCBindRoutine | ( | ) |
Definition at line 122 of file ACCBindRoutine.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCBindRoutine | ( | ACCBindRoutineOptions | options | ) |
Definition at line 126 of file ACCBindRoutine.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCCGToGPU | ( | ) |
Definition at line 237 of file ACCCGToGPU.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCCGToGPU | ( | ACCCGToGPUOptions | options | ) |
Definition at line 241 of file ACCCGToGPU.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCComputeLowering | ( | ) |
Definition at line 339 of file ACCComputeLowering.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCComputeLowering | ( | ACCComputeLoweringOptions | options | ) |
Definition at line 343 of file ACCComputeLowering.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareCtorDtorConversion | ( | ) |
Definition at line 449 of file ACCDeclareCtorDtorConversion.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareCtorDtorConversion | ( | ACCDeclareCtorDtorConversionOptions | options | ) |
Definition at line 453 of file ACCDeclareCtorDtorConversion.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareGPUModuleInsertion | ( | ) |
Definition at line 544 of file ACCDeclareGPUModuleInsertion.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCDeclareGPUModuleInsertion | ( | ACCDeclareGPUModuleInsertionOptions | options | ) |
Definition at line 548 of file ACCDeclareGPUModuleInsertion.cpp.
| 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 623 of file ACCEmitRemarksData.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCEmitRemarksLoop | ( | ) |
Definition at line 714 of file ACCEmitRemarksLoop.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCEmitRemarksLoop | ( | ACCEmitRemarksLoopOptions | options | ) |
Definition at line 718 of file ACCEmitRemarksLoop.cpp.
| 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 793 of file ACCEmitRemarksPrivate.cpp.
| 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 868 of file ACCIfClauseLowering.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitData | ( | ) |
Definition at line 963 of file ACCImplicitData.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitData | ( | ACCImplicitDataOptions | options | ) |
Definition at line 967 of file ACCImplicitData.cpp.
| 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 1042 of file ACCImplicitDeclare.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitRoutine | ( | ) |
Definition at line 1139 of file ACCImplicitRoutine.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCImplicitRoutine | ( | ACCImplicitRoutineOptions | options | ) |
Definition at line 1143 of file ACCImplicitRoutine.cpp.
| 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 1218 of file ACCLegalizeSerial.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCLoopTiling | ( | ) |
Definition at line 1309 of file ACCLoopTiling.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCLoopTiling | ( | ACCLoopTilingOptions | options | ) |
Definition at line 1313 of file ACCLoopTiling.cpp.
| 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 1388 of file ACCRecipeMaterialization.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineLowering | ( | ) |
Definition at line 1486 of file ACCRoutineLowering.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineLowering | ( | ACCRoutineLoweringOptions | options | ) |
Definition at line 1490 of file ACCRoutineLowering.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineToGPUFunc | ( | ) |
Definition at line 1587 of file ACCRoutineToGPUFunc.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCRoutineToGPUFunc | ( | ACCRoutineToGPUFuncOptions | options | ) |
Definition at line 1591 of file ACCRoutineToGPUFunc.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForDevice | ( | ) |
Definition at line 1682 of file ACCSpecializeForDevice.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForDevice | ( | ACCSpecializeForDeviceOptions | options | ) |
Definition at line 1686 of file ACCSpecializeForDevice.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForHost | ( | ) |
Definition at line 1777 of file ACCSpecializeForHost.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createACCSpecializeForHost | ( | ACCSpecializeForHostOptions | options | ) |
Definition at line 1781 of file ACCSpecializeForHost.cpp.
| 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(), and emitWaitCall().
| 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().
| std::unique_ptr<::mlir::Pass > mlir::acc::createLegalizeDataValuesInRegion | ( | ) |
Definition at line 1875 of file LegalizeDataValues.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createLegalizeDataValuesInRegion | ( | LegalizeDataValuesInRegionOptions | options | ) |
Definition at line 1879 of file LegalizeDataValues.cpp.
| 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 1954 of file OffloadLiveInValueCanonicalization.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createOffloadTargetVerifier | ( | ) |
Definition at line 2054 of file OffloadTargetVerifier.cpp.
| std::unique_ptr<::mlir::Pass > mlir::acc::createOffloadTargetVerifier | ( | OffloadTargetVerifierOptions | options | ) |
Definition at line 2058 of file OffloadTargetVerifier.cpp.
| 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 206 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 emitDataRuntimeCall(), emitGetDevicePtrCall(), and emitWaitCall().
| 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 287 of file ACCToLLVMUtils.cpp.
References createIdent(), createRuntimeCall(), mlir::Builder::getI64Type(), mlir::Operation::getLoc(), mlir::acc::ACCRuntimeCallConfig::getMapFlagsRuntimeValue(), and getParentFunctionName().
| 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 321 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().
|
inline |
Emit an OpenACC remark for the given operation with the given message.
| op | The operation to emit the remark for. |
| message | The remark message. |
| category | Optional category for the remark. Defaults to "openacc". |
Definition at line 151 of file OpenACCUtils.h.
References emitRemark().
| 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.
| 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". |
Definition at line 415 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().
| 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 346 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().
| LogicalResult mlir::acc::emitWaitCall | ( | Location | loc, |
| ValueRange | waitOperands, | ||
| Value | asyncQueue, | ||
| OpBuilder & | builder, | ||
| Region & | globalSymbolRegion, | ||
| SymbolTable & | symbolTable, | ||
| const ACCRuntimeCallConfig & | config ) |
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.
Definition at line 240 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().
| 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 258 of file ACCToLLVMUtils.h.
References mlir::acc::OpenACCSupport::emitNYI(), emitWaitCall(), getWaitClauseValues(), hasWaitDevnum(), and success().
Definition at line 535 of file OpenACCUtilsCG.cpp.
References getPairedDataExitOps(), and nullptr.
Referenced by computeDataClauseMapFlags(), and getMappingExitLoc().
|
static |
Definition at line 143 of file OpenACCUtilsCG.cpp.
Referenced by insertParDim(), and removeParDim().
| 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().
| 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 46 of file OpenACCUtils.cpp.
References getACCOperandForBlockArg(), and mlir::Value::getDefiningOp().
| 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 31 of file OpenACCUtils.cpp.
References mlir::Block::getParentOp().
Referenced by getACCDataClauseOpForBlockArg().
| 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 5387 of file OpenACC.cpp.
| 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 5400 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.
| ActiveParDimsAttr mlir::acc::getActiveParDimsAttr | ( | Operation * | op | ) |
Obtain the active parallel dimensions carried by op, if any.
Definition at line 232 of file OpenACCUtilsCG.cpp.
References mlir::Operation::getDiscardableAttrOfType().
Referenced by hasActiveParDimsAttr().
| 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 163 of file ACCToLLVMUtils.h.
References getDeviceTypesByPrecedence().
Referenced by getAsyncQueue().
| 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 5458 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.
| 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 5436 of file OpenACC.cpp.
| 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 5450 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS, and ACC_DATA_EXIT_OPS.
| 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 227 of file ACCToLLVMUtils.cpp.
References castToI64(), mlir::acc::ACCRuntimeCallConfig::getAsyncNoValueRuntimeValue(), mlir::acc::ACCRuntimeCallConfig::getAsyncSyncRuntimeValue(), and mlir::Builder::getI64Type().
Referenced by getAsyncQueue().
| 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 240 of file ACCToLLVMUtils.h.
References getAsyncClauseValue(), and getAsyncQueue().
| mlir::Value mlir::acc::getBaseEntity | ( | mlir::Value | val | ) |
Definition at line 189 of file OpenACCUtils.cpp.
References mlir::Value::getDefiningOp().
| 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 5419 of file OpenACC.cpp.
Referenced by mlir::createACCArgumentDescriptor(), and getMapSize().
| 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().
|
staticconstexpr |
Name for an attribute attached to a loop indicating the number of loops collapsed to create that loop.
Definition at line 212 of file OpenACC.h.
Referenced by getCollapseCount(), and setCollapseCountAttr().
|
staticconstexpr |
| 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 5477 of file OpenACC.cpp.
Referenced by checkDeclareOperands(), computeDataClauseMapFlags(), hasAttachPoint(), and hasCopyOutSibling().
| DataDescKind mlir::acc::getDataDescKind | ( | Operation * | mapEntryOp | ) |
Returns descriptor kind from acc.map_info, or none for other ops.
Definition at line 481 of file OpenACCUtilsCG.cpp.
Referenced by mlir::createACCArgumentDescriptor(), and getMapSize().
| int64_t mlir::acc::getDataDescriptorFieldIndex | ( | FieldEnum | field | ) |
The index an insert or extract of field addresses.
Definition at line 63 of file OpenACCRuntimeUtils.h.
Referenced by mlir::createACCArgumentDescriptor(), mlir::createACCDataDescriptor(), and mlir::createACCMemRefDescriptorWrapperArg().
| 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().
| StringRef mlir::acc::getDataDescriptorName | ( | DataDescriptor | desc | ) |
Returns the name of the runtime type desc materializes.
Definition at line 52 of file OpenACCRuntimeUtils.cpp.
| 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().
| 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.
| op | The operation to get the data layout for. |
| allowDefault | If true, allow returning a default data layout. |
Definition at line 34 of file OpenACCUtilsCG.cpp.
References mlir::Operation::getParentOfType(), and mlir::Operation::getParentOp().
Referenced by computeMapInfoSizeBytes(), and getTypeSizeAndAlignment().
| mlir::ValueRange mlir::acc::getDataOperands | ( | mlir::Operation * | accOp | ) |
Used to get an immutable range iterating over the data operands.
Definition at line 5497 of file OpenACC.cpp.
References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.
|
staticconstexpr |
|
staticconstexpr |
Used to obtain the attribute name for declare.
Definition at line 177 of file OpenACC.h.
Referenced by isDeviceValue(), and isValidSymbolUse().
| 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 78 of file OpenACCUtils.cpp.
References mlir::Operation::getParentOp().
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 487 of file OpenACCUtilsCG.cpp.
Referenced by mlir::createACCArgumentDescriptor(), and mlir::emitACCDataRuntimeArgs().
| 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 312 of file ACCToLLVMUtils.cpp.
Referenced by getAsyncClauseValue(), and getWaitClauseValues().
| 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:
| computeConstructOp | The compute construct operation |
| domInfo | Dominance information |
| postDomInfo | Post-dominance information |
Definition at line 356 of file OpenACCUtils.cpp.
References mlir::DominanceInfo::dominates(), mlir::Operation::getParentOfType(), and mlir::Operation::getParentOp().
| 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 26 of file OpenACCUtils.cpp.
References ACC_COMPUTE_CONSTRUCT_OPS, and mlir::Region::getParentOfType().
|
static |
Returns the module var lives in.
Definition at line 742 of file OpenACCUtilsCG.cpp.
References mlir::Value::getDefiningOp(), mlir::Operation::getParentOp(), and mlir::Value::getParentRegion().
Referenced by computeMapInfoSizeBytes(), and computeMapInfoSizeBytes().
|
static |
Definition at line 597 of file OpenACCUtilsCG.cpp.
References ACC_DATA_ENTRY_OPS.
Referenced by computeDataClauseMapFlags().
|
static |
Definition at line 556 of file OpenACCUtilsCG.cpp.
Referenced by computeDataClauseMapFlags().
| 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().
|
staticconstexpr |
|
inline |
Create a gang dim 1 GPUParallelDimsAttr based on the mapping policy.
Definition at line 123 of file OpenACCUtilsCG.h.
References mlir::acc::ACCParMappingPolicy< ParDimAttrT >::gangDim().
|
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.
| 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().
|
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().
| 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().
|
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().
| 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 5487 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS.
Referenced by computeDataClauseMapFlags().
| 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 createIdent(), and mlir::emitACCDataRuntimeArgs().
Returns element size in bytes from acc.map_info, if present.
Definition at line 500 of file OpenACCUtilsCG.cpp.
Referenced by getElementSize().
| std::optional< MapFlags > mlir::acc::getMapFlags | ( | Operation * | mapEntryOp | ) |
Returns offload map-type flags from acc.map_info, if present.
Definition at line 513 of file OpenACCUtilsCG.cpp.
Referenced by hasCopyOutSibling().
Definition at line 540 of file OpenACCUtilsCG.cpp.
References findCorrespondingDataExit(), and mlir::Value::getDefiningOp().
| 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 549 of file OpenACCUtilsCG.cpp.
References getMappingExitLoc().
Referenced by getMappingExitLoc().
Returns the optional size operand from acc.map_info, or null.
Definition at line 507 of file OpenACCUtilsCG.cpp.
Referenced by mlir::emitACCDataRuntimeArgs(), and getMapSize().
| mlir::MutableOperandRange mlir::acc::getMutableDataOperands | ( | mlir::Operation * | accOp | ) |
Used to get a mutable range iterating over the data operands.
Definition at line 5508 of file OpenACC.cpp.
References ACC_COMPUTE_AND_DATA_CONSTRUCT_OPS.
| 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::emitACCDataRuntimeArgs().
| 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.
| 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. |
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().
| 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 519 of file OpenACCUtilsCG.cpp.
References ACC_DATA_EXIT_OPS, and mlir::Value::getUses().
Referenced by findCorrespondingDataExit().
| 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().
| 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(), emitWaitCall(), getParentFunctionName(), and getParentFunctionName().
| 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().
| 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().
| 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 292 of file OpenACCUtilsCG.cpp.
Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().
| 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 433 of file OpenACCUtilsCG.cpp.
References collectPrivateLocalParDims(), getPrivateBaseMemRefType(), getTypeSizeAndAlignment(), getWorkerPrivateSharedMemoryNumCopies(), isPrivateLocalSharedMemoryCandidate(), and mlir::acc::ACCParMappingPolicy< ParDimAttrT >::isWorker().
| PrivatizeOp mlir::acc::getPrivatizeOp | ( | PrivateLocalOp | privateLocal, |
| ComputeRegionOp | computeRegion ) |
Resolve the acc.privatize operation associated with a private local.
Definition at line 273 of file OpenACCUtilsCG.cpp.
References mlir::Value::getDefiningOp().
Referenced by isPrivateLocalSharedMemoryCandidate().
| mlir::SymbolRefAttr mlir::acc::getRecipe | ( | mlir::Operation * | accOp | ) |
Used to get the recipe attribute from a data clause operation.
Definition at line 5518 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS.
| 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 151 of file OpenACCUtils.cpp.
References mlir::Type::print().
Referenced by mlir::acc::OpenACCSupport::getRecipeName().
| 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().
| 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().
|
static |
Definition at line 141 of file OpenACCUtilsReduction.cpp.
References mlir::emitError(), mlir::arith::getIdentityValueAttr(), and mlir::Type::isIntOrIndexOrFloat().
Referenced by createIdentityValue().
|
staticconstexpr |
Definition at line 185 of file OpenACC.h.
Referenced by mlir::acc::impl::ACCBindRoutineBase< DerivedT >::createACCBindRoutine, mlir::acc::impl::ACCCGToGPUBase< DerivedT >::getDependentDialects(), isAccRoutine(), and isValidSymbolUse().
| 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().
| 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().
|
inline |
Create a sequential GPUParallelDimsAttr based on the mapping policy.
Definition at line 129 of file OpenACCUtilsCG.h.
References mlir::acc::ACCParMappingPolicy< ParDimAttrT >::seqDim().
|
staticconstexpr |
Definition at line 189 of file OpenACC.h.
Referenced by mlir::acc::impl::ACCCGToGPUBase< DerivedT >::getDependentDialects(), isSpecializedAccRoutine(), and isValidSymbolUse().
| mlir::acc::VariableTypeCategory mlir::acc::getTypeCategory | ( | mlir::Value | var | ) |
Get the type category of an OpenACC variable.
Definition at line 99 of file OpenACCUtils.cpp.
References mlir::Value::getType().
| 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(), getMapSize(), getPrivateLocalSharedMemoryUpperBoundBytes(), mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::getTypeSizeAndAlignment(), getTypeSizeAndAlignment(), mlir::acc::OpenACCSupport::getTypeSizeAndAlignment(), and getTypeSizeAndAlignmentHelper().
| 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().
|
static |
Definition at line 24 of file OpenACCUtilsType.cpp.
References getTypeSizeAndAlignment(), and mlir::acc::OpenACCSupport::getTypeSizeAndAlignment().
Referenced by getTypeSizeAndAlignment().
| 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 5368 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS.
Referenced by checkDeclareOperands(), mlir::createACCArgumentDescriptor(), mlir::emitACCDataRuntimeArgs(), getMapSize(), and hasCopyOutSibling().
| 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 116 of file OpenACCUtils.cpp.
References getConstantIntValue(), mlir::Value::getDefiningOp(), getVarName(), and getVarNameAttrName().
Referenced by mlir::acc::OpenACCSupport::getVariableName().
| std::optional< llvm::StringRef > mlir::acc::getVarName | ( | mlir::Operation * | accOp | ) |
Used to obtain the name from an acc operation.
Definition at line 5465 of file OpenACC.cpp.
Referenced by getVariableName().
|
staticconstexpr |
Definition at line 216 of file OpenACC.h.
Referenced by getVariableName().
| 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 112 of file OpenACCUtils.cpp.
| 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 5354 of file OpenACC.cpp.
Referenced by hasCopyOutSibling().
| 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 5409 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS.
Referenced by mlir::emitACCDataRuntimeArgs(), and hasAttachPoint().
| 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 5376 of file OpenACC.cpp.
References ACC_DATA_ENTRY_OPS.
Referenced by getElementSize(), and getMapSize().
| 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 194 of file ACCToLLVMUtils.h.
References getDeviceTypesByPrecedence().
Referenced by emitWaitClause().
|
static |
Definition at line 343 of file OpenACCUtilsCG.cpp.
References mlir::ValueBoundsConstraintSet::computeConstantBound(), mlir::acc::OpenACCSupport::emitNYI(), and mlir::presburger::EQ.
Referenced by getPrivateLocalSharedMemoryUpperBoundBytes(), and isPrivateLocalSharedMemoryCandidate().
Return whether op carries active parallel dimensions.
Definition at line 237 of file OpenACCUtilsCG.cpp.
References getActiveParDimsAttr().
Returns true when mapEntryOp carries an attach point (varPtrPtr).
Definition at line 467 of file OpenACCUtilsCG.cpp.
References getDataClause(), and getVarPtrPtr().
Referenced by computePackedMapFlags().
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 562 of file OpenACCUtilsCG.cpp.
References getDataClause(), getMapFlags(), mlir::Operation::getResult(), mlir::Value::getUsers(), getVar(), and getVarPtr().
Referenced by computeDataClauseMapFlags().
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().
Return whether op carries parallel dimensions.
Definition at line 185 of file OpenACCUtilsCG.cpp.
References getParDimsAttr().
Referenced by copyParDimsAttr(), setParDimsAttr(), and updateParDimsAttr().
Return whether op carries sequential parallel dimensions.
Definition at line 187 of file OpenACCUtilsCG.cpp.
References getParDimsAttr().
| 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 228 of file ACCToLLVMUtils.h.
Referenced by emitWaitClause(), getWaitDevnumValue(), and getWaitValuesWithoutDevnum().
| 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().
|
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::hasDiscardableAttr().
Referenced by isValidValueUse().
| 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.
| val | The value to check |
Definition at line 279 of file OpenACCUtils.cpp.
References getDeclareAttrName(), mlir::Value::getDefiningOp(), mlir::Operation::getDiscardableAttrOfType(), mlir::Value::getType(), isDeviceValue(), and mlir::SymbolTable::lookupNearestSymbolFrom().
Referenced by isDeviceValue(), and isValidValueUse().
Definition at line 20 of file OpenACCUtilsReduction.cpp.
Referenced by translateACCReductionOperator().
| bool mlir::acc::isGangWorkerVectorAllOne | ( | ComputeOpT | op | ) |
Definition at line 243 of file OpenACC.h.
References mlir::isConstantIntValue().
Definition at line 374 of file OpenACCUtilsCG.cpp.
References mlir::Operation::getParentOfType(), and isSpecializedAccRoutine().
Referenced by isPrivateLocalSharedMemoryCandidate().
|
inline |
Used to check whether the provided type implements the MappableType interface.
Definition at line 172 of file OpenACC.h.
Referenced by checkRecipe().
| 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 67 of file OpenACCUtils.cpp.
References isOnlyUsedByOpClauses().
Referenced by isValidValueUse().
| 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 72 of file OpenACCUtils.cpp.
References isOnlyUsedByOpClauses().
|
inline |
| 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 379 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().
|
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::hasDiscardableAttr().
Referenced by mlir::acc::impl::ACCCGToGPUBase< DerivedT >::getDependentDialects(), isInsideACCSpecializedRoutine(), isValidValueUse(), and mlir::acc::impl::ACCCGToGPUBase< DerivedT >::~ACCCGToGPUBase().
|
static |
Definition at line 285 of file OpenACCUtilsCG.cpp.
Referenced by isPrivateLocalSharedMemoryCandidate().
| 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.
| user | The operation using the symbol |
| symbol | The symbol reference being used |
| definingOpPtr | Optional output parameter to receive the defining op |
Definition at line 215 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().
| 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:
| val | The value to check |
| region | The OpenACC region |
Definition at line 323 of file OpenACCUtils.cpp.
References mlir::Value::getDefiningOp(), mlir::Region::getParentOp(), mlir::Value::getType(), isAccRoutine(), isDeviceValue(), mlir::Type::isIntOrIndexOrFloat(), isOnlyUsedByPrivateClauses(), and isSpecializedAccRoutine().
Referenced by mlir::acc::detail::OpenACCSupportTraits::Model< ImplT >::isValidValueUse(), and mlir::acc::OpenACCSupport::isValidValueUse().
| 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().
| 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.
| enableLoopConversion | Whether to convert orphan acc.loop operations. |
Definition at line 408 of file ACCSpecializeForHost.cpp.
References mlir::RewritePatternSet::getContext(), and mlir::RewritePatternSet::insert().
| 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.
| enableLoopConversion | Whether to convert orphan acc.loop operations. |
Definition at line 371 of file ACCSpecializeForHost.cpp.
References mlir::RewritePatternSet::getContext(), and mlir::RewritePatternSet::insert().
| 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().
| 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 190 of file OpenACCRuntimeUtils.cpp.
References mlir::acc::ACCRuntimeCallConfig::setDeviceTypeRuntimeValue().
Referenced by mlir::acc::ACCRuntimeCallConfig::ACCRuntimeCallConfig().
| 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 197 of file OpenACCRuntimeUtils.cpp.
References mlir::acc::ACCRuntimeCallConfig::setMapFlagRuntimeValue().
Referenced by mlir::acc::ACCRuntimeCallConfig::ACCRuntimeCallConfig().
| 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 790 of file OpenACCUtilsCG.cpp.
References mlir::arith::ConstantIndexOp::create(), and mlir::OpBuilder::setInsertionPoint().
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
|
inline |
Definition at line 2556 of file Passes.h.
Referenced by mlir::registerAllPasses().
| 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().
Set active parallel dimensions on op.
Definition at line 241 of file OpenACCUtilsCG.cpp.
References mlir::Operation::setDiscardableAttr().
Referenced by setActiveParDimsAttr().
Set active parallel dimensions on op from a dimension list.
Definition at line 245 of file OpenACCUtilsCG.cpp.
References mlir::Operation::getContext(), and setActiveParDimsAttr().
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().
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().
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().
| 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 387 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().
Sum aligned static_upper_bound_bytes for all acc.gpu_shared_memory in region.
Definition at line 264 of file OpenACCUtilsCG.cpp.
References mlir::acc::SharedMemoryBudget::alignOffset(), and mlir::Region::walk().
| 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:
Before Tiling:
After Tiling (each group is one multi-IV loop over all tiled IVs):
Unknown tile sizes (represented as -1 in acc dialect for tile(*)) are resolved to the provided default tile size.
| 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. |
Definition at line 155 of file OpenACCUtilsTiling.cpp.
References createElementGroupLoop(), mlir::RewriterBase::finalizeOpModification(), mlir::Builder::getDenseBoolArrayAttr(), mlir::Builder::getIntegerAttr(), mlir::getType(), moveOpsAndReplaceIVs(), removeWorkerVectorFromLoop(), resolveAndCastTileSize(), mlir::OpBuilder::setInsertionPoint(), and mlir::RewriterBase::startOpModification().
| 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().
| 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.
Unfuses fused locations, returning the last sub-location.
Definition at line 25 of file ACCToLLVMUtils.cpp.
Referenced by getFileLineColLoc().
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().
| 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.
| 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. |
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().
|
constexpr |
Default GPU module name used by OpenACC.
Definition at line 27 of file OpenACCUtilsGPU.h.
Referenced by getOrCreateGPUModule().