|
MLIR
20.0.0git
|
#include "mlir/Dialect/GPU/Transforms/Passes.h"#include "mlir/AsmParser/AsmParser.h"#include "mlir/Dialect/Arith/IR/Arith.h"#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"#include "mlir/Dialect/DLTI/DLTI.h"#include "mlir/Dialect/Func/IR/FuncOps.h"#include "mlir/Dialect/GPU/IR/GPUDialect.h"#include "mlir/Dialect/GPU/Transforms/Utils.h"#include "mlir/Dialect/MemRef/IR/MemRef.h"#include "mlir/IR/Builders.h"#include "mlir/IR/BuiltinAttributes.h"#include "mlir/IR/IRMapping.h"#include "mlir/IR/Matchers.h"#include "mlir/IR/SymbolTable.h"#include "mlir/Support/LLVM.h"#include "mlir/Transforms/RegionUtils.h"#include <limits>#include "mlir/Dialect/GPU/Transforms/Passes.h.inc"Go to the source code of this file.
Namespaces | |
| mlir | |
| Include the generated interface declarations. | |
Macros | |
| #define | GEN_PASS_DEF_GPULAUNCHSINKINDEXCOMPUTATIONS |
| #define | GEN_PASS_DEF_GPUKERNELOUTLINING |
Functions | |
| template<typename OpTy > | |
| static void | createForAllDimensions (OpBuilder &builder, Location loc, SmallVectorImpl< Value > &values) |
| static void | injectGpuIndexOperations (Location loc, Region &launchFuncOpBody, Region &launchOpBody, IRMapping &map, bool hasCluster=false) |
Adds operations generating block/thread ids and grid/block dimensions at the beginning of the launchFuncOpBody region. More... | |
| static bool | isLikelyAnIndexComputation (Operation *op) |
| Identifies operations that are beneficial to sink into kernels. More... | |
| static bool | extractBeneficiaryOps (Operation *op, const SetVector< Value > &existingDependencies, SetVector< Operation * > &beneficiaryOps, llvm::SmallPtrSetImpl< Value > &availableValues, llvm::function_ref< bool(Operation *)> isSinkingBeneficiary) |
For a given operation op, computes whether it is beneficial to sink the operation into the kernel. More... | |
| static DenseI32ArrayAttr | maybeConstantDimsAttr (gpu::KernelDim3 dims) |
| Return the provided KernelDim3 as an array of i32 constants if possible. More... | |
| static gpu::GPUFuncOp | outlineKernelFuncImpl (gpu::LaunchOp launchOp, StringRef kernelFnName, SetVector< Value > &operands) |
Outline the gpu.launch operation body into a kernel function. More... | |
| static void | convertToLaunchFuncOp (gpu::LaunchOp launchOp, gpu::GPUFuncOp kernelFunc, ValueRange operands) |
Replace gpu.launch operations with an gpu.launch_func operation launching kernelFunc. More... | |
| #define GEN_PASS_DEF_GPUKERNELOUTLINING |
Definition at line 34 of file KernelOutlining.cpp.
| #define GEN_PASS_DEF_GPULAUNCHSINKINDEXCOMPUTATIONS |
Definition at line 33 of file KernelOutlining.cpp.
|
static |
Replace gpu.launch operations with an gpu.launch_func operation launching kernelFunc.
The kernel func contains the body of the gpu.launch with constant region arguments inlined.
Definition at line 282 of file KernelOutlining.cpp.
References mlir::OpBuilder::create(), mlir::Value::getType(), and mlir::Operation::replaceAllUsesWith().
|
static |
Definition at line 41 of file KernelOutlining.cpp.
References mlir::OpBuilder::create(), and mlir::Builder::getIndexType().
|
static |
For a given operation op, computes whether it is beneficial to sink the operation into the kernel.
An operation can be sunk if doing so does not introduce new kernel arguments. Whether a value is already available in the kernel (and hence does not introduce new arguments) is checked by querying existingDependencies and availableValues. If an operand is not yet available, we recursively check whether it can be made available by siking its defining op. Operations that are indentified for sinking are added to beneficiaryOps in the order they should appear in the kernel. Furthermore, availableValues is updated with results that will be available after sinking the identified ops.
Definition at line 92 of file KernelOutlining.cpp.
References mlir::Operation::getOperands(), and mlir::Operation::getResults().
Referenced by mlir::sinkOperationsIntoLaunchOp().
|
static |
Adds operations generating block/thread ids and grid/block dimensions at the beginning of the launchFuncOpBody region.
Add mapping from argument in entry block of launchOpBody, to the corresponding result value of the added operations.
Definition at line 51 of file KernelOutlining.cpp.
References mlir::detail::enumerate(), mlir::Region::front(), mlir::Block::getArgument(), mlir::Attribute::getContext(), mlir::IRMapping::map(), and mlir::OpBuilder::setInsertionPointToStart().
Referenced by outlineKernelFuncImpl().
|
static |
Identifies operations that are beneficial to sink into kernels.
These operations may not have side-effects, as otherwise sinking (and hence duplicating them) is not legal.
Definition at line 76 of file KernelOutlining.cpp.
References mlir::m_Constant(), and mlir::matchPattern().
|
static |
Return the provided KernelDim3 as an array of i32 constants if possible.
Definition at line 158 of file KernelOutlining.cpp.
References mlir::detail::DenseArrayAttrImpl< int32_t >::get(), mlir::Value::getContext(), mlir::m_ConstantInt(), mlir::matchPattern(), max(), mlir::gpu::KernelDim3::x, mlir::gpu::KernelDim3::y, and mlir::gpu::KernelDim3::z.
Referenced by outlineKernelFuncImpl().
|
static |
Outline the gpu.launch operation body into a kernel function.
Replace gpu.terminator operations by gpu.return in the generated function. Set block and grid size bounds if known.
Definition at line 178 of file KernelOutlining.cpp.
References mlir::Region::cloneInto(), mlir::OpBuilder::create(), mlir::detail::enumerate(), mlir::Block::erase(), mlir::Operation::erase(), mlir::Region::front(), mlir::get(), mlir::Block::getArgument(), mlir::Block::getOperations(), mlir::Block::getTerminator(), mlir::Builder::getUnitAttr(), mlir::getUsedValuesDefinedAbove(), injectGpuIndexOperations(), mlir::IRMapping::lookup(), mlir::IRMapping::map(), maybeConstantDimsAttr(), and mlir::Operation::setAttr().
Referenced by mlir::outlineKernelFunc().
1.9.1