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< Attribute >	getLibrariesToLink () const
	Returns the LLVM bitcode libraries to be linked. More...

LogicalResult	appendStandardLibs (AMDGCNLibraries libs)
	Appends standard ROCm device libraries to `fileList`. 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
	Determines required Device Libraries and adds `oclc` control variables to the LLVM Module if needed. 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, function_ref< void(llvm::Module &)> initialLlvmIRCallback={}, function_ref< void(llvm::Module &)> linkedLlvmIRCallback={}, function_ref< void(llvm::Module &)> optimizedLlvmIRCallback={}, function_ref< void(StringRef)> isaCallback={})

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
void	addControlVariables (llvm::Module &module, AMDGCNLibraries libs, bool wave64, bool daz, bool finiteOnly, bool unsafeMath, bool fastMath, bool correctSqrt, StringRef abiVer)
	Adds `oclc` control variables to the LLVM Module if needed. More...

virtual std::optional< SmallVector< char, 0 > >	compileToBinary (const std::string &serializedISA)
	Compiles assembly to a binary. More...

std::optional< SmallVector< char, 0 > >	moduleToObjectImpl (const gpu::TargetOptions &targetOptions, llvm::Module &llvmModule)
	Default implementation of `ModuleToObject::moduleToObject`. 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< Attribute > librariesToLink, 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< Attribute >	librariesToLink
	List of LLVM bitcode files to link to. More...

AMDGCNLibraries	deviceLibs = AMDGCNLibraries::None
	AMD GCN libraries to use when linking, the default is using none. 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...

function_ref< void(llvm::Module &)>	initialLlvmIRCallback
	Callback invoked with the initial LLVM IR for the device module. More...

function_ref< void(llvm::Module &)>	linkedLlvmIRCallback
	Callback invoked with LLVM IR for the device module after linking the device libraries. More...

function_ref< void(llvm::Module &)>	optimizedLlvmIRCallback
	Callback invoked with LLVM IR for the device module after LLVM optimizations but before codegen. More...

function_ref< void(StringRef)>	isaCallback
	Callback invoked with the target ISA for the device, for example PTX assembly. 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 48 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 94 of file Target.cpp.

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

Member Function Documentation

◆ addControlVariables()

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

protected

Adds oclc control variables to the LLVM Module if needed.

It also sets amdhsa_code_object_version module flag which is equal to ABI version and it uses "llvm::Module::Error" to set that flag.

Definition at line 228 of file Target.cpp.

References mlir::LLVM::ModuleToObject::chip, Error, mlir::get(), mlir::Operation::getContext(), mlir::LLVM::ModuleToObject::module, mlir::ROCDL::None, mlir::ROCDL::Ockl, and mlir::ROCDL::Ocml.

Referenced by handleModulePreLink().

◆ appendStandardLibs()

LogicalResult SerializeGPUModuleBase::appendStandardLibs ( AMDGCNLibraries libs )

Appends standard ROCm device libraries to fileList.

Definition at line 133 of file Target.cpp.

References mlir::Operation::emitError(), mlir::Operation::emitRemark(), mlir::get(), mlir::LLVM::ModuleToObject::getOperation(), getToolkitPath(), mlir::ROCDL::Hip, librariesToLink, mlir::ROCDL::None, mlir::ROCDL::Ockl, mlir::ROCDL::Ocml, mlir::ROCDL::OpenCL, and target.

Referenced by loadBitcodeFiles().

◆ assembleIsa()

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

protected

Returns the assembled ISA.

Definition at line 280 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.

Referenced by compileToBinary().

◆ compileToBinary()

std::optional< SmallVector< char, 0 > > SerializeGPUModuleBase::compileToBinary ( const std::string & serializedISA )

protectedvirtual

Compiles assembly to a binary.

Definition at line 344 of file Target.cpp.

References assembleIsa(), mlir::Operation::emitError(), mlir::LLVM::ModuleToObject::getOperation(), and toolkitPath.

Referenced by moduleToObjectImpl().

◆ getLibrariesToLink()

