MLIR  19.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:

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::OperationPass< gpu::GPUModuleOp >
 ~OperationPass ()=default
 
- 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 void getDependentDialects (DialectRegistry &registry) const
 Register dependent dialects for the current pass. 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...
 
std::optional< StringRef > getOpName () const
 Returns the name of the operation that this pass operates on, or std::nullopt 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

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
 
 OperationPass (OperationPass &&)=delete
 
OperationPassoperator= (const OperationPass &)=delete
 
OperationPassoperator= (OperationPass &&)=delete
 
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, std::optional< StringRef > opName=std::nullopt)
 
 Pass (const Pass &other)
 
Passoperator= (const Pass &)=delete
 
 Pass (Pass &&)=delete
 
Passoperator= (Pass &&)=delete
 
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 >
std::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 >
std::optional< std::reference_wrapper< AnalysisT > > getCachedParentAnalysis (Operation *parent)
 Returns the analysis for the given parent operation if it exists. More...
 
template<typename AnalysisT >
std::optional< std::reference_wrapper< AnalysisT > > getCachedParentAnalysis ()
 Returns the analysis for the parent operation if it exists. More...
 
template<typename AnalysisT >
std::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< int > optLevel
 
Option< std::string > gpuBinaryAnnotation
 
Option< bool > dumpPtx
 

Additional Inherited Members

- 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 97 of file Passes.h.

Constructor & Destructor Documentation

◆ SerializeToBlobPass() [1/2]

gpu::SerializeToBlobPass::SerializeToBlobPass ( TypeID  passID)

Definition at line 37 of file SerializeToBlob.cpp.

◆ SerializeToBlobPass() [2/2]

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

Definition at line 40 of file SerializeToBlob.cpp.

Member Function Documentation

◆ 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 106 of file SerializeToBlob.cpp.

References mlir::makeOptimizingTransformer(), and mlir::success().

◆ runOnOperation()

void gpu::SerializeToBlobPass::runOnOperation ( )
finalvirtual

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

Implements mlir::Pass.

Definition at line 67 of file SerializeToBlob.cpp.

References mlir::get(), and getContext().

◆ translateToLLVMIR()

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

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

Definition at line 150 of file SerializeToBlob.cpp.

References mlir::translateModuleToLLVMIR().

Member Data Documentation

◆ chip

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

Definition at line 129 of file Passes.h.

◆ dumpPtx

Option<bool> mlir::gpu::SerializeToBlobPass::dumpPtx
protected
Initial value:
{*this, "dump-ptx",
::llvm::cl::desc("Dump generated PTX"),
llvm::cl::init(false)}

Definition at line 140 of file Passes.h.

◆ features

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

Definition at line 131 of file Passes.h.

◆ 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())}
std::string getDefaultGpuBinaryAnnotation()
Returns the default annotation name for GPU binary blobs.

Definition at line 136 of file Passes.h.

◆ optLevel

Option<int> mlir::gpu::SerializeToBlobPass::optLevel
protected
Initial value:
{*this, "opt-level",
llvm::cl::desc("Optimization level for compilation"),
llvm::cl::init(2)}

Definition at line 133 of file Passes.h.

◆ triple

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

Definition at line 127 of file Passes.h.


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