Passes.h File Reference

#include "mlir/Dialect/AMDGPU/Utils/Chipset.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/GPU/Utils/GPUUtils.h"
#include "mlir/IR/PatternMatch.h"
#include "mlir/Pass/Pass.h"
#include <optional>
#include "mlir/Dialect/GPU/Transforms/Passes.h.inc"

Go to the source code of this file.

Namespaces
namespace	llvm
	The OpAsmOpInterface, see OpAsmInterface.td for more details.
namespace	mlir
	Include the generated interface declarations.
namespace	mlir::func
namespace	mlir::gpu

Macros
#define	GEN_PASS_DECL
#define	GEN_PASS_REGISTRATION
	Generate the code for registering passes.

Functions
void	mlir::populateGpuGlobalIdPatterns (RewritePatternSet &patterns)
	Collect a set of patterns to rewrite GlobalIdOp op within the GPU dialect.
void	mlir::populateGpuSubgroupIdPatterns (RewritePatternSet &patterns)
	Collect a set of patterns to rewrite SubgroupIdOp op within the GPU dialect.
void	mlir::populateGpuShufflePatterns (RewritePatternSet &patterns)
	Collect a set of patterns to rewrite shuffle ops within the GPU dialect.
void	mlir::populateGpuAllReducePatterns (RewritePatternSet &patterns)
	Collect a set of patterns to rewrite all-reduce ops within the GPU dialect.
void	mlir::populateGpuBreakDownSubgroupReducePatterns (RewritePatternSet &patterns, unsigned maxShuffleBitwidth=32, PatternBenefit benefit=1)
	Collect a set of patterns to break down subgroup_reduce ops into smaller ones supported by the target of size <= maxShuffleBitwidth, where size is the subgroup_reduce value bitwidth.
void	mlir::populateGpuLowerSubgroupReduceToShufflePatterns (RewritePatternSet &patterns, unsigned subgroupSize, unsigned shuffleBitwidth=32, PatternBenefit benefit=1)
	Collect a set of patterns to lower gpu.subgroup_reduce into gpu.shuffle ops over shuffleBitwidth scalar types.
void	mlir::populateGpuLowerClusteredSubgroupReduceToShufflePatterns (RewritePatternSet &patterns, unsigned subgroupSize, unsigned shuffleBitwidth=32, PatternBenefit benefit=1)
	Disjoint counterpart of populateGpuLowerSubgroupReduceToShufflePatterns that only matches gpu.subgroup_reduce ops with a cluster_size.
void	mlir::populateGpuLowerSubgroupReduceToDPPPatterns (RewritePatternSet &patterns, unsigned subgroupSize, amdgpu::Chipset chipset, PatternBenefit benefit=1)
	Collect a set of patterns to lower gpu.subgroup_reduce into amdgpu.dpp ops over scalar types.
void	mlir::populateGpuLowerClusteredSubgroupReduceToDPPPatterns (RewritePatternSet &patterns, unsigned subgroupSize, amdgpu::Chipset chipset, PatternBenefit benefit=1)
	Disjoint counterpart of populateGpuLowerSubgroupReduceToDPPPatterns that only matches gpu.subgroup_reduce ops with a cluster_size.
void	mlir::populateGpuRewritePatterns (RewritePatternSet &patterns)
	Collect all patterns to rewrite ops within the GPU dialect.
LogicalResult	mlir::gpu::transformGpuModulesToBinaries (Operation *op, OffloadingLLVMTranslationAttrInterface handler=nullptr, const gpu::TargetOptions &options={})
	Searches for all GPU modules in op and transforms them into GPU binary operations.
void	mlir::populateGpuDecomposeMemrefsPatterns (RewritePatternSet &patterns)
	Collect a set of patterns to decompose memrefs ops.
void	mlir::populateGpuEliminateBarriersPatterns (RewritePatternSet &patterns)
	Erase barriers that do not enforce conflicting memory side effects.
void	mlir::populateGpuPromoteShuffleToAMDGPUPatterns (RewritePatternSet &patterns, std::optional< amdgpu::Chipset > maybeChipset)
	Tries to promote gpu.shuffles to specialized AMDGPU intrinsics.

Macro Definition Documentation

GEN_PASS_DECL

#define GEN_PASS_DECL

Definition at line 36 of file Passes.h.

GEN_PASS_REGISTRATION

#define GEN_PASS_REGISTRATION

Generate the code for registering passes.

Definition at line 121 of file Passes.h.