mlir::amdgpu Namespace Reference

Namespaces
namespace	impl

Classes
struct	AmdgpuEmulateAtomicsPassOptions
struct	Chipset
	Represents the amdgpu gfx chipset version, e.g., gfx90a, gfx942, gfx1103. More...

Functions
void	populateCommonGPUTypeAndAttributeConversions (TypeConverter &typeConverter)
	Remap common GPU memory spaces (Workgroup, Private, etc) to LLVM address spaces.
ParseResult	parseMNKDimensionList (OpAsmParser &parser, IntegerAttr &m, IntegerAttr &n, IntegerAttr &k)
	Parser for the custom<MNKDimensionList> custom assembly format used by WMMAOp.
ParseResult	parseMNKDimensionList (OpAsmParser &parser, Operation *, IntegerAttr &m, IntegerAttr &n, IntegerAttr &k)
void	printMNKDimensionList (OpAsmPrinter &printer, IntegerAttr m, IntegerAttr n, IntegerAttr k)
	Printer for the custom<MNKDimensionList> custom assembly format used by WMMAOp.
void	printMNKDimensionList (OpAsmPrinter &printer, Operation *, IntegerAttr m, IntegerAttr n, IntegerAttr k)
int32_t	getGlobalPrefetchLLVMEncoding (amdgpu::LoadTemporalHint hint, amdgpu::Scope scope, bool isSpeculative)
void	registerMemoryAccessOpInterfacesExternalModels (DialectRegistry &registry)
std::unique_ptr<::mlir::Pass >	createAmdgpuEmulateAtomicsPass ()
std::unique_ptr<::mlir::Pass >	createAmdgpuEmulateAtomicsPass (AmdgpuEmulateAtomicsPassOptions options)
std::unique_ptr<::mlir::Pass >	createAmdgpuMaskedloadToLoadPass ()
std::unique_ptr<::mlir::Pass >	createAmdgpuResolveStridedMetadataPass ()
void	registerAmdgpuEmulateAtomicsPass ()
void	registerAmdgpuEmulateAtomicsPassPass ()
void	registerAmdgpuMaskedloadToLoadPass ()
void	registerAmdgpuMaskedloadToLoadPassPass ()
void	registerAmdgpuResolveStridedMetadataPass ()
void	registerAmdgpuResolveStridedMetadataPassPass ()
void	registerAMDGPUPasses ()
void	populateAmdgpuEmulateAtomicsPatterns (ConversionTarget &target, RewritePatternSet &patterns, const ROCDL::TargetInfo &targetInfo, PatternBenefit benefit=1)
void	populateAmdgpuResolveStridedMetadataPatterns (RewritePatternSet &patterns, PatternBenefit benefit=1)
void	populateAmdgpuMaskedloadToLoadPatterns (RewritePatternSet &patterns, PatternBenefit benefit=1)
	LLVM_DEPRECATED_WITH_FIXIT ("Chipset is deperecated", "test FEAT_OCP_FP8_CONVERSION_INSTS on a ROCDL::TargetInfo") inline bool hasOcpFp8(const Chipset &chipset)
bool	isGlobalMemorySpace (Attribute memorySpace, bool allowFlat)
bool	isWorkgroupMemorySpace (Attribute memorySpace)
bool	isFatRawBufferMemorySpace (Attribute memorySpace)

Function Documentation

createAmdgpuEmulateAtomicsPass() [1/2]

std::unique_ptr<::mlir::Pass > mlir::amdgpu::createAmdgpuEmulateAtomicsPass

(

)

Definition at line 102 of file EmulateAtomics.cpp.

References result.

createAmdgpuEmulateAtomicsPass() [2/2]

std::unique_ptr<::mlir::Pass > mlir::amdgpu::createAmdgpuEmulateAtomicsPass

(

AmdgpuEmulateAtomicsPassOptions

options

)

Definition at line 106 of file EmulateAtomics.cpp.

createAmdgpuMaskedloadToLoadPass()

std::unique_ptr<::mlir::Pass > mlir::amdgpu::createAmdgpuMaskedloadToLoadPass

(

)

We declare an explicit private instantiation because Pass classes should only be visible by the current library.

Definition at line 182 of file MaskedloadToLoad.cpp.

createAmdgpuResolveStridedMetadataPass()

std::unique_ptr<::mlir::Pass > mlir::amdgpu::createAmdgpuResolveStridedMetadataPass

(

)

We declare an explicit private instantiation because Pass classes should only be visible by the current library.

Definition at line 258 of file ResolveStridedMetadata.cpp.

getGlobalPrefetchLLVMEncoding()

int32_t mlir::amdgpu::getGlobalPrefetchLLVMEncoding ( amdgpu::LoadTemporalHint hint, amdgpu::Scope scope, bool isSpeculative )

inline

Definition at line 18 of file AMDGPUEnums.h.

isFatRawBufferMemorySpace()

bool mlir::amdgpu::isFatRawBufferMemorySpace ( Attribute memorySpace )

inline

Definition at line 49 of file MemorySpaceUtils.h.

isGlobalMemorySpace()

bool mlir::amdgpu::isGlobalMemorySpace ( Attribute memorySpace, bool allowFlat )

inline

Definition at line 23 of file MemorySpaceUtils.h.

Referenced by verifyBase(), and verifyRawBufferOp().

isWorkgroupMemorySpace()

bool mlir::amdgpu::isWorkgroupMemorySpace ( Attribute memorySpace )

inline

Definition at line 36 of file MemorySpaceUtils.h.

Referenced by verifyBase(), verifyDescriptorOp(), and verifyDsBarrierOpCommon().

LLVM_DEPRECATED_WITH_FIXIT()

