MLIR  16.0.0git
Public Member Functions | Protected Member Functions | Protected Attributes | List of all members
mlir::gpu::SerializeToBlobPass Class Referenceabstract

Base pass class to serialize kernel functions through LLVM into user-specified IR and add the resulting blob as module attribute. More...

#include "mlir/Dialect/GPU/Transforms/Passes.h"

+ Inheritance diagram for mlir::gpu::SerializeToBlobPass:
+ Collaboration diagram for mlir::gpu::SerializeToBlobPass:

Public Member Functions

 SerializeToBlobPass (TypeID passID)
 
 SerializeToBlobPass (const SerializeToBlobPass &other)
 
void runOnOperation () final
 The polymorphic API that runs the pass over the currently held operation. More...
 
- Public Member Functions inherited from mlir::Pass
virtual ~Pass ()=default
 
TypeID getTypeID () const
 Returns the unique identifier that corresponds to this pass. More...
 
const PassInfolookupPassInfo () const
 Returns the pass info for this pass, or null if unknown. More...
 
virtual StringRef getName () const =0
 Returns the derived pass name. More...
 
virtual StringRef getArgument () const
 Return the command line argument used when registering this pass. More...
 
virtual StringRef getDescription () const
 Return the command line description used when registering this pass. More...
 
Optional< StringRef > getOpName () const
 Returns the name of the operation that this pass operates on, or None if this is a generic OperationPass. More...
 
virtual LogicalResult initializeOptions (StringRef options)
 Attempt to initialize the options of this pass from the given string. More...
 
void printAsTextualPipeline (raw_ostream &os)
 Prints out the pass in the textual representation of pipelines. More...
 
ArrayRef< Statistic * > getStatistics () const
 Returns the main statistics for this pass instance. More...
 
MutableArrayRef< Statistic * > getStatistics ()
 
const PassgetThreadingSibling () const
 Returns the thread sibling of this pass. More...
 
const PassgetThreadingSiblingOrThis () const
 Returns the thread sibling of this pass, or the pass itself it has no sibling. More...
 

Protected Member Functions

void getDependentDialects (DialectRegistry &registry) const override
 Register dependent dialects for the current pass. More...
 
virtual LogicalResult optimizeLlvm (llvm::Module &llvmModule, llvm::TargetMachine &targetMachine)
 Hook allowing the application of optimizations before codegen By default, does nothing. More...
 
virtual std::unique_ptr< llvm::Module > translateToLLVMIR (llvm::LLVMContext &llvmContext)
 Translates the 'getOperation()' result to an LLVM module. More...
 
- Protected Member Functions inherited from mlir::OperationPass< gpu::GPUModuleOp >
 OperationPass (TypeID passID)
 
 OperationPass (const OperationPass &)=default
 
bool canScheduleOn (RegisteredOperationName opName) const final
 Indicate if the current pass can be scheduled on the given operation type. More...
 
gpu::GPUModuleOp getOperation ()
 Return the current operation being transformed. More...
 
AnalysisT & getAnalysis ()
 Query an analysis for the current operation of the specific derived operation type. More...
 
- Protected Member Functions inherited from mlir::Pass
 Pass (TypeID passID, Optional< StringRef > opName=llvm::None)
 
 Pass (const Pass &other)
 
detail::PassExecutionStategetPassState ()
 Returns the current pass state. More...
 
MLIRContextgetContext ()
 Return the MLIR context for the current operation being transformed. More...
 
virtual LogicalResult initialize (MLIRContext *context)
 Initialize any complex state necessary for running this pass. More...
 
LogicalResult runPipeline (OpPassManager &pipeline, Operation *op)
 Schedule an arbitrary pass pipeline on the provided operation. More...
 
std::unique_ptr< Passclone () const
 A clone method to create a copy of this pass. More...
 
OperationgetOperation ()
 Return the current operation being transformed. More...
 
void signalPassFailure ()
 Signal that some invariant was broken when running. More...
 
template<typename AnalysisT >
AnalysisT & getAnalysis ()
 Query an analysis for the current ir unit. More...
 
template<typename AnalysisT , typename OpT >
AnalysisT & getAnalysis ()
 Query an analysis for the current ir unit of a specific derived operation type. More...
 
template<typename AnalysisT >
Optional< std::reference_wrapper< AnalysisT > > getCachedAnalysis ()
 Query a cached instance of an analysis for the current ir unit if one exists. More...
 
void markAllAnalysesPreserved ()
 Mark all analyses as preserved. More...
 
template<typename... AnalysesT>
void markAnalysesPreserved ()
 Mark the provided analyses as preserved. More...
 
void markAnalysesPreserved (TypeID id)
 
template<typename AnalysisT >
Optional< std::reference_wrapper< AnalysisT > > getCachedParentAnalysis (Operation *parent)
 Returns the analysis for the given parent operation if it exists. More...
 
template<typename AnalysisT >
Optional< std::reference_wrapper< AnalysisT > > getCachedParentAnalysis ()
 Returns the analysis for the parent operation if it exists. More...
 
template<typename AnalysisT >
Optional< std::reference_wrapper< AnalysisT > > getCachedChildAnalysis (Operation *child)
 Returns the analysis for the given child operation if it exists. More...
 
