MLIR
22.0.0git
|
Classes | |
class | PtxBuilder |
A class to build PTX assembly automatically. More... | |
struct | NVVMCheckSMVersion |
class | SerializeGPUModuleBase |
Base class for all NVVM serializations from GPU modules into binary strings. More... | |
Typedefs | |
using | IDArgPair = std::pair< llvm::Intrinsic::ID, llvm::SmallVector< llvm::Value * > > |
A pair type of LLVM's Intrinsic ID and args (which are llvm values). More... | |
Enumerations | |
enum class | PTXRegisterMod { Read = 0 , Write = 2 , ReadWrite = 1 } |
Register read/write modifier to build constraint string for PTX inline https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html#parameters. More... | |
Functions | |
void | registerConvertGpuToNVVMInterface (DialectRegistry ®istry) |
Registers the ConvertToLLVMAttrInterface interface on the NVVM::NVVMTargetAttr attribute. More... | |
llvm::raw_ostream & | operator<< (llvm::raw_ostream &os, PTXRegisterMod mod) |
void | countPlaceholderNumbers (StringRef ptxCode, llvm::SmallDenseSet< unsigned > &seenRW, llvm::SmallDenseSet< unsigned > &seenW, llvm::SmallDenseSet< unsigned > &seenR, llvm::SmallVectorImpl< unsigned > &rwNums, llvm::SmallVectorImpl< unsigned > &wNums, llvm::SmallVectorImpl< unsigned > &rNums) |
Count the number of placeholder variables such as {$r}, {$w}, {$rw} in the PTX code. More... | |
bool | operator== (unsigned as, NVVMMemorySpace memSpace) |
Utility functions to compare NVVMMemorySpace with unsigned values. More... | |
bool | operator== (NVVMMemorySpace memSpace, unsigned as) |
bool | operator!= (unsigned as, NVVMMemorySpace memSpace) |
bool | operator!= (NVVMMemorySpace memSpace, unsigned as) |
std::pair< mlir::Type, unsigned > | inferMMAType (mlir::NVVM::MMATypes type, mlir::NVVM::MMAFrag frag, int nRow, int nCol, mlir::MLIRContext *context) |
Return the element type and number of elements associated with a wmma matrix of given chracteristics. More... | |
void | registerInlinerInterface (DialectRegistry ®istry) |
Register the NVVMInlinerInterface implementation of DialectInlinerInterface with the NVVM dialect. More... | |
void | registerNVVMTargetInterfaceExternalModels (DialectRegistry ®istry) |
Registers the TargetAttrInterface for the #nvvm.target attribute in the given registry. More... | |
void | registerNVVMTargetInterfaceExternalModels (MLIRContext &context) |
Registers the TargetAttrInterface for the #nvvm.target attribute in the registry associated with the given context. More... | |
StringRef | getCUDAToolkitPath () |
Searches & returns the path CUDA toolkit path, the search order is: More... | |
Variables | |
constexpr int | kSharedMemoryAlignmentBit = 128 |
using mlir::NVVM::IDArgPair = typedef std::pair<llvm::Intrinsic::ID, llvm::SmallVector<llvm::Value *> > |
A pair type of LLVM's Intrinsic ID and args (which are llvm values).
This type is returned by the getIntrinsicIDAndArgs() methods.
Definition at line 53 of file NVVMDialect.h.
|
strong |
Register read/write modifier to build constraint string for PTX inline https://docs.nvidia.com/cuda/inline-ptx-assembly/index.html#parameters.
Enumerator | |
---|---|
Read | Read register with no modifier. |
Write | Write register with '=' modifier. |
ReadWrite | ReadWrite register with '+' modifier. Note that, this is not natively supported by LLVM, the Interface does mapping |
Definition at line 27 of file BasicPtxBuilderInterface.h.
void mlir::NVVM::countPlaceholderNumbers | ( | StringRef | ptxCode, |
llvm::SmallDenseSet< unsigned > & | seenRW, | ||
llvm::SmallDenseSet< unsigned > & | seenW, | ||
llvm::SmallDenseSet< unsigned > & | seenR, | ||
llvm::SmallVectorImpl< unsigned > & | rwNums, | ||
llvm::SmallVectorImpl< unsigned > & | wNums, | ||
llvm::SmallVectorImpl< unsigned > & | rNums | ||
) |
Count the number of placeholder variables such as {$r}, {$w}, {$rw} in the PTX code.
Referenced by rewriteAsmPlaceholders().
StringRef mlir::NVVM::getCUDAToolkitPath | ( | ) |
Searches & returns the path CUDA toolkit path, the search order is:
CUDA_ROOT
environment variable.CUDA_HOME
environment variable.CUDA_PATH
environment variable.Definition at line 88 of file Target.cpp.
References __DEFAULT_CUDATOOLKIT_PATH__.
Referenced by mlir::NVVM::SerializeGPUModuleBase::SerializeGPUModuleBase().
std::pair<mlir::Type, unsigned> mlir::NVVM::inferMMAType | ( | mlir::NVVM::MMATypes | type, |
mlir::NVVM::MMAFrag | frag, | ||
int | nRow, | ||
int | nCol, | ||
mlir::MLIRContext * | context | ||
) |
Return the element type and number of elements associated with a wmma matrix of given chracteristics.
This matches the logic in IntrinsicsNVVM.td WMMA_REGS structure.
Referenced by mlir::convertMMAToLLVMType(), and inferMMATypeFromMNK().
|
inline |
Definition at line 44 of file NVVMDialect.h.
|
inline |
Definition at line 41 of file NVVMDialect.h.
|
inline |
Definition at line 38 of file BasicPtxBuilderInterface.h.
|
inline |
Definition at line 38 of file NVVMDialect.h.
|
inline |
Utility functions to compare NVVMMemorySpace with unsigned values.
Definition at line 35 of file NVVMDialect.h.
void mlir::NVVM::registerConvertGpuToNVVMInterface | ( | DialectRegistry & | registry | ) |
Registers the ConvertToLLVMAttrInterface
interface on the NVVM::NVVMTargetAttr
attribute.
This interface populates the conversion target, LLVM type converter, and pattern set for converting GPU operations to NVVM.
Definition at line 781 of file LowerGpuOpsToNVVMOps.cpp.
References mlir::DialectRegistry::addExtension().
Referenced by mlir::registerAllExtensions().
void mlir::NVVM::registerInlinerInterface | ( | DialectRegistry & | registry | ) |
Register the NVVMInlinerInterface
implementation of DialectInlinerInterface
with the NVVM dialect.
Definition at line 835 of file InlinerInterfaceImpl.cpp.
References mlir::DialectRegistry::addExtension().
void mlir::NVVM::registerNVVMTargetInterfaceExternalModels | ( | DialectRegistry & | registry | ) |
Registers the TargetAttrInterface
for the #nvvm.target
attribute in the given registry.
Definition at line 73 of file Target.cpp.
References mlir::DialectRegistry::addExtension().
Referenced by mlir::registerAllDialects(), and registerNVVMTargetInterfaceExternalModels().
void mlir::NVVM::registerNVVMTargetInterfaceExternalModels | ( | MLIRContext & | context | ) |
Registers the TargetAttrInterface
for the #nvvm.target
attribute in the registry associated with the given context.
Definition at line 80 of file Target.cpp.
References mlir::MLIRContext::appendDialectRegistry(), and registerNVVMTargetInterfaceExternalModels().
|
constexpr |
Definition at line 49 of file NVVMDialect.h.
Referenced by mlir::populateGpuToNVVMConversionPatterns().