Base class for all ROCDL serializations from GPU modules into binary strings. More...

#include "mlir/Target/LLVM/ROCDL/Utils.h"

Inheritance diagram for mlir::ROCDL::SerializeGPUModuleBase:

Public Member Functions
	SerializeGPUModuleBase (Operation &module, ROCDLTargetAttr target, const gpu::TargetOptions &targetOptions={})
	Initializes the `toolkitPath` with the path in `targetOptions` or if empty with the path in `getROCMPath`. More...

ROCDLTargetAttr	getTarget () const
	Returns the target attribute. More...

StringRef	getToolkitPath () const
	Returns the ROCM toolkit path. More...

ArrayRef< std::string >	getFileList () const
	Returns the bitcode files to be loaded. More...

LogicalResult	appendStandardLibs ()
	Appends standard ROCm device libraries like `ocml.bc`, `ockl.bc`, etc. More...

virtual std::optional< SmallVector< std::unique_ptr< llvm::Module > > >	loadBitcodeFiles (llvm::Module &module) override
	Loads the bitcode files in `fileList`. More...

void	handleModulePreLink (llvm::Module &module) override
	Adds `oclc` control variables to the LLVM module. More...

LogicalResult	handleBitcodeFile (llvm::Module &module) override
	Removes unnecessary metadata from the loaded bitcode files. More...

Public Member Functions inherited from mlir::LLVM::ModuleToObject
	ModuleToObject (Operation &module, StringRef triple, StringRef chip, StringRef features={}, int optLevel=3)

virtual	~ModuleToObject ()

Operation &	getOperation ()
	Returns the operation being serialized. More...

virtual std::optional< SmallVector< char, 0 > >	run ()
	Runs the serialization pipeline, returning `std::nullopt` on error. More...

Static Public Member Functions
static void	init ()
	Initializes the LLVM AMDGPU target by safely calling `LLVMInitializeAMDGPU*` methods if available. More...

Protected Member Functions
LogicalResult	getCommonBitcodeLibs (llvm::SmallVector< std::string > &libs, SmallVector< char, 256 > &libPath, StringRef isaVersion)
	Appends the paths of common ROCm device libraries to `libs`. More...

void	addControlVariables (llvm::Module &module, bool wave64, bool daz, bool finiteOnly, bool unsafeMath, bool fastMath, bool correctSqrt, StringRef abiVer)
	Adds `oclc` control variables to the LLVM module. More...

std::optional< SmallVector< char, 0 > >	assembleIsa (StringRef isa)
	Returns the assembled ISA. More...

Protected Member Functions inherited from mlir::LLVM::ModuleToObject
virtual void	setDataLayoutAndTriple (llvm::Module &module)
	Hook for computing the Datalayout. More...

virtual void	handleModulePostLink (llvm::Module &module)
	Hook for performing additional actions on the llvmModule post linking. More...

virtual std::optional< SmallVector< char, 0 > >	moduleToObject (llvm::Module &llvmModule)
	Serializes the LLVM IR bitcode to an object file, by default it serializes to LLVM bitcode. More...

std::optional< llvm::TargetMachine * >	getOrCreateTargetMachine ()
	Create the target machine based on the target triple and chip. More...

std::unique_ptr< llvm::Module >	loadBitcodeFile (llvm::LLVMContext &context, StringRef path)
	Loads a bitcode file from path. More...

LogicalResult	loadBitcodeFilesFromList (llvm::LLVMContext &context, ArrayRef< std::string > fileList, SmallVector< std::unique_ptr< llvm::Module >> &llvmModules, bool failureOnError=true)
	Loads multiple bitcode files. More...

std::unique_ptr< llvm::Module >	translateToLLVMIR (llvm::LLVMContext &llvmContext)
	Translates the operation to LLVM IR. More...

LogicalResult	linkFiles (llvm::Module &module, SmallVector< std::unique_ptr< llvm::Module >> &&libs)
	Link the llvmModule to other bitcode file. More...

