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 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 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 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
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::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 >
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< 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< std::string >	gpuBinaryAnnotation

Additional Inherited Members
Static Public Member Functions inherited from mlir::Pass
static const PassInfo *	lookupPassInfo (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:

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

◆ getDependentDialects()

◆ optimizeLlvm()

◆ runOnOperation()

◆ translateToLLVMIR()

Member Data Documentation

◆ chip

◆ features

◆ gpuBinaryAnnotation

◆ triple