MLIR
20.0.0git
|
#include "Utils.h"
#include "mlir/Dialect/GPU/IR/GPUDialect.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 | |
llvm | |
Include the generated interface declarations. | |
mlir | |
Include the generated interface declarations. | |
mlir::func | |
mlir::gpu | |
Macros | |
#define | GEN_PASS_DECL |
#define | GEN_PASS_REGISTRATION |
Generate the code for registering passes. More... | |
Functions | |
std::unique_ptr< Pass > | mlir::createGpuLauchSinkIndexComputationsPass () |
Pass that moves ops which are likely an index computation into gpu.launch body. More... | |
std::unique_ptr< OperationPass< ModuleOp > > | mlir::createGpuKernelOutliningPass (StringRef dataLayoutStr=StringRef()) |
Replaces gpu.launch with gpu.launch_func by moving the region into a separate kernel function. More... | |
std::unique_ptr< OperationPass< func::FuncOp > > | mlir::createGpuAsyncRegionPass () |
Rewrites a function region so that GPU ops execute asynchronously. More... | |
std::unique_ptr< OperationPass< func::FuncOp > > | mlir::createGpuMapParallelLoopsPass () |
Maps the parallel loops found in the given function to workgroups. More... | |
void | mlir::populateGpuGlobalIdPatterns (RewritePatternSet &patterns) |
Collect a set of patterns to rewrite GlobalIdOp op within the GPU dialect. More... | |
void | mlir::populateGpuShufflePatterns (RewritePatternSet &patterns) |
Collect a set of patterns to rewrite shuffle ops within the GPU dialect. More... | |
void | mlir::populateGpuAllReducePatterns (RewritePatternSet &patterns) |
Collect a set of patterns to rewrite all-reduce ops within the GPU dialect. More... | |
void | mlir::populateGpuBreakDownSubgrupReducePatterns (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. More... | |
void | mlir::populateGpuLowerSubgroupReduceToShufflePattenrs (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. More... | |
void | mlir::populateGpuRewritePatterns (RewritePatternSet &patterns) |
Collect all patterns to rewrite ops within the GPU dialect. More... | |
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. More... | |
void | mlir::populateGpuDecomposeMemrefsPatterns (RewritePatternSet &patterns) |
Collect a set of patterns to decompose memrefs ops. More... | |
std::unique_ptr< Pass > | mlir::createGpuDecomposeMemrefsPass () |
Pass decomposes memref ops inside gpu.launch body. More... | |
void | mlir::populateGpuEliminateBarriersPatterns (RewritePatternSet &patterns) |
Erase barriers that do not enforce conflicting memory side effects. More... | |