virtual LogicalResult	optimizeModule (llvm::Module &module, int optL)
	Optimize the module. More...

Protected Attributes
ROCDLTargetAttr	target
	ROCDL target attribute. More...

std::string	toolkitPath
	ROCM toolkit path. More...

SmallVector< std::string >	fileList
	List of LLVM bitcode files to link to. More...

Protected Attributes inherited from mlir::LLVM::ModuleToObject
Operation &	module
	Module to transform to a binary object. More...

StringRef	triple
	Target triple. More...

StringRef	chip
	Target chip. More...

StringRef	features
	Target features. More...

int	optLevel
	Optimization level. More...

Additional Inherited Members
Static Protected Member Functions inherited from mlir::LLVM::ModuleToObject
static std::optional< std::string >	translateToISA (llvm::Module &llvmModule, llvm::TargetMachine &targetMachine)
	Utility function for translating to ISA, returns `std::nullopt` on failure. More...

Detailed Description

Base class for all ROCDL serializations from GPU modules into binary strings.

By default this class serializes into LLVM bitcode.

Definition at line 32 of file Utils.h.

Constructor & Destructor Documentation

◆ SerializeGPUModuleBase()

SerializeGPUModuleBase::SerializeGPUModuleBase	(	Operation &	module,
		ROCDLTargetAttr	target,
		const gpu::TargetOptions &	targetOptions = `{}`
	)

Initializes the toolkitPath with the path in targetOptions or if empty with the path in getROCMPath.

Definition at line 97 of file Target.cpp.

References appendStandardLibs(), fileList, mlir::ROCDL::getROCMPath(), target, and toolkitPath.

Member Function Documentation

◆ addControlVariables()

void SerializeGPUModuleBase::addControlVariables	(	llvm::Module &	module,
		bool	wave64,
		bool	daz,
		bool	finiteOnly,
		bool	unsafeMath,
		bool	fastMath,
		bool	correctSqrt,
		StringRef	abiVer
	)

protected

Adds oclc control variables to the LLVM module.

Definition at line 219 of file Target.cpp.

References mlir::get(), mlir::Operation::getContext(), and mlir::LLVM::ModuleToObject::module.

Referenced by handleModulePreLink().

◆ appendStandardLibs()

LogicalResult SerializeGPUModuleBase::appendStandardLibs ( )

Appends standard ROCm device libraries like ocml.bc, ockl.bc, etc.

Definition at line 141 of file Target.cpp.

References mlir::LLVM::ModuleToObject::chip, mlir::Operation::emitRemark(), mlir::failure(), fileList, getCommonBitcodeLibs(), mlir::LLVM::ModuleToObject::getOperation(), getToolkitPath(), and mlir::success().

Referenced by SerializeGPUModuleBase().

◆ assembleIsa()

std::optional< SmallVector< char, 0 > > SerializeGPUModuleBase::assembleIsa ( StringRef isa )

protected

Returns the assembled ISA.

Definition at line 254 of file Target.cpp.

References mlir::LLVM::ModuleToObject::chip, mlir::emitError(), mlir::LLVM::ModuleToObject::features, mlir::Operation::getLoc(), mlir::LLVM::ModuleToObject::getOperation(), target, and mlir::LLVM::ModuleToObject::triple.

◆ getCommonBitcodeLibs()

LogicalResult SerializeGPUModuleBase::getCommonBitcodeLibs	(	llvm::SmallVector< std::string > &	libs,
		SmallVector< char, 256 > &	libPath,
		StringRef	isaVersion
	)

protected

Appends the paths of common ROCm device libraries to libs.

Definition at line 191 of file Target.cpp.

References mlir::Operation::emitRemark(), mlir::failure(), mlir::LLVM::ModuleToObject::getOperation(), and mlir::success().

Referenced by appendStandardLibs().

◆ getFileList()