template<typename AnalysisT >
AnalysisT & getChildAnalysis (Operation *child)
 Returns the analysis for the given child operation, or creates it if it doesn't exist. More...
 
template<typename AnalysisT , typename OpTy >
AnalysisT & getChildAnalysis (OpTy child)
 Returns the analysis for the given child operation of specific derived operation type, or creates it if it doesn't exist. More...
 
AnalysisManager getAnalysisManager ()
 Returns the current analysis manager. More...
 
virtual std::unique_ptr< PassclonePass () const =0
 Create a copy of this pass, ignoring statistics and options. More...
 
void copyOptionValuesFrom (const Pass *other)
 Copy the option values from 'other', which is another instance of this pass. More...
 

Protected Attributes

Option< std::string > triple
 
Option< std::string > chip
 
Option< std::string > features
 
Option< std::string > gpuBinaryAnnotation
 

Additional Inherited Members

- Static Public Member Functions inherited from mlir::Pass
static const PassInfolookupPassInfo (StringRef passArg)
 Returns the pass info for the specified pass class or null if unknown. More...
 
- Static Protected Member Functions inherited from mlir::OperationPass< gpu::GPUModuleOp >
static bool classof (const Pass *pass)
 Support isa/dyn_cast functionality. More...
 

Detailed Description

Base pass class to serialize kernel functions through LLVM into user-specified IR and add the resulting blob as module attribute.

Definition at line 63 of file Passes.h.

Constructor & Destructor Documentation

◆ SerializeToBlobPass() [1/2]

gpu::SerializeToBlobPass::SerializeToBlobPass ( TypeID  passID)

Definition at line 32 of file SerializeToBlob.cpp.

◆ SerializeToBlobPass() [2/2]

gpu::SerializeToBlobPass::SerializeToBlobPass ( const SerializeToBlobPass other)

Definition at line 35 of file SerializeToBlob.cpp.

References mlir::failed(), and optimizeLlvm().

Member Function Documentation

◆ getDependentDialects()

void gpu::SerializeToBlobPass::getDependentDialects ( DialectRegistry registry) const
overrideprotectedvirtual

Register dependent dialects for the current pass.

A pass is expected to register the dialects it will create entities for (Operations, Types, Attributes), other than dialect that exists in the input. For example, a pass that converts from Linalg to Affine would register the Affine dialect but does not need to register Linalg.

Reimplemented from mlir::Pass.

Definition at line 108 of file SerializeToBlob.cpp.

References chip, mlir::emitError(), features, mlir::Pass::getDependentDialects(), mlir::OperationPass< gpu::GPUModuleOp >::getOperation(), mlir::registerLLVMDialectTranslation(), and triple.

◆ optimizeLlvm()

LogicalResult gpu::SerializeToBlobPass::optimizeLlvm ( llvm::Module &  llvmModule,
llvm::TargetMachine &  targetMachine 
)
protectedvirtual

Hook allowing the application of optimizations before codegen By default, does nothing.

Definition at line 101 of file SerializeToBlob.cpp.

References mlir::success().

Referenced by SerializeToBlobPass().

◆ runOnOperation()

void gpu::SerializeToBlobPass::runOnOperation ( )
finalvirtual

The polymorphic API that runs the pass over the currently held operation.

Implements mlir::Pass.

Definition at line 62 of file SerializeToBlob.cpp.

References mlir::Pass::getContext(), mlir::OperationPass< gpu::GPUModuleOp >::getOperation(), gpuBinaryAnnotation, mlir::Pass::signalPassFailure(), and translateToLLVMIR().

◆ translateToLLVMIR()

std::unique_ptr< llvm::Module > gpu::SerializeToBlobPass::translateToLLVMIR ( llvm::LLVMContext &  llvmContext)
protectedvirtual

Translates the 'getOperation()' result to an LLVM module.

Definition at line 135 of file SerializeToBlob.cpp.

References mlir::OperationPass< gpu::GPUModuleOp >::getOperation(), and mlir::translateModuleToLLVMIR().

Referenced by runOnOperation().

Member Data Documentation

◆ chip

Option<std::string> mlir::gpu::SerializeToBlobPass::chip
protected
Initial value:
{*this, "chip",
::llvm::cl::desc("Target architecture")}

Definition at line 97 of file Passes.h.

Referenced by getDependentDialects().

◆ features

Option<std::string> mlir::gpu::SerializeToBlobPass::features
protected
Initial value:
{*this, "features",
::llvm::cl::desc("Target features")}

Definition at line 99 of file Passes.h.

Referenced by getDependentDialects().

◆ gpuBinaryAnnotation

Option<std::string> mlir::gpu::SerializeToBlobPass::gpuBinaryAnnotation
protected
Initial value:
{
*this, "gpu-binary-annotation",
llvm::cl::desc("Annotation attribute string for GPU binary"),
llvm::cl::init(getDefaultGpuBinaryAnnotation())}

Definition at line 101 of file Passes.h.

Referenced by runOnOperation().

◆ triple

Option<std::string> mlir::gpu::SerializeToBlobPass::triple
protected
Initial value:
{*this, "triple",
::llvm::cl::desc("Target triple")}

Definition at line 95 of file Passes.h.

Referenced by getDependentDialects().


The documentation for this class was generated from the following files: