Base class for all NVVM serializations from GPU modules into binary strings. More...
#include "mlir/Target/LLVM/NVVM/Utils.h"
Inheritance diagram for mlir::NVVM::SerializeGPUModuleBase:
Public Member Functions | |
| SerializeGPUModuleBase (Operation &module, NVVMTargetAttr target, const gpu::TargetOptions &targetOptions={}) | |
Initializes the toolkitPath with the path in targetOptions or if empty with the path in getCUDAToolkitPath. More... | |
| NVVMTargetAttr | getTarget () const |
| Returns the target attribute. More... | |
| StringRef | getToolkitPath () const |
| Returns the CUDA toolkit path. More... | |
| ArrayRef< Attribute > | getLibrariesToLink () const |
| Returns the bitcode libraries to be linked into the gpu module after translation to LLVM IR. More... | |
| LogicalResult | appendStandardLibs () |
Appends nvvm/libdevice.bc into librariesToLink. More... | |
| virtual std::optional< SmallVector< std::unique_ptr< llvm::Module > > > | loadBitcodeFiles (llvm::Module &module) override |
Loads the bitcode files in librariesToLink. 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 NVPTX target by safely calling LLVMInitializeNVPTX* methods if available. More... | |
Protected Attributes | |
| NVVMTargetAttr | target |
| NVVM target attribute. More... | |
| std::string | toolkitPath |
| CUDA toolkit path. More... | |
| SmallVector< Attribute > | librariesToLink |
| List of LLVM bitcode to link into after translation to LLVM IR. 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 | |
| Protected Member Functions inherited from mlir::LLVM::ModuleToObject | |
| virtual void | setDataLayoutAndTriple (llvm::Module &module) |
| Hook for computing the Datalayout. More... | |
| virtual LogicalResult | handleBitcodeFile (llvm::Module &module) |
| Hook for performing additional actions on a loaded bitcode file. More... | |
| virtual void | handleModulePreLink (llvm::Module &module) |
| Hook for performing additional actions on the llvmModule pre linking. 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... | |
| 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 NVVM serializations from GPU modules into binary strings.
By default this class serializes into LLVM bitcode.
Constructor & Destructor Documentation
◆ SerializeGPUModuleBase()
| SerializeGPUModuleBase::SerializeGPUModuleBase | ( | Operation & | module, |
| NVVMTargetAttr | target, | ||
| const gpu::TargetOptions & | targetOptions = {} |
||
| ) |
Initializes the toolkitPath with the path in targetOptions or if empty with the path in getCUDAToolkitPath.
Definition at line 94 of file Target.cpp.
References appendStandardLibs(), mlir::NVVM::getCUDAToolkitPath(), librariesToLink, target, and toolkitPath.
Member Function Documentation
◆ appendStandardLibs()
| LogicalResult SerializeGPUModuleBase::appendStandardLibs | ( | ) |
Appends nvvm/libdevice.bc into librariesToLink.
Returns failure if the library couldn't be found.
Definition at line 140 of file Target.cpp.
References _mlir_embedded_libdevice, _mlir_embedded_libdevice_size, mlir::UnmanagedAsmResourceBlob::allocateInferAlign(), mlir::Operation::emitError(), mlir::get(), mlir::MLIRContext::getLoadedDialect(), mlir::DialectResourceBlobHandle< DialectT >::getManagerInterface(), mlir::LLVM::ModuleToObject::getOperation(), getToolkitPath(), librariesToLink, and target.
Referenced by SerializeGPUModuleBase().
◆ getLibrariesToLink()
Returns the bitcode libraries to be linked into the gpu module after translation to LLVM IR.
Definition at line 135 of file Target.cpp.
References librariesToLink.
◆ getTarget()
| NVVMTargetAttr SerializeGPUModuleBase::getTarget | ( | ) | const |
◆ getToolkitPath()
| StringRef SerializeGPUModuleBase::getToolkitPath | ( | ) | const |
Returns the CUDA toolkit path.
Definition at line 133 of file Target.cpp.
References toolkitPath.
Referenced by appendStandardLibs().
◆ init()
|
static |
Initializes the LLVM NVPTX target by safely calling LLVMInitializeNVPTX* methods if available.
Definition at line 118 of file Target.cpp.
◆ loadBitcodeFiles()
|
overridevirtual |
Loads the bitcode files in librariesToLink.
Reimplemented from mlir::LLVM::ModuleToObject.
Definition at line 192 of file Target.cpp.
References mlir::Operation::getContext(), librariesToLink, mlir::LLVM::ModuleToObject::loadBitcodeFilesFromList(), and mlir::LLVM::ModuleToObject::module.
Member Data Documentation
◆ librariesToLink
|
protected |
List of LLVM bitcode to link into after translation to LLVM IR.
The attributes can be StringAttr pointing to a file path, or a Resource blob pointing to the LLVM bitcode in-memory.
Definition at line 71 of file Utils.h.
Referenced by appendStandardLibs(), getLibrariesToLink(), loadBitcodeFiles(), and SerializeGPUModuleBase().
◆ target
|
protected |
NVVM target attribute.
Definition at line 63 of file Utils.h.
Referenced by appendStandardLibs(), getTarget(), and SerializeGPUModuleBase().
◆ toolkitPath
|
protected |
CUDA toolkit path.
Definition at line 66 of file Utils.h.
Referenced by getToolkitPath(), and SerializeGPUModuleBase().
The documentation for this class was generated from the following files:
- include/mlir/Target/LLVM/NVVM/Utils.h
- lib/Target/LLVM/NVVM/Target.cpp