ArrayRef< std::string > SerializeGPUModuleBase::getFileList ( ) const

Returns the bitcode files to be loaded.

Definition at line 137 of file Target.cpp.

References fileList.

◆ getTarget()

ROCDLTargetAttr SerializeGPUModuleBase::getTarget ( ) const

Returns the target attribute.

Definition at line 133 of file Target.cpp.

References target.

◆ getToolkitPath()

StringRef SerializeGPUModuleBase::getToolkitPath ( ) const

Returns the ROCM toolkit path.

Definition at line 135 of file Target.cpp.

References toolkitPath.

Referenced by appendStandardLibs().

◆ handleBitcodeFile()

LogicalResult SerializeGPUModuleBase::handleBitcodeFile ( llvm::Module & module )

overridevirtual

Removes unnecessary metadata from the loaded bitcode files.

Reimplemented from mlir::LLVM::ModuleToObject.

Definition at line 170 of file Target.cpp.

References mlir::LLVM::ModuleToObject::module, and mlir::success().

◆ handleModulePreLink()

void SerializeGPUModuleBase::handleModulePreLink ( llvm::Module & module )

overridevirtual

Adds oclc control variables to the LLVM module.

Reimplemented from mlir::LLVM::ModuleToObject.

Definition at line 180 of file Target.cpp.

References addControlVariables(), mlir::LLVM::ModuleToObject::getOrCreateTargetMachine(), mlir::LLVM::ModuleToObject::module, and target.

◆ init()

void SerializeGPUModuleBase::init ( )

static

Initializes the LLVM AMDGPU target by safely calling LLVMInitializeAMDGPU* methods if available.

Definition at line 119 of file Target.cpp.

◆ loadBitcodeFiles()

std::optional< SmallVector< std::unique_ptr< llvm::Module > > > SerializeGPUModuleBase::loadBitcodeFiles ( llvm::Module & module )

overridevirtual

Loads the bitcode files in fileList.

Reimplemented from mlir::LLVM::ModuleToObject.

Definition at line 162 of file Target.cpp.

References mlir::failed(), fileList, mlir::Operation::getContext(), mlir::LLVM::ModuleToObject::loadBitcodeFilesFromList(), and mlir::LLVM::ModuleToObject::module.

Member Data Documentation

◆ fileList

SmallVector<std::string> mlir::ROCDL::SerializeGPUModuleBase::fileList

protected

List of LLVM bitcode files to link to.

Definition at line 86 of file Utils.h.

Referenced by appendStandardLibs(), getFileList(), loadBitcodeFiles(), and SerializeGPUModuleBase().

◆ target

ROCDLTargetAttr mlir::ROCDL::SerializeGPUModuleBase::target

protected

ROCDL target attribute.

Definition at line 80 of file Utils.h.

Referenced by assembleIsa(), getTarget(), handleModulePreLink(), and SerializeGPUModuleBase().

◆ toolkitPath

std::string mlir::ROCDL::SerializeGPUModuleBase::toolkitPath

protected

ROCM toolkit path.

Definition at line 83 of file Utils.h.

Referenced by getToolkitPath(), and SerializeGPUModuleBase().

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

include/mlir/Target/LLVM/ROCDL/Utils.h
lib/Target/LLVM/ROCDL/Target.cpp

Public Member Functions

Static Public Member Functions

Protected Member Functions

Protected Attributes

Additional Inherited Members

Detailed Description

Constructor & Destructor Documentation

◆ SerializeGPUModuleBase()

Member Function Documentation

◆ addControlVariables()

◆ appendStandardLibs()

◆ assembleIsa()

◆ getCommonBitcodeLibs()

◆ getFileList()

◆ getTarget()

◆ getToolkitPath()

◆ handleBitcodeFile()

◆ handleModulePreLink()

◆ init()

◆ loadBitcodeFiles()

Member Data Documentation

◆ fileList

◆ target

◆ toolkitPath