mlir::amdgpu::LLVM_DEPRECATED_WITH_FIXIT	(	"Chipset is deperecated"	,
		"test FEAT_OCP_FP8_CONVERSION_INSTS on a ROCDL::TargetInfo"	) const &

Deprecated

Test llvm::AMDGPU::FEAT_OCP_FP8_CONVERSION_INSTS on a ROCDL::TargetInfo instead. This misses gfx11.7, which does have the OCP fp8 conversions.

Definition at line 63 of file Chipset.h.

parseMNKDimensionList() [1/2]

ParseResult mlir::amdgpu::parseMNKDimensionList	(	OpAsmParser &	parser,
		IntegerAttr &	m,
		IntegerAttr &	n,
		IntegerAttr &	k )

Parser for the custom<MNKDimensionList> custom assembly format used by WMMAOp.

Definition at line 355 of file AMDGPUOps.cpp.

References mlir::AsmParser::emitError(), mlir::AsmParser::getBuilder(), mlir::AsmParser::getCurrentLocation(), mlir::Builder::getI32IntegerAttr(), mlir::AsmParser::parseDimensionList(), and success().

Referenced by parseMNKDimensionList().

parseMNKDimensionList() [2/2]

ParseResult mlir::amdgpu::parseMNKDimensionList ( OpAsmParser & parser, Operation * , IntegerAttr & m, IntegerAttr & n, IntegerAttr & k )

inline

Definition at line 37 of file AMDGPUDialect.h.

References parseMNKDimensionList().

populateAmdgpuEmulateAtomicsPatterns()

void mlir::amdgpu::populateAmdgpuEmulateAtomicsPatterns	(	ConversionTarget &	target,
		RewritePatternSet &	patterns,
		const ROCDL::TargetInfo &	targetInfo,
		PatternBenefit	benefit = 1 )

Definition at line 219 of file EmulateAtomics.cpp.

References mlir::RewritePatternSet::add(), mlir::RewritePatternSet::getContext(), mlir::ROCDL::TargetInfo::has(), isFaddNativelySupported(), mlir::ROCDL::TargetInfo::isGeneration(), and target.

populateAmdgpuMaskedloadToLoadPatterns()

void mlir::amdgpu::populateAmdgpuMaskedloadToLoadPatterns	(	RewritePatternSet &	patterns,
		PatternBenefit	benefit = 1 )

Definition at line 239 of file MaskedloadToLoad.cpp.

References mlir::RewritePatternSet::add(), and mlir::RewritePatternSet::getContext().

Referenced by AmdgpuMaskedloadToLoadPass::runOnOperation().

populateAmdgpuResolveStridedMetadataPatterns()

void mlir::amdgpu::populateAmdgpuResolveStridedMetadataPatterns	(	RewritePatternSet &	patterns,
		PatternBenefit	benefit = 1 )

Definition at line 68 of file ResolveStridedMetadata.cpp.

References mlir::RewritePatternSet::add(), and mlir::RewritePatternSet::getContext().

populateCommonGPUTypeAndAttributeConversions()

void mlir::amdgpu::populateCommonGPUTypeAndAttributeConversions

(

TypeConverter &

typeConverter

)

Remap common GPU memory spaces (Workgroup, Private, etc) to LLVM address spaces.

Definition at line 4559 of file AMDGPUToROCDL.cpp.

References mlir::populateGpuMemorySpaceAttributeConversions().

printMNKDimensionList() [1/2]

void mlir::amdgpu::printMNKDimensionList ( OpAsmPrinter & printer, IntegerAttr m, IntegerAttr n, IntegerAttr k )

inline

Printer for the custom<MNKDimensionList> custom assembly format used by WMMAOp.

Definition at line 45 of file AMDGPUDialect.h.

References mlir::AsmPrinter::printDimensionList().

Referenced by printMNKDimensionList().

printMNKDimensionList() [2/2]

void mlir::amdgpu::printMNKDimensionList ( OpAsmPrinter & printer, Operation * , IntegerAttr m, IntegerAttr n, IntegerAttr k )

inline

Definition at line 49 of file AMDGPUDialect.h.

References printMNKDimensionList().

registerAmdgpuEmulateAtomicsPass()

void mlir::amdgpu::registerAmdgpuEmulateAtomicsPass ( )

inline

Definition at line 275 of file Passes.h.

registerAmdgpuEmulateAtomicsPassPass()

void mlir::amdgpu::registerAmdgpuEmulateAtomicsPassPass ( )

inline

Definition at line 282 of file Passes.h.

registerAmdgpuMaskedloadToLoadPass()

void mlir::amdgpu::registerAmdgpuMaskedloadToLoadPass ( )

inline

Definition at line 296 of file Passes.h.

registerAmdgpuMaskedloadToLoadPassPass()

void mlir::amdgpu::registerAmdgpuMaskedloadToLoadPassPass ( )

inline

Definition at line 303 of file Passes.h.

registerAMDGPUPasses()

void mlir::amdgpu::registerAMDGPUPasses ( )

inline

Definition at line 338 of file Passes.h.

Referenced by mlir::registerAllPasses().

registerAmdgpuResolveStridedMetadataPass()

void mlir::amdgpu::registerAmdgpuResolveStridedMetadataPass ( )

inline

Definition at line 317 of file Passes.h.

registerAmdgpuResolveStridedMetadataPassPass()

void mlir::amdgpu::registerAmdgpuResolveStridedMetadataPassPass ( )

inline

Definition at line 324 of file Passes.h.

registerMemoryAccessOpInterfacesExternalModels()

void mlir::amdgpu::registerMemoryAccessOpInterfacesExternalModels

(

DialectRegistry &

registry

)

Definition at line 233 of file MemoryAccessOpInterfacesImpl.cpp.

References mlir::DialectRegistry::addExtension().

Referenced by mlir::registerAllDialects().