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::Pass
virtual	~Pass ()=default

TypeID	getTypeID () const
	Returns the unique identifier that corresponds to this pass. More...

const PassInfo *	lookupPassInfo () 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 Pass *	getThreadingSibling () const
	Returns the thread sibling of this pass. More...

const Pass *	getThreadingSiblingOrThis () 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

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)

detail::PassExecutionState &	getPassState ()
	Returns the current pass state. More...

MLIRContext &	getContext ()
	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< Pass >	clone () const
	A clone method to create a copy of this pass. More...

Operation *	getOperation ()
	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< Pass >	clonePass () 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 82 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 114 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 125 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 116 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())}

Definition at line 121 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 118 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 112 of file Passes.h.

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

include/mlir/Dialect/GPU/Transforms/Passes.h
lib/Dialect/GPU/Transforms/SerializeToBlob.cpp

Public Member Functions

Protected Member Functions

Protected Attributes

Additional Inherited Members

Detailed Description

Constructor & Destructor Documentation

◆ SerializeToBlobPass() [1/2]

◆ SerializeToBlobPass() [2/2]

Member Function Documentation

◆ optimizeLlvm()

◆ runOnOperation()

◆ translateToLLVMIR()

Member Data Documentation

◆ chip

◆ dumpPtx

◆ features

◆ gpuBinaryAnnotation

◆ optLevel

◆ triple