MLIR  22.0.0git
Classes | Typedefs | Enumerations | Functions | Variables
mlir::NVVM Namespace Reference

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 &registry)
 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 &registry)
 Register the NVVMInlinerInterface implementation of DialectInlinerInterface with the NVVM dialect. More...
 
void registerNVVMTargetInterfaceExternalModels (DialectRegistry &registry)
 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
 

Typedef Documentation

◆ IDArgPair

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.

Enumeration Type Documentation

◆ PTXRegisterMod

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.

Function Documentation

◆ countPlaceholderNumbers()

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().

◆ getCUDAToolkitPath()

StringRef mlir::NVVM::getCUDAToolkitPath ( )

Searches & returns the path CUDA toolkit path, the search order is:

  1. The CUDA_ROOT environment variable.
  2. The CUDA_HOME environment variable.
  3. The CUDA_PATH environment variable.
  4. The CUDA toolkit path detected by CMake.
  5. Returns an empty string.

Definition at line 88 of file Target.cpp.

References __DEFAULT_CUDATOOLKIT_PATH__.

Referenced by mlir::NVVM::SerializeGPUModuleBase::SerializeGPUModuleBase().

◆ inferMMAType()

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().

◆ operator!=() [1/2]

bool mlir::NVVM::operator!= ( NVVMMemorySpace  memSpace,
unsigned  as 
)
inline

Definition at line 44 of file NVVMDialect.h.

◆ operator!=() [2/2]

bool mlir::NVVM::operator!= ( unsigned  as,
NVVMMemorySpace  memSpace 
)
inline

Definition at line 41 of file NVVMDialect.h.

◆ operator<<()

llvm::raw_ostream& mlir::NVVM::operator<< ( llvm::raw_ostream &  os,
PTXRegisterMod  mod 
)
inline

Definition at line 38 of file BasicPtxBuilderInterface.h.

References Read, ReadWrite, and Write.

◆ operator==() [1/2]

bool mlir::NVVM::operator== ( NVVMMemorySpace  memSpace,
unsigned  as 
)
inline

Definition at line 38 of file NVVMDialect.h.

◆ operator==() [2/2]

bool mlir::NVVM::operator== ( unsigned  as,
NVVMMemorySpace  memSpace 
)
inline

Utility functions to compare NVVMMemorySpace with unsigned values.

Definition at line 35 of file NVVMDialect.h.

◆ registerConvertGpuToNVVMInterface()

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().

◆ registerInlinerInterface()

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().

◆ registerNVVMTargetInterfaceExternalModels() [1/2]

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().

◆ registerNVVMTargetInterfaceExternalModels() [2/2]

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().

Variable Documentation

◆ kSharedMemoryAlignmentBit

constexpr int mlir::NVVM::kSharedMemoryAlignmentBit = 128
constexpr

Definition at line 49 of file NVVMDialect.h.

Referenced by mlir::populateGpuToNVVMConversionPatterns().