ArrayRef< Attribute > SerializeGPUModuleBase::getLibrariesToLink ( ) const

Returns the LLVM bitcode libraries to be linked.

Definition at line 129 of file Target.cpp.

References librariesToLink.

◆ getTarget()

ROCDLTargetAttr SerializeGPUModuleBase::getTarget ( ) const

Returns the target attribute.

Definition at line 125 of file Target.cpp.

References target.

◆ getToolkitPath()

StringRef SerializeGPUModuleBase::getToolkitPath ( ) const

Returns the ROCM toolkit path.

Definition at line 127 of file Target.cpp.

References toolkitPath.

Referenced by appendStandardLibs(), and moduleToObjectImpl().

◆ handleBitcodeFile()

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

overridevirtual

Removes unnecessary metadata from the loaded bitcode files.

Reimplemented from mlir::LLVM::ModuleToObject.

Definition at line 190 of file Target.cpp.

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

◆ handleModulePreLink()

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

overridevirtual

Determines required Device Libraries and adds oclc control variables to the LLVM Module if needed.

Also sets amdhsa_code_object_version module flag.

Reimplemented from mlir::LLVM::ModuleToObject.

Definition at line 203 of file Target.cpp.

References addControlVariables(), mlir::ROCDL::All, deviceLibs, mlir::ROCDL::Hip, mlir::LLVM::ModuleToObject::module, mlir::ROCDL::Ockl, mlir::ROCDL::Ocml, mlir::ROCDL::OpenCL, and target.

◆ init()

void SerializeGPUModuleBase::init ( )

static

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

Definition at line 111 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 177 of file Target.cpp.

References appendStandardLibs(), deviceLibs, mlir::Operation::getContext(), librariesToLink, mlir::LLVM::ModuleToObject::loadBitcodeFilesFromList(), mlir::LLVM::ModuleToObject::module, and mlir::ROCDL::None.

◆ moduleToObjectImpl()

std::optional< SmallVector< char, 0 > > SerializeGPUModuleBase::moduleToObjectImpl	(	const gpu::TargetOptions &	targetOptions,
		llvm::Module &	llvmModule
	)

protected

Default implementation of ModuleToObject::moduleToObject.

Definition at line 403 of file Target.cpp.

References compileToBinary(), mlir::Operation::emitError(), mlir::gpu::TargetOptions::getCompilationTarget(), mlir::LLVM::ModuleToObject::getOperation(), mlir::LLVM::ModuleToObject::getOrCreateTargetMachine(), getToolkitPath(), mlir::LLVM::ModuleToObject::moduleToObject(), mlir::LLVM::ModuleToObject::translateToISA(), and mlir::LLVM::ModuleToObject::triple.

Member Data Documentation

◆ deviceLibs

AMDGCNLibraries mlir::ROCDL::SerializeGPUModuleBase::deviceLibs = AMDGCNLibraries::None

protected

AMD GCN libraries to use when linking, the default is using none.

Definition at line 114 of file Utils.h.

Referenced by handleModulePreLink(), and loadBitcodeFiles().

◆ librariesToLink

SmallVector<Attribute> mlir::ROCDL::SerializeGPUModuleBase::librariesToLink

protected

List of LLVM bitcode files to link to.

Definition at line 111 of file Utils.h.

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

◆ target

ROCDLTargetAttr mlir::ROCDL::SerializeGPUModuleBase::target

protected

ROCDL target attribute.

Definition at line 105 of file Utils.h.

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

◆ toolkitPath

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

protected

ROCM toolkit path.

Definition at line 108 of file Utils.h.

Referenced by compileToBinary(), 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()

◆ compileToBinary()

◆ getLibrariesToLink()

◆ getTarget()

◆ getToolkitPath()

◆ handleBitcodeFile()

◆ handleModulePreLink()

◆ init()

◆ loadBitcodeFiles()

◆ moduleToObjectImpl()

Member Data Documentation

◆ deviceLibs

◆ librariesToLink

◆ target

◆ toolkitPath